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

Presentation on theme: "PGIコンパイラーによる ディレクティブベースのGPGPU"— Presentation transcript:

1 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 では 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 事例１：姫野ベンチマーク 連立一次方程式をヤコビ法で解く メモリ性能が現われる

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

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

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 スレッドごとに別々のインスタンスを持つと指 定
省略しても、コンパイラーが自動的に判断 ストライド 0 構造体の配列と等価なデータ構造 CPUでは、キャッシュのヒット率向上に効果 GPUでは、遅くなる

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

17 事例２：行列積 SGEMM Netlibでソースコード公開 BLASに含まれるサブルーチン S --- 単精度実数

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

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

20 ループの回し方を変更

21 不可解な現象 vector(16) ディレクティブがあると 1 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 リリース
アクセラレータコンパイラーは枯れていない ライセンスファイルとライセンスサーバーが新 しければ、古いコンパイラーも動く 複数のバージョンのコンパイラーをインストー ルして、パスの設定で選択可能

30 お問い合わせは 価格表、オンライン見積もりによる割引 ご注文窓口 技術ご質問、ライセンス発行窓口
ご注文窓口 技術ご質問、ライセンス発行窓口

31 Happy hacking!