THE PARALLEL

UNIVERSE

データ並列 C++ のリダクション操作
oneAPI Great Cross-Architecture Challenge の受賞者を発表oneAPI によるフーリエ相関アルゴリズムの実装

00001101
00001010
00001101
00001010
01001100
01101111
Issue
01110010
01100101
01101101
2021
00100000
01101000
01110001
01110011
01110101

44

目次

The Parallel Universe2

編集者からのメッセージインテル® Embree レイトレーシング ・ ライブラリーがアカデミー賞を受賞

3

oneAPI Great Cross-Architecture Challenge の受賞者を発表

5

データ並列 C++ のリダクション操作

9

oneAPI によるフーリエ相関アルゴリズムの実装

17

oneAPI を使用して有限差分法を高速化

25

oneAPI 仕様で提案されている高度なレイトレーシング API

33

インテル® oneTBB フローグラフと OpenVINO™ 推論エンジン

35

明示的なベクトル化を使用したスキャン操作の最適化 Inference

46

注目記事

Henry A. Gabb インテル コーポレーション シニア主席エンジニア

一般的なリダクション並列パターンの実装とチューニング

数行の DPC++ とインテル® oneMKL で複雑な数学演算を実行する

CUDA* ベースのステンシル計算を DPC++ へ移行

一度記述すれば複数のベンダーのアクセラレーターで動作する高忠実度のレイトレーシング

C++ でのディープラーニング ・ モデル間の依存関係の表現

インテル® AVX-512 SIMD 命令を活用してプリフィクス ・ サム計算を高速化

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

3

編集者からのメッセージ
Henry A. Gabb インテル コーポレーション シニア主席エンジニア
HPC と並列コンピューティング分野で長年の経験があり、 並列プログラミングに関する記事を多数出版しています。 『Developing Multithreaded Applications: A Platform Consistent Approach』の編集者 / 共著者で、インテルと Microsoft* による Universal Parallel Computing Research Centers のプログラム・マネージャーを務めました。

インテル® Embree レイトレーシング ・ライブラリーがアカデミー賞を受賞 !
インテル® oneAPI レンダリング・ツールキット （英語） の主要コンポーネントであるインテル® Embree レイトレー
シング ・ ライブラリーが映画芸術科学アカデミー （AMPAS） の科学技術賞 （英語） を受賞しました。

「過去 10 年間、 インテル® Embree レイトレーシング ・ ライブラリーは、 広範な研究論文によって裏付けされた
洗練されたオープンソース・コードを通じて、 業界をリードするハイパフォーマンスな CPU ベースのレイジオメトリー
交差フレームワークを提供してきました。 インテル® Embree レイトレーシング ・ ライブラリーは、 映画制作のレンダリングには欠かせないリソースとなっています。」

本号では、 インテル コーポレーションのアドバンスド ・ レンダリング & ビジュアライゼーション担当シニア ・ ディレク
ターの Jim Jeffers が「oneAPI 仕様で提案されている高度なレイトレーシング API」の概要を説明します。 また、インテルが欧州原子核研究機構 （CERN） およびアルゴンヌ国立研究所と共同で主催した oneAPI コーディング ・コンテスト 「The Great Cross-Architecture Challenge」 の受賞者を紹介します。
昨年から oneAPI （英語） 業界イニシアチブの勢いが増している中、 本号では oneAPI コードの開発とチュー
ニングに関する記事もいくつか掲載しています。 本号の注目記事 「データ並列 C++ のリダクション操作」 は、一般的なリダクション並列パターンの最適化について紹介する全 2 パートのシリーズの第 1 弾です。「oneAPIによるフーリエ相関アルゴリズムの実装」 は、 私が最近データ並列 C++ （DPC++） とインテル® oneMKL を実際に使用した経験を基に記事を執筆しました。 「oneAPI を使用して有限差分法を高速化」 は、 oneAPI
ツールを使用して CUDA* ベースのステンシルコードを DPC++ へ移行した、 ブラジルの SENAI/CIMATEC
Supercomputing Center による寄稿記事です。 「インテル® oneTBB フローグラフと OpenVINO™ 推論エンジン」 では、 軽量な C++ API を使用して、 複数のマシンラーニング ・ モデルを同じパイプラインで連携させる方法を紹介します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

4

本号の最後を締めくくるのは、 コードの現代化に関する記事です。 「明示的なベクトル化を使用したスキャン操作の最適化」 では、 OpenMP* ディレクティブとベクトル組込み関数を使用して、 プリフィクス ・ サム （prefix sum）の計算をチューニングします。
コードの現代化、 ビジュアル・コンピューティング、 データセンターとクラウド・コンピューティング、 データサイエンス、
システムと IoT 開発、 oneAPI を利用したヘテロジニアス並列コンピューティング向けのインテルのソリューションの
詳細は、 Tech.Decoded （英語） を参照してください。
Henry A. Gabb
2021 年 4 月

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

5

oneAPI Great Cross-Architecture
Challenge の受賞者を発表
インテルは、 欧州原子核研究機構 （CERN） およびアルゴンヌ国立研究所と共同で主催した oneAPI Great
Cross-Architecture Challenge（英語）の受賞者を発表しました。 このコンテストには、クロスアーキテクチャー、
マルチベンダー、 オープンアプローチである oneAPI の勢いが増していることを示すように、 5 大陸 52 カ国から応募がありました。 oneAPI とデータ並列 C++ （DPC++） を使用した、バイオインフォマティクス、暗号、データ分析、教育、 金融サービス、 ゲノム、 ヘルスケア、 画像処理、 数学、 分子動力学、 素粒子物理学、 レイトレーシングなどの分野のさまざまなアプリケーションが集まりました。
参加者には、 oneAPI を使用して、 CPU、 GPU、 FPGA、 その他のアクセラレーターを含むさまざまな XPU を最大限に活用する、 高速で効率良い将来性のあるヘテロジニアス ・アプリケーションを作成することが求められました。 インテル® oneAPI ツールキットと、 インテル® XPU でコードやワークロードをテストできるインテル®
DevCloud （英語） への無料アクセスが提供され、 参加者はインテル® DPC++ 互換性ツールを使用して既存の C/C++ または CUDA* アプリケーションを移行するか、 全く新しい oneAPI アプリケーションを作成するかを選択しました。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

6「このコンテストは、 oneAPI のオープンでクロスアーキテクチャーなプログラミング ・ モデルの使いやすさと選択の自由を如実に示しました。 oneAPI を使用することで、 参加者はさまざまな分野で実世界に影響を与えるアプリ
ケーションを素早く移行したり、 ゼロから開発できました。 世界中から寄せられた革新的で創造的な応募作品と、
oneAPI に対する好意的なフィードバックと採用の拡大に非常に感銘を受けています。」
インテル コーポレーション データセンター XPU 製品 & ソリューション担当副社長 Jeff McVeigh「Great Cross-Architecture Challenge の参加者は、 oneAPI の可能性を示しました。 oneAPI を使用して、 ヘテロ
ジニアス ・ ハードウェア ・アーキテクチャーに対応した多様なアプリケーションのコードを記述しました。 インテルの開発環境をオープンにすることで、 参加者に必要な環境を提供し、 このコンテストを通じて、 世界中の人々が最先端の技術に触れることができました。 コンテストの受賞者を CERN に迎えることを楽しみにしています。」
CERN openlab 最高技術責任者 Maria Girone 氏

インターンシップ賞を受賞した学生の応募作品は以下のとおりです。
•• Rafael Campos （英語） （ポルトガル）：最新の構文と最小限のプログラミング作業により、 OpenCL* アプリケー
ションを適応させることで、 oneAPI の高速で効率良い開発を実証しました。 その結果、 バイオインフォマティクス・
アプリケーション （特にエピスタシス検出） のパフォーマンスと電力効率が向上しました。
•• Andrew Pastrello（英語）（オーストラリア）：連星ブラックホールのインスパイラル・マージャー・リングダウン・シミュ
レーションによって生成される重力波からオーディオを合成するように音楽制作ツールを最適化して、 CUDA* コー
ドから DPC++ へ容易に移行できることを示しました。

CERN の特別ツアーに参加する機会を得たプロの開発者は以下のとおりです。
•• Ricardo Nobre （英語） （ポルトガル） ： インテル® DPC++ 互換性ツールを使用して、 CPU と GPU デバイスを活用して遺伝子型と表現型の間の新たな関連性を見つける CUDA* ベースのアプリケーションをシームレスに移行しました。 手作業でチューニングした CUDA* コードの 95% 以上が自動的に移行されました。
•• Zhen Ju （英語） （中国） ： CUDA* ベースのアプリケーションの移行と、 すべてのアーキテクチャーに対応するオー
プンなプログラミング ・ モデルの利点を紹介しました。 移行したアプリケーションは、 遺伝子データの冗長な配列をフィルタリングする、 効率良く正確なソリューションを提供しました。
•• Eugenio Marinelli （英語） （フランス） ： oneAPI のクロスアーキテクチャー ・ ライブラリーとツールを活用して、合成 DNA に格納されたデジタルデータを迅速かつ正確に解読する、 スケーラブルなヘテロジニアス並列処理ア
ルゴリズムを実装する新しいアプリケーションを効率良く開発しました。

参加者には、 サンプルコード、 デベロッパー ・ ガイド、 ウェビナー、 インテル® DevMesh （英語） コラボレーション ・
ポータルなどの追加の無料リソースへのアクセスも提供されました。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

7

このコンテストでは、 4 万ドル以上の賞金のほか、 CERN のツアー、 CERN openlab でのインターンシップ、 アルゴ
ンヌ国立研究所でのプロジェクトに参加する機会など、 一生に一度のチャンスが提供されました。 以下の 6 名の著名な審査員 （英語） により、 5 名の大賞受賞者が選ばれました。
•• CERN openlab 最高技術責任者 Maria Girone 氏•• ストックホルム大学物理学教授 Erik Lindahl 氏•• ブリストル大学ハイパフォーマンス ・コンピューティング教授 Simon McIntosh-Smith 氏•• Hewlett Packard Enterprise ディスティングイッシュド ・ テクノロジスト Heidi Poxon 氏•• アルゴンヌ国立研究所サイエンス ・ ディレクター Katherine Riley 氏•• Codeplay Software ディスティングイッシュド ・ エンジニア Michael Wong 氏

応募作品は、 革新性、 人類への影響、 クロスアーキテクチャー・コンピューティングの利用、 コーディングの専門性、プロジェクトの説明の質などが評価されました。 上位 5 名には、 以下の大賞賞品のいずれかが贈られました。
•• CERN の特別ツアー1、 または現金 5,000 ドル （3 名）•• CERN openlab での夏季インターンシップ （直接またはバーチャル）、 または現金 8,000 ドル （1 名）•• アルゴンヌ国立研究所での oneAPI 関連プロジェクトへの参加 （直接またはバーチャル）、 または現金 8,000 ドル （1 名）また、 大賞に加えて、 質の高い作品を応募した 20 名の参加者に 500 ドルの賞金が贈られました。
2019 年以降、 oneAPI エコシステムのサポートは着実に拡大しています。 60 以上の主要な研究機関、 企業、大学が oneAPI イニシアチブをサポートしており、 インテル® oneAPI ツールキットの成功事例は、 oneAPI エコシ
ステム ・ サポート （英語） と賛同者の声 （英語） サイトで紹介されています。 最新の oneAPI アプリケーション ・
カタログ （英語） には、 oneAPI を利用した 230 以上のアプリケーションの詳細が記載されています。 oneAPI
およびインテル® oneAPI ツールキットを使用したハイパフォーマンスなクロスアーキテクチャー ・アプリケーションの開発を支援するため、 以下のリソースが提供されています。
•• 重要な情報：ドキュメント （英語）、oneAPI プログラミング・ガイド、サンプルコード （英語）、無料のトレーニング （英語）。
• • サポート ： 無料のコミュニティー ・ フォーラム （英語） を利用して、 インテル® oneAPI ツールキットに関する
テクニカルサポートが得られます。 商用版 （英語） では、 直接インテルのエンジニアに問い合わせることができる優先サポート （英語） を利用できます。
•• トレーニング ： ウェビナー、 ワークショップ、 その他のラーニングパスなど、 無料のトレーニング （英語） が用意されています。 oneAPI を使用した HPC および AI ソリューションの高速化について支援が必要なお客様には、 特別な訓練を受けた企業グループが、 新しいインテル® oneAPI Technology Partner プログラム （英語） を介してコンサルティングを提供しています。

1

COVID-19 の渡航制限によります。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

インテル® DevCloud
デー タセンターか ら エッジ ま でさ まざ ま な
ワークロードに対応した開発サンドボックス
最新のインテル製ハードウェアとソフトウェアを搭載したクラスター上でワークロードを開発、 テスト、 実行
統合されたインテルの最適化フレームワーク、 ツール、 ライブラリー
を利用することで、 プロジェクトに必要なものがすべて揃います。

インテルのさまざまなハードウェアを評価
自然言語処理や時系列解析などの機能を備えた最先端のクラスターと、 エッジ・アクセラレーション・ハードウェアを利用して、 スキルを高め、さまざまな実験を行うことができます。

インテル® ソフトウェア開発ツールで開発
インテルから最新のコンパイラー、 パフォーマンス ・ ライブラリー、
ツールをダウンロード、 設定、 およびインストールすることなく、すぐにプロジェクトを開始できます。

ポピュラーな AI フレームワークを使用
トレ ーニングと推 論をすぐに開始できるインテル® Optimized AI
Frameworks を利用してアルゴリズムとアプリケーションを高速化できます。

今すぐ開始 （英語） >コンパイラーの最適化に関する詳細は、最適化に関する注意事項 (software.intel.com/articles/optimization-notice#opt-jp) を参照してください。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。
Intel、インテル、Intel
ロゴ、その他のインテルの名称やロゴは、Intel
Corporation またはその子会社の商標です。
© Intel Corporation

8

The Parallel Universe9

データ並列 C++ のリダクション操作
一般的なリダクション並列パターンの実装とチューニング

Ramesh Peri インテル ・ コーポレーション シニア主席エンジニア

リダクション操作の概要
リダクションは、 並列プログラミングにおける一般的な操作で、 配列の要素を 1 つの結果にレデュースします。
例えば、 大きな配列の要素の総和を求める場合 （図 1）、 操作を並列に実行するには、 部分和を計算し、 それを順次組み合わせて最終的な結果を算出する必要があります。 リダクション演算子 （総和、 最小値、 最大値、
最小位置、 最大位置など） は結合的であり、 多くの場合、 可換的です。 リダクションの実装方法は多数あり、そのパフォーマンスはプロセッサー ・アーキテクチャーに依存します。 この記事では、 データ並列 C++ （DPC++） でリダクションを表現するいくつかの方法を紹介し、 それぞれのパフォーマンスへの影響について考察します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

10

図 1. 総和リダクションの図

グローバルアトミックを使用した DPC++ のリダクション
以下の実装では、DPC++ カーネルの各ワークアイテムが入力配列の要素を処理して、グローバル変数をアトミックに更新します。
void reductionAtomics1(sycl::queue &q,
sycl::buffer<int> inbuf,
int &res,int size) {const size_t data_size = inbuf.get_size() / sizeof(int);
int num_work_items = data_size;
sycl::buffer<int> sum_buf(&res, 1);
q.submit([&](auto &h) {sycl::accessor buf_acc(inbuf, h, sycl::read_only);
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::noinit);});h.parallel_for(num_work_items, [=](auto index) {size_t glob_id = index[0];auto v = sycl::ONEAPI::atomic_ref<int,sycl::ONEAPI::memory_order::relaxed,
sycl::ONEAPI::memory_scope::device,
sycl::access::address_space::global_space>(sum_acc[0]);v.fetch_add(buf_acc[glob_id]);});

コンパイラーが生成するスレッドの数 （コンパイラーが選択するターゲットデバイスのデフォルトのワークグループとサブグループのサイズに依存） によっては、 単一のグローバル変数である sum_buf へのアクセス競合が頻繁に発生します。 一般的に、 このようなソリューションのパフォーマンスはあまり高くありません。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

11

グローバルアトミックの競合の軽減
グローバル変数の更新時の競合を軽減する方法として、 この変数にアクセスするスレッドの数を減らすことが挙げられます。 これは、 各ワークアイテムが配列の複数の要素を処理し、 要素のチャンクに対してローカルなリダクションを実行した後、 グローバルアトミック更新を実行することで実現できます （図 2）。

図 2. 各ワークアイテムは部分和を計算

以下は、 この実装のコードです。
void reductionAtomics2(sycl::queue &q,
sycl::buffer<int> inbuf,
int &res) {const size_t data_size = inbuf.get_size() / sizeof(int);
sycl::buffer<int> sum_buf(&res, 1);
int num_work_items =q.get_device().get_info<sycl::info::device::max_compute_units>();int BATCH = (data_size + num_work_items - 1) / num_work_items;
q.submit([&](auto &h) {sycl::accessor buf_acc(inbuf, h, sycl::read_only);
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::noinit);}h.parallel_for(num_processing_elements, [=](auto index) {size_t glob_id = index[0];size_t start = glob_id * BATCH;
size_t end = (glob_id + 1) * BATCH;
if (end > N)
end = N;int sum = 0;
for (size_t i = start; i < end; i++)sum += buf_acc[i];auto v = sycl::ONEAPI::atomic_ref<int,sycl::ONEAPI::memory_order::relaxed,
sycl::ONEAPI::memory_scope::device,
sycl::access::address_space::global_space>(sum_acc[0]);v.fetch_add(sum);});});

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe12

この実装は、 各ワークアイテムが連続するメモリー位置にアクセスするためあまり効率的ではなく、 コンパイラーは非効率なコードを生成します。 DPC++ コンパイラーは各ワークアイテムをベクトルレーンのように扱うため、 連続するメモリー位置にアクセスするワークアイテムは非効率的になります。

ワークアイテムによる効率良いメモリーアクセス
DPC++ コンパイラーは、 各ワークアイテムをプロセッサーのベクトルレーンにマップするため、 ワークアイテムがプロセッサーのベクトル長よりも大きなストライドでメモリー位置にアクセスする場合、 コンパイラーは効率的なコードを生成できます。 このアクセスパターンを図 3 に示します。 ラベル wi-1 … wi-n は n 個のワークアイテムを示します。

図 3. 各ワークアイテムは連続しないメモリー位置から部分和を計算

以下は、 各ワークアイテムが入力ベクトルの複数の要素をインターリーブして処理するカーネルのコードです。
void reductionAtomics3(sycl::queue &q,
sycl::buffer<int> inbuf,
int &res) {const size_t data_size = inbuf.get_size() / sizeof(int);
sycl::buffer<int> sum_buf(&res, 1);
int num_work_items =q.get_device().get_info<sycl::info::device::max_compute_units>() *q.get_device().get_info<sycl::info::device::native_vector_width_int>();q.submit([&](auto &h) {sycl::accessor buf_acc(inbuf, h, sycl::read_only);
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::noinit);}});h.parallel_for(num_work_items, [=](auto index) {size_t glob_id = index[0];int sum = 0;
for (size_t i = glob_id; i < data_size; i += num_work_items)sum += buf_acc[i];auto v = sycl::ONEAPI::atomic_ref<int,sycl::ONEAPI::memory_order::relaxed,
sycl::ONEAPI::memory_scope::device,
sycl::access::address_space::global_space>(sum_acc[0]);v.fetch_add(sum);});

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

13

ワークアイテム数の選択は重要です。 処理要素の数と処理要素に最適なベクトル幅の積になるように選択します。
あるプラットフォームでは十分であっても、 サポートするスレッド数が処理要素の数よりもはるかに多いプラットフォームでは十分ではない場合があるため、 ワークアイテム数はターゲット ・プラットフォームに応じて慎重に選択する必要があります。

ツリー ・ リダクション
ツリー ・ リダクションは、 カーネル内の各ワークアイテムが、 隣接する要素にリダクション演算子を適用し、 複数のレベルの中間結果を生成するポピュラーな手法です。 これは、 図 4 に示すように、 スレッドのスケジュールと同期が高効率であることから、 ハードウェア ・ サポートを利用してワークグループ内で適用できます。

図 4. ツリー ・ リダクション

ワークグループのサイズはかなり小さく （約 256 または 512）、 ハードウェア ・ デバイスに依存しますが、 レデュー
スする要素の数は、 ワークグループのサイズよりも 1 ケタ大きくなります。 つまり、 各ワークグループで実行される
リダクション操作は、 さらにリダクションが必要な中間結果を生成します。 これに対応するには、 2 つの方法があります。
1. 以下に示すように、 各ワークグループにおいて中間結果のグローバルアトミック加算を呼び出します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

void buf2ﬁnalReduction(sycl::queue &q,
sycl::buffer<int> inbuf,
int &res) {const size_t num_work_items = inbuf.get_size() / sizeof(int);
int work_group_size =q.get_device().get_info<sycl::info::device::max_work_group_size>();sycl::buffer<int> sum_buf(&res, sizeof(int));
q.submit([&](auto &h) {sycl::accessor buf_acc(inbuf, h, sycl::read_only);
sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::noinit);
sycl::accessor<int, 1, sycl::access::mode::read_write,
sycl::access::target::local> scratch(work_group_size, h);h.parallel_for(sycl::nd_range<1>(num_work_items, work_group_size),[=](sycl::nd_item<1> item) {size_t global_id = item.get_global_id(0);int local_id = item.get_local_id(0);int group_id = item.get_group(0);int sum = 0;
if (global_id < data_size)
scratch[local_id] = buf_acc[global_id];
elsescratch[local_id] = 0;
for (int i = work_group_size / 2; i > 0; i >>= 1)
item.barrier(sycl::access::fence_space::local_space);if (local_id < i)scratch[local_id] += scratch[local_id + i];}
if (local_id == 0) {auto v = sycl::ONEAPI::atomic_ref<int,sycl::ONEAPI::memory_order::relaxed,sycl::ONEAPI::memory_scope::device,
sycl::access::address_space::global_space>(sum_acc[0]);v.fetch_add(scratch[0]);
}
});

});

2. ワークグループによって生成された中間値のリストに対して、 同じリダクション ・ カーネルを再度呼び出して、別の中間値のセットを生成します。 そして、 中間結果の数が十分に少なくなった時点で、 最終結果を得るため
カーネルを呼び出す代わりに、 グローバルアトミック更新を使用できます。
この手法を実装するカーネルを以下に示します。 ここでは、 カーネルの出力は各ワークグループが出力バッファーに格納する別の値のセットです。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

14

The Parallel Universe

15

void buf2bufReduction(sycl::queue &q,
sycl::buffer<int> inbuf,
sycl::buffer<int> outbuf) {const size_t data_size = inbuf.get_size() / sizeof(int);
int num_work_items = data_size;
int work_group_size =q.get_device().get_info<sycl::info::device::max_work_group_size>();q.submit([&](auto &h) {sycl::accessor inbuf_acc(inbuf, h, sycl::read_only);
sycl::accessor outbuf_acc(outbuf, h, sycl::write_only, sycl::noinit);
sycl::accessor<int, 1, sycl::access::mode::read_write,
sycl::access::target::local> scratch(work_group_size, h);h.parallel_for(sycl::nd_range<1>(num_work_items, work_group_size),[=](sycl::nd_item<1> item) {size_t global_id = item.get_global_id(0);int local_id = item.get_local_id(0);int group_id = item.get_group(0);int sum = 0;
if (global_id < data_size)
scratch[local_id] = buf_acc[global_id];
elsescratch[local_id] = 0;
for (int i = work_group_size / 2; i > 0; i >>= 1)
item.barrier(sycl::access::fence_space::local_space);if (local_id < i)scratch[local_id] += scratch[local_id + i];

}

});

}
if (local_id == 0)
outbuf_acc[group_id] = scratch[0];});前の 2 つのカーネルを以下のように呼び出すことで、 最終的な結果が得られます。 ここで、 X は最終的なアトミッ
クベースのリダクションが呼び出される中間結果のサイズであり、 この値はグローバル ・アトミック ・ リダクションの効率に基づいて選択されます。
int work_group_size =q.get_device().get_info<sycl::info::device::max_work_group_size>();int ibufsize = (size + work_group_size - 1) / work_group_size;
sycl::buffer<int> *prev = &buf;
while (ibufsize > X) {int *bufptr=(int *)malloc(sizeof(int) * ibufsize);
sycl::buffer<int> *cur = new sycl::buffer<int>(bufptr, ibufsize);
ibufsize = (ibufsize + work_group_size-1) / work_group_size;
buf2bufReduction(q, *prev, *cur);prev = cur;}buf2ﬁnalReduction(q, *prev, res);

DPC++ のビルトイン ・ リダクション演算子
DPC++ は、 以下のカーネルで使用されている総和リダクション演算子を提供しています。 この場合、 コンパイラー
はプラットフォームにとって最も効率良い実装を選択します。 通常は、 ビルトイン ・ リダクション演算子を使用することを推奨します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

16

void builtinReduction(sycl::queue &q,
sycl::buffer<int> inbuf,
int &res) {sycl::buffer<int> sum_buf(&res,1);q.submit([&](auto &h) {sycl::accessor buf_acc(inbuf, h, sycl::read_only);
sycl::accessor sum_acc(sum_buf, h, sycl::read_write);
auto sumr = sycl::ONEAPI::reduction(sum_acc, sycl::ONEAPI::plus<>());
h.parallel_for(sycl::nd_range<1>{data_size, 256}, sumr,[=](sycl::nd_item<1> item, auto &sumr_arg) {int glob_id = item.get_global_id(0);sumr_arg += buf_acc[glob_id];
});
});
}

まとめ
リダクションは、 多くのアプリケーションで使用される一般的な並列プログラミング・パターンです。 この記事では、
DPC++ でこの操作を実装するいくつかの方法を紹介しました。 これらの実装のパフォーマンスは、 コンパイラーやターゲット ・プラットフォームによって大きく異なります。 次回の記事では、 これらのカーネルの CPU および GPU
上でのパフォーマンスを調査する予定です。

関連情報•• oneAPI （英語）•• インテル® oneAPI ツールキット•• 『Data Parallel C++: Mastering DPC++ for Programming of Heterogeneous Systems using C++ andSYCL（データ並列 C++: C++ と SYCL* を使用したヘテロジニアス・システムのプログラミングで DPC++ を使いこなす）』（英語）『oneAPI GPU 最 適 化 ガ イド』 を入手して、 並列化からカーネル、 メモリーに至るまで最高のパフォーマンスを実現してください。
インテル® oneAPI ベース・ツールキットは、
さまざ まな アー キテクチャーに 対応した
ハイパフォーマンスなアプリケーションやソリューションの開発を可能にします。

詳細 （英語）

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

Github* に ア ク セ ス し てoneAPI サ ン プ ル コ ー ド を取得しましょう。
詳細

Github* にアクセス （英語）The Parallel Universe

17

oneAPI によるフーリエ相関アルゴリズムの実装
数行の DPC++ とインテル® oneMKL で複雑な数学演算を実行する

Henry A Gabb インテル コーポレーション シニア主席エンジニア兼 The Parallel Universe 編集長[ 編集者注 : この記事は、 『oneAPI GPU 最適化ガイド』 のために作成した例を参考にしています。 完全なソースコードとビルドおよび実行手順は、 oneAPI-samples リポジトリー （英語） で入手できます。 このガイドとリポジトリーは、どちらも oneAPI 開発者にとって優れたリソースです。 ]相互相関の概要
個々のインテル® oneAPI マス ・ カーネル ・ ライブラリー （インテル® oneMKL） カーネル関数をアクセラレーターにオフロードするのは簡単ですが、 ここでは複数のカーネル関数を必要とするより複雑な数学演算である相互相

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

18

関を見てみましょう。 相互相関には、 2 つの 1 次元信号の類似性を測定したり、 類似する画像の最適なオーバーレイ変換を見つけたり、 ボリューム医療画像をセグメント化するなど、 多くの用途があります。 1 と 0 のベクトルで表される次のような単純な信号について考えてみます。
信号 1: 0 0 0 0 0 1 1 0
信号 2: 0 0 1 1 0 0 0 0

信号は互いに循環シフトしたものとして扱われるので、 2 番目の信号を 1 番目の信号に対して 3 要素シフトさせ
ると、 最大相関スコアは 2 になります。
信号 1: 0 0 0 0 0 1 1 0
信号 2:
0 0 1 1 0 0 0 0
相関 : (1 * 1) + (1 * 1) = 2

2 または 4 要素シフトした場合、 相関スコアは 1 になります。 それ以外のシフトでは相関スコアは 0 です。
これは次のように計算されます。

ここで、 N は信号ベクトルの要素数、 i は sig1 に対する sig2 のシフト量です。
実際の信号にはより多くのデータ （およびノイズ） が含まれていますが、 1 次元の信号を一致させる場合でも、2 次元の画像をオーバーレイする場合でも、 3 次元のボリューム画像のレジストレーションを行う場合でも原理は同じです。 目的は、 相関を最大にする変換を見つけることです。 しかし、 上記の総当たり法の総和は、 N シフト
ごとに N 個の乗算と加算を必要とします。 1 次元、 2 次元、 3 次元では、 この問題はそれぞれ O(N2)、 O(N3)、O(N4) となります。
フーリエ相関アルゴリズムは、 フーリエ変換の複雑さが O(N log N) であることを利用して、 この計算を効率良く実行できます。
corr = IDFT(DFT(sig1) * CONJG(DFT(sig2)))

ここで、 DFT は離散フーリエ変換、 IDFT は逆 DFT、 CONJG は複素共役です。 フーリエ相関アルゴリズムは、最適化された順変換 / 逆変換と複素共役乗算関数を含むインテル® oneMKL を使用して構成できます。そのため、すべての計算をアクセラレーター ・ デバイス上で実行できます。

インテル® oneMKL によるテストデータの生成
多くのアプリケーションでは、 最終的な相関結果のみが重要となるため、 デバイスからホストへの転送はこの結果のみとなります。 この例では、 デバイス上で 2 つの人工信号を作成してインプレース変換し、 相関させます。
ホストは最終結果を取得し、 最適な変換結果と相関スコアを報告します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

19

一般通念では、 バッファリングはホストとデバイス間のデータ移動を明示的に制御できるため、 最高のパフォーマンスを達成できると考えられています。 この仮説を検証するため、 2 つの入力信号を生成してみます。
// 信号データ用のバッファーを作成。 これはデバイス上でのみ使用される。
sycl::buffer<ﬂoat> sig1_buf{N + 2};
sycl::buffer<ﬂoat> sig2_buf{N + 2};// ( デバイスで計算され、 ホストで使用される ) 相関結果を格納するコンテナーを宣言
std::vector<ﬂoat> corr(N + 2);

ニューラル ・ ネットワークのトレーニング中に過剰適合を防ぐため、 画像に視覚的効果を加えるため、 あるいは最適ではない検出器から得た信号の検出性を向上させるためなど、 信号にランダムノイズを加えることはよくあります。
バッファーは、 インテル® oneMKL の基本的な乱数ジェネレーターを使用して、 ランダムノイズで初期化されます。
// SYCL* キューを初期化
sycl::queue Q(sycl::default_selector{});// 相関結果の更新をトリガーする新しいスコープを開く{
sycl::buffer<ﬂoat> corr_buf(corr);// 入力信号を人工データで初期化
std::uint32_t seed = (unsigned)time(NULL); // RNG シード値を取得
oneapi::mkl::rng::mcg31m1 engine(Q, seed); // RNG エンジンを初期化// RNG 分布を設定
oneapi::mkl::rng::uniform<ﬂoat, oneapi::mkl::rng::uniform_method::standard>
rng_distribution(-0.00005, 0.00005);
oneapi::mkl::rng::generate(rng_distribution, engine, N, sig1_buf); // ノイズ
oneapi::mkl::rng::generate(rng_distribution, engine, N, sig2_buf);

新しいスコープが開かれ、 相関結果用に corr_buf というバッファーが宣言されています。 このバッファーがスコープ外になると、 ホスト上で corr が更新されます。
上記の例と同様に、 各バッファーの反対側の端に人工信号が配置されます。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

20

Q.submit([&](sycl::handler &h) {sycl::accessor sig1_acc{sig1_buf, h, sycl::write_only};
sycl::accessor sig2_acc{sig2_buf, h, sycl::write_only};h.single_task<>([=]() {sig1_acc[N - N / 4 - 1] = 1.0;
sig1_acc[N - N / 4] = 1.0;
sig1_acc[N - N / 4 + 1] = 1.0; // 信号
sig2_acc[N / 4 - 1] = 1.0;
sig2_acc[N / 4] = 1.0;
sig2_acc[N / 4 + 1] = 1.0;});}); // 信号の初期化終了

明示的バッファリングによる 1 次元フーリエ相関の実装
信号の準備ができたので、 インテル® oneMKL の DFT 機能を使用して信号を変換してみましょう。
// FFT 記述子を初期化
oneapi::mkl::dft::descriptor<oneapi::mkl::dft::precision::SINGLE,
oneapi::mkl::dft::domain::REAL>
transform_plan(N);transform_plan.commit(Q);// 実数配列の順変換を実行
oneapi::mkl::dft::compute_forward(transform_plan, sig1_buf);
oneapi::mkl::dft::compute_forward(transform_plan, sig2_buf);

単精度の実数から複素数への順変換が SYCL* キューにコミットされ、 その後、 両バッファーのデータに対して
インプレース DFT が実行されます。DFT(sig1) の結果は、CONJG(DFT(sig2)) と掛け合わせる必要があります。
これは、 ハンドコーディングしたデータ並列 C++ （DPC++） で行います。
Q.submit([&](sycl::handler &h){
sycl::accessor sig1_acc{sig1_buf, h, sycl::read_only};
sycl::accessor sig2_acc{sig2_buf, h, sycl::read_only};
sycl::accessor corr_acc{corr_buf, h, sycl::write_only};h.parallel_for<>(sycl::range<1>{N/2}, [=](auto idx){
corr_acc[idx*2+0] = sig1_acc[idx*2+0] * sig2_acc[idx*2+0] +sig1_acc[idx*2+1] * sig2_acc[idx*2+1];
corr_acc[idx*2+1] = sig1_acc[idx*2+1] * sig2_acc[idx*2+0] sig1_acc[idx*2+0] * sig2_acc[idx*2+1];});});

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

21

しかし、 この基本的な実装では、 最適なクロスアーキテクチャー ・ パフォーマンスを達成できません。 幸い、インテル® oneMKL は、 このステップに使用できる便利な関数、 oneapi::mkl::vm::mulbyconj を提供しています。 mulbyconj 関数 （英語） は std::complex<float> 入力を想定していますが、 バッファーは float
データ型として初期化されています。 順変換後に複素数データが入っている場合でも、 バッファーを再キャストする
必要があります。
// 計算 : DFT(sig1) * CONJG(DFT(sig2))
auto sig1_buf_cplx =sig1_buf.template reinterpret<std::complex<ﬂoat>, 1>((N + 2) / 2);
auto sig2_buf_cplx =sig2_buf.template reinterpret<std::complex<ﬂoat>, 1>((N + 2) / 2);
auto corr_buf_cplx =corr_buf.template reinterpret<std::complex<ﬂoat>, 1>((N + 2) / 2);
oneapi::mkl::vm::mulbyconj(Q, N / 2, sig1_buf_cplx, sig2_buf_cplx,
corr_buf_cplx);

IDFT ステップで計算が完了します。
// 複素相関配列の逆変換を実行
oneapi::mkl::dft::compute_backward(transform_plan, corr_buf);} // 相関結果を保持するバッファーがスコープ外になったため、 ホストコンテナーの更新が必要

この例の冒頭で開いていたスコープを閉じると、 相関結果を保持するバッファーがスコープ外となり、 ホスト
コンテナーの更新が必要になります。 これが、 ホストとデバイス間の唯一のデータ転送です。
ソースコード一式 （fcorr_1d_buffers.cpp） は、 oneAPI-samples リポジトリーにあります。

USM による 1 次元フーリエ相関の実装
フーリエ相関アルゴリズムを USM （統合共有メモリー） を使用して再実装し、 明示的なバッファリングと比較してみましょう。 ここでは、 2 つの実装の違いのみを紹介します。 まず、 信号配列と相関配列を USM に割り当て、
人工データで初期化します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

22// 信号配列と相関配列を初期化
auto sig1 = sycl::malloc_shared<ﬂoat>(N + 2, sycl_device, sycl_context);
auto sig2 = sycl::malloc_shared<ﬂoat>(N + 2, sycl_device, sycl_context);
auto corr = sycl::malloc_shared<ﬂoat>(N + 2, sycl_device, sycl_context);// 入力信号を人工データで初期化
std::uint32_t seed = (unsigned)time(NULL); // RNG シード値を取得
oneapi::mkl::rng::mcg31m1 engine(Q, seed); // RNG エンジンを初期化// RNG 分布を設定
oneapi::mkl::rng::uniform<ﬂoat, oneapi::mkl::rng::uniform_method::standard>
rng_distribution(-0.00005, 0.00005);// 警告 : これらの文はデバイス上で実行
auto evt1 =oneapi::mkl::rng::generate(rng_distribution, engine, N, sig1); // ノイズ
auto evt2 = oneapi::mkl::rng::generate(rng_distribution, engine, N, sig2);
evt1.wait();evt2.wait();// 警告 : これらの文はホスト上で実行されるため、// デバイス上で sig1 と sig2 の更新が必要
sig1[N - N / 4 - 1] = 1.0;
sig1[N - N / 4] = 1.0;
sig1[N - N / 4 + 1] = 1.0; // 信号
sig2[N / 4 - 1] = 1.0;
sig2[N / 4] = 1.0;
sig2[N / 4 + 1] = 1.0;

残りの実装はほとんど同じですが、 SYCL* バッファーの代わりに USM へのポインターがインテル® oneMKL 関数に渡されます。
// 実数配列の順変換を実行
evt1 = oneapi::mkl::dft::compute_forward(transform_plan, sig1);
evt2 = oneapi::mkl::dft::compute_forward(transform_plan, sig2);// 計算 : DFT(sig1) * CONJG(DFT(sig2))
oneapi::mkl::vm::mulbyconj(
Q, N / 2, reinterpret_cast<std::complex<ﬂoat> *>(sig1),
reinterpret_cast<std::complex<ﬂoat> *>(sig2),
reinterpret_cast<std::complex<ﬂoat> *>(corr), {evt1, evt2}).wait();// 複素相関配列の逆変換を実行
oneapi::mkl::dft::compute_backward(transform_plan, corr).wait();

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

23

また、 割り当てたメモリーを解放する必要があります。
sycl::free(sig1, sycl_context);
sycl::free(sig2, sycl_context);
sycl::free(corr, sycl_context);

USM 実装は、 より親しみやすい構文になっています。 また、 DPC++ ランタイムによって処理される暗黙のデータ転送に依存しているため、 概念的にもよりシンプルです。 しかし、 プログラマーのミスがパフォーマンスを低下させることは避けられません。
前述のコード例にある警告メッセージに注目してください。 インテル® oneMKL の乱数生成エンジンはデバイス上で初期化されるため、 sig1 と sig2 はデバイス上のランダムノイズで初期化されます。 しかし、 人工信号を追加するコードはホスト上で実行されるため、 DPC++ ランタイムはホストとデバイスのデータの整合性を取る必要があります。 フーリエ相関で使用される信号は通常大きいため、 特に 3D イメージング・アプリケーションでは、不要なデータ転送はパフォーマンスを低下させます。
信号データをデバイス上で直接更新することで、 データの一貫性を保持し、 不要なデータ転送を回避できます。
Q.single_task<>([=]() {sig1[N - N / 4 - 1] = 1.0;
sig1[N - N / 4] = 1.0;
sig1[N - N / 4 + 1] = 1.0; // 信号
sig2[N / 4 - 1] = 1.0;
sig2[N / 4] = 1.0;
sig2[N / 4 + 1] = 1.0;}).wait();

明示的バッファリング実装と USM 実装のパフォーマンスは同等であり、 DPC++ ランタイムは （プログラマーがデータの一貫性に注意を払えば） 不必要なデータ転送を回避できることを示しています。
ソースコード一式 （fcorr_1d_usm.cpp） は、 oneAPI-samples リポジトリーにあります。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

24

まとめ
最終的に相関値が最大となる場所を探すワークはホスト側で行います。
// 相関値が最大となるシフトを見つける
ﬂoat max_corr = 0.0;
int shift = 0;
for (unsigned int idx = 0; idx < N; idx++) {if (corr[idx] > max_corr) {max_corr = corr[idx];
shift = idx;}}
shift =
(shift > N / 2) ? shift - N : shift;// 信号は互いに循環シフトしたものとして扱う
std::cout << "Shift the second signal " << shift<< " elements relative to the ﬁrst signal to get a maximum, ""normalized correlation score of "<< max_corr / N << "."<< std::endl;

特に入力データが大きい場合は、 この計算をデバイス上で行った方が良いでしょう。 幸い、 MAXLOC リダクショ
ンは DPC++ で実装できる一般的な並列パターンです。 これは皆さんの練習問題として残します。 データ並列 C++（英語） の図 14-11 に開始点として役立つ例があるので参考にしてください。 また、oneAPI-samples リポジトリーにある USM の例では、 MAXLOC リダクションがすでに実装されており、 すべての計算がデバイス上で行われるようになっています。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe25

oneAPI を使用して有限差分法を高速化
CUDA* ベースのステンシル計算を
DPC++ へ移行

Clícia Pinto インテル コーポレーション テクニカルリーダー兼パフォーマンス ・ エンジニア
Lucas Batista SENAI CIMATEC Supercomputing Center HPC 開発者
Pedro de Santana SENAI CIMATEC Supercomputing Center HPC 開発者
Georgina González SENAI CIMATEC Supercomputing Center HPC 開発者

リバース ・ タイム ・ マイグレーション （RTM） では、 有限差分法 （FD） を使用して音響波動方程式の近似値を計算します。 これは、 RTM アプリケーションの計算上のボトルネックとなっており、 炭化水素探査のリソースを割り当てる際に、 タイムリーな結果と効率を保証するため最適化する必要があります。 この記事では、 インテル®
DPC++ 互換性ツールを使用して、 CUDA* ベースの RTM コードをデータ並列 C++ （DPC++） に移行した経験について説明します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

26

RTM の概要と入力データ
いくつかの地震探査法では、 FD スキームにおいて、 波動方程式の数値解としてステンシルが適用されています。
石油ガス産業で地下構造の画像を生成するために広く使用されている RTM でもこの方法を採用しています。
この方法には利点がありますが、 2 つの大きな計算上のボトルネックがあります。 伝播ステップで浮動小数点演算が多いことと、 メインメモリーへの波動場の保存が困難なことです。 これらのボトルネックの影響を軽減するため、 エンジニアはタスクの並列性と計算リソースの最適化の両方を追求し、 さまざまなアクセラレーター
で実行可能なソリューションを設計しています。 この方法の最適化は、 井戸掘削におけるエラーの可能性を軽減させるため、 探査地球物理学に大きな経済的利点をもたらします。 Claerbout1 が提案したように、 RTM
アルゴリズムは通常、 前方時間伝播、 後方伝播、 および画像条件の相互相関を行います。 RTM アルゴリズ
ムのフローチャートでは、 これらのステップをホストとデバイスの通信とともに示しています （図 1）。

図 1. 簡易 2 次元 RTM フローチャート

図 2 は、 2 次元の CUDA* ベースの実装におけるステンシル計算を示しています。 order は FD スキームの次数、nx と nz は 2 次元音響速度モデルを表す入力行列のサイズ、 p はソース / レシーバーの波動場、 cx と cy はそれぞれ FD 係数の x 軸と y 軸、 l は外挿された波動場を表しています。 RTM アルゴリズム全体では、 ベクト
ル P は異なる時間ステップにおける圧力ポイントの状態を格納しています。 RTM は、 計算量とデータ量が多い
ため、 専用の処理ユニットによる高速化に適しています。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

27

Input: order, nx, nz,*p, *l, *cx, *cy
Output: *lAssignments: h_order ← order/2i ← h_order + blockId.x * blockDim.x + threadId.x
j ← h_order + blockId.y * blockDim.y + threadId.y
mult ← i * nz
if < nx – h_order thenif < nz - h_order thenfor k = 0 to h_order do
aux = k - h_orderaux += p[mult + j+aux] * cz[k]aux += p[(i+aux)*nz + j] * cx[k]end forl[mult + j] = accz + accx
accz = 0
accx = 0
end ifend if
図 2. RTM ステンシル計算アルゴリズム

リファレンス RTM から oneAPI への移行
インテル® DPC++ 互換性ツールは、 CUDA* 言語カーネルやライブラリーの API 呼び出しを DPC++ に移行するのを支援します。 通常、 CUDA* ソースコードの 80 ～ 90% は自動的に移行されるため、 このプロセスは 「準備」、「移行」、 「レビュー」 の 3 つの段階で構成されます。 準備段階では、 ソースコードを移行ツールに適応させることを目指します。 この段階では、 すべての CUDA* ヘッダーファイルがデフォルトの場所、 または --cuda-includepath=<path/to/cuda/include> オプションで指定したカスタムの場所にあることを確認する必要があります。 移行段階では、 インテル® DPC++ 互換性ツールがオリジナルのアプリケーションを入力として受け取り、注釈付きの DPC++ コードを生成します。 レビュー段階では、 自動変換されたコードを確認し、 移行されなかった
コードを手動で変換するための注釈をレビューし、 コードの改善の可能性を探ります。
最初の移行では、 図 3 に示すように、 インテル® DPC++ 互換性ツールによって、 CUDA* のメモリーコピー API
呼び出しが sycl::queue.memcpy() に移行されることを確認しました。 正常に動作し、 エラーのない DPC++
ソースコードが得られても、 明示的なメモリー管理では最高のパフォーマンスが得られない可能性があります。
メモリー管理の改善について調査するため、 移行後のソースコードを手動で変更し、 各データ ・ オブジェクトで
SYCL* のバッファーとアクセサーを使用するようにしました。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

28

dpct::device_ext &dev_ct1 = dpct::get_current_device();sycl::queue &q_ct1 = dev_ct1.default_queue();q_ct1.memcpy(d_p, p[0], mtxBufferLength).wait();q_ct1.memcpy(d_pp, pp[0], mtxBufferLength).wait();q_ct1.memcpy(d_v2, v2[0], mtxBufferLength).wait();q_ct1.memcpy(d_coefs_x, coefs_x, coefsBufferLength).wait();q_ct1.memcpy(d_coefs_z, coefs_z, coefsBufferLength).wait();q_ct1.memcpy(d_taperx, taperx, brdBufferLength).wait();q_ct1.memcpy(d_taperz, taperz, brdBufferLength).wait();
図 3. 移行後のメモリー管理

void fd_forward(int order, ﬂoat **p, ﬂoat **pp, ﬂoat **v2, int nz, int nx, int nt, int
is, int sz, int *sx, ﬂoat *srce, int propag){
dpct::device_ext &dev_ct1 = dpct::get_current_device();sycl::queue &q_ct1 = dev_ct1.default_queue();sycl::range<3> dimGrid(1, gridz, gridx);
sycl::range<3> dimGridTaper(1, gridBorder_z, gridx);
sycl::range<3> dimGridSingle(1, 1, 1);
sycl::range<3> dimGridUpb(1, 1, gridx);
sycl::range<3> dimBlock(1, sizeblock, sizeblock);{sycl::buffer<ﬂoat, 1> *b_p = new sycl::buffer<ﬂoat, 1>(p[0], sycl::range<1>(nxe*nze));sycl::buffer<ﬂoat, 1> *b_pp = new sycl::buffer<ﬂoat, 1>(pp[0],sycl::range<1>(nxe*nze));sycl::buffer<ﬂoat, 1> b_v2(v2[0], sycl::range<1>(nxe*nze));sycl::buffer<ﬂoat, 1> b_coefs_x(coefs_x, sycl::range<1>(order+1));
sycl::buffer<ﬂoat, 1> b_coefs_z(coefs_z, sycl::range<1>(order+1));
sycl::buffer<ﬂoat, 1> b_taperx(taper_x, sycl::range<1>(nxb));
sycl::buffer<ﬂoat, 1> b_taperz(taper_z, sycl::range<1>(nxb));
sycl::buffer<ﬂoat, 1> *b_swap;
for (int it = 0; it < nt; it++){b_swap = b_pp;
b_pp = b_p;b_p = b_swap;
kernel_tapper()kernel_lap()
kernel_time()
kernel_scr()}}}
図 4. DPC++ 前方伝播 (簡易バージョン)

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe29

図 4 は、 DPC++ に移行された前方時間伝播を行う関数です。 ここで、 nt は波動方程式のモデル化に必要な時間ステップの数です。 このアルゴリズムでは、 データ ・ オブジェクトが、 デバイスのメモリーの制御と変更に使用されるバッファーとして定義されています。 後方伝播も同じ構造になっています。 前方関数と後方関数はどちらも、
GPU 実行を処理するためカーネルを呼び出します。

q_ct1.submit([&](sycl::handler &cgh) {auto acc_pr = b_pr->get_access<sycl::access::mode::read_write>(cgh);auto d_laplace_ct4 = d_laplace;
auto acc_coefs_x = b_coefs_x.get_access<sycl::access::mode::read>(cgh);auto acc_coefs_z = b_coefs_z.get_access<sycl::access::mode::read>(cgh);
cgh.parallel_for(sycl::nd_range<3>(dimGrid * dimBlock, dimBlock),[=](sycl::nd_item<3> item_ct1) {kernel_lap(order, nx, nz, acc_pr, d_laplace_ct4,
acc_coefs_x, acc_coefs_z, item_ct1);});});

図 5. DPC++ の kernel_lap 関数呼び出し

図 5 は、 kernel_lap 呼び出しを使用した DPC++ カーネルの呼び出しを示しています。 これは、 ステンシル計
算のメイン・プロシージャーです。 移行したアプリケーションの各カーネルは、特定のデバイス用のキューに投入され、並列カーネルが起動する前にデータアクセス要件を満たす必要があります。
最終的な RTM 画像を使用して、 オリジナルの CUDA* 実装と移行した DPC++ コードを比較します。 地震画像の生成には、 表 1 に示す入力パラメーターと、 図 6 に示す速度モデルを使用します。 CUDA* 実装と DPC++ 実装のそれぞれの出力画像を図 7 に示します。 DPC++ 実装の最終画像は、 リファレンスと比較して十分な精度を達成しています。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe30

表 1. ここで紹介したモデル固有のパラメーターの説明

図 6. Koslov 速度モデル

図 7. オリジナルの CUDA* ベースの RTM ソースコード （左） と移行後の DPC++ コード （右） によって生成された地震画像

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

メモリー管理の改善

図 8. 第 9 世代インテル® HD グラフィックス NEO 上で明示的なメモリー割り当て （上） とバッファー /アクセサーによるメモリー管理 （下） を使用して実行した DPC++ バージョンのルーフライン ・ ビュー

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

31

The Parallel Universe

32

DPC++ のメモリー管理オプションを調査するため、 移行したアプリケーションからバッファー / アクセサーベースのバージョンを開発します。 これより、 デバイス上で明示的にメモリーを割り当てたり、 解放したりする必要がな
くなります。 また、 異なる処理ユニットとの間のデータ転送を管理する必要もありません。 これを実現するには、デバイスへ明示的にコピーしていた構造体ごとに、 対応するバッファー / アクセサーを作成します。 バッファーは計算終了後に破棄され、 データはホストにコピーバックされます （メモリー同期）。 波括弧は、 バッファー定義の周囲にデータ ・ オブジェクトを共有できるスコープを作成します。 実行がこのスコープ外になると、 実行フロー間の同期が行われ、 バッファーが破棄されます。 オリジナルと変更した移行ソースコードを比較するため、 インテル®
Advisor でルーフライン解析を行い、 ハードウェアの制限とシステム上の異なるメモリー階層間のデータ ・ トラン
ザクションを評価することで、 パフォーマンスを推定します。
図 8 は、 シングル ・ タイム ・ステップのみの単純化された RTM 実行のルーフライン ・ グラフです。 浮動小数点演算の数が減っているので、 低パフォーマンス ・ メトリックが想定されます。 上のグラフは、 明示的なデータ管理を使用する移行後のソースコードのルーフラインを示しています。 6.052 GFLOPS のパフォーマンスと 3.617
FLOP/ バイトの演算強度を達成しています。 演算強度とは、 移動されたデータ量 （メモリー ・ トラフィック） に対する浮動小数点演算の総数の割合です。 下のグラフは、 バッファー／アクセサーバージョンのルーフラインを示しており、 パフォーマンスは 2 倍の 12.246 GFLOPS、 演算強度は 17.896 FLOP/ バイトを達成しています。

まとめ
この記事では、 インテル® DPC++ 互換性ツールを使用して RTM コードを CUDA* から DPC++ に移行し、インテル® Advisor でチューニングして、 oneAPI の概念を実証しました。 移行後のソースコードは、 アプリケー
ションのアルゴリズム実行フローが独自の構造に統一されているため、より読みやすく、保守も容易になりました。
また、 移行後にバッファーやアクセサーを適用してメモリー管理を行うことで、 パフォーマンスと演算強度をさらに向上できました。 この記事で紹介したソースコードと手順2 は公開リポジトリーにあります。

関連情報
1. Claerbout, JF. Toward a unified theory of reflector mapping. Geophysics, v. 36, n. 3, p. 467-481, 1971.
2. https://github.com/cs2isenaicimatec/finite-difference-computation/ （英語）インテル® DPC++ 互換性ツール

CUDA* アプリケーションを標準ベースのデータ並列 C++ コードに移行

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

詳細 （英語）The Parallel Universe

33

oneAPI 仕様で提案されている高度なレイトレーシング
API
一度記述すれば複数のベンダーのアクセラレーターで動作する高忠実度のレイトレーシング

Jim Jeffers インテル コーポレーション シニア主席エンジニア兼アドバンスド・レンダリング & ビジュアライゼーション・
アーキテクチャー ・ シニア ・ ディレクター

画像提供 - Bentley Motors Limited: インテル® OSPRay Studio、 インテル® Embree、 インテル® Open Image Denoise を活用した高度なレイトレーシング

このたび、一連の高度なレイトレーシング API （英語） が oneAPI 仕様 （英語） に追加されたことをお知らせします。
映画、 科学的ビジュアライゼーション、 設計、 ゲームの世界におけるレイトレーシング計算は急激に増加しており、これらの API を XPU アーキテクチャー向けの oneAPI 仕様に追加することで、 この分野でのロバストで効率的な開発が進むことが期待されます。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

34

oneAPI 仕様にレイトレーシング機能を追加することで、 業界全体のソフトウェア開発者が、 複数のベンダーのシステムやアクセラレーターにまたがる高忠実度のレイトレーシング計算を 「一度記述する」 だけで済むようになります。 これらのインターフェイスを標準化することにより、 幅広い計算とレンダリング ・ インフラストラクチャー開発のため、 適切に設計され、 十分に試行された API とオプションが提供されます。 機能はいくつかの分野に細分化されます。
•• 幾何学的レイトレーシング計算•• ボリューム計算とレンダリング•• 画像のノイズ除去•• スケーラブル ・ レンダリングとビジュアライゼーション ・ インフラストラクチャー

レイトレーシング API のオープンソース実装は、 インテル® oneAPI レンダリング ・ ツールキット （英語） で活用されています。 ライブラリーは、 映画やテレビの写真のようなビジュアル効果やアニメーション、 科学的ビジュア
ライゼーション、 ハイパフォーマンス ・ コンピューティング計算、 CAD、 建築工学、 ゲームなど、 さまざまなソフト
ウェア ・アプリケーションの 3D グラフィックス計算で使用されます。
oneAPI 仕様では、 次のレイトレーシング ・ ライブラリーが推奨されています。
•• アカデミー賞を受賞したインテル® Embree （英語） — 幾何学的レイトレーシング•• インテル® オープン ・ ボリューム ・ カーネル ・ ライブラリー （インテル® Open VKL） （英語） —ボリューム処理•• インテル® Open Image Denoise （英語） —AI ベースのノイズ除去•• インテル® OSPRay （英語） —スケーラブルなミドルウェア ・ レンダリング API

これらのレイトレーシング API の機能に関する仕様作成プロセスと活発なコミュニティーからのフィードバック、Apple やその他の ARM ベースのプロセッサーでの 採用、 将来のインテル® Iris® Xe グラフィックスに向けた
インテル独自の開発体制は、 これらのレイトレーシング API がクロスアーキテクチャーで業界標準に対応していることを裏付けています。 API が開発者に提供する統合の容易性と利点は、 オープンソース ・コミュニティーのパート
ナーにより広く知られています。これを、オープンでコミュニティー主導の、クロスプラットフォームのレイトレーシング・
エコシステムの利点を提供する次のステップとします。
現在、 これらの API のさらなる拡張を行っています。 oneAPI 仕様に対する開発者の皆さんからのフィードバックとコラボレーションを GitHub* リポジトリー （英語） までお寄せください。

ポッドキャスト

災害管理の使命と科学的発見

アニメーター、 デジタルコンテンツ制作者、 建築エンジニア、 熟練したゲーマーの皆さん、インテル® oneAPI ベース & レンダリング ・ ツールキットでビジュアライゼーションの限界に挑戦してみてください。 詳細 （英語） >コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

聴取する （英語）The Parallel Universe

35

インテル® oneTBB フロー
グラフと OpenVINO™ 推論
エンジン
C++ でのディープラーニング ・ モデル間の依存関係の表現

Elvis G Fefey インテル コーポレーション ソフトウェア開発エンジニア
Michael J Voss インテル コーポレーション シニア主任ソフトウェア ・ エンジニア
Maxim Y Shevtsov インテル コーポレーション ディープラーニング ・ ソフトウェア ・ エンジニア

OpenVINO™ と調整フレームワークの必要性
データの可用性が高まるにつれて、 問題解決における従来のアプローチは、 データから学習するマシンラーニ
ング （ML） の手法に置き換えられています。 エンドツーエンドの ML パイプラインの 2 つの主要なステージは、
トレーニングと推論です。 トレーニングでは、 データを使用してモデルを作成します。 推論では、 モデルが新しいデータから出力を生成します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

36

OpenVINO ™ ツールキットは、 インテル® アーキテクチャーでハイパフォーマンスなディープラーニング推論を行う
開発者向けのツールスイートです。 専用の推論アクセラレーターも登場しつつありますが、 普及しているマルチコ
ア CPU での推論のパフォーマンスが向上すれば、 多くのユーザーに大きなメリットを提供できることになります。
OpenVINO™ のマルチスレッド・モデルは移植性に優れており、低レベルの制御を行う必要はありません。 ユーザーがスレッドを明示的に開始および終了する必要はなく、使用するプロセッサー数やコア数を知る必要もありません。
これにより、 簡単に展開できる最適化されたパフォーマンスが得られます。
ML が複雑になるにつれて、 推論アプリケーションで複数のモデルが必要になり、 一部のモデルはほかのモデルの出力に依存するようになります。 そのようなアプリケーションでは、 実行中にモデル間の依存関係を強制しつつ、独立して処理できるモデルを同時に実行できるように調整が必要です。
OpenVINO™ 推論エンジン自体は、 異なるモデルを同時に処理する方法は提供しません。 推論エンジンの主な役割は、 ハイパフォーマンス ・ モデルを調整してインテル® アーキテクチャーに展開するメカニズムを提供することです。 ただし、 OpenVINO™ ツールキットには、 広く利用されているオープンソースの GStreamer フレームワーク
を拡張した DL ストリーマーが含まれています。 GStreamer は、 モデル間の依存関係を強制する複雑なメディア
解析パイプラインを作成するフレームワークです。 DL ストリーマーは、 GStreamer を拡張したもので、 インテル®
アーキテクチャーにパイプラインの相互運用性と最適化された推論を提供します。 DL ストリーマーと GStreamer
はどちらも複雑なメディア ・ パイプラインを構築するための優れた選択肢ですが、 すべての開発者がこれらの大規模なフレームワークを使用するとは限りません。
この記事では、 OpenVINO™ ツールキットに含まれるインテル® oneAPI スレッディング ・ ビルディング ・ブロック（インテル® oneTBB） ライブラリーを利用して、 GStreamer や DL ストリーマーの代わりに、 軽量の C++ を使用して推論モデルを調整する方法を説明します。

インテル® oneTBB フローグラフ
インテル® oneTBB は、 CPU 上での並列プログラミング向けの汎用 C++ ライブラリーです。 2006 年から利用可能なインテル® スレッディング ・ ビルディング ・ブロック （インテル® TBB） ライブラリーが進化したもので、 長い歴史を持ちます。 このライブラリーは、 並列プログラミングを容易にする、 汎用の並列アルゴリズム、 フローグラフ・
インターフェイス、 コンカレント・コンテナー、 タスクベースのワーク・スケジューラー、 スケーラブルなメモリー・マネー
ジャー、 補助機能を提供します。 フローグラフ機能は、 計算グラフとして表現できる、 アプリケーションの関数とクラス （tbb::flow 名前空間） を提供します。 フローグラフは、 コードで実行の依存関係を表現する場合、または単純な線形パイプライン以上のものが必要なストリーミング ・アプリケーションがある場合に使用します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

37

フローグラフ ・ インターフェイスは、 3 つの主要なコンポーネントから構成されます。
1. 計算グラフ全体を表すグラフ ・ オブジェクト
2. ユーザーが提供するラムダ式の実行、 データストリームの結合、 データの分割とブロードキャストを行うノード
3. ノード間の依存関係または通信チャネルを表すエッジ

推論エンジンモデルのグラフ作成に役立つノードのサマリーを次に示します。
•• source_node （または input_node） ： source_node は、 グラフの残りの部分に供給されるデータを生成する機能を提供します。 このノードは、 ビデオストリームからフレームを読み取って利用するために使用できます。 インテル® oneTBB では、 source_node が input_node に置き換えられていることに注意してください。 ただし、記事執筆時点の最新バージョンの OpenVINO™ には、 source_node が含まれています。

•• function_node ： function_node 本体は、 ノードに入力された各データ項目に対してユーザーが提供するコードを実行して、 ノードから出力するデータを生成します。 このノードタイプは、 推論計算の実行に使用できます。
互いに直接的または間接的に依存しない （グラフのエッジで表される） 関数ノードは、 同時に実行できます。 グラ
フのエッジがノード間の依存関係を表す場合、 インテル® oneTBB ランタイム ・ ライブラリーは依存関係を強制し、ノードが正しい順序で実行されるようにします。 function_node は、ノードを流れる異なるデータ項目でユーザーが提供するコード本体の同時実行を許可または禁止するように設定できるため、 異なるデータ項目を並列で操作するのみならず、 シリアル化する必要がある操作 （処理されたビデオフレームを正しい順序で表示するなど） を実行することにも適しています。
•• sequencer_node ： 出力が特定の順序で必要な場合は、 シーケンサー ・ノードを使用して順序を管理します。 例
えば、 順不同で計算されたビデオフレームの順序付けに使用できます。
•• join_node ： 結合ノードは、 複数のデータストリームをまとめます。 独立して計算された出力を 1 つのユニットとしてまとめる必要がある場合は、 join_node を使用できます。

サンプル・アプリケーション
インテル® oneTBB を利用して OpenVINO™ 推論を調整する方法を示すため、 OpenVINO™ ツールキットに含まれるセキュリティー ・ バリアのデモを使用しました。 ビデオフレームから車両を検出する車両検出器、 検出された車両の色と種類を分類する車両分類器、 ナンバープレートのテキストを抽出するナンバープレート認識モデルの 3
つのトレーニング済みモデルにより推論を実装します。 車両分類器モデルとナンバープレート認識モデルは、 車両検出器モデルに依存します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

38

グラフを設定するには、 source_node を使用してビデオフレームを読み取った後、 function_node を使用して車両検出を行います。 ほかの 2 つの関数ノードを使用して、 車両分類器とナンバープレート認識を実行します。
join_node を使用して各フレームの結果を集計した後、 sequencer_node を使用して推論されたビデオフレー
ムの順序を変更します。 最後のノードは、 必要に応じて結果を表示または保存する別の function_node です。
図 1 にグラフを示します。

図 1. セキュリティー ・ バリア ・ デモのフローグラフ

入力ビデオフレームの読み取り
source_node を使用して、 ビデオから推論する画像を生成します。 終了条件が発生するまで、 本体を繰り返し実行してフレームを生成します。 このケースでは、 終了条件はビデオファイルの最後にある空のフレームです。
source_node のテンプレート ・ パラメーターは、 出力のデータ型です。 ノードは、 画像と対応する画像番号で構成される std::tuple を返します。 ダウンストリームでの集約およびフレーム順序の再確立で必要になるため、 フレーム番号が含まれています。 不要な場合は、 フレーム番号を省略できます。
まず、 フローグラフのノードについてヘッダーファイルをインクルードします。
#include "tbb/ﬂow_graph.h"

グラフを流れるメッセージを保持するため、 いくつかのデータ型を宣言します。 f はメッセージにフレームが含まれることを、 n はメッセージにフレーム番号が含まれることを、v はメッセージに車両の境界長方形が含まれることを、
p はメッセージにナンバープレートの境界長方形が含まれることを、 s はメッセージに vector<std::string>が含まれることをそれぞれ示します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

39

using fn_t = std::tuple<cv::Mat, size_t>;using fnvp_t = std::tuple<cv::Mat, size_t,
std::vector<cv::Rect> /* vehicles */,std::vector<cv::Rect> /* plates */ >;using sn_t = std::tuple<std::vector<std::string>, size_t>;

フレームカウンターとフレームを読み取るキャプチャー ・ オブジェクトを初期化します。
size_t frame_nu = 0;
VideoCapture my_video( “video_ﬁle_name” );次に、 フローグラフ ・ オブジェクトとフレームを読み取る source_node を作成します。
tbb::ﬂow::graph g;
tbb::ﬂow::source_node<fn_t> src_node(g,[&](fn_t &fn)->bool {cv::Mat f;
my_video.read(f);if (f.empty()) { return false; }std::get<0>(fn) = f;
std::get<1>(fn) = frame_nu++;
return true;}, false);

推論計算
セキュリティー ・ バリア ・ デモには、 車両の検出、 分類、 ナンバープレート認識の 3 つのクラスが含まれています。
これらのクラスで、 推論要求を作成し、 推論計算を送り、 結果を収集する機能を提供します。 detector （検出）、
vclassifier （分類）、 preader （ナンバープレート認識） の 3 つのクラスをインスタンス化します。

車両の検出
車 両 検出 の 推 論 計算には、 function_node を 使 用し ま す。 関 数 ノード の テンプレ ート ・ パラメー ター
は、 関数ノードが受け取る入力データ型と関数ノードが出力するデータ型の 2 つです。 車両検出器の入力は、
source_node から受け取った画像と対応する番号のペアです。 ノードの本体で、 推論要求を作成し、 推論のブロブを準備し、 推論計算を開始して、 計算結果が生成されるのを待ちます。 検出器の結果は、 信頼水準、車両なのかナンバープレートなのかを示すラベル、 検出されたオブジェクトの座標 （x、 y、 幅、 高さ） です。
各フレームで検出されたオブジェクトのリストをベクトルに設定し、 さらなる処理のためにダウンストリームに渡します。 出力タイプ fnvp_t には、 フレーム、 フレーム番号、 検出された車両の位置のリストが含まれます。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

40

tbb::ﬂow::function_node<fn_t, fnvp_t> detector_node(g,tbb::ﬂow::unlimited, [&](const fn_t &fn)->fnvp_t{
std::vector<cv::Rect> vehicles, plates;
auto req = detector.createInferRequest();auto& f = std::get<0>(fn);
auto& n = std::get<1>(fn);detector.setImage(req, f);req.StartAsync();
req.Wait(IInferRequest::WaitMode::RESULT_READY);
auto results = detector.getResults(req, f.size());
for (auto& r : results) {if (r.label == 1) {if (r.conﬁdence > 0.5) {vehicles.push_back(r.location &cv::Rect{cv::Point(0, 0), f.size()});}}
if (r.label == 2) {if (r.conﬁdence > 0.5) {plates.push_back(r.location &cv::Rect{cv::Point(0, 0), f.size()});}}}
return std::make_tuple(f, n, vehicles, plates);});

検出した車両の分類
車両分類器は、 フレーム、 フレーム番号、 検出された車両位置のリストを受け取ります。 各車両について、 推論
要求を作成し、 ブロブを準備し、 推論を開始して車両を分類し、 準備が整ったら結果を収集します。 分類結果は車両の色と種類で、 さらなる処理のためにダウンストリームに渡されます。
分類器の出力のデータ型は、 文字列のリストです。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

41

tbb::ﬂow::function_node<fnvp_t, sn_t> vclassiﬁer_node(g,tbb::ﬂow::unlimited, [&](const fnvp_t &fnvp)->sn_t{
auto& f = std::get<0>(fnvp);
auto& n = std::get<1>(fnvp);
auto& v = std::get<2>(fnvp);
std::vector<std::string> rlist;});

if (!FLAGS_m_va.empty() && !v.empty()) {for( auto& x : v) {auto req = vclassiﬁer.createInferRequest();vclassiﬁer.setImage(req, f, x);req.StartAsync();
req.Wait(IInferRequest::WaitMode::RESULT_READY);
auto r = vclassiﬁer.getResults(req);
rlist.push_back((r.ﬁrst + " " + r.second));}}return std::make_tuple(rlist, n);検出した車両のナンバープレートの認識分類器と同様に、 ナンバープレート認識は、 フレーム、 フレーム番号、 検出された車両位置のリストを受け取ります。
同様に、 推論要求を作成して送ります。 結果は、 ナンバープレートの文字列です。
tbb::ﬂow::function_node<fnvp_t, sn_t> preader_node(g,tbb::ﬂow::unlimited, [&](fnvp_t fnvp)->sn_t{
auto& f = std::get<0>(fnvp);
auto& n = std::get<1>(fnvp);
auto& p = std::get<3>(fnvp);
std::vector<std::string> rlist;});

if (!FLAGS_m_lpr.empty() && !p.empty()) {for( auto& x : p ) {auto req = preader.createInferRequest();preader.setImage(req, f, x);req.StartAsync();
req.Wait(IInferRequest::WaitMode::RESULT_READY);
auto r = preader.getResults(req);
rlist.push_back(r);}}return std::make_tuple( rlist, n);コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

42

生の推論結果のみが必要な場合は、 すでに説明したノードで十分です。 しかし、 後処理が必要な場合は、 ノードを追加します。

集約
ここで、 各フレームの計算を集約します。 フレーム番号をキーとして使用し、 タグ一致 join_node を使用して、各フレームのすべての結果をまとめます。 各フレームには、 フレーム番号、 検出された車両の分類結果、 検出されたナンバープレートの文字列が含まれます。
tbb::ﬂow::join_node<std::tuple<fnvp_t, sn_t, sn_t,
tbb::ﬂow::tag_matching >agg_node(g, [](const fnvp_t &x) { return std::get<1>(x); },[](const sn_t &x){ return std::get<1>(x); },[](const sn_t &x){ return std::get<1>(x); });順序設定
計算は並列で行われるため、 フレームの推論結果は順不同で完了する可能性があります。 sequencer_node を使用すると、 ビデオの再生が必要な場合にフレームの順序を変更できます。
tbb::ﬂow::sequencer_node<std::tuple<fnvp_t, sn_t, sn_t >>seq_node(g, [] (const std::tuple<fnvp_t, sn_t, sn_t > &x) -> size_t {return std::get<1>(std::get<0>(x));});

結果のスーパーインポーズ
最後に、 検出された車両の位置とナンバープレートの文字列をフレームに重ね合わせる function_node を実装します。 結果は再生したり、 出力ビデオに保存できます。 計算が再生よりも速く完了する場合は、 スロットリング ・メカニズムを使用して処理されるフレームの増加を防ぎます。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

43

tbb::ﬂow::function_node<std::tuple<fnvp_t, sn_t, sn_t>, cv::Mat>
output_node(g, tbb::ﬂow::serial,[&](std::tuple<fnvp_t, sn_t, sn_t> agg_results){
auto& f = std::get<0>(std::get<0>(agg_results));
auto& v = std::get<2>(std::get<0>(agg_results));
auto& p = std::get<3>(std::get<0>(agg_results));
auto& vs = std::get<0>(std::get<1>(agg_results));
auto& ps = std::get<0>(std::get<2>(agg_results));
auto frame_nu = std::get<1>(std::get<0>(agg_results));for (auto &r : v) {cv::rectangle(f, r, {0, 255, 0}, 4);}for (auto &r : p) {cv::rectangle(f, r, {0, 255, 0}, 4);}if(!v.empty()) {for (size_t i=0; i< v.size(); ++i) {cv::putText( f, vs[i], cv::Point{v[i].x, v[i].y + 35},
cv::FONT_HERSHEY_COMPLEX, 1.3, cv::Scalar(0, 255, 0), 4);}}
if (!p.empty()) {for (size_t i=0; i< p.size(); ++i) {slog::info << "Frame Number: " << frame_nu << slog::endl;
slog::info << "License Plate: " << ps[i] << slog::endl;
slog::info << "====" << slog::endl;
cv::putText( f, ps[i], cv::Point{p[i].x, p[i].y + 35},
cv::FONT_HERSHEY_COMPLEX, 1.3, cv::Scalar(0, 255, 0), 4);

});

}
}
return f;

グラフの作成と依存関係の表現
グラフを作成して、 依存関係をインテル® oneTBB 関数 tbb::flow::make_edge で表現します。 先行ノードと後続ノードを使用して、 任意のグラフトポロジーを作成できます。 唯一必要なのは、 エッジのデータ型が一致していることを確認することです。 つまり、 出力ノードのデータ型は入力ノードのデータ型と一致する必要があります。
フレーム生成から車両検出までエッジを作成してグラフを作成します。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe44

make_edge(src_node, detector_node);

車両検出器は、 車両分類器とナンバープレート認識に並列で接続します。
make_edge(detector_node, vclassiﬁer_node);
make_edge(detector_node, preader_node);

検出器、 分類器、 ナンバープレート認識をアグリゲーターに接続します。 検出器はフレームを提供し、 分類器とナンバープレート認識は集約のため、 そのフレームの推論結果を提供します。
make_edge(detector_node, tbb::ﬂow::input_port<0>(agg_node));make_edge(vclassiﬁer_node, tbb::ﬂow::input_port<1>(agg_node));make_edge(preader_node, tbb::ﬂow::input_port<2>(agg_node));アグリゲーターをフレームを並べ替えるシーケンサーに接続します。
make_edge(agg_node, seq_node);

シーケンサーを結果を出力する出力ノードに接続します。
make_edge(seq_node, output_node);

グラフのアクティベート
最後に、グラフをアクティベートして実行できるようにします。 この処理には、ソースノードの activate() メンバーを使用します。 g.wait_for_all() を呼び出して、 グラフが終了する前にグラフのすべての計算が完了するようにします。
src_node.activate();g.wait_for_all();

結果
実装の評価には、 Open Model Zoo （英語） リポジトリーの 3 つの事前トレーニング済みモデル （vehiclelicense-plate-detection-barrier-0106_fp16、 vehicle-attributes-recognition-barrier-0039_fp16、license-plate-recognition-barrier-0001） を使用しました。 インテル® Iris® Pro グラフィックス 580 が統合されたインテル® Core™ i7-6770HQ プロセッサーを搭載したシステムでアプリケーションを実行しました。 ビデオ
から 32,030 フレームのすべての推論を実行する経過時間を測定しました。 手動でスレッド化した実装に対するインテル® oneTBB フローグラフのパフォーマンス ・ ゲインを図 2 に示します。 最大 11% のパフォーマンス ・ ゲインが得られました。 CPU と統合 GPU に異なる方法でモデルを割り当てた、 5 つの異なる設定を評価しました。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

45

図 2. セキュリティー ・ バリア ・ デモのパフォーマンス結果。
D-CPU、 VC-CPU、 LPR-CPU: 3 つのモデルをすべて CPU で実行。
D-GPU、 VC-GPU、 LPR-GPU: 3 つのモデルをすべて GPU で実行。
D-GPU、 VC-GPU、 LPR-CPU: 検出と分類を GPU で実行し、 認識を CPU で実行。
D-GPU、 VC-CPU、 LPR-GPU: 検出と認識を GPU で実行し、 分類を CPU で実行。
D-CPU、 VC-GPU、 LPR-GPU: 分類と認識を GPU で実行し、 検出を CPU で実行。

インテル® oneTBB フローグラフの入手方法
OpenVINO™ ツールキットにはすでにインテル® oneTBB が含まれているため、 フローグラフを使用しても追加の依存関係は発生しません。 ビルド ・ インフラストラクチャー （CMakeLists.txt） を更新して、 利用できるようにするだけです。
include_directories(/path/to/openvino/inference_engine/external/tbb/include)
link_libraries(-L/path/to/openvino/inference_engine/external/tbb/lib -ltbb)

まとめ
インテル® oneTBB フローグラフを使用して、 推論に使用したモデル間の依存関係を表現する方法を示しました。
フローグラフを使用すると、 依存関係を簡単に表現できるだけでなく、 モデル間の並列処理を C++ で直接表現できます。 GStreamer や DL ストリーマーのような大規模なフレームワークを使用しない開発者は、 インテル®
oneTBB フローグラフを使用することを検討してください。 OpenVINO™ セキュリティー・バリア・デモの実装では、
インテル® oneTBB フローグラフを使用した場合、 手動でスレッド化した実装に対して最大 11 ％のパフォーマン
ス・ゲインが得られました。フローグラフは C++ フレームワークであり、OpenVINO™ ツールキットに含まれていて、スレッド化をサポートしていることから、 OpenVINO™ で推論の ML モデルのグラフを作成するのは自然な選択と言えるでしょう。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

46

明示的なベクトル化を使用したスキャン操作の最適化
インテル ® AVX-512 SIMD 命令を活用してプリフィクス ・ サム計算を高速化

Vamsi Sripathi インテル コーポレーション シニア HPC アプリケーション ・ エンジニア
Ruchira Sasanka インテル コーポレーション シニア HPC アプリケーション ・ エンジニア

はじめに
最近のインテル® プロセッサーは、 命令レベル、 データレベル、 スレッドレベルの並列処理を提供しています。
SIMD （Single Instruction Multiple Data） オペランドを同時に実行する機能により、 プロセッサーの演算実行ユニットの利用率が最大化されます。 この記事では、 ベクトルスキャン操作に適用する SIMD の最適化について
取り上げます。
スキャン （包括的 / 排他的スキャン、 プリフィクス・サム、 または累積和とも呼ばれる） は、 多くのアプリケーション・
ドメインで一般的な操作であり1、 C++、 OpenMP* ランタイム、 Python* NumPy* パッケージの標準ライブラリー

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe47

関数として定義されています2,3。 ベクトルのスキャンは、ソースベクトルからインデックス i （包括的スキャンの場合）
および i-1 （排他的スキャンの場合） までのすべての値を合計してインデックス i の結果を得るもう 1 つのベクトル （配列） です。 例えば、 要素の数が n のベクトル x の包括的スキャンは、 次の方法で取得します。

一般的には、 次のように定義されます。

包括的スキャン （以降、 単にスキャンと呼びます） について説明します。 SIMD の概要 （ベクトル化） から始め、C および OpenMP* ディレクティブを使用したスキャン操作のベースラインを実装します。 次に、 512 ビットのベクトルレジスターで動作する SIMD に最適化されたスキャン ・アルゴリズムを実装します。 最後に、 ベースライン
実装と最適化された実装のパフォーマンスを比較します。

SIMD の概要
SIMD 命令は、 複数のデータオペランドを保持できる広いベクトルレジスターを取り扱います。 ベクトルレジスターの幅と数はアーキテクチャーに依存し、 最新世代のインテル® プロセッサーは 32 の 512 ビット幅のベクトルレジ
スターをサポートします。 インテルは、 x86 ISA 拡張命令をインテル ® アドバンスト ・ ベクトル ・ エクステンション（インテル® AVX）、 インテル® AVX2、 インテル® AVX-512 形式で提供し、 さまざまな算術演算と論理演算をそれ
ぞれ 256 ビット幅と 512 ビット幅のレジスターで実行できます。 ベクトルレジスターの幅に応じて、 XMM、 YMM、ZMM という名前が付けられています （表 1）。

レジスター幅レジスター名64b オペランド /レジスター数x86 ISASIMD 加算あたりの 64b ops 数

128 ビット

XMM

2

インテル® SSE

2256 ビット

YMM

4インテル® AVX/AVX2

4

512 ビット

ZMM

8インテル® AVX-5128

表 1. SIMD レジスターの種類

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

48

SIMD 実行は、 主に次の 2 つの方法で行われます。
1. 暗黙的なベクトル化 ： コンパイラーに依存し、 スカラー計算ループをベクトル化 /SIMD ブロックに自動的に変換します。 コンパイラーは、 コストモデルに基づいて、 ベクトル化する計算の特定、 アンロールするループの決定などの手間のかかる作業を行います。 また、 生成されたオブジェクト ・コードで必要なループのプロローグセクションとエピローグセクション （ピールとリマインダーとも呼ばれます） を生成することにより、 SIMD 命令のデータ ・
アライメント要件を処理します。 プログラマーが行う作業は、 ターゲット CPU アーキテクチャー向けに適切なコンパイラー最適化オプションを指定することだけです。 インテル® コンパイラーは、 SIMD 命令の生成を制御する
明示的なヒントを提供するプラグマ 4 も提供しています。
2. 明示的なベクトル化 ：a. ベクトル組込み関数 ： ハードウェア SIMD 命令にマップする、 コンパイラーによって提供される C API です5。
ベクトル組込み関数を使用すると、 プログラミングで、 コードで生成される SIMD 命令の種類 （インテル®
AVX、 インテル® AVX2、 インテル® AVX-512） と、 それらをどの程度積極的に使用するかを制御できます。
レジスターの選択や命令スケジューリングなどのチューニングの重要な部分はコンパイラーに任せて、 スキャンなどの重要なデータ依存操作のコード生成はプログラマーが制御する、 ハイブリッド ・ モードと見なすことができます。
b. アセンブリー言語 ： コンパイラーの最適化を完全に避けて、 マイクロアーキテクチャーの専門知識が必要なアセンブリー言語でコーディングします。 このアプローチは推奨しません。

表 2 は、 これらのアプローチのトレードオフを要約したものです。
モード

移植性

労力

暗黙的なベクトル化

高い

低い

ベクトル組込み関数

中中

アセンブリー命令低い

高い

表 2. SIMD プログラミング ・ アプローチの比較

ベースライン ・ コード
ベクトルのスキャンを計算するベースライン・コードと、インテル® コンパイラーの最適化レポート（-qopt-report=5
コンパイラー ・ オプションを使用して生成） を次に示します。 スキャン操作には固有のループ伝搬依存関係がある
ため、 コンパイラーは計算を完全にはベクトル化できません。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

49

OpenMP* ディレクティブを追加した同じ操作とコンパイラーの最適化レポートを次に示します。 OpenMP* は、
SIMD サポートを含むポータブルなディレクティブ ・ ベースの並列プログラミング ・ モデルです。 （reduction 節
および scan ディレクティブを組み合わせた） SIMD 構文を使用して、 ベクトルのスキャン操作を実行します（詳細は、 OpenMP* 仕様を参照してください6,7）。 SIMD がワークロードのパフォーマンスに与える影響を理解する
ため、 ループを並列化しないで、 OpenMP* ディレクティブをベクトル化のみに制限します。 コンパイラーの最適化
レポートは、 ベクトル長 8 （つまり、 インテル® AVX-512） でベクトル化され、 係数 16 でループがアンロールされたことを示しています。

ベースラインと自動ベクトル化バージョンを、 次のセクションで説明する明示的なインテル® AVX-512 SIMD 実装と比較します。

明示的なベクトル化
このセクションでは、 インテル® AVX-512 実行ユニットをサポートしているインテル® プロセッサーでのベクトルス
キャン操作を最適化する SIMD の手法を示します。 入力ベクトルと出力ベクトルには倍精度浮動小数点 （FP64）
データ型を使用します。 実装で使用したインテル® AVX-512 SIMD 命令を表 3 に示します。
実装の主な考え方は、 必要なシャッフル操作のシーケンスを適用した後、512 ビット・レジスターの下位 256 ビット・
レーンと上位 256 ビット ・ レーンで一連の加算操作を同時に実行することです。 8 つの入力要素ごとに、 合計 5 つコンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

50の加算と 5 つの並べ替えを実行します。 結果を 1 つのレジスターに累計しているため、 すべての加算操作は依存
関係チェーンを形成します。 しかし、 加算から独立している一部の並べ替えは同時に実行できます。 図 1 は実装をビジュアルに表現したもので、 反復ごとに 16 の要素を処理しているメインブロックのコードが続きます。

名前インテル® AVX-512 ベクトル組込み関数 API（x86 命令）

説明

SIMD ロード__m512d _mm512_loadu_pd(void *mem_address);

8 つの FP64 要素のセットをメモリーからインテル® AVX-512 ベクトルレジス
ターにロードします。

SIMD 加算__m512d _mm512_add_pd(__m512d src1, __m512d src2);

src1 および src2 の FP64 要素を加算
して結果を返します。

imm8 の 8 ビット ・コントロールを使用
SIMD
して 128 ビット ・ レーン内の src の要素
並べ替え__m512d _mm512_maskz_permute_pd
をシャッフルし、 ゼロマスク k を使用しマスクあり (__mmask8 k, __m512d src, const int imm8); て結果をストアします （対応するマスク（128b レーン）
ビットが設定されていない場合、 要素
は 0 に設定されます）。
imm8 の 8 ビット ・コントロールを使用
SIMD
して 256 ビット ・ レーン内の src の要素
並べ替え__m512d _mm512_maskz_permutex_pd
をシャッフルし、 ゼロマスク k を使用しマスクあり (__mmask8 k, __m512d src, const int imm8); て結果をストアします （対応するマスク（256b レーン）
ビットが設定されていない場合、 要素
は 0 に設定されます）。
SIMD ストア

void _mm512_storeu_pd(void *mem_address, _m512d src);

8 つの FP64 要素のセットをインテル®
AVX-512 ベクトルレジスターからメモ
リーにストアします。

表 3. インテル® AVX-512 SIMD 命令

図 1. 明示的なインテル® AVX-512 SIMD スキャン操作のシーケンス

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

51

The Parallel Universe

52

パフォーマンスの評価
インテル® Xeon® Platinum 8260L プロセッサーの 1 つのコアで、 スキャン実装の 3 つのバージョンのパフォーマンスを評価しました。 GCC （9.3.0）、Clang （10.0.1）、ICC （インテル® C++ コンパイラー 19.1.3.304） コンパイラーで、
ベースライン、 OpenMP* SIMD、 明示的なインテル® AVX-512 SIMD 実装を比較しました。 表 4 は、 ICC で生成された各実装のメインブロックのアセンブリー ・コードです。

表 4. ICC で生成された 3 つのスキャン実装のアセンブリー ・ コード

ベースライン・コードでは、 ループは 2 つの要素のみでアンロールされ、 両方のスカラー加算操作が依存関係チェーンを形成していることがわかります。 生成されたコードは広いインテル® AVX-512 実行ユニットを利用していない
ため、 これは非効率的です。
OpenMP* SIMD 実装では、 生成されたコードはインテル® AVX-512 の加算と並べ替え操作を使用し、 16 の要
素でアンロールされます。 16 の入力要素を処理するために、合計 8 つの加算、8 つの並べ替え、2 つのブロードキャ
スト命令が生成されています。 8 つの要素のセットについて、加算と並べ替えが依存関係チェーン （赤でハイライト）を形成します。 このバージョンのコードで生成されたすべての並べ替えは、256b レーン内で要素をシャッフルします。
128b レーン内で要素をシャッフルする並べ替え（1 サイクル）は、256b レーン内でシャッフルする並べ替え（3 サイクル）
よりも低レイテンシーです。

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

53

明示的なインテル® AVX-512 SIMD 実装は 16 の要素でアンロールされ、 11 の加算と 10 の並べ替え命令を使用
して 16 の要素を処理します。 OpenMP* SIMD スキャンよりも多くの加算および並べ替え命令を使用しますが、2 つの利点があります。 まず、 10 の並べ替えの 4 つが 128b レーン （レイテンシーが 1 サイクル） 内にあるため、より効率的です。 これらの 4 つの並べ替え命令は緑でハイライトされています。 次に、 一部の並べ替えは加算命
令から独立しています。 これは、 多くの命令レベルの並列処理とインテル® AVX-512 実行ユニットの効果的な利用に役立ちます。 できるだけ早く累計結果を計算するため、 追加の加算を行います。 これにより、 多くの同時命令が実行され、 将来の反復での実行パイプラインのストールが軽減されます。
図 2 は、 64 から 1,024 要素の範囲のベクトルサイズについて、 すべてのベクトルを L1 キャッシュに配置した場合のパフォーマンスを示しています。 パフォーマンス ・ データから次のことが分かります。
1. 明示的なインテル® AVX-512 SIMD 実装は、 ベースラインおよび OpenMP* SIMD 実装よりも優れています。
2. GCC と Clang は、 スキャン 計算をベクトル 化できません。 OpenMP* SIMD ディレ クティブ を 使 用しても、パフォーマンスは変わりません。
3. ICC は、 OpenMP* SIMD ディレクティブを使用すると優れた自動ベクトル化を行います。
4. ベースラインおよび OpenMP* SIMD スキャンに対する明示的な SIMD スキャン実装の平均スピードアップは、それぞれ 4.6 倍 （GCC と Clang） および 1.6 倍 （ICC） でした。

図 2. 明示的なインテル® AVX-512 スキャンのパフォーマンスの比較

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

The Parallel Universe

54

まとめ
明示的な SIMD プログラミングを簡単に紹介して、 ベクトルスキャン計算に適用しました。 多くの場合、 最適化
コンパイラーを使用すると優れたパフォーマンスが得られますが、 この記事で示されているように、 改善の余地があるケースもあります。 そのため、 開発者が最大のパフォーマンスを達成するための明示的な SIMD プログラ
ミングを理解することは有用です。

関連情報 （英語）
1. Guy E Blelloch (2018) Prefix sums and their applications
2. https://en.cppreference.com/w/cpp/algorithm/inclusive_scan
3. https://numpy.org/doc/stable/reference/generated/numpy.cumsum.html
4. https://sof tware.intel.com/content/w w w/us/en/develop/documentation/cpp-compilerdeveloper-guide-and-reference/top/compiler-reference/pragmas/intel-specific-pragmareference.html#intel-specific-pragma-reference
5. https://sof tware.intel.com/content/w w w/us/en/develop/documentation/cpp-compilerdeveloper-guide-and-reference/top/compiler-reference/intrinsics/details-about-intrinsics.
html
6. https://www.openmp.org/spec-html/5.0/openmpsu42.html#x65-1400002.9.3.1
7. https://www.openmp.org/spec-html/5.0/openmpsu45.html#x68-1940002.9.6

コンパイラーの最適化に関する詳細は、 最適化に関する注意事項を参照してください。

THE PARALLEL

UNIVERSE

インテルのテクノロジーを使用するには、 対応したハードウェア、 ソフトウェア、 またはサービスの有効化が必要となる場合があります。 詳細については、 OEM または販売店にお問い合わせいただくか、
http://www.intel.co.jp/ を参照してください。
実際の費用と結果は異なる場合があります。
インテルは、 サードパーティーのデータについて管理や監査を行っていません。 ほかの情報も参考にして、 正確かどうかを評価してください。
最適化に関する注意事項 : インテル® コンパイラーでは、 インテル® マイクロプロセッサーに限定されない最適化に関して、 他社製マイクロプロセッサー用に同等の最適化を行えないことがあります。
これには、 インテル® ストリーミング SIMD 拡張命令 2、 インテル® ストリーミング SIMD 拡張命令 3、 インテル® ストリーミング SIMD 拡張命令 3 補足命令などの最適化が該当します。 インテルは、他社製マイクロプロセッサーに関して、 いかなる最適化の利用、 機能、 または効果も保証いたしません。 本製品のマイクロプロセッサー依存の最適化は、 インテル® マイクロプロセッサーでの使用を前提としています。 インテル® マイクロアーキテクチャーに限定されない最適化のなかにも、 インテル® マイクロプロセッサー用のものがあります。 この注意事項で言及した命令セットの詳細については、該当する製品のユーザー ・ リファレンス ・ ガイドを参照してください。 注意事項の改訂 #20110804。 https://software.intel.com/en-us/articles/optimization-notice#ja-jp
性能に関するテストに使用されるソフトウェアとワークロードは、 性能がインテル® マイクロプロセッサー用に最適化されていることがあります。
SYSmark* や MobileMark* などの性能テストは、 特定のコンピューター ・ システム、 コンポーネント、 ソフトウェア、 操作、 機能に基づいて行ったものです。 結果はこれらの要因によって異なります。
製品の購入を検討される場合は、 他の製品と組み合わせた場合の本製品の性能など、 ほかの情報や性能テストも参考にして、 パフォーマンスを総合的に評価することをお勧めします。 構成の詳細は、補足資料を参照してください。 性能やベンチマーク結果について、 さらに詳しい情報をお知りになりたい場合は、 http://www.intel.com/benchmarks/ （英語） を参照してください。
性能の測定結果はシステム構成の日付時点のテストに基づいています。 また、 現在公開中のすべてのセキュリティー・アップデートが適用されているとは限りません。 詳細は、 システム構成を参照してください。
絶対的なセキュリティーを提供できる製品またはコンポーネントはありません。
本資料は、 （明示されているか否かにかかわらず、 また禁反言によるとよらずにかかわらず） いかなる知的財産権のライセンスも許諾するものではありません。
インテルは、 明示されているか否かにかかわらず、 いかなる保証もいたしません。 ここにいう保証には、 商品適格性、 特定目的への適合性、 および非侵害性の黙示の保証、 ならびに履行の過程、取引の過程、 または取引での使用から生じるあらゆる保証を含みますが、 これらに限定されるわけではありません。
© Intel Corporation. Intel、 インテル、 Intel ロゴ、 その他のインテルの名称やロゴは、 Intel Corporation またはその子会社の商標です。
* その他の社名、 製品名などは、 一般に各社の表示、 商標または登録商標です。
JPN/2106/PDF/XL/SPI/ND