PGIコンパイラーによる ディレクティブベースのGPGPU 株式会社 ベストシステムズ 石川 直太 [email protected] 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 --- スレッドごとにインスタンスを持つ 変数 多重ループのベクトル化 !$acc to vector(4) do k = 1, km !$acc do vector(4) do j = 1, jm !$acc do vector(16) do i = 1, im 多重ループをベクト ル化できる それぞれのループの 並列度を定数で指定 する 並列度の積は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の可能性? 事例1:姫野ベンチマーク 連立一次方程式をヤコビ法で解く メモリ性能が現われる http://accc.riken.jp/HPC/HimenoBMT.html GPU版姫野ベンチのコード主要部 !$acc region & !$acc copyin(aa(1:mimax,1:12,1:mjmax,1:mkmax)) & !$acc copyin( p(1:mimax,1:mjmax,1:mkmax)) & !$acc copyout(p(1:mimax,1:mjmax,1:mkmax)) & #ifdef USE_WW !$acc copyout(ww(1:mimax,1:2,1:mjmax,1:mkmax)) & #else !$acc copyout(wrk2(1:mimax,1:mjmax,1:mkmax)) & !$acc copyout(gosatmp(1:mimax,1:mjmax,1:mkmax)) & #endif ! USE_WW !$acc local(k,j,i,s0,ss) & !$acc copyin(kmax,jmax,imax,omega) !$acc do host do loop=1,nn !$acc do parallel, vector(4) ! Cut and try. do k=2,kmax-1 !$acc do parallel !!!, vector(2) ! Cut and try. do j=2,jmax-1 !$acc do vector(64) ! Cut and try. do i=2,imax-1 姫野ベンチの配列のパディング Portland Groupによる改良 – 配列の第一次元の大きさを調節する。 – #ifdef PAD – mimax = 272 ! GPUに適するマジックナンバー? – #else – mimax = 257 ! 姫野ベンチオリジナルの値 – #endif CPUによる計算:850 MFLOPS GPUによる計算:20292 MFLOPS --- 23.8倍 マクロによる添え字のすり替え !!!! BEGIN BSI HACK !!!! #define a1(i,j,k) aa(i,1,j,k) #define a2(i,j,k) aa(i,2,j,k) #define a3(i,j,k) aa(i,3,j,k) #define b1(i,j,k) aa(i,4,j,k) #define b2(i,j,k) aa(i,5,j,k) #define b3(i,j,k) aa(i,6,j,k) #define c1(i,j,k) aa(i,7,j,k) #define c2(i,j,k) aa(i,8,j,k) #define c3(i,j,k) aa(i,9,j,k) #define a4(i,j,k) aa(i,10,j,k) #define bnd(i,j,k) aa(i,11,j,k) #define wrk1(i,j,k) aa(i,12,j,k) #ifdef USE_WW #define gosatmp(i,j,k) ww(i,1,j,k) #define wrk2(i,j,k) ww(i,2,j,k) #endif ! USE_WW !!!! END BSI HACK !!!! Fortranのマクロ C/C++のマクロと同様 配列構造の試行錯誤に便利 大文字小文字の区別に注意 implicit none との併用をお勧め PGIではソースの拡張子が「F」または「F90」 (大文字)の場合と、オプション「-Mpreprocess 」で有効 Intel コンパイラーではオプション「-fpp」で有 効 姫野ベンチで解ったこと 配列の最も左側の添え字(Fortranの場合)を、最 も内側のループで、1づつ増やすとよい。 コンパイル時に「Non-stride-1 accesses」と表 示された場合には、性能が出にくい。 copyin, copyout ディレクティブが重要。 配列の構造を変えると、性能が上がる可能性が あるが、若干工数を要する。 vector(64) のパラメーターは試行錯誤。 姫野ベンチで効果がなかったこと • • private ディレクティブ – スレッドごとに別々のインスタンスを持つと指 定 – 省略しても、コンパイラーが自動的に判断 ストライド 0 – 構造体の配列と等価なデータ構造 – CPUでは、キャッシュのヒット率向上に効果 – GPUでは、遅くなる 姫野ベンチの他の研究との比較 富士通研究所 (情報処理学会HPC研究会) CUDAプログラミングで、69.7 GFLOPS メモリ転送速度がピーク性能の80%を超えるチュー ニング これと比較して本実験は 0.28倍 ソフテック – PGIコンパイラーを使って、20457.79 MFLOPS NEC (2009年9月2日 セミナー資料) – PGIコンパイラーを使って、18477.78 MFLOPS 事例2:行列積 SGEMM BLASに含まれるサブルーチン S --- 単精度実数 GE --- 一般的な行列 MM --- 行列×行列 Netlibでソースコード公開 http://www.netlib.org/blas/index.html オリジナルコード主要部 * * Form C := alpha*A*B + beta*C. * !$acc region DO 90, J = 1, N IF( BETA.EQ.ZERO )THEN DO 50, I = 1, M C( I, J ) = ZERO 50 CONTINUE ELSE IF( BETA.NE.ONE )THEN DO 60, I = 1, M C( I, J ) = BETA*C( I, J ) 60 CONTINUE END IF DO 80, L = 1, K IF( B( L, J ).NE.ZERO )THEN TEMP = ALPHA*B( L, J ) DO 70, I = 1, M C( I, J ) = C( I, J ) + TEMP*A( I, L ) 70 CONTINUE END IF 80 CONTINUE 90 CONTINUE !$acc end region 1570 MFLOPS (CPU よ り 遅い ) ないほうがよい条件分枝を除去 * * Form C := alpha*A*B + beta*C. * !$acc region DO 90, J = 1, N CBSI2010 IF( BETA.EQ.ZERO )THEN CBSI2010 DO 50, I = 1, M CBSI2010 C( I, J ) = ZERO CBSI2010 50 CONTINUE CBSI2010 ELSE IF( BETA.NE.ONE )THEN DO 60, I = 1, M C( I, J ) = BETA*C( I, J ) 60 CONTINUE CBSI2010 END IF DO 80, L = 1, K CBSI2010 IF( B( L, J ).NE.ZERO )THEN TEMP = ALPHA*B( L, J ) DO 70, I = 1, M C( I, J ) = C( I, J ) + TEMP*A( I, L ) 70 CONTINUE CBSI2010 END IF 80 CONTINUE 90 CONTINUE !$acc end region 1583 MFLOPS ループの回し方を変更 * * Form C := alpha*A*B + beta*C. * !$acc region local(temp,i,j,l) !$acc do vector(16) DO J = 1, N !$acc do vector(16) private(temp) DO I = 1, M ! C( I, J ) = BETA*C( I, J ) temp = 0.0 DO L = 1, K ! C( I, J ) = C( I, J ) + ALPHA * B(L,J) *A( I, L ) temp = temp + B(L,J) *A( I, L ) enddo c(i,j) = c(i,j) * beta + temp * alpha enddo enddo !$acc end region 不可解な現象 vector(16) ディレクティブがあると 1 GFLOPS vector(16) ディレクティブがないと 23 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 リリース PGI 10.3 でサンプルコードがコンパイルエラー アクセラレータコンパイラーは枯れていない ライセンスファイルとライセンスサーバーが新 しければ、古いコンパイラーも動く 複数のバージョンのコンパイラーをインストー ルして、パスの設定で選択可能 お問い合わせは 価格表、オンライン見積もりによる割引 ご注文窓口 http://www.bestsystems.co.jp/ [email protected] 技術ご質問、ライセンス発行窓口 [email protected] Happy hacking!
© Copyright 2024 ExpyDoc