どのように計算集約型コードをIntel GPUにオフロードするか
本記事は以下の翻訳です。
techdecoded.intel.io
Intelプロセッサーのグラフィックスアーキテクチャーは多くのIntel SoC製品においてグラフィックス、コンピュート、メディアとディスプレイ能力を与えるIntelのテクノロジーです。
Intelプロセッサーのグラフィックスアーキテクチャーはgeneration（世代）を短くしたGenとして、非公式な呼び方ですが知られています。
アーキテクチャーのリリースごとにGenの後に対応するバージョンが記載されます。
例えば、Intelグラフィックスアーキテクチャーの最新リリースはGen11です。
何年にもわたって洗練されたグラフィックス（3Dレンダリングとメディアパフォーマンス）と最大1TFLOPSの性能を持つ汎用計算能力を持つまでに進化しています。
この記事では、IntelプロセッサーのGen9とGen11グラフィックスアーキテクチャーの汎用計算能力とIntel oneAPI Base ToolkitにあるData Parallel C++(DPC++)を使ってどのようにプログラムするのかについて探っていきます。
特に、DPC++を使って2つのGenアーキテクチャのプログラミングと性能の側面を示すケーススタディを見ていきます。
Intelプロセッサーのグラフィックスは、Intel CPUのダイに統合された高効率で高性能なグラフィックスとメディアアクセラレーターです。
統合GPUはCPUとラストレベルキャッシュ(LLC)を共有しており、細粒度で低いレイテンシーと高い帯域幅でのコヒレンシーなデータ共有が可能です。
図１はGen11グラフィックスがあるSoCを示しています。オンダイの統合はディスクリートのグラフィックスカードより低い電力消費を可能にします。
図１：IntelプロセッサーのGen11グラフィックスSoC（CPU SoCより大きい部分）
図２はGen9GPUのアーキテクチャーブロックダイアグラムを示します。GPUは多くの実行ユニット(EU)を含み、それぞれはSingle Instruction Multiple Data(SIMD)計算を行う能力があります。
8つのEUの集まりはサブスライスとなります。
サブスライスは以下を持ちます。
- 命令キャッシュ
- L1/L2サンプラーキャッシュ
- メモリーのload/storeユニットのポート
これらのサブスライスが集まってスライスを形成し、スライスは（CPUとコヒレントな）共有L3キャッシュとバンク構成の共有ローカルメモリー(SLM)から構成されます。
Intelの統合GPUは1つかそれ以上のスライスを持つかもしれません。
このようなコンフィギュレーションでは、L3はインターコネクトファブリックを通して複数のスライスと接続されています。
図２：Intel Gen9 GPUアーキテクチャー
図３はGen9アーキテクチャーにおけるEUのある程度の詳細を示しています。
各EUごとに最大7つのスレッドを扱えるマルチスレッディングをサポートし、各スレッドは128本のSIMD-8 32bitレジスターを持っています。
EUは1サイクルに最大4命令を発行できます（Intel GPUのアーキテクチャーの詳細とベンチマークのより詳しい内容はここにあります）。例えば、ハードウェアの理論上のピークGFLOPSは以下のように計算できます。
(EUs) * (SIMD units/EU) * (FLOPS per cycle/SIMD unit) * (Freq GHz)
図３：サブスライスとEUアーキテクチャー詳細
GPUのようなデバイスをプログラミングするとき、最も良い性能を得るには利用可能なハードウェア機能を上手くマッピングするプログラミング言語構成が必要です。いくつかのAPIが利用できますが、ここではoneAPIを掘り下げていきましょう。
oneAPIとDPC++
oneAPIはオープンかつ無償で、そしてアクセラレーターや複数世代のハードウェアにわたって可搬性と性能を提供する標準化ベースのプログラミングモデルです。
oneAPIは様々なハードウェアターゲットにわたってコードを再利用するために、核となるプログラミング言語のDPC++を含みます。
以前の記事、oneAPIを使ったヘトロジニアスプログラミング(The Parallel Universe, 39号)で詳細を見つけることができます。
DPC++は以下の特徴を含んでいます。
- A Unified Shared Memory：ホスト・デバイス間の簡便なメモリマネジメントが可能になります
- OpenCLスタイルのNDRange subgroups : ベクトル化を補助するための機能です
- 一般的なポインタや関数ポインタのサポート
- その他、数多くの特徴
この記事でCUDAコードをDPC++に変換するケーススタディを示していきます。
Hogbom Clean 画像アルゴリズムについて見ていきましょう。これは電波天文学での画像で広く使われているアルゴリズムです。この画像アルゴリズムは2つのホットスポット（処理が重たい部分）があります。
簡潔のため、Find Peakの性能面についてフォーカスしていきます。オリジナルの実装はC++やOpenMP, CUDA, そしてOpenCLがあります。ホストCPUは利用可能なときGPU上で動作するCUDA/OpenCLカーネルにオフロードします（CUDAはNVidia GPUにのみ計算をオフロードするプロプライエタリなアプローチです）。図４と５はそれぞれホストおよびデバイスコードのスニペットを示しています。
図４：Find Peak のホストコード：C++とCUDA
図５：Find Peakのデバイスコード：CUDA
CUDAコードを手動でDPC++コードに置き換えることもできますし、DPC++互換性ツール(DPCT)を使うこともできます。DPCTはCUDAプログラムをDPC++（図６と７）へ移行するのを補佐します。移行にはIntel oneAPI Base ToolkitとNVIDIA CUDAヘッダーのみ必要です。DPCTを起動してexample.cuファイルを移行するのは以下のようにシンプルです。
dpct example.cu
多くのCUDAファイルがあるアプリケーションの移行では、プログラムソースの場所を指定するDPCTのオプション-in-rootと移行したコードを書き込むための-out-rootが利用できます。アプリケーションがmakeやcmakeを使っているのなら、intercept-buildを使って移行を行うのをお勧めします。これはコンパイラ呼び出しとともにコンパイル処理のデータベースファイルを作成します（ホストC++コードとデバイスCUDAコード両方の入力ファイル名と、付随するコンパイラオプションです）。
intercept-build make
dpct -p=<path to .json file> --out-root=dpct_output ...
特にHogbom Clean CUDAコードをDPC++に移行するのには、CUDAカーネルを持つHogbomCuda.cuファイルにDPCTツールを直接かける方法、intercept-buidを使う方法、どちらも利用可能です。デフォルトでは移行したコードはdp.cppという拡張子名のファイルが得られます。
以下に移行したDPC++コードとオリジナルのCUDAコードの比較を見ることができます。
図６：DPCTを使って移行したFind Peak DPC++のホストコード
図７：CUDAホストコードと移行したDPC++ホストコードの比較
図８：DPCTを使って移行したFind Peak DPC++のデバイスコード
図９：Find Peak CUDAカーネルと移行したDPC++デバイスコードの比較
DPC++コードで鍵となる側面がいくつかあります。
- SYCLキューを使ったデバイスコードの呼び出し
- デバイスコードを実行するためのラムダ関数ハンドラー
- マルチスレッド実行のための
parallel_for構文（オプション）
ここで、移行後のDPC++コードは単一共有メモリ(USM)プログラミングモデルを使い、デバイスカーネルによって読み書きされるデータをデバイスメモリに確保します。（ホストではなく）デバイス上のメモリ確保なので、ホストとデバイス相互の明示的なデータコピーが必要となります。また、共有としてメモリを確保し、ホストとデバイス両方がアクセス、更新することもできます。ここで示していないこととして非USMモードがあり、これはSYCLバッファーとアクセサーを使ってデータ転送を行います。
DPCTで移行したコードは実行するデバイスを決め、デバイスのためのキューを作成します（get_current_device()とget_default_queue()の呼び出し）。GPUへDPC++コードをオフロードするために、sycl::gpu_selectorパラメーターを持ったキューを作成する必要があります。処理されるデータはデバイスと（そのデバイスである）GPU上で実行するカーネルで利用できるようにすべきです。GPUへ/またはGPUからコピーされるデータのサイズと次元はsycl::range, sycl::nd_rangeで指定します。DPCTを使うとき、CUDAコードの各ソース行は等価のDPC++コードに移行されます。Find Peakのデバイスカーネルコード(d_findPeak)では、（CUDAコードから）生成されたDPC++コードはほぼ1対1で等価の移行になります。なので、DPCTは素早い移植とプロトタイピングのための非常に強力なツールになります。移行したDPC++コードとCUDAコードの比較を図７と９に示します。
DPCTを使ってDPC++に移行したコードを得て、次はその正確性と効率性を検証してみます。いくつかのケースでは、DPCTツールはプリプロセッサのディレクティブ変数を対応する値に置き換えるかもしれません。この置換を手動で戻す必要があるかもしれません。また、移行したコードに修正が示されるコンパイルエラー（CUDAのthreadId.xを等価であるnd_rangeアクセサーに置き換えることなど）が起きるかもしれません。Hogbom Cleanアプリケーションコードには、移行後のDPC++コードが生成した結果の検証を手助けする正確性チェッカーがあります。正確性チェックはGPU上でのDPC++コード実行の結果と元となるホストCPU上でのC++実装の結果を比較することで行われます。
そして、移行したDPC++コードのGPU上での利用率（EU占有率やキャッシュ使用状況、単精度や倍精度のFLOPS）やホスト・デバイス間のデータ転送を解析することでその効率性を測定することができます。GPU利用率にインパクトを与えるパラメーターの1つとしてワークグループのサイズとそのrange dimensionがあります。Hogbom Cleanアプリケーションでは、Find Peak処理においてnBlocksとfindPeakWidthにあたります。
性能へのインパクトとチューニングのこの機会を説明するために、nBlocksの値を24と4に設定して測定した性能プロファイルを図１０に示します。
findPeakWidthは256に設定しています。このプロファイルはGPUプロファイリング機能をサポートするIntel Vtune Profilerで測定しました。CUDAを使ったNVidia GPUでの効率的なパラメータがDPC++コードを実行するIntel GPUにおいて効率的ではないかもしれないので、DPCTを使った場合チューニングはより明確に必要となります。表１はGen9(48EU)で測定した結果を示します。
図１０：nBlocksの2つの値 (a) 24, (b) 4 におけるGen9でのHogbom Cleanのプロファイル結果
表１：処理量の多い Find Peak におけるGen9 GPUでの性能測定
GPUの利用率や効率性の最適化に加えて、ホスト・デバイス間のデータ転送もまたチューニングすべき点です。Hogbom CleanアプリケーションにはFind Peak, Subtract PSFカーネルの複数回の呼び出しがあり、これらのカーネルで使われるデータはデバイス上に常駐させることができます。それゆえ、メモリの再確保やホスト・デバイス間のコピーは必要ありません（今後の記事でデータ転送と単一共有メモリに関するこれらの最適化についてお話しする予定です）。
IntelプロセッサーのグラフィックスアーキテクチャとDPC++の特徴を理解することは、より良いアルゴリズムと移植可能な実装を書くことの手助けになります。この記事では、いくつかのアーキテクチャの詳細をレビューし、DPC++を構成するものとDPCTを使ったケーススタディを見てきました。Intel GPUでの最高性能を得るために、特にDPCTを使うときはカーネルパラメーターを調整することが重要です。最新のIntelハードウェアとソフトウェアでアプリケーションの開発、テスト、実行するのにIntel DevCloudをお試しいただくことをお勧めします。
