OpenACC vs OpenMP GPU Offloading (環境構築編)

アルバイトの大友です。
NVIDIA GPUではCUDAというアーキテクチャにより容易に汎用計算を実行できます。
しかし、既にあるCPU用のソースコードをCUDA/C,C++によりGPU上で動作するソースコードに移植することは一般にコストが高いという問題があります。
この問題の解決策として、pragmaディレクティブを追加するだけで容易にGPU上で動作するプログラムを記述することができるOpenMP GPU OffloadingとOpenACCがありますが、
今回はこれらの環境構築の方法の調査と簡単な実行速度比較を行いました。

OpenACC

OpenACCではpragmaディレクティブを用いてGPU上での並列処理を記述することができます。
OpenACCを用いるためには、

• OpenACCに対応したコンパイラのビルド
• 対象のソースコードでOpenACCを有効化

を行う必要があります。

OpenACCに対応したコンパイラのビルド

OpenACCに対応したコンパイラにはpgiやgcc 9.1以降などがあります。
今回はpgiのインストールについてのみ書きます。

pgi編

pgiコンパイラ（PGI Community Edition）のインストールはPGI Installation Guide – PGI Documentに従えばできるかと思います。
対話型のインストーラなので、質問に答えるだけでインストールができます。
environment-modulesを予めインストールしておくことでmoduleファイルのインストールまで行ってくれます。

ソースコードでのOpenACCの有効化

例えば

#pragma acc kernels loop independent
for (...){}

とすることで用いることができます。

OpenMP GPU Offloading

CPU上でマルチスレッドにより取り敢えずプログラムの高速化をしたければ、OpenMPを用いて

#pragma omp parallel for
for (...){}

とすると思うのですが、このOpenMPにはバージョン4.0からプログラムをGPUなどのアクセラレータ上で実行するOffloading機能があります。
Offloading機能を用いるためには

• gccやclangなどのコンパイラでOpenMP Offloaginを有効化
• 対象のソースコードでpragmaディレクティブを用いてOffloadingを有効化

を行う必要があります。

コンパイラのOffloading有効化

clang編

NVIDIA GPUへのOffloadingを有効化したclangのビルドは
Building LLVM/Clang with OpenMP Offloading to NVIDIA GPUs – HPCWIKI
を参考にしていただければできるかと思います。
大まかな手順としては

1. llvm/clang/openmpをビルド
2. ビルドしたclangを用いてopenmpを再ビルド

の2段階となっています。
2段階目のopenmpの再ビルドの際にはCMAKE_BUILD_TYPE=Debugとしておくことをおすすめします。
こうするとOffloadingが期待通り動作しない場合などに実行時ログを出力できるようになります。

gcc編

NVIDIA GPUへのOffloadingを有効化したgccのビルドは
Offloading Support in GCC – GCC Wiki
を参考にしていただければできるかと思います。
大まかな手順としては

1. GPU用のコードのためのコンパイラをビルド（NVIDIA GPUの場合は上記リンクの”For Nvidia PTX”）
2. 1でビルドしたコンパイラをリンクしてホスト用のgccをビルド（上記リンクの”Building host compiler”）

の2段階となっています。

ソースコードでのOffloading有効化

例えばOpenMPのOffloading機能はソースコードとしては

#pragma omp target teams distribute parallel for
for (...){}

とすることで用いることができます。

Offloadingが実際に期待通り行われているかは、環境変数OMP_TARGET_OFFLOADをMANDATORYとDISABLEにして実行時間が異なることを確認したり、NVIDIA GPUであればnvprofなどでカーネル関数が実行されているかを確認することができます。
clangの場合、もしOffloadingが行われていなさそうでしたら、

LIBOMPTARGET_DEBUG=1 ./a.out

として実行時ログを出力し、何が悪いかを調査することができます。
このログ出力を行うためにはclangのビルドで書いた通りCMAKE_BUILD_TYPEをDebugとしておく必要があります。

カーネル関数のイメージの作成について

通常CUDAではカーネル関数をイメージとして実行ファイル等にリンクさせたり、前方互換性のためにPTXと呼ばれるアセンブリとして実行ファイル等に保存しておき実行時にコンパイルして実行したりします。
ではOpenACCやOpenMP GPU Offloadingではどの様にカーネル関数を保持しているのでしょうか？
これはobjdumpやcuobjdumpを用いることで調べることができます。

• OpenACC/pgiの場合は、ソースコードのコンパイル時にカーネル関数のイメージやPTXを実行ファイルにリンク・埋め込みます。
• OpenMP GPU Offloading/gccの場合は、実行ファイルのELFのrodataセクションにカーネル関数のPTXが書き込まれており、実行時にアセンブル・リンクされます。
• OpenMP GPU Offloading/clangの場合は、コンパイル時にカーネルのイメージが作成されリンクされます。

スレッド階層について

OpenMP Offloading、OpenACCではCUDAのgrid/blockのように階層的にスレッドを管理しています。
規格として定まっているわけではなく実際の実装によって変わるようですが、おおよそこの様になっているそうです[1,2]。

OpenMP Offloadingでは多重ループを用いてこれらの大きさを指定できます。
OpenACCでは多重ループとpragmaディレクティブでこれらの大きさを指定できます。

計算性能比較

今回はvectoradd計算を用いて簡単な計算性能比較を行います。
並列化を行わないCPU版のコードは以下の通りです。

const size_t N = 1lu << 28;
float* a = malloc(sizeof(float) * N);
float* b = malloc(sizeof(float) * N);
float* c = malloc(sizeof(float) * N);

for (size_t i = 0; i < N; i++) {
        c[i] = a[i] + b[i];
}

性能評価では、OpenMP GPU Offloading/gcc、clangとOpenMP/pgi、CUDA、CPUでのOpenMPによるマルチスレッド（gccとclangの2種類）の計6種類を比較しました（図1）。
評価に用いたGPUはNVIDIA GeForce GTX1060、CPUはIntel Core i9-7900Xです。
性能のチューニングはteam/gang/block sizeの調整のみ行います。
評価に用いたコードはこちらです。

そもそも要求されるB/f比が大きな計算なこともありGeForce GTX1060の単精度計算性能の理論ピークである4.4TFlop/sと比較すると性能がでていませんが、
CUDAとOpenACC/pgiの計算性能はほぼ同じとなりました。
一方でOpenMP GPU OffloadingではCPUと比較しても性能が低いという結果となりました。

なぜOpenMPでは遅いのか？

各GPU実装で最も計算性能が高かったパラメータでの実行時の状況をnv-nsight-cu-cliで見てみます。
OpenMP/clangは原因は分かりませんがnv-nsight-cu-cliで解析が始まらないため結果を得られませんでした。

OpenMP/gccではTheoretical occupancyが低いのがわかります。
Occupancyが低い理由としてはレジスタの使いすぎが考えられます。
ではなぜOpenMP/gccでは多くのレジスタを使っているかということですが、CUDAの中間アセンブリであるPTXを見たところteamsの制御などのコードがカーネル関数に実装されており、これに多数のレジスタやSharedメモリを用いているようでした。
OpenMP/gccが遅いもう一つの理由として、grid sizeが他と比べて小さいことが挙げられそうですが、team sizeを制御するompディレクティブ中のnum_teamsを使っても31以上には上げられなかったため検証を断念しました。

おわりに

個人的な感想ですが、環境構築が一番楽なのはOpenACC/pgiでした。かつvectoraddの性能評価を見る限り性能も納得できるものかなと思います。
OpenMPはgccかclangに依らず環境構築は比較的大変でした。ビルドは指示に従えばできるのですが、実行するとOffloadingが行われないなどの問題が起き、それの対応に結構時間がかかりました。
今回比較に用いた計算はメモリアクセスの比率が大きいため、次回はできれば計算の割合が大きくメモリやキャッシュの階層性を意識して工夫のしようがあるベンチマークで比較を行いたいものです。

参考ページ

• [1] Introduction to OpenACC Course Lecture #2 Q&A
• [2] Assessing One-to-One Parallelism LevelsMapping for OpenMP Offloading to GPUs, Chen Shen, et al, PMAM’17

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください。