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

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

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

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

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

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 では性能向上なし

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

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

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

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

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

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

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

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

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

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

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

オリジナルコード主要部

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

ループの回し方を変更

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

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

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

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

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

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

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

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

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

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

Happy hacking!