OpenACCによる並列化

(1)

実習

OpenACC

による

ICCG

ソルバーの並列化

(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

(3)

Makefile for OpenACC

• OpenACC用のMakefileの作成

CC = icc OPTFLAGS= -O3 -qopenmp -ipo -xCORE-AVX2 -align TARGET = ../run/L3-solx

C/L3/srcx_acc/Makefile

F90 = ifort

F90OPTFLAGS= -O3 -qopenmp -ipo -xCORE-AVX2 -align array32byte

…

TARGET = ../run/L3-solx

F/L3/srcx_acc/Makefile

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)で並列化されたのか知るため}

(4)

ジョブスクリプト

for OpenACC

• $ cd ../run

• $ cp gor.sh gor-acc.sh

–

or $ cp go.sh gor-acc.sh

• gor-acc.sh の編集

#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

C/L3/run/gor-acc.sh

F/L3/run/gor-acc.sh

#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

PGI_ACC_TIME=1 で、GPU実行情

報のサマリを標準エラーに出力

(5)

GPU用のインプットファイルの作成

• INPUT.dat の編集

128 128 128 NX/NY/NZ 1.00e-0 1.00e-0 1.00e0 DX/DY/DZ 1.0e-08 OMEGA, EPSICCG 18 PEsmpTOT -30 NCOLORtot 0 0

C/L3/run/INPUT.dat

F/L3/run/INPUT.dat

128 128 128 NX/NY/NZ 1.00e-0 1.00e-0 1.00e0 DX/DY/DZ 1.0e-08 OMEGA, EPSICCG 1 PEsmpTOT -30 NCOLORtot 0

GPUのスレッド数は数百万に及ぶの

で、スレッド数を

_{PEsmpTOTで制御す}

るのは現実的ではない

(6)

CPU版の結果の取得

• 結果の正しさを検証するために、

1 CPUコアで実行した結果

を取っておく

• 下の結果は

1CPUコアで実行した際のもの

–

PGIコンパイラ使用（OpenACCなし）

–

NCOLORtot = -30 # INPUT.dat

–

自分で実行してみてください

### CMRCM ### FINAL COLOR NUMBER 30 3.546910e-01 sec. (assemble) 1 1.437100e+01 101 4.996898e-01 201 2.067294e-04 301 3.749532e-08 313 9.310708e-09 N= 2097152 2.163149e+01 sec. (solver)

C/L3/run/gor-acc.sh.oXXXXX

### CM-RCM ### FINAL COLOR NUMBER 30 69889 3.818271E-01 sec. (assemble) 1 1.437100E+01 101 4.996898E-01 201 2.067294E-04 301 3.749532E-08 313 9.310708E-09 8.931846E+00 9.503229E+00 N= 2097152 2.442678E+01 sec. (solver)

F/L3/run/gor-acc.sh.oXXXXX

← Quick check Is the number of iteration same as baseline?

(7)

OpenACCでの並列化戦略

• OpenACCではどのループを並列化するべきか

#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 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; } }

C/L3/srcx_acc/solver_ICCG_mc.c

!$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

F/L3/srcx_acc/solver_ICCG_mc.f

← PEsmpTOT = 1 ip1 = ic ip2 = ic+1

OpenACCでのターゲットループ

←ここで同期が必要！

前の色が終わったところで同期

が必要となる。

_{OpenMPでは暗}

黙的に同期が入る。

OpenACCでは同期を取るため

にカーネルを閉じるしかない！

(8)

OpenACC 指示文の挿入

#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; }

C/L3/srcx_acc/solver_ICCG_mc.c

!$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

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; } !$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

C では配列のサイズ情報必須！

Fortranでは配列がサイズ情報も持っているた

め必要ない

(9)

PGIコンパイラのメッセージ確認

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

C/L3/srcx_acc/

$make

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

F/L3/srcx_acc/

$ make

並列化されていない！

XとWがaliasを持ってる可能性があるため、

コンパイラは並列化をしない

並列化されてる！

gang, vector レベルで並列化されている

(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]) #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 が必要になる！

nothing to do

#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; }

(11)

PGIコンパイラのメッセージ確認

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 */

C/L3/srcx_acc/

$make

F/L3/srcx_acc/

$ make

ようやく並列化

(12)

OpenACC 指示文の挿入

#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]; #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; } }

C/L3/srcx_acc/solver_ICCG_mc.c

!$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

F/L3/srcx_acc/solver_ICCG_mc.f

← PEsmpTOT = 1 ip1 = ic ip2 = ic+1

配列の確保は

_{poi_gen.cでなされている。}

AL, itemL の長さNPLはこのソースからは見えな

い！

配列のサイズ情報を書く必要はない

↓このループは短いので逐 次計算(長さ3 or 6)

(13)

関数の引数の書き換え

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)

C/L3/srcx_acc/solver_ICCG_mc.c

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

(14)

main.c のトラップ

int main() { double *WK; int NPL, NPU; int ISET, ITR, IER; int icel, ic0, i; double xN, xL, xU; double Stime, Etime;

C/L3/srcx_acc/main.c

F/L3/srcx_acc/main.f

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;

(15)

OpenACC 版の結果チェック

15 ### CMRCM ### FINAL COLOR NUMBER 30 3.461821e-01 sec. (assemble) 1 1.437100e+01 101 4.996898e-01 201 2.067294e-04 301 3.749532e-08 313 9.310708e-09 N= 2097152 2.437671e+01 sec. (solver)

C/L3/run/gor-acc.sh.oXXXXX

### CM-RCM ### FINAL COLOR NUMBER 30 69889 3.770702E-01 sec. (assemble) 1 1.437100E+01 101 4.996898E-01 201 2.067294E-04 301 3.749532E-08 313 9.310708E-09 9.091849E+00 9.593184E+00 N= 2097152 2.574174E+01 sec. (solver)

F/L3/run/gor-acc.sh.oXXXXX

← ベースラインと一致！ 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=166

C/L3/run/gor-acc.sh.eXXXXX

← この時点ではCPUより遅 いが、気にしない

←標準エラー出力側に、PGI_ACC_TIMEによる

サマリが出力される

(Fortran版も同様)

gridの値がgang(スレッドブロック)数、

blockの値がgangあたりのvector(スレッド)数

(16)

データ転送の最適化

#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; } } }

C/L3/srcx_acc/solver_ICCG_mc.c

!$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 !$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

F/L3/srcx_acc/solver_ICCG_mc.f

!$acc data 指示文で確保済みの配列に対して、それより内側で_{copy指示子を適用しても} 無視される。 →冗長な転送を除去！

(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

Data transfer

Device

Host

(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

Device

Host

Data transfer

(19)

実習

• EX1：Makefile、ジョブスクリプト、INPUT.datを編集し、CPU実行時の

出力を得てください

• EX2：全並列化ループ(OMP DO で並列化されているところ)を

OpenACCで並列化してください

–

コンパイラメッセージに注意！

–

常に

CPUでの結果と突き合わせ！

• 計算順序が変わるため、必ずしも一致しない

–

この時点で遅くても気にしない！

• EX3：データ転送を最適化してください

• EX4：INPUT.datの色数(NCOLORtot)を変更し、速度への影響をみて

ください

–

PGI_ACC_TIMEは若干速度に影響を与えるので、計測時は

PGI_ACC_TIME=0 (ジョブスクリプト)としてください

マルチコア・メニィコア並列プログラミング入門

(20)

参考：

Reedbush での nvvp の使い方

• _{NVVP：NVIDIA Visual Profiler}

(

https://developer.nvidia.com/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プログラミング入門

(21)

GPUプログラミング入門 21

①

Fileをクリック

(22)

②

Importをクリック

③

Nvprof を選択

(23)

⑤

Multiple processesに変更

⑥

⑧ ファイル選択画面が開くので、

_{nvvpで作成したプロファ}

イリングデータを選択し、

_OK

ファイルシステム

_{→上の「場所」欄に}

/lustre/gt00h/ユーザー名/OpenACC_samples/…

とするのが速いか

_…？

⑨

diffusion.nvpを選択できたら Finish

(25)

この辺がメインの計算部分

(26)

計算部分をクリックすると、

カーネルの情報が得られる

ここをクリックするとプロ

ファイラがヒントをくれる

(27)

まとめ

• _{GPUとOpenACCの基礎について}

• ICCGソルバーを題材とし、OpenACC化を行なった

• ICCGソルバーのOpenACC化を通じ、以下を実習した

–

kernels指示文による並列化

–

コンパイラメッセージの見方

–

PGI_ACC_TIMEの出力

–

データ転送の最適化

• さらなる最適化

–

NVVPを使ったプロファイリング

–

Async 指示子

–

カーネルフュージョン

マルチコア・メニィコア並列プログラミング入門

(28)

Q & A

• アカウントは

1ヶ月

有効です。

–

ぜひご自分のアプリで試してください

• 資料の

PDF版はWEBページにございます。

–

http://www.cc.u-tokyo.ac.jp/support/kosyu/87/

• アンケートへの協力をお願いします。

GPUプログラミング入門