OpenACCによる並列化

Transcription

1 実習 OpenACC による ICCG ソルバーの並列化 1

2 ログイン Reedbush へのログイン $ ssh reedbush.cc.u-tokyo.ac.jp l txxxxx Module のロード $ module load pgi/17.3 cuda ログインするたびに必要です! ワークディレクトリに移動 $ cdw ターゲットプログラム <$O-L3>/srcx OpenACC 用のディレクトリの作成 $ cd (C or F)/L3 $ cp r srcx srcx_acc $ cd srcx_acc $ make clean 2

3 Makefile for OpenACC OpenACC 用の Makefile の作成 C/L3/srcx_acc/Makefile CC = icc OPTFLAGS= -O3 -qopenmp -ipo -xcore-avx2 -align TARGET =../run/l3-solx F/L3/srcx_acc/Makefile F90 = ifort F90OPTFLAGS= -O3 -qopenmp -ipo -xcore-avx2 - align array32byte TARGET =../run/l3-solx CC = pgcc OPTFLAGS= -O3 -acc -Minfo=accel -ta=tesla:cc60 TARGET =../run/l3-solx-acc F90 = pgfortran F90OPTFLAGS= -O3 -acc -Minfo=accel -ta=tesla:cc60 TARGET =../run/l3-solx-acc -Minfo=accel でコンパイラメッセージの出力 OpenACC ではコンパイラメッセージの確認が極めて重要保守的に並列化しないことがあるため 並列化されているかどうかの確認や どのループがどのレベル (gang, worker, vector) で並列化されたのか知るため 3

4 ジョブスクリプト for OpenACC $ cd../run $ cp gor.sh gor-acc.sh or $ cp go.sh gor-acc.sh gor-acc.sh の編集 C/L3/run/gor-acc.sh F/L3/run/gor-acc.sh PGI_ACC_TIME=1 で GPU 実行情報のサマリを標準エラーに出力 #PBS -q u-lecture #PBS -N test #PBS -l select=1:ncpus=18 #PBS -Wgroup_list=gt00 #PBS -l walltime=00:10:00 #PBS -e test2.err #PBS -o test2.lst cd $PBS_O_WORKDIR. /etc/profile.d/modules.sh export OMP_NUM_THREADS=18 export KMP_AFFINITY=granularity=fine,compact,1,0 numactl./l3-rsol0 numactl./l3-rsol0 numactl./l3-rsol0 #PBS -q h-lecture #PBS -N test #PBS -l select=1 #PBS -Wgroup_list=gt00 #PBS -l walltime=00:10:00 #PBS -e test2.err #PBS -o test2.lst cd $PBS_O_WORKDIR. /etc/profile.d/modules.sh module load pgi/17.3 cuda export PGI_ACC_TIME=1 numactl./l3-solx-acc numactl./l3-solx-acc numactl./l3-solx-acc 4

5 GPU 用のインプットファイルの作成 INPUT.dat の編集 C/L3/run/INPUT.dat F/L3/run/INPUT.dat NX/NY/NZ 1.00e e e0 DX/DY/DZ 1.0e-08 OMEGA, EPSICCG 18 PEsmpTOT -30 NCOLORtot NX/NY/NZ 1.00e e e0 DX/DY/DZ 1.0e-08 OMEGA, EPSICCG 1 PEsmpTOT -30 NCOLORtot 0 GPU のスレッド数は数百万に及ぶので スレッド数を PEsmpTOT で制御するのは現実的ではない 5

6 CPU 版の結果の取得 結果の正しさを検証するために 1 CPU コアで実行した結果を取っておく 下の結果は 1CPU コアで実行した際のもの PGI コンパイラ使用 (OpenACC なし ) NCOLORtot = -30 # INPUT.dat 自分で実行してみてください C/L3/run/gor-acc.sh.oXXXXX ### CMRCM ### FINAL COLOR NUMBER e-01 sec. (assemble) e e e e e-09 N= e+01 sec. (solver) Quick check Is the number of iteration same as baseline? F/L3/run/gor-acc.sh.oXXXXX ### CM-RCM ### FINAL COLOR NUMBER E-01 sec. (assemble) E E E E E E E+00 N= E+01 sec. (solver) 6

7 OpenACC での並列化戦略 OpenACC ではどのループを並列化するべきか C/L3/srcx_acc/solver_ICCG_mc.c #pragma omp parallel private (ic, ip1, ip2, i, WVAL, j) for(ic=0; ic<ncolortot; ic++) { ip1 = ic * PEsmpTOT; PEsmpTOT = 1 ip1 = ic ip2 = ic * PEsmpTOT + PEsmpTOT; ip2 = ic+1 #pragma omp for for(i=smpindex[ip1]; i<smpindex[ip2]; i++) { VAL = D[i]; for(j=indexl[i]; j<indexl[i+1]; j++) { VAL = VAL - AL[j]*AL[j] * W[DD][itemL[j] - 1]; W[DD][i] = 1.0 / VAL; OpenACC でのターゲットループ F/L3/srcx_acc/solver_ICCG_mc.f!$omp parallel private(ic,ip1,ip2,i,val,k) do ic= 1, NCOLORtot ip1= SMPindex((ic-1)*PEsmpTOT) + 1 ip2= SMPindex((ic-1)*PEsmpTOT + PEsmpTOT)!$omp do do i= ip1, ip2 VAL= D(i) do k= indexl(i-1)+1, indexl(i) VAL= VAL - (AL(k)**2) * W(itemL(k),DD) enddo W(i,DD)= 1.d0/VAL enddo enddo!$omp end parallel ここで同期が必要! 前の色が終わったところで同期が必要となる OpenMPでは暗黙的に同期が入る OpenACCでは同期を取るためにカーネルを閉じるしかない! 7

8 OpenACC 指示文の挿入 C/L3/srcx_acc/solver_ICCG_mc.c #pragma omp parallel for private (i) for(i=0; i<n; i++) { X[i] = 0.0; W[1][i] = 0.0; W[2][i] = 0.0; W[3][i] = 0.0; F/L3/srcx_acc/solver_ICCG_mc.f!$omp parallel do private(i) do i= 1, N X(i) = 0.d0 W(i,2)= 0.0D0 W(i,3)= 0.0D0 W(i,4)= 0.0D0 W(i,R)= B(i) enddo #pragma omp parallel for private (i) #pragma acc kernels copyout(x[0:n],w[1:3][0:n]) for(i=0; i<n; i++) { X[i] = 0.0; W[1][i] = 0.0; W[2][i] = 0.0; W[3][i] = 0.0; C では配列のサイズ情報必須!!$omp parallel do private(i)!$acc kernels copyout(x,w) copyin(b) do i= 1, N X(i) = 0.d0 W(i,2)= 0.0D0 W(i,3)= 0.0D0 W(i,4)= 0.0D0 W(i,R)= B(i) enddo!$acc end kernels Fortran では配列がサイズ情報も持っているため必要ない 8

9 PGI コンパイラのメッセージ確認 C/L3/srcx_acc/ $make F/L3/srcx_acc/ $ make solve_iccg_mc: 34, Generating copyout(x[:n],w[1:3][:n]) 35, Complex loop carried dependence of X-> prevents parallelization Loop carried dependence of W->-> prevents parallelization Loop carried backward dependence of W->-> prevents vectorization Accelerator scalar kernel generated Accelerator kernel generated Generating Tesla code 35, #pragma acc loop seq solve_iccg_mc: 53, Generating copyout(w(:,:),x(:)) Generating copyin(b(:)) 54, Loop is parallelizable Accelerator kernel generated Generating Tesla code 54,!$acc loop gang, vector(128)! blockidx%x threadidx%x 並列化されていない! X と W が alias を持ってる可能性があるため コンパイラは並列化をしない 並列化されてる! gang, vector レベルで並列化されている 9

10 !$acc loop independent C/L3/srcx_acc/solver_ICCG_mc.c F/L3/srcx_acc/solver_ICCG_mc.f #pragma omp parallel for private (i) #pragma acc kernels copyout(x[0:n],w[1:3][0:n]) for(i=0; i<n; i++) { X[i] = 0.0; W[1][i] = 0.0; W[2][i] = 0.0; W[3][i] = 0.0; nothing to do #pragma omp parallel for private (i) #pragma acc kernels copyout(x[0:n],w[1:3][0:n]) #pragma acc loop independent for(i=0; i<n; i++) { X[i] = 0.0; W[1][i] = 0.0; W[2][i] = 0.0; W[3][i] = 0.0; C では loop independent が必要になる! 10

11 PGI コンパイラのメッセージ確認 C/L3/srcx_acc/ $make F/L3/srcx_acc/ $ make solve_iccg_mc: 34, Generating copyout(x[:n],w[1:3][:n]) 36, Loop is parallelizable Accelerator kernel generated Generating Tesla code 36, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ nothing to do ようやく並列化 11

12 OpenACC 指示文の挿入 C/L3/srcx_acc/solver_ICCG_mc.c #pragma omp parallel private (ic, ip1, ip2, i, WVAL, j) for(ic=0; ic<ncolortot; ic++) { ip1 = ic * PEsmpTOT; PEsmpTOT = 1 ip1 = ic ip2 = ic * PEsmpTOT + PEsmpTOT; ip2 = ic+1 #pragma omp for #pragma acc kernels copyin(d[0:n], indexl[0:n+1], AL[0:NPL], iteml[0:npl], SMPindex[0:NCOLORtot*PEsmpTOT] ) copy(w[0:4][0:n]) #pragma acc loop independent for(i=smpindex[ip1]; i<smpindex[ip2]; i++) { VAL = D[i]; #pramga acc loop seq for(j=indexl[i]; j<indexl[i+1]; j++) { VAL = VAL - AL[j]*AL[j] * W[DD][itemL[j] - 1]; W[DD][i] = 1.0 / VAL; 配列の確保は poi_gen.c でなされている AL, iteml の長さ NPL はこのソースからは見えない! F/L3/srcx_acc/solver_ICCG_mc.f!$omp parallel private(ic,ip1,ip2,i,val,k) do ic= 1, NCOLORtot ip1= SMPindex((ic-1)*PEsmpTOT) + 1 ip2= SMPindex((ic-1)*PEsmpTOT + PEsmpTOT)!$omp do!$acc kernels copyin(d,indexl,al,iteml) copy(w)!$acc loop independent do i= ip1, ip2 VAL= D(i)!$acc loop seq do k= indexl(i-1)+1, indexl(i) このループは短いので逐次計算 ( 長さ 3 or 6) VAL= VAL - (AL(k)**2) * W(itemL(k),DD) enddo W(i,DD)= 1.d0/VAL enddo!$acc end kernels enddo!$omp end parallel 配列のサイズ情報を書く必要はない 12

13 関数の引数の書き換え C/L3/srcx_acc/solver_ICCG_mc.c solve_iccg_mc(int N, int NL, int NU, int *indexl, int *iteml, int *indexu, int *itemu, double *D, double *B, double *X, double *AL, double *AU, int NCOLORtot, int PEsmpTOT, int *SMPindex, int *SMPindexG, double EPS, int *ITR, int *IER) F/L3/srcx_acc/solver_ICCG_mc.f nothing to do solve_iccg_mc(int N, int NL, int NU, int NPL, int NPU, int *indexl, int *iteml, int *indexu, int *itemu, double *D, double *B, double *X, double *AL, double *AU, int NCOLORtot, int PEsmpTOT, int *SMPindex, int *SMPindexG, double EPS, int *ITR, int *IER) 以下の書き換えも必要 : C/L3/srcx_acc/solver_ICCG_mc.h C/L3/srcx_acc/main.c 13

14 main.c のトラップ C/L3/srcx_acc/main.c F/L3/srcx_acc/main.f int main() { double *WK; int NPL, NPU; int ISET, ITR, IER; int icel, ic0, i; double xn, xl, xu; double Stime, Etime; nothing to do 偽物 NPL, NPU は poi_gen.h で宣言されているものが本物 以下の関数呼び出しの引数として本物を使うため 使われてない偽物はコメントアウト int main() { double *WK; // int NPL, NPU; int ISET, ITR, IER; int icel, ic0, i; double xn, xl, xu; double Stime, Etime; 14

15 OpenACC 版の結果チェック C/L3/run/gor-acc.sh.oXXXXX ### CMRCM ### FINAL COLOR NUMBER 30 F/L3/run/gor-acc.sh.oXXXXX ### CM-RCM ### FINAL COLOR NUMBER e-01 sec. (assemble) e e e e e-09 N= e+01 sec. (solver) C/L3/run/gor-acc.sh.eXXXXX ベースラインと一致! この時点では CPU より遅いが 気にしない Accelerator Kernel Timing data /lustre/gt25/z30108/c/l3/srcx/solver_iccg_mc.c solve_iccg_mc NVIDIA devicenum=0 time(us): 586,647 34: compute region reached 1 time 36: kernel launched 1 time grid: [16384] block: [128] device time(us): total=114 max=114 min=114 avg=114 elapsed time(us): total=138 max=138 min=138 avg=138 34: data region reached 2 times 34: kernel launched 3 times grid: [1] block: [128] device time(us): total=8 max=4 min=2 avg=2 elapsed time(us): total=499 max=430 min=32 avg= E-01 sec. (assemble) E E E E E E E+00 N= E+01 sec. (solver) 標準エラー出力側に PGI_ACC_TIME によるサマリが出力される (Fortran 版も同様 ) grid の値が gang( スレッドブロック ) 数 block の値が gang あたりの vector( スレッド ) 数 15

16 データ転送の最適化 C/L3/srcx_acc/solver_ICCG_mc.c #pragma acc data copyin(d[:n], indexl[:n+1], AL[:NPL], iteml[:npl]) copyin(smpindex[0:ncolortot*pesmptot]) copyout(x[:n], W[:4][:N]) { #pragma omp parallel for private (i) #pragma acc kernels copyout(x[0:n],w[1:3][0:n]) #pragma acc loop independent for(i=0; i<n; i++) { X[i] = 0.0; W[1][i] = 0.0; W[2][i] = 0.0; W[3][i] = 0.0; #pragma omp parallel private (ic, ip1, ip2, i, WVAL, j) for(ic=0; ic<ncolortot; ic++) { ip1 = ic * PEsmpTOT; ip2 = ic * PEsmpTOT + PEsmpTOT; #pragma omp for #pragma acc kernels copyin(d[0:n], indexl[0:n+1], AL[0:NPL], iteml[0:npl], SMPindex[0:NCOLORtot*PEsmpTOT]) copy(w[0:4][0:n]) #pragma acc loop independent for(i=smpindex[ip1]; i<smpindex[ip2]; i++) { VAL = D[i]; #pragma acc loop seq for(j=indexl[i]; j<indexl[i+1]; j++) { VAL = VAL - AL[j]*AL[j] * W[DD][itemL[j] - 1]; W[DD][i] = 1.0 / VAL; F/L3/srcx_acc/solver_ICCG_mc.f!$acc data copyin(b,d,indexl,al,iteml) copyout(x,w)!$omp parallel do private(i)!$acc kernels copyin(b) copyout(x,w) do i= 1, N X(i) = 0.d0 W(i,2)= 0.0D0 W(i,3)= 0.0D0 W(i,4)= 0.0D0 W(i,R)= B(i) enddo!$acc end kernels!$acc data 指示文で確保済みの配列に対して それより内側で copy 指示子を適用しても無視される 冗長な転送を除去!!$omp parallel private(ic,ip1,ip2,i,val,k) do ic= 1, NCOLORtot ip1= SMPindex((ic-1)*PEsmpTOT) + 1 ip2= SMPindex((ic-1)*PEsmpTOT + PEsmpTOT)!$omp do!$acc kernels copyin(d,indexl,al,iteml) copy(w)!$acc loop independent do i= ip1, ip2 VAL= D(i)!$acc loop seq do k= indexl(i-1)+1, indexl(i) VAL= VAL - (AL(k)**2) * W(itemL(k),DD) enddo W(i,DD)= 1.d0/VAL enddo!$acc end kernels enddo!$omp end parallel!$acc end data 16

17 データ転送最適化のイメージ (1/2) 初期実装 Compute r (0) = b-[a]x (0) for i= 1, 2, solve [M]z (i-1) = r (i-1) r i-1 = r (i-1) z (i-1) if i=1 p (1) = z (0) else b i-1 = r i-1 /r i-2 p (i) = z (i-1) + b i-1 p (i-1) endif q (i) = [A]p (i) a i = r i-1 /p (i) q (i) x (i) = x (i-1) + a i p (i) r (i) = r (i-1) - a i q (i) check convergence r end Host Device Data transfer 17

18 データ転送最適化のイメージ (2/2) 最適化実装 Compute r (0) = b-[a]x (0) for i= 1, 2, solve [M]z (i-1) = r (i-1) r i-1 = r (i-1) z (i-1) if i=1 p (1) = z (0) else b i-1 = r i-1 /r i-2 p (i) = z (i-1) + b i-1 p (i-1) endif q (i) = [A]p (i) a i = r i-1 /p (i) q (i) x (i) = x (i-1) + a i p (i) r (i) = r (i-1) - a i q (i) check convergence r end Host Device Data transfer 18

19 実習 EX1:Makefile ジョブスクリプト INPUT.dat を編集し CPU 実行時の出力を得てください EX2: 全並列化ループ (OMP DO で並列化されているところ ) を OpenACC で並列化してください コンパイラメッセージに注意! 常に CPU での結果と突き合わせ! 計算順序が変わるため 必ずしも一致しない この時点で遅くても気にしない! EX3: データ転送を最適化してください EX4:INPUT.dat の色数 (NCOLORtot) を変更し 速度への影響をみてください PGI_ACC_TIME は若干速度に影響を与えるので 計測時は PGI_ACC_TIME=0 ( ジョブスクリプト ) としてください マルチコア メニィコア並列プログラミング入門 19

20 参考 :Reedbush での nvvp の使い方 NVVP:NVIDIA Visual Profiler ( ) ログインして module をロード $ ssh -Y reedbush.cc.u-tokyo.ac.jp -l txxxxx $ module load cuda pgi/17.3 # cuda module が必須 サンプルプログラムのコピー $ cp /lustre/gt00h/share/openacc_samples.tar.gz. $ tar zxvf OpenACC_samples.tar.gz nvprof コマンドを使ってデータを収集 ジョブスクリプトの中に書く nvvp.sh 参照 nvvp の起動 $ nvvp GPU プログラミング入門 20

21 1File をクリック GPU プログラミング入門 21

22 2Import をクリック 3 Nvprof を選択 4 Next GPU プログラミング入門 22

23 5 Multiple processes に変更 7 Browse 6 Next GPU プログラミング入門 23

24 8 ファイル選択画面が開くので nvvp で作成したプロファイリングデータを選択し OK ファイルシステム 上の 場所 欄に /lustre/gt00h/ ユーザー名 /OpenACC_samples/ とするのが速いか? 9 diffusion.nvp を選択できたら Finish GPU プログラミング入門 24

25 ここで拡大 この辺がメインの計算部分 GPU プログラミング入門 25

26 計算部分をクリックすると カーネルの情報が得られる ここをクリックするとプロファイラがヒントをくれる GPU プログラミング入門 26

27 まとめ GPUとOpenACCの基礎について ICCGソルバーを題材とし OpenACC 化を行なった ICCGソルバーのOpenACC 化を通じ 以下を実習した kernels 指示文による並列化 コンパイラメッセージの見方 PGI_ACC_TIMEの出力 データ転送の最適化 さらなる最適化 NVVPを使ったプロファイリング Async 指示子 カーネルフュージョン マルチコア メニィコア並列プログラミング入門 27

28 Q & A アカウントは1ヶ月有効です ぜひご自分のアプリで試してください 資料のPDF 版はWEBページにございます アンケートへの協力をお願いします GPU プログラミング入門 28