PGIコンパイラーによる ディレクティブベースのGPGPU

1. 株式会社 ベストシステムズ 石川 直太 naota@bestsystems.co.jp PGIコンパイラーによるディレクティブベースのGPGPU

2. PGIによるGPGPUの概要 • アクセラレータ対応ライセンスが必要 • Fortran 95 または C99 のコードにディレクティブ • !$acc .. • !$acc& … 継続行 • #pragma acc … • オプションを付けてコンパイル • -ta=nvidia • -ta=nvidia,time 実行時に統計情報を表示 • -Minfo,accel コンパイル時に詳細な情報を表示

3. 主要なディレクティブ • acc region --- GPUによる処理の開始 • acc end region --- GPUによる処理の終了 • copyin, copy, copyout --- まとめてコピー • local --- GPU側だけで使う変数 • acc do vector --- GPUによる並列処理 • private --- スレッドごとにインスタンスを持つ変数

4. 多重ループのベクトル化 • 多重ループをベクトル化できる • それぞれのループの並列度を定数で指定する • 並列度の積は256以下 • Cut and try • 明示的に指定しなくても、自動ベクトル化

5. PGI 9.0 から 10.0 への改良 • 総和演算が可能 • 姫野ベンチマークのコードより • GOSA = GOSA + SS * SS • すべてのスレッドに渡る総和演算 • 9.0 ではコンパイルエラー • 10.0 ではコンパイル、計算可能

6. PGI 9.0から10.0への改良 • コンパイラーまかせでの、性能向上 • 姫野ベンチマークに、「!$acc region」と「!$acc end region」だけを追加したコード • PGI 9.0-4 では 0.229 GFLOPS (CPUより遅い) • PGI 10.0 では 15.3 GFLOPS • チューニングすると、9.0でも10.0でも、20.5 GFLOPS • 10.1、10.2 では性能向上なし

7. GPGPU化可能なコード • 並列化可能が大前提 • リストベクトルでなく多重ループ • OpenMP から GPGPU への移行は容易 • 倍精度演算よりも単精度演算が高速

8. ディレクティブベースGPGPUの特徴 • CUDAプログラミングよりも容易 • ディレクティブを無視すれば、通常のC/Fortran • GPUによる計算とCPUによる計算の比較が容易 • オリジナルコードの変更への対処が比較的容易 • ハードウェアに依存しない • 理論的には、AMDのGPUや将来のCPUに対応可能 • ヘテロジニアスマルチコアCPUの可能性？

9. 事例１：姫野ベンチマーク • 連立一次方程式をヤコビ法で解く • メモリ性能が現われる • http://accc.riken.jp/HPC/HimenoBMT.html

10. GPU版姫野ベンチのコード主要部

11. 姫野ベンチの配列のパディング • Portland Groupによる改良 • 配列の第一次元の大きさを調節する。 • #ifdef PAD • mimax = 272 ! GPUに適するマジックナンバー? • #else • mimax = 257 ! 姫野ベンチオリジナルの値 • #endif • CPUによる計算：850 MFLOPS • ＧＰＵによる計算：20292 MFLOPS --- 23.8倍

12. マクロによる添え字のすり替え

13. Fortranのマクロ • C/C++のマクロと同様 • 配列構造の試行錯誤に便利 • 大文字小文字の区別に注意 • implicit none との併用をお勧め • PGIではソースの拡張子が「F」または「F90」(大文字)の場合と、オプション「-Mpreprocess」で有効 • Intel コンパイラーではオプション「-fpp」で有効

14. 姫野ベンチで解ったこと • 配列の最も左側の添え字(Fortranの場合)を、最も内側のループで、1づつ増やすとよい。 • コンパイル時に「Non-stride-1 accesses」と表示された場合には、性能が出にくい。 • copyin, copyout ディレクティブが重要。 • 配列の構造を変えると、性能が上がる可能性があるが、若干工数を要する。 • vector(64) のパラメーターは試行錯誤。

15. 姫野ベンチで効果がなかったこと • private ディレクティブ • スレッドごとに別々のインスタンスを持つと指定 • 省略しても、コンパイラーが自動的に判断 • ストライド 0 • 構造体の配列と等価なデータ構造 • CPUでは、キャッシュのヒット率向上に効果 • GPUでは、遅くなる

16. 姫野ベンチの他の研究との比較 • 富士通研究所 (情報処理学会HPC研究会) • CUDAプログラミングで、69.7 GFLOPS • メモリ転送速度がピーク性能の80%を超えるチューニング • これと比較して本実験は 0.28倍 • ソフテック • PGIコンパイラーを使って、20457.79 MFLOPS • NEC (2009年9月2日 セミナー資料) • PGIコンパイラーを使って、18477.78 MFLOPS

17. 事例２：行列積 • SGEMM • BLASに含まれるサブルーチン • S --- 単精度実数 • GE --- 一般的な行列 • MM --- 行列×行列 • Netlibでソースコード公開 • http://www.netlib.org/blas/index.html

18. オリジナルコード主要部

19. ないほうがよい条件分枝を除去

20. ループの回し方を変更

21. 不可解な現象 • vector(16) ディレクティブがあると 1 GFLOPS • vector(16) ディレクティブがないと 23 GFLOPS • 診断メッセージによると、どちらも16x16ブロック • 試行錯誤が必要 • PGI コンパイラーは、まだ発展途上か？

22. 行列積のコードをGPUで計算しましょう 実習

23. 用意してあるファイル • sample.f --- GPU化していないサブルーチンコード • sgemm-4.f --- GPU化の例 • Makefile • test-sgemm.f --- 評価用メインプログラム

24. まずはCPUで実行 • make sample • ./sample

25. 最初の一歩 • 高速化したいブロックの最初に !$acc region • 終わりに !$acc end region • 転置行列の積を計算するブロックもあるが、とりあえずは、最初のブロックだけ • make sample • ./sample

26. チューニング(1) • オリジナルコードには、0による乗算を避けるための条件分枝があります。 • なくてもよい条件分枝を削除しましょう。

27. チューニング(2) • ループの回し方を変えてみましょう。 • ヒントは「sgemm-4.f」

28. 試行錯誤の例 • 明示的な、「local」、「private」 • !$acc do vector(並列度) • 「copyin」、「copy」、「copyout」 • コンパイルオプション「-ta=nvidia,mul24」

29. 終わりに、最新 Bad know how • 2010年3月5日に、PGI 10.3 リリース • PGI 10.3 でサンプルコードがコンパイルエラー • アクセラレータコンパイラーは枯れていない • ライセンスファイルとライセンスサーバーが新しければ、古いコンパイラーも動く • 複数のバージョンのコンパイラーをインストールして、パスの設定で選択可能

30. お問い合わせは • 価格表、オンライン見積もりによる割引 • http://www.bestsystems.co.jp/ • ご注文窓口 • sales@bestsystems.co.jp • 技術ご質問、ライセンス発行窓口 • license@bestsystems.co.jp

31. Happy hacking!