1 / 31

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

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

mercury
Download Presentation

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

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  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!

More Related