SYCL* の事例
C++ Thrust アプリケーションを SYCL* とoneAPI DPC++ ライブラリー （oneDPL） へ移行する
DPEcho ： 2020 年代以降の SYCL* による一般相対論

Issue

51
2023

目次

注目記事

The Parallel Universe2

編集者からのメッセージ

3

SYCL* の事例

6

C++ Thrust アプリケーションを SYCL* と oneAPI DPC++ ライブラリー（oneDPL） へ移行する

11

DPEcho ： 2020 年代以降の SYCL* による一般相対論

14

Fortran、 oneMKL、 OpenMP* を使用して LU 因数分解を高速化

20

ISO C++ がヘテロジニアス ・コンピューティングに十分でない理由

oneAPI で C++ アプリケーションを現代化および高速化し、 ベンダーへの依存を排除する

科学計算における SYCL* の優位性

一般的な数学演算を高速化する標準ベースのオープンなアプローチ

ACM RecSys Challenge 2022 のインテル® SIHG4SR ソリューション

26

薬の良し悪しを見分ける

33

インテルの推薦システムが上位入賞

製薬業界におけるコンピューター ・ ビジョン

インテル® Arc™ GPU で TensorFlow* Stable Diffusion を実行する
PC で高速なディープラーニングを実現© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

41

The Parallel Universe

3

編集者からのメッセージ
Henry A. Gabb インテル コーポレーション シニア主席エンジニア

HPC と並列コンピューティング分野で長年の経験があり、 並列プログラミングに関する記事を多数出版しています。 『Developing Multithreaded Applications: A Platform Consistent Approach』の 編 集 者 / 共 著 者 で、 イ ン テル と Microsoft* に よる Universal Parallel Computing ResearchCenters のプログラム ・ マネージャーを務めました。

oneAPI 誕生から 4 年
2018 年の Intel Architecture Day でインテル コーポレー
ションのチーフ ・アーキテクトである Raja Koduri が oneAPI
業界イニシアチブを発表してから、 4 年が経ちました。 彼は次のようにツイート （英語） しています。

「2018 年の Intel Architecture Day で oneAPI を発表してから 4 年。
長い道のりを歩んできました。 2020 年のベータ版リリース以来、 主要な
HPC、 AI、 レンダリング ・ フレームワークから支持され、 開発者による採用はほぼ 10 倍に拡大しました。 さらなる進化への素晴らしい始まりです。」
oneAPI の 4 周年を記念して、本号では SYCL* と oneMKL や oneDPL などの oneAPI ライブラリーに注目します。
注目記事 「SYCL* の事例」 では、現在の ISO C++ がヘテロジニアス並列処理には不十分である理由を概説します。
一言で言えば、 C++ にはアクセラレーター ・ デバイスの概念がないため、 デバイスの検出、 不連続メモリー、
デバイス境界を越えた例外処理などの概念がありません。 しかし、 将来の C++ バージョンでは、 ヘテロジニアス
なハードウェアに対応していくと考えられるため、 今後変更される可能性があります。 ホストとデバイス間のデータ転送については、 Parallel Universe 48 号の記事 「oneAPI の maxloc リダクション」 で取り上げました。 SYCL*
デバイス検出 API については、 今後の記事で取り上げる予定です。 「C++ Thrust アプリケーションを SYCL* とoneAPI DPC++ ライブラリー （oneDPL） へ移行する」 では、 オープンでベンダーに依存しないライブラリーを使用して、 C++ コードを現代化および高速化する方法を説明します。 「DPEcho ： 2020 年代以降の SYCL* による一般相対論」 では、 Salvatore Cielo 氏 （ライプニッツ・スーパーコンピューティング・センター）、 Alexander Pöppl（インテル コーポレーション）、 Luca Del Zanna 氏 （フィレンツェ大学）、 および Matteo Bugli 氏 （トリノ大学） が重要な科学アプリケーションを SYCL* へ移行した方法を紹介します。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

4

最近、 Codeplay が NVIDIA* GPU と AMD* GPU 向
けの oneAPI プラグインをリリースしました。 また、
oneAPI によるオープンソースの高速コンピューティン
グという共有ビジョンを実現するため、ArrayFire* チー
ムがインテルに加わりました （英語）。 ArrayFire* オー
プンソース ・プロジェクトは、 The ArrayFire Missionに従って継続され、 Google、 Twitter、 VoltronData、
そしてインテルを含むさまざまな企業がスポンサーとなっているメンテナーによって管理されます。 Parallel
Universe の 以 前 の 記 事 「ArrayFire* と oneAPI、各 種 ラ イ ブ ラ リ ー、 OpenCL* の 相 互 運 用 性 」 や「ArrayFire* と oneAPI による 2 次元フーリエ相関
アルゴリズムの高速化」 から、 ArrayFire* は 1 つのコードでどこでも実行できる、 ヘテロジニアス並列処理を実現する高水準の抽象化であることを記憶されている方もいらっしゃるでしょう。
「Fortran、 oneMKL、 OpenMP* を使用して LU 因数分解を高速化」 では、 初心に帰り OpenMP* ヘテロジニ
アス並列ディレクティブを使用して、 数学ライブラリー関数をアクセラレーターにオフロードし、 ホストとデバイスのデータ転送を効率良く管理する方法を示します。
oneAPI を利用した AI アプリケーションにも目を向けます。 「ACM RecSys Challenge 2022 の インテル®
SIHG4SR ソリューション」 では、 この推薦システムコンテストで上位入賞を果たしたインテルのソリューションについて説明します。 「薬の良し悪しを見分ける：製薬業界におけるコンピューター・ビジョン」 では、 AI リファレンス・
キットの Visual Quality Inspection （英語） を紹介します。 このエンドツーエンドのコンピューター ・ ビジョンとディープラーニング ・ パイプラインは、 医薬品製造工程で不良錠剤を識別します。 Stable Diffusion （文章による説明を画像に変換する、 楽しくも本格的なディープラーニング ・ モデル） を PC で実行することに興味がある場合、「インテル® Arc™ GPU で TensorFlow* Stable Diffusion を実行する」 は必読です。 セットアップから実行までの手順が分かります。
コードの現代化、ビジュアル・コンピューティング、 データセンターとクラウド・コンピューティング、 データサイエンス、
システムと IoT 開発、 oneAPI を利用したヘテロジニアス並列コンピューティング向けのインテル ・ ソリューション
の詳細は、 Tech.Decoded （英語） を参照してください。
Henry A. Gabb
2023 年 1 月© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

将来に通用するコード
プロプライエタリーの壁を超えて成長

複数の言語をサポートし、 ヘテロジニアス計算パフォーマンスを実現する、単一のオープンなプログラミング ・ モデルで、 コードの幅を広げましょう。
オープンスタンダードに根ざしたインテル® oneAPI は、 クロスアーキテクチャーのライブラリー、コンパイラー、 ツールを提供し、 コードをより多くのハードウェアに対応させ、比類ないパフォーマンスを実現します。
インテル® oneAPI の詳細

The Parallel Universe

6

SYCL* の事例

ISO C++ がヘテロジニアス ・ コンピューティングに十分でない理由

Henry A Gabb インテル コーポレーション シニア主席エンジニア兼 The Parallel Universe 編集長

計算科学に長く携わっていると、 プログラミング言語やツールは移り変わるものだということが分かります。
1 つのプログラミング言語だけで仕事ができる人は、 もうほとんどいないでしょう。 2000 年にインテルに入社したとき、 私のプロジェクトでは、 Fortran、 C、 そして残念なことに C++ が混在していました。 当時、 私は C++が好きではありませんでした。 コンパイルに時間がかかり、 オブジェクトやオブジェクト指向プログラミングに価値があるとは思えなかったのです。 C++ の構成要素は言語に豊かな表現力をもたらす一方、 処理速度が遅くなる1
ことも、 魅力を感じなかった要因です。
そのため、 oneAPI を使うようになり、 C++ を勉強し直さなければならなくなったときには、 気が重くなりました。
oneAPI のダイレクト ・プログラミング手法は、 ISO C++17 をベースとした Khronos SYCL* （英語） を採用
しています。 しかし、 2000 年頃の C++ と比較すると、 C++17 の見た目は大きく異なります。 特に C++ 標準
テンプレート ・ ライブラリー （STL） や oneAPI ライブラリーと組み合わせた場合、 非常に表現力が豊かで、

1

C++ コンパイラーはこの 20年で大幅に改善されたので、もはやこれは真実ではありません。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

7

生産性に優れていると言わざるを得ません。 これが、 本記事のトピックである SYCL*、 および SYCL* と ISO C++の関係につながります。
C++ は、 （Rust* 支持者が何と言おうと2） 時の試練に耐えてきた優れたプログラミング言語です。 多くの C++
プログラマーがおり、 何十億行もの C++ コードが存在します。 しかし、 「フリーランチ」 （英語） はずいぶん前に終わり、 Herb Sutter は現在 「ハードウェア ・ ジャングル」3 （英語） について語っています。 ヘテロジニアス並列処理時代にありながら、 C++ にはアクセラレーター ・ デバイスを活用するのに必要な主要機能が欠けています。
まず、 最も重要なことですが、 C++ にはアクセラレーター ・ デバイスという概念がないため、 利用可能なデバイスを検出したり、 それらのデバイスにワークをオフロードしたり、 デバイス上で実行中のコードで発生する例外を処理する言語構造がありません。 コードがホストとさまざまなデバイス上で非同期に実行される場合、 例外処理は特に複雑になります。 SYCL* は、 利用可能なデバイスを検出し、 デバイスのハードウェアとバックエンドの特性を照会する platform （英語）、 context （英語）、 および device （英語） の C++ クラスを提供しています。 queue （英語）クラスと event （英語） クラスは、 プログラマーがデバイスにワークをサブミットし、 非同期タスクの完了を監視できるようにします。 デバイス上で実行されているコードでエラーが発生した場合、queue::wait_and_throw()、queue::throw_asynchronous()、 または event::wait_and_throw() メソッドによって、 あるいはキューやコンテキストを破棄するときに自動的に非同期のエラーハンドラーをトリガーできます。
また、 C++ には不連続メモリーという概念がありません。 アクセラレーター ・ デバイスは通常、 ホストシステムのメインメモリーとは別に独自のメモリーを持っています。 ヘテロジニアス並列処理を行うには、 ホストメモリーから
デバイスメモリーにデータを転送する必要があります。 C++ にはそのようなメカニズムはありませんが、 SYCL* では バッファー （英語） / アクセサー （英語） とポインターベースの統合共有メモリー (USM) （英語） への割り当てという形で提供されています。4 ホストとデバイス間のデータ転送は、 C++ クラスまたは使い慣れた動的メモリー割り当て構文を使って行われます。

2「C++ vs. Rust」 で検索すると、 議論を見ることができます。
「フリーランチ」 とは、 周波数スケーリングからコア ・ スケーリングへの移行を指します。
「ハードウェア ・ ジャングル」 とは、 専用プロセッサーの普及を指します。
4
S YCL* は、 画像データに特化したクラスも提供していますが、 画像の取り扱いは本記事の範囲外です。
3© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

8

SYCL* は、 アクセラレーター ・ デバイスをサポートする拡張を備えた C++ であり、 C++ テンプレート ・ ライブラリーの新しいセットを追加するだけです （図 1）。 これは、 一般的な SAXPY （単精度の A × X ＋ Y） 計算を C++ で実装した例 （左） と SYCL* で実装した例 （右） を並べるとよく分かります。
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24

#include <vector>#include <CL/sycl.hpp>

int main(){size_t N{...};
ﬂoat A{2.0};

int main(){size_t N{...};
ﬂoat A{2.0};

std::vector<ﬂoat> X(N);
std::vector<ﬂoat> Y(N);
std::vector<ﬂoat> Z(N);sycl::queue Q;
auto X = sycl::malloc_shared<ﬂoat>(N, Q);
auto Y = sycl::malloc_shared<ﬂoat>(N, Q);
auto Z = sycl::malloc_shared<ﬂoat>(N, Q);

for (int i = 0; i < N; i++){// X、 Y、 Z ワーク配列を初期化

for (int i = 0; i < N; i++){// X、 Y、 Z ワーク配列を初期化}}

for (int i = 0; i < N; i++){Z[i] += A * X[i] + Y[i];}Q.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> i){Z[i] += A * X[i] + Y[i];});Q.wait();}}

2 つのコードはよく似ています。 C++ コードは、 SAXPY 計算をアクセラレーターにオフロードするように変更されています。
• 行 1 ： SYCL* ヘッダーをインクルードします。
• 行 8 ： デフォルトのアクセラレーター用の SYCL* キューを作成します。 アクセラレーターが利用できない場合、キューにサブミットされたワークはホストで実行されます。
• 行 9–11 ： ワーク配列を USM に割り当てます。 SYCL* ランタイムは、 ホストとキュー ・ オブジェクトで指定されたアクセラレーターの間のデータ転送を処理します。
• 行 19 ： SAXPY 計算をキューの引数で指定したアクセラレーター上で並列に実行する C++ ラムダ関数としてサブミットします。
• 行 23 ： ヘテロジニアス並列プログラミングに対する一部の独自アプローチとは異なり、 SYCL* は非同期のイベン
トベースの言語です。 SYCL* キューにワークをサブミットしても、 ホスト （またはほかのアクセラレーター） が有用なワークを続行するのを妨げることはありません。 wait() メソッドは、 アクセラレーターが SAXPY 計算を完了する前にホストが終了するのを防ぎます。

oneAPI は生産的な高速化アプローチであると前述しましたが、 それを証明する例を示します。 STL は C++
プログラマーの生産性を大幅に向上します。 oneAPI データ並列 C++ ライブラリー （oneDPL） は、 C++ STL のoneAPI により高速化された実装です。 次の例は、C++ STL （左） と oneDPL （右） を使って maxloc リダクションを実行します。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

9

#include <iostream>#include <vector>#include <algorithm>#include <iostream>#include <oneapi/dpl/algorithm>#include <oneapi/dpl/execution>#include <oneapi/dpl/iterator>

int main(){
std::vector<int> data{2, 2, 4, 1, 1};

int main(){
std::vector<int> data{2, 2, 4, 1, 1};

auto maxloc = max_element(

auto maxloc = oneapi::dpl::max_element(
oneapi::dpl::execution::dpcpp_default,
data.cbegin(),
data.cend());

data.cbegin(),
data.cend());
std::cout << "Maximum at element "<< distance(data.cbegin(),maxloc)<< std::endl;}

std::cout << "Maximum at element "<< oneapi::dpl::distance(data.cbegin(),maxloc)<< std::endl;}

この 2 つのコードも似ています。
• 行 2–4 ： 必要な oneDPL ヘッダーをインクルードします。
• 行 10 と 16 ： STL の max_element() 関数と distance() 関数の oneDPL 実装を呼び出します。
• 行 11 ： 実 行ポ リシー を 指 定し ま す。 ここでは、 デ フォルト のアク セラレ ー ター のポ リシー を 指 定し ま す。
アクセラレーターが利用できない場合、 キューにサブミットされたワークはホストで実行されます。

oneDPL コードには、 明示的なホストとデバイス間のデータ転送はありません。 ランタイムによって暗黙的に処理されます。 これは、 oneAPI がヘテロジニアス・アクセラレーションに生産的で移植性の高いアプローチを提供するという私の主張を支持するものです。 SYCL* プログラムは、 SYCL* コンパイラーとランタイムがあればどのシステムでも動作します （図 2）。

図 1. SYCL* はヘテロジニアス ・ アクセラレーターをサポートする
追加のテンプレート ・ ライブラリーを含む C++ （出典 ： Khronos.org/sycl）© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

10

図 2. SYCL* コンパイラー ・ エコシステムは拡大している （出典 ： Khronos.org/sycl）

C++ はプログラマーの期待と安定性への要求が定着している言語であるため、 標準への変更はゆっくりと慎重に行われます。 ヘテロジニアス並列処理は、 将来の C++ バージョンがサポートしなければならない多くの新機能の 1 つに過ぎません。 P2300R5 std::execution の提案 （英語） は、 「汎用実行コンテキストでの非同期実行を管理するフレームワーク」 と 「スケジューラー、 送信者、 受信者に基づく非同期モデル」 を定義することで、 ヘテロ
ジニアス性に関する C++ の現在の制限を解決するのに役立つ可能性があります。 この提案は、 C++26 の Library
Working Group に送付されました。 一方、 図 2 の右上のコメントにあるように、 SYCL* は、 ISO C++ 標準が進化するまでの間、 ヘテロジニアス並列処理を表現する柔軟で非独占的な方法を提供してくれます。
SYCL* プログラミングについてもっと学びたい方は、 以下のリソースが役立つでしょう。
• SYCL* 2020 仕様 （英語 | 日本語） と SYCL* 2020 API リファレンス ・ ガイド （英語 | 日本語） は SYCL* コードを記述する際に役立ちます。
• oneAPI は生産的なプログラミング手法であると前述しましたが、 「oneAPI によるフーリエ相関アルゴリズムの実装」 （記事 | ウェビナー （英語） | ソースコード （英語）） を参照するとその理由が分かります。 数行の SYCL* と
oneMKL コードで複雑なアルゴリズムを高速化する方法を紹介しています。
• SYCL* バッファーとアクセサーを使用したホストとデバイス間のデータ転送については、 「oneAPI の maxloc リダクション」 （記事 | ウェビナー （英語） | ソースコード （英語）） が参考になります。 oneDPL のコンテキストで、暗黙のバッファリング、 明示のバッファリング、 USM を使ったホストとデバイス間のデータ転送を説明しています。
• 非同期のヘテロジニアス並列 SYCL* コードの例外処理については、 『Data Parallel C++ ： Mastering DPC++ for
Programming of Heterogeneous Systems using C++ and SYCL』（英語） の 「第 5 章：エラー処理」（英語）が参考になります。 この本は、 SYCL* に興味のあるすべての方にお勧めです。
• James Reinders の HPCwire の記事 「SYCL* でヘテロジニアス ・プログラミングの課題を解決する」。
• 「SYCL* の基本」 （英語） オンライン ・ チュートリアル。
• oneAPI サンプル （英語） リポジトリーには、 SYCL* と oneAPI プログラミングを説明する数多くのサンプル
コードがあります。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

11

C++ Thrust アプリケー
ションを SYCL* と oneAPI
DPC++ ライブラリー（oneDPL） へ移行する
oneAPI で C++ アプリケーションを現代化および高速化し、ベンダーへの依存を排除する

Pablo Reble インテル コーポレーション ソフトウェア ・ エンジニア

プログラマーはライブラリーが大好きです。 単に怠け者だからというだけでなく、 ライブラリーが生産性とパフォー
マンスを向上するからです。 よくあるパターンを最適化するのに時間を浪費し、 ターゲット ・ ハードウェアが変更
されるたびにまた同じことをするのは無意味です。 パターンを特定する作業はすでに終わっているため、 ISO C++
標準テンプレート・ライブラリー （STL） のような事前定義の抽象化機能を使うべきです。 これらはすぐに使えて、専門家によって最適化されています。
STL は CPU だけをターゲットにしているのであれば最適ですが、コンピューティングの世界はヘテロジニアスです。
CPU だけでなく、 それ以外の選択肢も必要です。 NVIDIA* Thrust ライブラリーと oneAPI データ並列 C++
ライブラリー （oneDPL） は CPU 以外にも対応しています。 この 2 つを比較してみましょう。 どちらもオープン
ソース ・プロジェクト1、2 ですが、 Thrust のデバイスサポートはプロプライエタリーなソフトウェア ・ コンポーネントに依存しています （CUB*/CUDA* を 通じての み GPU をサポートしています）。 つまり、 ベンダー 依存 で す。
一方、 oneDPL は SYCL* をベースにしており、 複数のベンダーのアクセラレーターをサポートするように設計されています。31
2
3oneDPL リポジトリー ： https://github.com/oneapi-src/oneDPL （英語）
Thrust リポジトリー ： https://github.com/NVIDIA/thrust （英語）「SYCL* の事例 ： ISO C++ がヘテロジニアス ・ コンピューティングに十分でない理由」© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

12

Thrust からベンダーに依存しない oneDPL への移行を検討するのは当然と言えます。 ここでは、 2 つの簡単な例
を使用して、 ツールを使用するアプローチと手動でコーディングし直すアプローチの 2 つの補完的な移行戦略について見ていきます。 CUDA* を SYCL* に移行する SYCLomatic ツールについては、 皆さんすでにご存知かもしれません。4 ツールを使用した移行には限界があり、 特にテンプレート化されたコードでは、 必ずしも機能的ではなく、 生成されるコードの品質や保守性には限界があります。 例えば、 ラムダ関数や自動型推論のような新しい
C++ の機能を導入する場合、 手作業による編集やコードの現代化が通常必要になります。 移行後のコードは、
オープン・スタンダードに準拠した、最新の C++ に沿ったものになります。 Thrust は登場から 10 年が経っており、 そろそろコードをアップグレードする時期と言えます。 いくつかのコード例を見てみましょう。
最初の 例は、 シーケンス中のいくつかの 値をインプレース変換します。 変換する値はマスクで定義します。
この関数は、 値を格納するストリームと、 ステンシルシーケンスを格納するストリームの 2 つの入力ストリームを受け取ります。 最初のシーケンスの各値は、 2 番目の入力シーケンスのステンシル値が特定の条件 （コード例
では非ゼロのプレディケート） を満たす場合に変換 （コード例では符号反転） されます。
// Thrust コード例 ( ステージ 0)thrust::transform_if(dev_inp,
dev_inp + 10,
dev_stencil,
dev_inp,
thrust::negate<int>(),thrust::identity<int>());

oneDPL は標準 C++ アルゴリズムと密接に連携しています。 STL はステンシルのオーバーロードをサポートしていないため、 oneDPL も同様にサポートしていません。 このパターンは別の方法で表現する必要があるため、
SYCLomatic は、 ドロップイン置換として代替案を提供しています。
// 移行後のコード ( ステージ 1)
dpct::transform_if(oneapi::dpl::execution::make_device_policy(syclQueue),
dev_inp,
dev_inp + 10,
dev_stencil,
dev_inp,
std::negate<int>(),dpl::identity<int>()); // identity は C++20 の機能

ラムダ関数をカスタム ・ファンクターとして使用し、 標準 C++ コードを作成することで、 さらに改善できます。
この方法では、 ステンシルシーケンスが一般的な入力シーケンスとは異なる方法で処理されます。 プレディケートの評価はカスタム ・ファンクターの内部で行われ、 ライブラリーの内部に隠されることはありません。
// 手動編集後 ( ステージ 2)
dpl::transform(dpl::execution::dpcpp_default,
dev_inp,
dev_inp + 10,
dev_stencil,
dev_inp,[&](const auto& input, const auto& mask){
return mask ? std::negate<>()(input) : input;});

4

SYCLomatic リポジトリー ： https://github.com/oneapi-src/SYCLomatic （英語）。 この記事で紹介した SYCLomatic の例は、説明を目的としたものであることに注意してください。 インテル® DPC++ 互換性ツールなどの実装では、 実際の結果は異なる場合があります。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

13

これで、 SYCL* デバイスをターゲットにするために必要なカスタム実行ポリシーを除いて、 STL の動作のみを使用してパターンを表現できました。 これは、 Thrust と oneDPL を区別する重要な点です。 Thrust には、
入力イテレーターがアルゴリズムの実行場所を定義するデフォルトモードがあります。 これは、 実行パラメーター
がイテレーターではなく、 実行ポリシーによって定義される ISO C++ とは異なります。 oneDPL は標準 C++
アルゴリズムと密接に連携し、 カスタムポリシーを提供します。 それを使って、 アルゴリズムの実行場所を定義できます。 便宜上、 デフォルトの SYCL* デバイスをターゲットとする定義済みポリシーも用意されています。
どちらのポリシーも、 上記のステージ 1 およびステージ 2 のサンプルコードで使用されています。
次の移行例では、 別のよくあるパターンであるソートについて見てみましょう。 Thrust バージョンはキーでソートします。 このアルゴリズムは 2 つの入力シーケンスを受け取ります。 1 つは値を、 もう 1 つはキーを保持しています。
// Thrust コード例 ( ステージ 0)thrust::stable_sort_by_key(dev_values, dev_values + 10, dev_keys);

値の入力シーケンスは、 キーのシーケンスの要素を比較してソートされます。 このパターンは既存の ISO C++関数にはないため、 SYCLomatic は機能的な代替案を提供します。
// 移行後のコード ( ステージ 1)
dpct::stable_sort(oneapi::dpl::execution::make_device_policy(syclQue),
dev_values,
dev_values + 10,
dev_keys);

カスタムのイテレーターとファンクターを使用することで、 さらに改善できます。 ZIP イテレーターを使用してキーと値のシーケンスを結びつけ、 単一の反復処理空間を作成します。 カスタム ・ファンクターは、 値とキーのペアでキー要素のみを比較します。 その結果、 若干冗長になりますが、 カスタム ・ イテレーターのような汎用的な拡張を含む
ISO C++ 関数のみを使用します。
// 手動編集後 ( ステージ 2)
dpl::stable_sort(dpl::execution::dpcpp_default,
make_zip_iterator(dev_values, dev_keys),
make_zip_iterator(dev_values, dev_keys) + 10,[](const auto& a, const auto& b){
return get<1>(a) < get<1>(b);});

高速化された C++ STL アルゴリズムを実装したライブラリーは、 開発者の生産性を大幅に向上できます。 ツールによる移行ステップを手動コード編集で補完することで、 保守性とコード品質を向上できます。 将来の開発に向けて、 範囲ベースの API と P2300R5 std::execution の提案 （英語） は、 現在の C++ の制限に対応し、 この記事で説明したような並列パターンを表現する有力な代替手段です。 ISO C++ でこれらの機能を実現するには、 既存のアプリケーションを大幅に作り直す必要があります。今のところ、プロプライエタリーな言語やベンダー依存を避け、
マルチベンダーのハードウェアに対応するには、 可能な限り ISO C++ に従い、 拡張機能を一般的なものにするのが最善です。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

14

DPEcho: 2020 年代以降の
SYCL* による一般相対論
科学計算における SYCL* の優位性

Salvatore Cielo ライプニッツ ・ スーパーコンピューティング ・ センター 宇宙プラズマ物理学アプリケーション ・
サポート担当 HPC エキスパート
Alexander Pöppl インテル コーポレーション ソフトウェア ・ イネーブリング & 最適化エンジニア
Luca Del Zanna フィレンツェ大学 物理および天文学部准教授
Matteo Bugli トリノ大学 天体物理学研究者兼 MSCA フェロー

コードは複雑であるものの、 定量的かつ定性的に大きな問題クラスのシミュレーションを可能にするヘテロジニア
ス ・ コンピューティングのおかげで、 数値科学が見直されています。 SYCL* （英語） プログラミング言語は、 拡張性、 移植性、 オープン性を備えたヘテロジニアスへの標準アプローチを提供し、 ワークシェアリングやスケジューリ
ング、 シミュレーション領域を実行空間に直接マッピングする直感的な C++ API により、 アクセラレーターの能力を引き出します。 後者は、 空間と時間の性質が運動方程式と緊密に結合し、 計算量とメモリー消費量が多い数値
相対論において、 特に便利です。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

一般相対論的磁気流体力学（GR-MHD）の ECHO コード1 には、
最近の HPC 向けに最適化されたバージョン （ECHO-3DHPC2）がすでにあります。 これは Fortran で記述されており、 MPI +OpenMP* でハイブリッド並列処理を実装しています。 この記事
では、 このコードに手を加える代わりに SYCL* で書き換えた、
GR-MHD の MPI + SYCL* バージョンである DPEcho を紹介します。 DPEcho は、 不安定性、 乱流、 波動の伝播、 恒星風と磁気圏、 およびブラックホール周辺の天体物理学的プロセスのモデル化に使用され、 ミンコフスキーやコード化された一般
相対性理論 （GR） メトリックで従来の MHD と相対論的 MHDをサポートしています。 ここでは、 DPEcho の公開バージョンを紹介し、 SYCL* で複雑な現実のアプリケーションのアルゴリズムを表現する方法を説明し、 このアプローチによって達成される
パフォーマンスを定量化します。

DPEcho を使用する

15

CMAKE_BUILD_TYPE
SYCL
oneAPIGPUON
ENABLE_MPISINGLE_PRECISIONFILE_IO
VISIT_BOVPHYSICSGRMHDMETRICCARTESIAN
FD
4
NRK
2
REC_ORDER

Release

OFF

OFF

5

SYCL: Select the SYCL target
architecture
Keys: [enter] Edit an entry [d] Delete
an entry[l] Show log output[c]
Configure[h] Help[q] Quit
without generating[t] Toggle advanced mode
(currently off)

DPEcho の公開バージョン （英語） は、Apache* 2.0 ライセンスの下に GitHub* で公開されています。 このコードは CMake をサポートしており、 ccmake のインタラクティブなターミナル ・ユーザー ・ インターフェイスを意識して設計されています （図 1）。
DPEcho は、 このインターフェイスを介していくつかの SYCL*
コンパイラーをサポートしています。 ccmake 画面にコードの図 1. DPEcho のインタラクティブな内容 （SYCL* キューセレクターと MPI による GPU サポート、 ccmake ターミナル ・ ユーザー ・ インターフェイス
単精度、I/O トグル）が表示されます。 次に、物理と数値のオプションがあります。 ソルバーは従来の MHD と GR-MHD の切り替えが可能で、 関心のある問題に応じて、 サポートされている GR メトリックの 1 つを選ぶことができます （公開バージョンではデカルトまたはミンコフスキーのみサポート）。 数値オプションには、オリジナルの ECHO コードと同様に、数値導出 （FD）、可変グリッド再構成 （REC）、
ルンゲ ・ クッタ ・ ソルバー ・ステップ （NRK: 1-3） の異なるオーダーが含まれています。

1

2"ECHO: An Eulerian conservative high order scheme for general relativistic magnetohydrodynamics andmagnetodynamics," Apr. 2007, doi: 10.1051/0004-6361:20077093.
"ECHO-3DHPC: Advance the performance of astrophysics simulations with code modernization," Oct. 2018,
doi: 10.48550/arxiv.1810.04597.

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

16

DPEcho の開発、 解析、 インテルの CPU および GPU での実行には、 インテル® oneAPI ベース ・ ツールキットとインテル® oneAPI HPC ツールキットを使用しました。 ユーザーは、MPI ランクを CPU コアに固定するのと同様に、
MPI ランクと GPU リソース （タイルなど） のマッピングを指定できます。3 これは実行時に環境変数を通じて行われます。 パフォーマンス解析には、 インテル® VTune™ プロファイラーを使用しました。 最後に、 インテル® MPI
ライブラリーは、 統合共有メモリー （USM） に配置されたメモリーセグメントなどを直接操作することができます。
DPEcho の公開バージョンに存在する物理および数値手法により、 図 2 に示す平面アルヴェーン波のような相対論的磁化波のモデル化が可能です。
異なる問題 （GR-MHD 爆風など） をモデリング
す る イ ンターフェイスも 用 意 さ れて い ま す が、
ユーザー定義の実装が必要です。 DPEcho の開発バージョンには、 回転する天体物理学のブラック
ホールの周囲のモデリングを目的としたさまざまな GR メトリックや、 非平面時空間を扱う一般的なインフラストラクチャーが含まれる予定です。
これらの 機能は一般公開の予定はありません。
開発協力のリクエストには、 個別に対応します。
この点を除いて 2 つのバージョンは類似しており、互換性のあるインターフェイスを持ちます。
図 2. DPEcho 公開バージョンでモデル化されたGR （x 速度成分） での平面アルヴェーン波のレンダリング。
矢印は磁場 （大きさは一様） をトレースします。

DPEcho の主な構造
DPEcho 開発中に最も苦労したのは、 SYCL* のデバイス中心の構造をサポートする物理アルゴリズムと数値
アルゴリズムの再設計でした。 図 3 は、DPEcho のホストとデバイスのコード、 通信、およびフロー制御を示します。
慎重な計画と USM の採用により、 実行時間全体にわたって主要な変数をデバイスメモリーに保持することができました。 これにより、（MPI デカルトグリッドでの分解に続く） 各 MPI タスクのドメインの周りのゴーストレイヤーが、最小サイズのバッファーを介してコピーされ、 データ転送が大幅に軽減されます。 さらに処理するため、 デバイス上のデータをロックしないように、 出力ダンプも同様の原則に従います。 単純な SYCL* 制御構造 （例えば、 ワーク
グループ、 共有ローカルメモリー、 およびキューの依存関係を除外する） のみを使用しても、 データ転送ロックをバイパスし、読みやすいソースコードで高いパフォーマンスを実現できます。 将来、さらに制御が必要になった場合、これらの高度な機能は、 大きな問題なしに、 効率良くベースラインに組み込まれます。 これも SYCL* の主な利点です。

3"Intel MPI Library Developer Reference - GPU Support." https://www.intel.com/content/www/us/en/develop/documentation/
mpi-developer-reference-linux/top/environment-variable-reference/gpu-support.html?wapkw=GPU%20pinning （英語）© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

17

図 3. DPEcho のロジックの概要

USM の使用法を図 4 に示します。 まず、 計算をオフロードするターゲットデバイスを選択します。 ここでは、
SYCL* ランタイムによって選択されるデフォルトの GPU を使用します。 未定ベクトルの配列構造体として、 データ
をデバイスメモリーに直接割り当てます。 各ベクトル ・コンポーネントを個別の配列に割り当てることで、 ランタイム
がタイル化された GPU アーキテクチャーでデータを正しく分割できるようになります。
アプリケーションの計算の大部分は、 磁束計算カー
ネルで実行されます （図 5）。 磁束とは別に、 反復処理の以降のステップでは、 安全なタイムステップの増分を決定するため、 最大波動伝播速度の計算も必要です。 SYCL* はリダクション ・ カーネルによってこれをサポートします。 parallel_for は、 計算の形状に加えて、 共有リダクション ・ オブジェクトを受け取ります。 デバイスの実行コンテキストでは、各ワークアイテムがローカル波速をサブミットします。 最後に、 SYCL* ランタイムがグローバルの最大値を決定します。

// SYCL* デバイスとキューを初期化
sycl::gpu_selector sDev; sycl::queue qDev(sDev);// USM で主要な変数を割り当て
double *v[VAR_NUM], *f[VAR_NUM], [...];for (int i=0; i < VAR_NUM; ++i) {// 各変数のプリミティブと磁束v[i] = malloc_device<double>(grid.numCells, qDev);
f[i] = malloc_device<double>(grid.numCells, qDev);[...]}

図 4. 直接デバイスアクセスのため SYCL* USM で
DPEcho の主要な変数を割り当て© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe18//-- SYCL* の ranges: グリッドサイズとローカル ・ ワークグループ ・ サイズ
range<3> r = range(grid.nx, grid.ny, grid.nz), rLoc = range(8, 8, 8);[..]//-- メインのコードループ
auto maxReduction = sycl::reduction(aMax+directionIndex, sycl::maximum<ﬁeld>());
qDev.parallel_for(nd_range<3>(r, rLoc), maxReduction, [=](nd_item<3> it, auto &max) {id<3> id
= it.get_global_id();// SYCL* のインデックスとオフセット
int myId= globLinId(id, grid.nh, dOﬀset); // ゴーストレイヤーのカスタム ・ インデックス
int dStride = stride(id, myDir, grid.nh);// GPU レジスターでローカル変数を宣言
double vRecR[VAR_NUM], vRecL[VAR_NUM] , ...;
Metric g(xCenter, yCenter, zCenter);// GR の心臓部 : 各グリッドポイントのメトリック// 各変数を補間する空間を意識した関数 -> 並列化を考慮 !
holibRec(myId, v, dStride, vRecL, vRecR);
physicalFlux(directionIndex, g, vRecL, vRecR, ...);[ ... ] // 磁束から特性を計算
max.combine(localMax); // タイムステップ

図 5. DPEcho のデバイスコード例 : 磁束計算を行う parallel_for 構文 （分かりやすくするため一部省略）

カーネルでは、 各ワークアイテムが 1 つのグリッドセルの磁束を計算します。 ワークアイテムのグローバル ID を
グリッドセルにマッピングするため、 カスタム 「計算機クラス」 の grid を使用し、 ゴーストレイヤーをオフセットします。 最初に Metric オブジェクトを作成します。 Metric オブジェクトのコンポーネントは各セルでローカルに定義され、 解析的に計算されるため、 複数の追加変数を保存する必要はありません。 次に、 磁束計算のため、カーネルはまずセル境界で未定値を補間する必要があります。 ここでは、 コンパイル時に設定される、 異なる
オーダーの再構成をサポートしています。 この補間は holibRec で実行されます。 この関数は、 ワークアイテム
ID を使用して、 シミュレーション ・ グリッド内の位置を決定します。 現在は実装されていませんが、 この機能は隣接するワークアイテム間でワークシェアリングを可能にします。例えば、セル中央の値をセル間でブロードキャスト
して、 グローバルメモリーからの冗長なロードを回避できます。 補間後の磁束計算は純粋にローカルであり、 驚異的な並列性をもたらします。

パフォーマンスの比較
同じ問題設定を実行して、 DPEcho とベースライン ECHO のスケーリングを比較します （図 6）。 最初に、ライプニッツ ・ スーパーコンピューティング ・ センターにあ る SuperMUC-NG （英 語） の インテル® Xeon®
Platinum 8174 プロセッサー ・ ベースのノードで、 弱いスケーリング ・ テストを実 施します （図 6 の左側）。
16 ノードまでのスケーリングは完璧であり、 MPI + OpenMP* のハイブリッドなベースライン ・ バージョンは最大 4 倍のパフォーマンス （1 ウォールクロック秒に 1 ステップで処理される 100 万単位のセル更新数で測定） を達成していることが分かります。 DPEcho のメモリー ・フットプリントが軽減されたことで、 MPI タスク
ごとに 1923 個ではなく 3843 個のセルを配置できました。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

19

図 6. ECHO （Fortran バージョン） と DPEcho の比較。
左 ： HPC ハードウェア （ライプニッツ ・ スーパーコンピューティング ・ センターの SuperMUC-NG のインテル® Xeon® Platinum 8174 プロセッサー ・ ベースのノード） での倍精度、 CPU のみの弱スケーリング ・
テスト。 4 倍のスピードアップに加え、 メモリー消費量の減少により、 シミュレーションを最大 8 倍まで拡大可能。
右 ： コンシューマー向けハードウェア （DG1 搭載の第 9 世代インテル® Core™ i9 プロセッサー） でグリッドサイズを大きくして単精度テストを実行。 ディスクリート ・ グラフィックスは十分な大きさのワークロードを必要とするが、高速化されていない HPC ハードウェアに匹敵するパフォーマンスを達成。

HPC ハードウェアのほかに、 コンシューマー向けのハードウェア （インテル® Iris® Xe MAX グラフィックス （開発コード名 DG-1） 搭載のインテル® Core™ i9-9900K CPU） でもテストしています。 このテストでは、 グリッドサイズのスケーリングを行い （図 6 の右側）、 常にフルノードまたはデバイスを使用します （ハードウェアの制限により単精度に切り替えています）。 CPU でのベースライン ・ パフォーマンスは、 HPC の結果と一致しています。 DG-1 によりパフォーマンスは最大 7 倍向上し、 高速化されていない HPC ハードウェアを上回っています。 これは、 SYCL*
コードの優れた移植性と効率性を証明するものであり、 あらゆるクラスのハードウェアを最も効率良く利用することができます。

今後の展望と考察
DPEcho には、磁場のソレノイド状態を維持する方法はまだ実装されていません。 オリジナルの ECHO コードでは、磁場成分のスタッガリングと磁束の多次元近似リーマンソルバーに基づく戦略が採用されています。 これはかなり複雑であるため、 コミュニティーには SYCL* が基本的な役割を果たす、 シンプルで拡張性のあるコードを提供しています。 理想的でない効果 （例えば、 誘導方程式における物理的な粘性、 抵抗、 ダイナモの項） などのより複雑な物理オプションについても、 同様の理由から実装されていません。 開発バージョンでは、 円盤からカー ・
ブラックホール解への降着シミュレーションがまもなくテストされる予定です。
個々のアップグレードがすべて実装されれば、 DPEcho の可能性は確実に高まるでしょう。 超巨大ブラックホールや
恒星サイズのブラックホールからの降着流入と平行噴射の 3D 高解像度シミュレーションを詳細に研究し、衝撃波捕獲コードにつきものの数値拡散の影響を最小限に抑えながら、 理想的な磁気回転不安定性と小さな再結合サイトの非理想的なプラズモイド不安定性の両方を解決できる可能性があります。 当面は、 公開バージョン
で実験することを歓迎します。 開発協力に興味のある方は、 著者にご連絡ください。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

20

Fortran、 oneMKL、
OpenMP* を使用してLU 因数分解を高速化
一般的な数学演算を高速化する標準ベースのオープンなアプローチ

Henry A Gabb インテル コーポレーション シニア主席エンジニア兼 The Parallel Universe 編集長

BASIC、 Pascal、 C は私が最初に学んだプログラミング言語ですが、 初めて本格的な （教室外での） コーディングをしたのは Fortran でした。 インテルに入社して C と C++ に乗り換えるまでの 12 年間は、ほとんど Fortran しか使いませんでした。 この 20 年間、 Fortran を 1 行も書いていないことを、 私はいくつかの著作で嘆いています。
旧友の Jim Cownie が 「Fortran は過去の言語なのか ?」 （英語） において、 そうではないという説得力のある答えを出すまで、 Fortran がまだ使われているかどうか疑問に思っていたくらいです。 そのため、 同僚から最新のFortran コンパイラーの OpenMP* アクセラレーター ・ オフロード ・ サポートを試してみないかと誘われたとき、私はそのチャンスに飛びつきました。
私は、 Fortran と同じくらい OpenMP* が好きです。 OpenMP* は常にコードを並列化する簡単な方法であり、1997 年に並列コンピューティング分野に登場して爆発的に広がって以来、 OpenMP* 仕様は長い道のりを歩んできました。 現在では、 アクセラレーターへのオフロードもサポートしています。 そこで、 いつも使っている SYCL* ではなく、 並列プログラミングの原点に戻って、 Fortran、 OpenMP*、 インテル® oneAPI マス・カーネル・ライブラリー（インテル® oneMKL） を使ってアクセラレーターに計算をオフロードする方法を説明したいと思います。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

21以前の記事 （記事 （英語）、 ウェビナー （英語）、 サンプルコード （英語）） を読んでいただいた方は、 私がフーリエ相関アルゴリズムを気に入っていることをご存じだと思います。 私はこれを生体分子のドッキングの研究に使い、最近では 2D 画像と 3D 画像の調整に使用しています。 これまで、 このアルゴリズムにはうんざりするくらい取り組んできたので、 ここでは線形代数計算 （特に LU 分解） をアクセラレーターにオフロードすることがいかに簡単か見てみましょう。 LU 分解について詳しく説明することがこの記事の目的ではないので、 線形方程式を解くために使われるものだと理解いただければ十分です。 詳しく知りたい場合は、 インテル® MKL クックブック （英語） の「LU 因数分解されたブロック三重対角係数行列を含む連立線形方程式を解く」 （英語） と 「一般的なブロック三重対角行列の因数分解」 （英語） を参照してください。 ここでは、 Fortran、 インテル® oneMKL、 および OpenMP*の実装を通して、 オフロード構文を説明し、 ヘテロジニアス並列処 理におけるホストとデバイス間のデータ転送に注目します。 まず、インテル® oneMKL に含まれている dgetri_oop_batch_strided.f90 の例から開始し （「サンプルコードの使用」 （英語））、 効率を改善するためこれを変更します。
このプログラムは、 まず OpenMP* オフロードをサポートする Fortran インターフェイスをインクルードし、 次に
32 ビット整数型と 64 ビット整数型のどちらを使用するかを決定します （詳細は、 「ILP64 インターフェイスと LP64 インターフェイスの使用」 （英語） を参照） （図 1）。 この判断は、 この記事の最後にあるコンパイラー ・コマンドで示すように、 コンパイル時に行います。 セットアップを完了するため、 因数分解する行列の数とそのサイズを設定し、 インテル® oneMKL サブルーチンが必要とする次元とストライドのパラメーターを設定し、 ワーク
配列を割り当てます。 これで計算を開始する準備が整いました。
include "mkl_omp_oﬄoad.f90"
program sgetri_oop_batch_strided_example#if deﬁned(MKL_ILP64)
use onemkl_lapack_omp_oﬄoad_ilp64#else
use onemkl_lapack_omp_oﬄoad_lp64#endif! 64-bit integer type! 32-bit integer type

integer,parameter :: n = 10, batch_size = 4
integer:: lda, ldainv, stride_a, stride_ainv, stride_ipiv
real,allocatable :: a(:,:), ainv(:,:)
integer, allocatable :: ipiv(:,:), info(:)
lda= nldainv= n
stride_a= lda* nstride_ainv = ldainv * n
stride_ipiv = n
allocate(a(stride_a, batch_size), ainv(stride_ainv, batch_size), &ipiv(stride_ipiv, batch_size), info(batch_size))

図 1. サンプルプログラムのセットアップ。 OpenMP* ターゲットオフロードをサポートするインテル® oneMKL ヘッダーとモジュールは、 青色で表示しています。 ハードコードされた行列の次元とバッチサイズのパラメーターは、
テストには適していますが、 高速化には小さすぎるということに注意してください。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

22

手元にある GPU は単精度の計算しかサポートしていないので、 Fortran REAL 型と適切な LAPACK （英語）サ ブ ル ー チ ン を 使 用 し て、 LU 因 子 分 解 (sgetrf_batch_strided （ 英 語 ） と LU 因 子 分 解 し た 行 列 のアウトオ ブ プ レ ース 反 転 （sgetri_oop_batch_strided （ 英 語）） を 計 算し ま す （倍 精 度に 変 換 す るには、
Fortran DOUBLE PRECISION 型 に 切 り 替 え、 LAPACK の dgetrf_batch_strided サ ブ ル ー チ ン とdgetri_oop_batch_strided サブル ー チン を 使 用するだ け で す）。 オリジナル のコード では、 2 つ のOpenMP* target data 領域を使用しています （図 2）。 ターゲット構造は、 デバイスに制御を転送します。
最初の領域は、 入力行列をデバイスメモリーに転送し ［map(tofrom:a)］、 インプレース LU 分解 （sgetrf_
batch_strided） を デバイスにディス パッチし、 LU 分 解した行列 （a に格 納）、 ピボット ・ インデックス[map(from:ipiv)］、 ステータス ［map(from:info)］ をデバイスメモリーから取得します。 因数分解が正常に実行されると、 プログラムは次の OpenMP* 領域に進みます。 LU 因数分解された行列とピボット・インデックス
はデバイスメモリーに戻され ［map(to:a) と map(to:ipiv)］、 アウトオブプレース反転はデバイス上で計算され （sgetri_oop_batch_strided）、 逆行列 ［map(from:ainv)］ とステータス ［map(from:info)］
はデバイスメモリーから取得されます。
! OpenMP* オフロードで LU 因数分解を計算! 開始時に A は入力行列を格納! 終了時に A は LU 因数分解された行列を格納!$omp target data map(tofrom:a) map(from:ipiv) map(from:info)!$omp dispatch
call sgetrf_batch_strided(n, n, a, lda, stride_a, ipiv, stride_ipiv, &batch_size, info)!$omp end target data
if (any(info .ne. 0)) thenprint *, 'Error: sgetrf_batch_strided returned with errors.'
else! OpenMP* オフロードで行列の反転を計算し、 終了時に逆行列は Ainv に格納!$omp target data map(to:a) map(to:ipiv) map(from:ainv) map(from:info)!$omp dispatch
call sgetri_oop_batch_strided(n, a, lda, stride_a, ipiv, stride_ipiv, &ainv, ldainv, stride_ainv, batch_size, info)!$omp end target data
endif
if (any(info .ne. 0)) thenprint *, 'Error: sgetri_oop_batch_strided returned with errors.'
else
print '("The calculations executed successfully.")'endif

図 2. 2 つの LAPACK サブルーチンそれぞれに OpenMP* ターゲット領域を使用して、 LAPACK 計算をアクセラレーターにオフロード。
OpenMP* ディレクティブは青で表示しています。 LAPACK の呼び出しは緑色で表示しています。

このサンプルプログラムの 作者は、 反転に進む前に因数 分解 の 結果 を 確 認するため、 2 つ の OpenMP*
ターゲット領域を使用しているのでしょう。 このプログラマーの注意深さは賞賛に値しますが、 追加のホストとデバイス間のデータ転送が発生するため、 ヘテロジニアス並列プログラミングでは避けるべきです。 図 3 の矢印
はすべて、 不連続メモリー間のデータ移動を表しています。 例えば、 入力行列はデバイスに転送され、 因数は© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

23

ホストに戻されます ［map(tofrom:a)］。 そして、 同じ因数が次の OpenMP* target data 領域でデバイス［map(to:a)］ に再度転送されます。 同様に、 ピボット・インデックスは、 sgetrf_batch_strided がリターンした後、デバイス ［map(from:ipiv)］ から取得され、sgetri_oop_batch_strided への入力としてデバイス［map(to:ipiv)］ に戻されます。 ヘテロジニアス ・ コンピューティングを行うということは、 これらの行列はおそらく大きいため、 ホストからデバイスへ、 またはその逆に転送すると、 計算と電力のオーバーヘッドが大きくなる可能性があります。 1 次元のステータス配列はそれほど大きくないので、 map(from:info) の転送にそれほど悩まされることはないです。

図 3. 図 2 に示した 2 つの OpenMP* ターゲットオフロード領域のホストとデバイス間のデータ転送。
各矢印は、 ホストとデバイスのメモリー間のデータ転送を示します。

理想的には、 OpenMP* ランタイムが、 必要なデータがすでにデバイスメモリーにあるかどうかをチェックすることです。 しかし、 OpenMP* のターゲット領域間でそれらのデータがフラッシュされず、 上書きもされないという保証はありません。 また、 2 つの OpenMP* オフロード領域を作成することによるオーバーヘッドを考慮する必要があります。 このため、 1 つのオフロード領域ですべての計算を行うほうがよいでしょう （図 4）。
!$omp target data map(to:a) map(from:ainv) map(from:info_rf) map(from:info_ri)!$omp dispatch
call sgetrf_batch_strided(n, n, a, lda, stride_a, ipiv, stride_ipiv, &batch_size, info_rf)!$omp dispatch
call sgetri_oop_batch_strided(n, a, lda, stride_a, ipiv, stride_ipiv, &ainv, ldainv, stride_ainv, batch_size, info_ri)!$omp end target data
if (any(info_rf .ne. 0) thenprint *, 'Error: getrf_batch_strided returned with errors.'
elseif (any(info_ri .ne. 0) thenprint *, 'Error: getri_oop_batch_strided returned with errors.'
else
print '("Computation executed successfully")'endif

図 4. 1 つの OpenMP* ターゲット領域を使用して、 LAPACK 計算をアクセラレーターにオフロード。
OpenMP* ディレクティブは青で表示しています。 LAPACK の呼び出しは緑色で表示しています。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

24

2 つ の OpenMP* ター ゲット 領 域 を 1 つに 融 合すると、 ホ ストとデバイス 間 の デー タ転 送 が少なくなる
こと が分 か り ま す （ 図 5）。 入 力 行 列 （a） は デバイス に 転 送 さ れ、 イ ン プ レ ースで 因 数 分 解 さ れ ま す。
LU 因 子 と ピ ボ ット （a、 ipiv） は デバイ ス 上 で sgetrf_batch_strided に よ って 計 算 さ れ る の で、
sgetri_oop_batch_strided が呼び出されたときにはすでにデバイスメモリーにあります。 後で連立 1 次方程式を解くために使用する予定がない限り、それらをホストに転送する必要はありません。 ただし、その場合は、必要な入力データがデバイスメモリーにある間に、 同じ OpenMP* target data 領域で sgetrs_batch_
strided （英語） を呼び出すことができます。 これは読者の演習として残しておきます。

図 5. 図 4 の OpenMP* ターゲットオフロード領域におけるホストとデバイス間のデータ転送。
各矢印 （アクセラレーター ・ デバイス内を除く） は、 ホストとデバイスのメモリー間のデータ転送を示します。

この記事で説明するサンプルプログラムは、このリポジトリー （英語） から入手できます。 サンプルプログラムには、
入力行列を初期化し、 計算の各ステップの出力を表示するコードが含まれていますが、 この記事には含めませんでした。 インテル® Fortran コンパイラーとインテル® oneMKL を使用してサンプルプログラムをビルドする
コマンドは、 OpenMP* を用いる場合とそうでない場合のそれぞれで、 次のとおりです。
$ ifx -i8 -fpp -DMKL_ILP64 -qopenmp -fopenmp-targets=spir64 -fsycl \lu_fact_omp_oﬄoad_example.f90 -o lu_fact_omp_oﬄoad_example.exe \-L${MKLROOT}/lib/intel64 -lmkl_sycl -lmkl_intel_ilp64 -lmkl_intel_thread \-lmkl_core -liomp5 -lpthread -ldl
$ ifx -i8 -fpp -DMKL_ILP64 \lu_fact_omp_oﬄoad_example.f90 -o lu_fact_example.exe \-L${MKLROOT}/lib/intel64 -lmkl_intel_ilp64 -lmkl_intel_thread \-lmkl_core -liomp5 -lpthread -ldl

コードはどちらの場合も同じであり、 これは OpenMP* の主な 利点の 1 つです。 コンパイラーが OpenMP*
ディレクティブを無視しても、 シリアルコードはそのまま残ります。
インテル® oneMKL 関数は SYCL* メソッドを使用してアクセラレーター ・ オフロードを実行するため、 インテル®
Fortran コンパイラーおよびインテル® oneMKL とともに、 デバイスドライバーとデータ並列 C++ （英語） コンパイ
ラーをインストールして設定する必要があります。 最も簡単な方法は、 インテル® oneAPI ベース ・ ツールキットとインテル® oneAPI HPC ツールキットの両方をインストールすることです。 インテル® oneAPI を初めて使用する場合、またはアップデートを探している場合は、 インテル® oneAPI ツールキットを参照してください。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe25

図 1 に示す配列の次元とバッチサイズは、 テストには適していますが、 アクセラレーションの恩恵を受けるには小さすぎることに注意してください。 計算量は、 制御フローをターゲットデバイスに転送し、 不連続メモリー間でデータを転送する際のオーバーヘッドを相殺できる程度に大きくする必要があります。 このサンプルコードのパフォーマンス特性については、 今後の記事で取り上げますが、 配列の次元やバッチサイズを自由に試すことで、計算をアクセラレーターにオフロードせずにホスト・プロセッサーに残したほうが良い場合について、 考察できます。
ここでは、 OpenMP* のアクセラレーター・オフロード機能のほんの一部を紹介しました。 OpenMP* 仕様 （英語）には、 ほかにも多くのオフロード構造が記載されています。 「OpenMP* API を使用した GPU プログラミングの概要」 （英語） で、 初心者向けのチュートリアルを提供しています。 最新のインテルのハードウェアとソフトウェア
がインストールされている無料のインテル® デベロッパー ・ クラウド （英語） では、 OpenMP* アクセラレーター ・
オフロードの実験を行うことができます。

ポッドキャスト

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

アニメーター、 デジタルコンテンツ制作者、 建築エンジニア、 熟練ゲーマーの皆さん、インテル® oneAPI ベース & レンダリング ・ ツール キットでビジュアライゼーションの限界に挑戦してみてください。 詳細 >© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

聴く （英語）The Parallel Universe

26

ACM RecSys Challenge
2022 のインテル® SIHG4SR
ソリューション
インテルの推薦システムが上位入賞

Chendi Xue インテル コーポレーション マシンラーニング ・ エンジニア
Xinyao Wang インテル コーポレーション AI フレームワーク ・ エンジニア
Zhou Yu インテル コーポレーション AI フレームワーク ・ エンジニア
Jian Zhang インテル コーポレーション AI ソフトウェア ・ エンジニアリング ・ マネージャー
Ke Ding インテル コーポレーション主任エンジニア AI 応用研究
Eric Anderson インテル コーポレーション AI フレームワーク ・ エンジニア
Rita Brugarolas Brufau インテル コーポレーション ソフトウェア ・ ソリューション ・ エンジニア

課題
米計算機学会 （ACM: Association for Computing Machinery） の推薦システム ・ カンファレンスでは、 毎年
RecSys Challenge と呼ばれるデータサイエンス ・ コンテストを開催しています。 インテルは、 このコンテストに2 年前から 参加しています。 2022 年、 インテルは SIHG4SR: Side Information Heterogeneous Graph for
Session Recommender （英語） ソリューション （『ACM Digital Library』 に掲載） を発表し、最終リーダーボード
で 4 位に入賞しました （チーム名 「MooreWins」）。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

27

RecSyc Challenge 2022 （英語） は Dressipi （英語） によって企画され、 課題はファッション推薦でした。 提供
されたデータセット （英語） には、 18 か月間にわたる 110 万件の匿名化されたオンライン小売セッションが含まれていました （図 1 と図 2）。 各セッションには、 購入されたアイテムをラベルとする、 1 日の顧客の行動が含まれており、 記述的なファッション特性のセットに基づいて、 顧客がどのアイテムを購入するかを予測することが課題でした。 コンテストに参加したソリューションは、 MRR （Mean Reciprocal Rank） スコアによって順位が付けられました。 具体的には、 各セッションにおいて、 購入される可能性のあるアイテムのリストに対して、 実際に購入されたアイテムがより上位に表示されれば、 MRR は高くなります。 逆に、 アイテムが予測リストの下位に表示された
場合、 MRR は低くなるか、 ゼロになります。 例えば、 予測リストが [item2, item3, item1, item0] で、 item3 が実際11に購入されたアイテムの場合、 MRR= item3 のランク = 2 になります。

図 1. オンライン ・ ショッピング ・ セッションの例

図 2. アイテム特徴の例© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

28

インテルのソリューション
SIHG4SR ソリューションは、 グラフ ・ニューラル ・ ネットワーク （GNN） を使用して各セッションから学習し、 閲覧
したアイテムと最終的に購入したアイテムの間のつながりを明らかにします（図 3）。まず、各セッションを DGL（DeepGraph Library） （英語） を使用して実装したグラフで表現し、 それをネットワークに投入します。 このモデルは3 つの段階に分けることができます。 第 1 段階では、 入力グラフをアイテム、 特徴、 カテゴリーの 3 つの学習可能な埋め込みテーブルにエンコードします。 第 2 段階では、特徴とカテゴリーをアイテムの埋め込みと同じ形状に再整形し、 アテンション読み出し層を適用して特徴をデコードします。 そして、 第 3 段階では、 アイテムとサイド情報を基に 2 つの独立した予測を 1 つに融合して、 アイテムの最終予測を行います。

図 3. SIHG4SR アーキテクチャー

この新しいモデルは、 ヘテロジニアス ・ グラフを利用して各セッションを表現し、 新しい集約手法でサイド情報とシーケンス特徴を融合します。 サイド情報としてアイテムの特徴を組み込みユーザーの意図を有効活用するこのソリューションは、 再帰的特徴量エンジニアリング、 2 段階トレーニング手法、マルチレベルの組み合わせ、 ハイパーパラメーター ・ チューニングにより、 MRR 0.20762 を達成し、 総合 4 位を獲得しました。
このソリューションでは、 表形式のセッションデータをグラフ化します （図 4）。 各セッションは 1 つのグラフに変換され、 合計 110 万個のグラフが作成されます。 各グラフには 3 種類のエッジが存在します。 閲覧順にアイテムをつなぐエッジ、 アイテムとその特徴をつなぐエッジ、 そして、 特徴と対応するカテゴリーをつなぐエッジ
です。 アイテムの順序は推薦において重要です。 閲覧履歴を学習することで、グラフ・ネットワークは異なるパターンを検出し、 最終的に次にクリックされるアイテムや購入されるアイテムを予測できます。 さらに、 各アイテムの記述的特徴と対応するカテゴリーにより、 この豊富なコンテキストを利用してユーザーの意図を有効活用できます。
ここでは、 これらのアイテムの特徴やアイテムの特徴カテゴリーをサイド情報と呼びます。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

29

図 4. 表形式データからヘテロジニアス ・ グラフへの変換

モデル
このソリューションは、 アイテム情報とサイド情報を集約し、 異なる予測を融合し、 重み付き損失関数で精度を向上する新しい方法を提案します。
アイテム情報とサイド情報の集約 ： 特徴とカテゴリーの埋め込みの形状を統一するため、 2 つの方法を採用しています。 特徴の埋め込みについては、 同じアイテムにつながる特徴の埋め込みを、 新しい埋め込みベクトルとして合計する 「シャッフル」 手法を使用します。 カテゴリーの埋め込みについては、 まず特徴カテゴリーの埋め込みの平均を計算し、 次にブロードキャストしてアイテムの埋め込みに再整形する 「平均 + ブロードキャスト」 手法を使用します。
アイテム予測とサイド情報予測の融合 ： ユーザーの嗜好を表す局所表現 Zl と大域表現 Zg を生成し、 この 2 つの表現を固定比率 Zs= α Zl+ β Zg. で結合します。
重み付き損失 ： 重み付き負対数尤度 （NLL） 損失関数を使用して、 古いセッションよりも新しいセッションに注目します。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

30

特徴量エンジニアリングによる精度向上
すべての特徴がモデルに貢献するわけではないので、 グラフの構築に有用な特徴を選択することは非常に重要です。 提供されたアイテム特徴には、74 のカテゴリーと 904 の一意の 「カテゴリー：値」 のペアが含まれています。
例えば、 特徴カテゴリー 3 が 「色」 の場合、 カテゴリー 3 に属する特徴の値は、 「赤」 や 「白」、 あるいは別の色である可能性があります。 もし、 あるカテゴリーが履歴で一貫性のない情報を提供したり、 ほかのカテゴリーと矛盾するメッセージを提供する場合は、 グラフから除外する必要があります。 さらに、 提供された情報を基に、新たな有用な特徴を生成することもできます。 このソリューションでは、 個々の特徴カテゴリーを除外するテストを 74 回行ってから、 過去に検出された有害な特徴に応じて、 特徴カテゴリーの組み合わせを除外するテストを8 回行いました。 さらに、 閲覧と購入のタイムスタンプ、 過去のアイテム閲覧履歴、 購入回数を解析することで、予測に有益な 3 つの新しい特徴を作成しました。

インテル® End-to-End AI Optimization Kit
インテル® End-to-End AI Optimization Kit （英語） は、 一般的なハードウェア上でハイパフォーマンスかつ軽量なモデルを実現する、エンドツーエンドの AI 最適化を行う構成可能なツールキットです。 データ処理、トレーニング、
モデル最適化、推論など、AI パイプラインの各段階を最適化することを目的に開発されました。 いくつかのコンポー
ネントが含まれています。 RecDP は大規模データ処理を行うシンプルな API を提供し、 SDA （Smart model
Democratization Advisor） は SigOpt （英語） でモデルを最適化してチューニングする組み込みの知識を提供
し、 トレーニングなしのアーキテクチャー検索キット （DE-NAS） はコンパクトなニューラル・ネットワークを構築し、Model Adapter はトレーニング前の知識を簡単に新しいモデルへ転移させます。

SDA によるハイパーパラメーター ・ チューニング
最適なモデル ・アーキテクチャーを得るため、 ハイパーパラメーター ・ チューニングを行い最適なパラメーターを選択しました。 手動のハイパーパラメーター ・ チューニング （HPO） には専門知識が必要であり、 通常、 大規模
かつ再帰的な試行の間にエラーが発生します。 そこで、 SIHG4SR では SDA を利用してハイパーパラメーター ・チューニングを行いました （図 5）。

図 5. インテル® End-to-End AI Optimization Kit によるハイパーパラメーター ・ チューニング

Model Adapter による 2 段階トレーニング
各セッションには 1 つの購入アイテムと複数の閲覧アイテムが含まれるため、 Model Adapter を使用して転移学習ベースの 2 段階トレーニングを行い、 データを深く掘り下げます （図 6）。 第 1 段階では、 ネクストクリックの問題に注目します。 このネットワークは次に閲覧されるアイテムを正しく予測しようとします。 これにより、 データ
サイズを 5 倍に増強し、 モデルを大幅に強化できます。 第 2 段階では、 ネクストクリック問題から、 小さなデータセットで十分な予測精度を持つモデルを生成できる購入アイテム予測問題へと知識を転移させます。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

31

図 6. Model Adapter による転移学習ベースの 2 段階トレーニング

複数のモデルを組み合わせる
予測能力を向上させるため、 SIHG4SR アーキテクチャーをベースに複数の派生モデルを構築します （図 7）。
トレーニング手法、 データ ・ サンプリング、 ノード設定、 エッジドロップ手法、 読み出しと融合の設定を調整して、サブモデルを作成します。 最終的に、SIHG4SR_V1、SIHG4SR_V2、SIHG4SR_V3、SIHG4SR_V3_enhanced と名付けた 4 つのサブモデルを作成しました。 2 レベルの組み合わせを使用して、 4 つのモデルからの予測値を加重和することで、 単一モデルの一般化誤差を大幅に低減できます。

図 7. モデルの組み合わせ

まとめ
RecSys Challenge 2022 はセッションベースのレコメンダー問題で、 参加者には匿名化されたデータセットが提供されました。 セッションデータは表でもグラフでも表現可能であり、 ディープラーニング ・ モデルのほうがマシンラーニング手法よりも優れた予測能力を示しました。 インテルのソリューションは、4 つの重要な要因により、MRR 0.20762 を達成し、 総合 4 位を獲得しました。
• SIHG4SR ： セッションベース ・ レコメンダー問題のサイド情報を融合するヘテロジニアス GNN ベースのモデル• グラフ内のノード / エッジを追加 / 削除して精度を向上する特徴量エンジニアリング手法• ハイパーパラメーター ・ チューニングと転移学習を行うインテル® End to End AI Optimization Kit （英語）• 複数のモデルを活用するマルチレベルの組み合わせ© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

コードの境界を乗り越える
oneAPI 向けインテル® デベロッパー ・ クラウドでクロスアーキテクチャー ・プログラミングのパワーを体験してください。

デモ

Mandelbrot デモを異なるアーキテクチャーで実行して、クロスアーキテクチャーのパフォーマンスを確認できます。

学習

サンプルコードを含む 25 の Jupyter* Notebook で
SYCL* を実際に体験できます。

開発

最新のインテルの CPU、 GPU、 FPGA で将来に対応した
アプリケーションを計画してテストできます。

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

The Parallel Universe

33

薬の良し悪しを見分ける
製薬業界におけるコンピューター ・ ビジョン

Eduardo Alvarez インテル コーポレーション シニア AI ソリューション ・ エンジニア

このチュートリアルでは、 AI リファレンス・キットの Visual Quality Inspection （英語） を使用して、 コンピューター・
ビジョン ・ ソリューションを構築する方法を紹介します。 このリファレンス ・ キットは、 製造工程における不良品の視覚的検査の概念を説明するさまざまなデータセット （英語） へのリンクを提供します。 製造工程で発生する
さまざまな不良品で試すことができますが、 このチュートリアルでは錠剤の品質に注目します。 このデータセット
では、 市販のサプリメントが 「良」 または 「不良」 のカテゴリーに分類されています。 ここでは、 このデータセットを使用して、 事前にトレーニングした VGG-16 モデルを転移学習し、 錠剤の自動品質管理ツールを作成します。
VGG-16 は、 16 層の深さを持つ畳み込みニューラル ・ ネットワークです。 ILSVRC 2014 で最も成績の良い
アーキテクチャーの 1 つでした。 クラス分類タスクでは、（分類エラー率 6.66% の GoogleNet に次ぐ）2 位となり、トップ 5 分類エラー率は 7.32% でした。 また、位置特定タスクでは優勝し、位置特定エラー率は 25.32% でした。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

34

しかし、 ゼロからトレーニングすると非常に時間がかかります。 VGG-16 でトレーニングした ImageNet の重みのサイズは 528MB あるため、 かなりのディスク容量と帯域幅を必要とし、 非効率的です。 1 億 3800 万個のパラメーターは、 勾配の爆発的な増加につながります。 そこで、 AWS* EC2* m6i.4xlarge インスタンス （第 3 世代
インテル® Xeon® Platinum 8375C プロセッサー 2.90GHz） 上で稼働する PyTorch* 向けインテル® エクステンション （IPEX） を利用して、 錠剤データセット上で事前にトレーニングした VGGNet 分類アーキテクチャーを転移学習させます。
IPEX には、 インテル® ハードウェア上で PyTorch* のパフォーマンスを向上する最適化が含まれています （図 1）。
また、 Python* の API も含まれており、 ユーザーは 2、 3 行のコードを変更するだけで、 これらの最適化を利用できます。 IPEX の最適化のほとんどは、 最終的に stock PyTorch* のリリースに含まれる予定です。

図 1. PyTorch* 向けインテル® エクステンション （IPEX） により、 stock PyTorch* と比較して予測時間が向上。
この比較は v1.8.0 で行っていますが、 掲載時の最新バージョンは v1.12.0 です。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

35

第 3 世代のインテル® Xeon® プロセッサーは、 インテル® アドバンスト ・ ベクトル ・ エクステンション （インテル®
AVX-512） で低精度の BFloat16 をネイティブにサポートし、 将来的にはインテル® アドバンスト ・ マトリクス ・
エクステンション （インテル® AMX） で混合精度をサポートする予定です。 インテル® AMX を使用すれば、単精度でネットワーク精度を維持したまま、 半精度でトレーニングできます。
IPEX は、 Conv2D + ReLU、 Linear + ReLU など、 頻繁に使用される演算子パターンの融合を透過的にサポート
し、 TorchScript でさらにパフォーマンスを最適化できます。 また、 演算子を最適化し、 いくつかのカスタマイズされた演算子を実装します。 一部の ATen 演算子は、 ATen の登録メカニズムにより IPEX で最適化された演算子に置き換えられます。 さらに、一般的なトポロジー向けにいくつかのカスタマイズされた演算子が実装されています。
例えば、ROIAlign と NMS は Mask R-CNN で定義されています。 IPEX はカスタマイズされた演算子も最適化し、これらのトポロジーのパフォーマンスを向上します。

探索的データ分析
製薬業界において品質管理は非常に重要です。 良 （図 2） と不良 （図 3） の錠剤のさまざまな側面を見てみましょう。
不良品の例としては、 色、 斑点 / 汚れ、 割れ / 欠け、 刻印の不具合、 錠剤の種類違い、 傷などが挙げられます。

図 2. 良 / 許容範囲の錠剤の例

図 3. 不良 / 欠陥錠剤の例© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

36

平均的な良と不良の錠剤の差分分析では、 錠剤の右上に欠陥があり、 錠剤の右側に変色のようなものがある
ことがわかります （図 4）。 これは、 一部の錠剤の右側面にダメージを与える製造工程の問題を示している可能性があります。

図 4. 平均的な許容範囲の錠剤と平均的な不良の錠剤の相違点

カスタム VGG-16 モデル定義
ImageNet で事前学習された VGG-16 特徴抽出器と、 カスタム分類ヘッドを使用して、 カスタム ・ マルチクラス
分類モデルを定義します。 転移学習を可能にするため、 最初の畳み込みブロックのパラメーターはフリーズ （固定）されます。 この関数は、 クラススコア （トレーニング・モードの場合）、 クラス確率、 そして正規化された特徴マップ（評価モードの場合） を返します。
PyTorch* モデルのトレーニング ・スキーマを初めて目にした方は、 いくつかの定型コードを記述する必要があることに気付くでしょう。 PyTorch* のモデルは通常、 以下のコンポーネントを必要とします。
• データの前処理：データの読み込みとカスタマイズのため、PyTorch* の 「Dataset クラス」 を拡張します。 これは、
フォルダーから画像、 それぞれのラベル、 その他必要なメタデータを抽出するのに便利です。
• データの読み込み：PyTorch* は「DataLoader クラス」も提供しており、モデルのトレーニングと評価時に「Dataset
クラス」 の操作を支援します。 この関数は、 トレーニングと推論中にモデルにデータを提供する役割を担うジェネレーターです。
• モデルの定義 ： PyTorch* モデルを定義するには、 「Module クラス」 を拡張するクラスを定義します。 コンストラ
クター （ __init__） はモデルの層を決定し、 forward() 関数は定義されたモデルの層を通じて入力データを前方伝播させる方法を定義する責任があります。
• トレーニング関数 ： トレーニング ・コンポーネントの損失関数と最適化アルゴリズムを定義する必要があります。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe37

以下のコード例は、 この例のモデル定義を表しています （図 5）。 コンストラクターでは、 拡張する VGG-16モデルと、 既存のアーキテクチャーに追加するカスタム層を定義しています。 _freeze_params メソッドは
requires_grad を False に設定することで、 VGG-16 モデルの浅い層をフリーズ （転移学習のため層の重みを変更できないように） します。 forward() メソッドは、 重みを適用し、 モデルを通じてデータを移動します。
class CustomVGG(nn.Module):
def __init__(self, n_classes=2):
super().__init__()self.feature_extractor = models.vgg16(pretrained=True).features[:-1]
self.classiﬁcation_head = nn.Sequential(
nn.MaxPool2d(kernel_size=2, stride=2),
nn.AvgPool2d(
kernel_size=(INPUT_IMG_SIZE[0] // 2 ** 5, INPUT_IMG_SIZE[1] // 2 ** 5)),
nn.Flatten(),nn.Linear(in_features=self.feature_extractor[-2].out_channels,
out_features=n_classes,),)# self._freeze_params()
def _freeze_params(self):
for param in self.feature_extractor[:23].parameters():
param.requires_grad = False
def forward(self, x_in):"""forward"""feature_maps = self.feature_extractor(x_in)scores = self.classiﬁcation_head(feature_maps)
if self.training:
return scoresprobs = nn.functional.softmax(scores, dim=-1)
weights = self.classiﬁcation_head[3].weightweights = (weights.unsqueeze(-1).unsqueeze(-1).unsqueeze(0).repeat((x_in.size(0),1,1,INPUT_IMG_SIZE[0] // 2 ** 4,
INPUT_IMG_SIZE[0] // 2 ** 4,)))feature_maps = feature_maps.unsqueeze(1).repeat((1, probs.size(1), 1, 1, 1))
location = torch.mul(weights, feature_maps).sum(axis=2)
location = F.interpolate(location, size=INPUT_IMG_SIZE, mode="bilinear")maxs, _ = location.max(dim=-1, keepdim=True)maxs, _ = maxs.max(dim=-2, keepdim=True)mins, _ = location.min(dim=-1, keepdim=True)mins, _ = mins.min(dim=-2, keepdim=True)
norm_location = (location - mins) / (maxs - mins)
return probs, norm_location

図 5. カスタム VGG-16 モデルクラス© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

38

PyTorch* 向けインテル® エクステンションによる転移学習
前述したように、 IPEX を使用して、 錠剤データセット上で事前にトレーニングした VGGNet 分類モデルを転移学習させます。 以下のコード例の 11 行目は、 モデルにさまざまな最適化を適用する IPEX optimize メソッドの呼び出しを示しています（図 6）。 IPEX のモデルへの適用については、公式ドキュメント（英語）を参照してください。
# モデルのトレーニングを開始# モデルのトレーニング# DL アーキテクチャー、 オプティマイザー、 損失関数の初期化
model = CustomVGG()
class_weight = torch.tensor(class_weight).type(torch.FloatTensor).to(DEVICE)
criterion = nn.CrossEntropyLoss(weight=class_weight)
optimizer = optim.Adam(model.parameters(), lr=LR)# IPEX 最適化
model, optimizer = ipex.optimize(model=model, optimizer=optimizer, dtype=torch.ﬂoat32)# モジュールのトレーニング
start_time = time.time()
trained_model = train(train_loader, model=model, optimizer=optimizer, criterion=criterion, epochs=EPOCHS,
device=DEVICE, target_accuracy=TARGET_TRAINING_ACCURACY)
train_time = time.time()-start_time# hdf5 PyTorch* モデルの保存
model_path = f"{subset_name}.h5"
torch.save(trained_model, model_path)

図 6. モデルのトレーニング

トレーニング済みモデルのアーキテクチャーを評価し、 トレーニング済みパラメーターの総数を確認します （図 7）。
IPEX でモデルを最適化すると、 適切な層に _IPEX というプリフィクスが追加されます。 これは、 IPEX が実装されていることを確認する良い方法です。
from torchsummary import summary
summary(trained_model, (3, 224, 224))© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

図 7. PyTorch* 向けインテル® エクステンション（IPEX）モデルのアーキテクチャー© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

39

The Parallel Universe40

ホールドアウト ・ テスト ・ データに対する予測では、 モデルは写真の 2/3 が 「不良」 錠剤であると判断しています（図 8）。 境界ボックスとヒートマップは、 錠剤で検出された最も顕著な欠陥を強調しています。 予想どおり、 「FF」
刻印の問題、 欠け、 変色が見られます。

図 8. 錠剤データセットの 3 つのテスト画像に対する予測値。 赤い境界ボックスで囲まれた領域は、 欠陥を示しています。
ヒートマップは、 重要な欠陥の濃度を示しています。

まとめ
コードを少し追加するだけで、インテル® ハードウェア上でカスタム VGG-16 バイナリー分類器を最適化できました。
この効率良い転移学習により、 事前にトレーニングした Torch VGG-16 モデルを、 錠剤データセットの画像を使用して転移学習 （チューニング） し、 医薬品の品質管理を行う効率良い分類ツールに変えることができました。

インテル® DPC++ 互換性ツール

CUDA* アプリケーションを標準ベースの SYCL* コードに移行© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

詳細 （英語）The Parallel Universe

41

インテル® Arc™ GPU で
TensorFlow* Stable
Diffusion を実行する
PC で高速なディープラーニングを実現

Chris Lishka、 Vivek Kumar、 Geetanjali Krishna、 および AG Ramesh インテル コーポレーション

インテルは、 ハイパフォーマンス GPU の 「インテル® Arc™ A770」 を含む、 新しいインテル® Arc™ A シリーズ ・グラフィックス ・ ハードウェアを発表しました。 A770 は、 ゲーム、 デジタルコンテンツ制作、 ストリーミングに優れていますが、 PC 上でディープラーニング （DL） ワークロードを実行することも可能です。
TensorFlow* （TF） は、 確立された DL フレームワークです。 インテルは、 TF が GPU やその他のオフロード
デバイスを含むさまざまなアクセラレーターをターゲットにできるように、 Google* とプラグイン ・ インターフェイス（英語） を設 計しました。 TensorFlow* 向けインテル® エクステンション （英語） は、 ネイティブ Linux* とWindows* Subsystem for Linux* 2 （WSL2） の両方で動作するインテル® Arc™ A シリーズ GPU の実験的サポートを含む、 インテル® GPU 上で TF DL ワークロードを実行できるようにするプラグインです。 ネイティブ
Windows* 上での実行には、 TensorFlow* DirectML プラグイン （英語） を使用できます。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

42

WSL2 上の Ubuntu* 20.04 に TensorFlow* 向けインテル® エクステンションをインストールするのは簡単です。
Windows* と WSL2 にドライバー、 ランタイムパッケージ、 TF、 拡張機能をセットアップする手順を説明します。
インテル® Arc™ A770 GPU でテキストから画像生成するため、 Jupyter* Notebook で TF を使用して人気のある
KerasCV* Stable Diffusion 機能を簡単に実行する方法を紹介します。

WSL2 での Ubuntu* 20.04 のセットアップ
本記事では、 第 12 世代インテル® Core™ i9-12900 プロセッサー （開発コード名 Alder Lake） 搭載の PC とインテル® Arc™ A770 ディスクリート ・ グラフィックス ・ カード 16GB を 使 用していま す。 前 提 条 件 として、Windows* 11 に WSL2 と Ubuntu* 20.04 コンテナがインストールされている必要があります。 WSL2 を使用
して Ubuntu* 20.04 コンテナをセットアップする手順は、 こちらを参照してください。 WSL2 は、 Windows*
ホストで GPU にアクセスするため、 DirectX* を含む特別な Linux* カーネルビルドを使用します。 これにより、
汎用の GPU 計算機能を Windows* と Linux* で共有できます。 GPU は Windows* ホストで管理されるため、GPU ドライバーは Windows* にインストールする必要があります。 GPU 用の低レベルドライバーをインストール
するには、 インテル® Arc™ グラフィックス Windows* DCH ドライバーの標準インストーラーを使用する必要があります。 ここでは、 自己インストール型の実行ファイルを使用して、 最新のドライバー （v31.0.101.3490） がインストールされました。

Ubuntu* Linux* で の イ ン テル® グ ラフィッ クス 計 算ランタイム とoneAPI パッケージのセットアップ
Windows* にドライバーをインストールし、 WSL2 上の Ubuntu* 20.04 コンテナを起動したら、 次のステップでは
apt-get コマンドを使用していくつかの Ubuntu*/Debian* パッケージをインストールします。 まず、 apt システムにインテル® グラフィックス ・ パッケージの専用リポジトリーを追加する必要があります。
$ sudo apt-get install -y gpg-agent wget
$ wget -qO - https://repositories.intel.com/graphics/intel-graphics.key | \sudo gpg --dearmor --output /usr/share/keyrings/intel-graphics.gpg$ echo 'deb [arch=amd64 signed-by=/usr/share/keyrings/intel-graphics.gpg]
https://repositories.intel.com/graphics/ubuntu focal main' | \sudo tee /etc/apt/sources.list.d/intel.gpu.focal.list$ sudo apt-get update

リポジトリーが登録されると、 以下のコマンドを使用して、 インテルが提供するランタイムパッケージをインストール
できます。
$ sudo apt-get install \intel-opencl-icd=22.28.23726.1+i419~u20.04 \intel-level-zero-gpu=1.3.23726.1+i419~u20.04 \level-zero=1.8.1+i419~u20.04

これにより、 上記のパッケージと暗黙の依存関係が Ubuntu* 20.04 （WSL2 で実行） にインストールされます。
TensorFlow* 向けイ ンテル® エクス テン ションはイ ンテル® oneAPI ツール キット を 使 用して ビ ルド さ れて いま す。 そ のため、 いくつ か の oneAPI ランタイム ・ ライブ ラリー も インスト ールする 必 要が ありま す。
まず、 インテルの oneAPI リポジトリーを apt システムに登録します。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

43$ wget -O- \https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | \sudo gpg --dearmor --output /usr/share/keyrings/oneapi-archive-keyring.gpg$ echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] \https://apt.repos.intel.com/oneapi all main" | \sudo tee /etc/apt/sources.list.d/oneAPI.list$ sudo apt-get update

次に、 DPC++ と oneMKL ランタイムパッケージをインストールします。
$ sudo apt-get install intel-oneapi-runtime-dpcpp-cpp intel-oneapi-runtime-mkl（この 記 事 で イ ンスト ール し た バージョンは、 intel oneapi runtime dpcpp cpp=2022.2.1 16953 と intel
oneapi runtime mkl=2022.2.1 16993 です。） これは、 TensorFlow* 向けインテル® エクステンションを実行するためのライブラリーのみをインストールし、 oneAPI の完全な開発パッケージ （より大規模なもの） は含まれません。 oneAPI の開発パッケージ全体をインストールする場合は、 oneAPI のインストール ・ドキュメント （英語）を参照してください。

Python* 仮想環境の作成
TF を GPU で 使 用する場 合は、 Miniconda* を推 奨しま す。 Linux* 用の conda 環 境を 作成 する完全なドキュメントは、 こちら （英語） を参照してください。 最初のステップとして、 Ubuntu* 20.04 に Miniconda*
をダウンロードしてインストールします。
$ curl https://repo.anaconda.com/miniconda/Miniconda3-latest-Linux-x86_64.sh \-o Miniconda3-latest-Linux-x86_64.sh$ bash Miniconda3-latest-Linux-x86_64.sh

次に、 TF と TensorFlow* 向けインテル® エクステンションを新しい環境にインストールし、 pip を最新バージョンに更新します。
$ conda create --name keras-cv python=3.9$ conda deactivate$ conda activate keras-cv(keras-cv)$ pip install --upgrade pip

これで、 TF （v2.10.0） と TensorFlow* 向けインテル® エクステンション （v1.0.0） をインストールできるようになりました。
(keras-cv)$ pip install tensorﬂow==2.10.0(keras-cv)$ pip install --upgrade intel-extension-for-tensorﬂow[gpu]パッケージ名の最後にある [gpu] は、 拡張機能の GPU バックエンドを選択するように指定するものです。

TensorFlow*/Keras の Stable Diffusion のインストールと実行
さて、 ここからが本題です。 一般的で実行しやすい Stable Diffusion の TF Keras モデルは、 KerasCV （英語）パッケージで利用可能です。 この実装を使用して、 Jupyter* Notebook 内でテキストから画像への生成を実行します。 まず、 定番の Jupyter* Notebook パッケージをインストールします。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

44(keras-cv)$ pip install notebook

次に、 KerasCV とそれが使用するパッケージをインストールし、 Notebook サーバーを起動します。
(keras-cv)$ pip install keras-cv tensorﬂow-datasets matplotlib ipywidgets(keras-cv)$ jupyter notebook

すると、 Notebook サーバーに接続するための URL が表示されます。 Windows* の場合、 ブラウザーを開き、
アドレス欄に URL をコピー & ペーストしてください。 これで、Jupyter* Notebook のページが表示されるはずです。
それでは、 テキストから画像への変換を試してみましょう。 新しいノートブックを作成し、 以下のコマンドを実行
してください。 コマンドの出力のほとんどは簡潔にするため省略されていますが、 この実行で生成された画像は例として残しています。

model.text_to_image(...) の最初の実行では、 起動時のオーバーヘッドが発生しますが、 それ以降の実行は高速になります。モデルが実行されると、プログレスバーが表示され、いくつのステップが完了したかが示されます。
ワークロードが終了すると、 このプログレスバーは、 各反復にかかった秒数と画像の生成にかかった全体的な時間を表示します。 エルヴィスの乗り物を変えるなど、 異なるプロンプトて実験すると、 生成される画像も変わります。
この Stable Diffusion モデルはランダム化を採用しているため、 同じプロンプトでも次回以降の実行で異なる画像が生成されます。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

The Parallel Universe

45

インテル® Arc™ GPU が使用されたかどうかの判定
Stable Diffusion モデルの TF 出力には、 インテル® Arc™ GPU が使用されているかどうかが表示されます。
TensorFlow* 向けインテル® エクステンションは 「XPU」 デバイスを使用しており、 インテル® Arc™ GPU を利用できる場合は自動的に XPU が選択されます。 TF の出力には、 以下のような内容が表示されます。
2022-11-10 19:06:45.243358: I
tensorﬂow/core/grappler/optimizers/custom_graph_optimizer_registry.cc:114]
Plugin optimizer for device_type XPU is enabled.

インテル® Arc™ GPU の使用状況をリアルタイムで監視するもう 1 つの方法は、 Windows* タスク ・ マネージャーを使用することです。 タスク ・ マネージャーを起動し、 ［パフォーマンス］ タブを選択して監視対象デバイスのリストを表示し、 インテル® Arc™ A770 グラフィックスに対応する GPU デバイスを選択するだけです。 GPU ページには、 上部に 4 つのパネルがあり、 それぞれに表示を変更するためのドロップダウン ・ メニューが用意されています。 これらのパネルの 1 つで、 ドロップダウン ・メニューから ［Compute］ を選択します。 WSL2 の Ubuntu*
20.04 コンテナで Stable Diffusion ワークロードを実行し、 model.text_to_image(...) が実行されると、Windows* のタスク ・ マネージャーの ［Compute］ パネルにインテル® Arc™ GPU の計算が表示されます。 また、［Copy］ や ［Dedicated GPU Memory Usage］ などのパネルにアクティビティーが表示されるはずです。 これは、ワークロードが （WSL2* の） Linux* で実行されている間に、 計算が GPU で行われていることを示しています。

まとめ
TensorFlow* 向けインテル® エクステンションは、 一般的な DL ワークロードをインテル® GPU 上で実行できるようにします。 インテル® Arc™ GPU のサポートは現在実験的なものですが、 Stable Diffusion などの人気のあるモデルはすでにうまく動作しています。 この記事で紹介した例では、WSL2 で Ubuntu* 20.04 コンテナをセットアップして、 インテル® GPU 上で DL ワークロードを素早く実行できるようにする方法を示しました。 ネイティブ
Linux* （Ubuntu* 20.04） で動作するインテル® Arc™ GPU で同じことを行うには、 GitHub* リポジトリーのTensorFlow* 向けインテル® エクステンションのインストール ・ ドキュメントを参照してください。

関連情報• TensorFlow* 向けインテル® エクステンションの GitHub* リポジトリー （英語）• TensorFlow* 向けインテル® エクステンションの pypi.org ページ （英語）• TensorFlow* 向けインテル® エクステンションのドキュメント （英語）• TensorFlow* （英語）• KerasCV の Stable Diffusion を使用したハイパフォーマンスな画像生成 （英語）© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、Intel Corporation またはその子会社の商標です。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

インテルのテクノロジーを使用するには、 対応したハードウェア、 ソフトウェア、 またはサービスの有効化が必要となる場合があります。 詳細については、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 またはその子会社の商標です。
* その他の社名、 製品名などは、 一般に各社の表示、 商標または登録商標です。