OpenACC

Size: px

Start display at page:

Download "OpenACC"

ゆたかかやぬま
5 years ago
Views:

1 109 OpenMP/OpenACC, cc.u-tokyo.ac.jp cc.u-tokyo.ac.jp 1

2 n Reedbush n $ ssh -Y reedbush.cc.u-tokyo.ac.jp l txxxxx n module n $ module load pgi/18.7 # n n $ cdw n OpenACC_samples n $ cp /lustre/gt00/share/openmp_openacc.tar.gz. n $ tar zxvf OpenMP_OpenACC.tar.gz u 自宅で学習したい人へ u PGI compiler の無償版をダウンロードしましょう -> PGI の代理店 (SofTek さん ) のページ u 本実習のプログラムは以下 -> 2

3 GPU 3

4 What s GPU? Graphics Processing Unit n もともと PC の3D描画専用の装置 n n パソコンの部品として量産されてる = 非常に安価 3D Game Computer Graphics GPU 4

5 GPUコンピューティング n n n n n GPUはグラフィックスやゲームの画像計算のために進化を続けている CPUがコア数が2-12個程度に対し GPUは1000以上のコアがある GPUを一般のアプリケーションの高速化に利用することを GPUコンピューティング GPGPU (General Purpose computation on GPU) などという 2007年にNVIDIA社のCUDA言語がリリースされて大きく発展ここ数年ディープラーニング深層学習機械学習 AI 人工知能などでも注目を浴びている 5

6 GPU n n n GPU ü ü NVIDIA P100 (Reedbush-H) 5,304 GFlops ü Intel Xeon Phi (Oakforest-PACS) 3,046.4 GFLops ü ü ü CPU ü CPU GPU 1000 ü CPU ü GPU CPU 6

7 NVIDIA Tesla P100 n 56 SMs, 3584 CUDA, 16 GByte Tesla P100 whitepaper 7

8 n T p T / p T p T/p n ü ü n ü 8

9 n n ü ü ü ü ü n 9

10 n ü n ü = = = = = = = = 4-6 = = = = = = = 3 11 = = = = = = = 1-10 = = 10

11 GPU 1. 必要なデータを送るノードの外へバス (PCIe など ) ~20GB/s CPU OS が動いている ~32GB/s GPU OS は存在しない 3. 計算結果を返す 2. 計算を行う ~200GB/s ~1,000GB/s メインメモリデバイスメモリ n OS CPU n 11

12 GPU n CPU GPU ü CPU GPU CPU GPU ü OpenACC n >> P ü ü ü Intel CPU Hyperthread x 2 12

13 GPU n ü 32 Warp ü Warp ü branch divergence, divergent branch if ( 奇数番スレッド ) { 処理 A; else { 処理 B; 13

14 GPU n Warp ü coalesced access ü 128 Byte 128 Byte Byte 128 byte x 1 Address Thread 128 byte x Address Thread

15 OPENACC 15

16 GPU n CUFFT, CUBLAS ü GPU ü n OpenACC ü ü n CUDA OpenCL ü GPU ü GPGPU 16

17 OpenACC n OpenACC ü ü OpenMP ü C /C++, Fortran ü 2011 OpenACC ü PGI, Cray, GCC PGI ü WEB n ü ü ü CPU C 言語 #pragma acc directive-name [clause, ] { // C code Fortran!$acc directive-name [clause, ]! Fortran code!$acc end directive-name 17

18 OpenMP OpenACC 1 OpenMP OpenACC int main(){... #pragma acc... for (i=0; i<n; i++) { CPU CPU CPU GPU) 18

19 OpenACC C openacc_hello/01_hello_acc CPU GPU int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; #pragma acc data copyin(a[0:n]), copyout(b[0:n]) #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; double sum = 0; for (int i=0; i<n; i++) { sum += b[i]; fprintf(stdout, "%f n", sum/n); free(a); free(b); return 0; a b GPU へ copyin a b GPU から copyout 19

20 OpenACC C openacc_hello/01_hello_acc CPU GPU int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; #pragma acc data copyin(a[0:n]), copyout(b[0:n]) #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; double sum = 0; for (int i=0; i<n; i++) { sum += b[i]; fprintf(stdout, "%f n", sum/n); free(a); free(b); return 0; a b GPU へ copyin a b GPU から copyout 20

21 OpenACC F openacc_hello/01_hello_acc CPU GPU program main implicit none! 変数宣言 allocate(a(n),b(n)) c = 2.0 do i = 1, n a(i) = 10.0 end do!$acc data copyin(a) copyout(b)!$acc kernels!$acc loop independent do i = 1, n b(i) = a(i) + c end do!$acc end kernels!$acc end data sum = 0.d0 do i = 1, n sum = sum + b(i) end do print *, sum/n deallocate(a,b) end program main a b GPU へ copyin a b GPU から copyout 21

22 OpenACC n ü kernels, parallel n ü data, enter data, exit data, update n ü loop n ü host_data, atomic, routine, declare 22

23 OpenACC n kernels ü ü int main() { #pragma acc kernels { for (int i=0; i<n; i++) { A; 1 program main!$ acc kernels do i = 1, n A; end do for (int i=0; i<n; i++) { B; 2 do i = 1, n B; end do!$acc end kernels ü parallel 23

24 CPU OpenACC C openacc_hello/01_hello_acc int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; n OpenACC ü kernels, loop for (int i=0; i<n; i++) { b[i] = a[i] + c; double sum = 0; for (int i=0; i<n; i++) { sum += b[i]; fprintf(stdout, "%f n", sum/n); free(a); free(b); return 0; 24

25 CPU OpenACC C openacc_hello/01_hello_acc int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; n OpenACC ü kernels, loop #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; double sum = 0; for (int i=0; i<n; i++) { sum += b[i]; fprintf(stdout, "%f n", sum/n); free(a); free(b); return 0; 25

26 CPU OpenACC C openacc_hello/01_hello_acc int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; n OpenACC ü kernels, loop #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; カーネルとしてコンパイルされ GPU 上で実行される配列の 1 要素が 1 スレッドで処理されるイメージ double sum = 0; for (int i=0; i<n; i++) { sum += b[i]; fprintf(stdout, "%f n", sum/n); free(a); free(b); return 0; 26

27 CPU OpenACC F openacc_hello/01_hello_acc program main implicit none! 変数宣言 allocate(a(n),b(n)) c = 2.0 n OpenACC ü kernels, loop do i = 1, n a(i) = 10.0 end do!$acc kernels!$acc loop independent do i = 1, n b(i) = a(i) + c end do!$acc end kernels Fortran も同じ sum = 0.d0 do i = 1, n sum = sum + b(i) end do print *, sum/n deallocate(a,b) end program main 27

28 OpenACC n ü firstprivate private ü OpenMP shared ü n ü shared ü ü n kernels ü OpenACC ü shared ü data 28

29 C n data ü (GPU) (CPU) (GPU) kernels data ü CUDA cudamalloc, cudamemcpy int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; #pragma acc data copyin(a[0:n]), copyout(b[0:n]) #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; openacc_hello/01_hello_acc 変数 c はスカラ変数のため自動的にデバイスへコピーされプライベート変数となる 29

0 do i = 1, n a(i) = 10.0 end do!$acc data copyin(a) copyout(b)!$acc kernels!

30 F n data ü (GPU) (CPU) (GPU) kernels data ü CUDA cudamalloc, cudamemcpy program main implicit none integer,parameter :: n = 1000 real(kind=4),allocatable,dimension(:) :: a,b real(kind=4) :: c integer :: i real(kind=8) :: sum allocate(a(n),b(n)) c = 2.0 do i = 1, n a(i) = 10.0 end do!$acc data copyin(a) copyout(b)!$acc kernels!$acc loop independent do i = 1, n b(i) = a(i) + c end do!$acc end kernels!$acc end data openacc_hello/01_hello_acc Fortran では配列サイズ情報が変数に付随するため (lbound,ubound,size などの組み込み関数をサポートしている ) 基本的にサイズを書く必要がない 30

31 data n n n n n copy ü allocate, memcpy(h->d), memcpy(d->h), deallocate copyin ü allocate, memcpy(h->d), deallocate ü copyout ü allocate, memcpy(d->h), deallocate ü create ü allocate, deallocate ü present ü n copy/copyin/copyout/create present OpenACC2.5 31

32 enter data, exit data n data #pragma acc data copyin(a[0:n]) copyout(b[0:n]) { goto hoge; hoge: NG!$acc data copyin(a), copyout(b) goto 1000!$acc end data 1000 NG ü goto n enter data exit data data #pragma acc enter data copyin(a[0:n]) create(b[0:n]) { goto hoge: hoge: #pragma acc exit data delete(a[0:n]) copyout(b[0:n]) OK!$acc enter data!$acc& copyin(a) create(b) goto !$acc exit data!$acc& delete(a) copyout(b) NG ü enter data copyout create, exit data copyin delete 32

33 n n n Fortran C n A n Fortran:!$acc data copy(a(lower1:upper1, lower2:upper2) )...!$acc end data n C : #pragma acc data copy(a[begin1:length1][begin2:length2])... 33

34 C openacc_hello/01_hello_acc CPU GPU int main(){ const int n = 1000; float *a = malloc(n*sizeof(float)); float *b = malloc(n*sizeof(float)); float c = 2.0; for (int i=0; i<n; i++) { a[i] = 10.0; #pragma acc data copyin(a[0:n]), copyout(b[0:n]) #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; double sum = 0; for (int i=0; i<n; i++) { sum += b[i]; fprintf(stdout, "%f n", sum/n); free(a); free(b); return 0; loop a b GPU へ copyin a b GPU から copyout 34

35 F openacc_hello/01_hello_acc CPU GPU program main implicit none! 変数宣言 allocate(a(n),b(n)) c = 2.0 a b do i = 1, n a(i) = 10.0 end do!$acc data copyin(a) copyout(b)!$acc kernels!$acc loop independent do i = 1, n b(i) = a(i) + c end do!$acc end kernels!$acc end data sum = 0.d0 do i = 1, n sum = sum + b(i) end do print *, sum/n deallocate(a,b) end program main loop GPU へ copyin a b GPU から copyout 35

36 loop n loop ü CUDA gang, worker, vector gang: CUDA thread block vector: CUDA block threads ü (independent clause) C ü (reduction clause) ü (seq clause) 36

37 n independent ü ü #pragma acc kernels #pragma acc loop independent for (int i=0; i<n; i++) { b[i] = a[i] + c; independent n // これは正しくない #pragma acc kernels #pragma acc loop independent for (int i=1; i<n; i++) { d[i] = d[i-1]; 37

38 OpenACC CUDA // OpenACC void calc(int n, const float *a, const float *b, float c, float *d) { #pragma acc kernels present(a, b, d) #pragma acc loop independent for (int i=0; i<n; i++) { d[i] = a[i] + c*b[i]; kernel // CUDA global void calc_kernel(int n, const float *a, const float *b, float c, float *d) { const int i = blockidx.x * blockdim.x + threadidx.x; if (i < n) { d[i] = a[i] + c*b[i]; int main() {... #pragma acc data copyin(a[0:n], b[0:n]) copyout(d[0:n]) { calc(n, a, b, c, d);... ü ü ü kernels GPU loop data kernels void calc(int n, const float *a, const float *b, float c, float *d) { dim3 threads(128); dim3 blocks((n + threads.x - 1) / threads.x); calc_kernel<<<blocks, threads>>>(n, a, b, c, d); cudathreadsynchronize(); int main() {... float *a_d, *b_d, *d_d; cudamalloc(&a_d, n*sizeof(float)); cudamalloc(&b_d, n*sizeof(float)); cudamalloc(&d_d, n*sizeof(float)); cudamemcpy(a_d, a, n*sizeof(float), cudamemcpydefault); cudamemcpy(b_d, b, n*sizeof(float), cudamemcpydefault); cudamemcpy(d_d, d, n*sizeof(float), cudamemcpydefault); calc(n, a_d, b_d, c, d_d); cudamemcpy(d, d_d, n*sizeof(float), cudamemcpydefault);... 38

39 OpenACC n PGI ü Reedbush OpenACC PGI $ module load pgi/18.7 $ pgcc -O3 -acc -Minfo=accel -ta=tesla,cc60 -c main.c -acc: OpenACC -Minfo=accel: OpenACC GPU OpenACC -ta=tesla,cc60: NVIDIA GPU Tesla compute capability 6.0 (cc60) n Makefile Makefile $ module load pgi/18.7 $ make 39

40 OpenACC n : openacc_basic/ C ü OpenACC kernels, data, loop ü F for (unsigned int j=0; j<ny; j++) { for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; do j = 1,ny do i = 1,nx c(i,j) = a(i,j) + b(i,j) end do end do ü openacc_basic/01_original openacc_basic/02_kernels openacc_basic/03_kernels_copy openacc_basic/04_loop openacc_basic/05_data openacc_basic/06_present openacc_basic/07_reduction CPUコード OpenACCコード上にkernels 指示文のみ追加 OpenACCコード上にcopy 指示節追加 OpenACCコード上にloop 指示文を追加 OpenACCコード上にdata 指示文を明示的に追加 OpenACCコード上でpresent 指示節を使用 OpenACC コード上に reduction 指示節を使用 40

41 C n : openacc_basic/ ü OpenACC kernels, data, loop ü void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ for (unsigned int j=0; j<ny; j++) { for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; ix = j*nx + i ny j i nx 41

42 F n : openacc_basic/ ü OpenACC kernels, data, loop ü subroutine calc(nx, ny, a, b, c) implicit none integer,intent(in) :: nx,ny real(kind=4),dimension(:,:),intent(in) :: a,b real(kind=4),dimension(:,:),intent(out) :: c integer :: i,j do j = 1,ny do i = 1,nx c(i,j) = a(i,j) + b(i,j) end do end do end subroutine calc Fortran 版では多次元配列を利用 42

43 OpenACC: CPU n CPU ü $ cd openacc_basic/01_original $ make $ qsub./run.sh $ cat run.sh.o?????? mean = Time = [sec] n ü a b c 1.0, 2.0, 0.0 ü calc c += a *b nt(=1000) ü openacc_basic/01_original 43

44 C F OpenACC: kernels n 02_kernels : calc ü CPU kernels openacc_basic/02_kernels void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ const unsigned int n = nx * ny; #pragma acc kernels for (unsigned int j=0; j<ny; j++) { for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; subroutine calc(nx, ny, a, b, c) implicit none integer,intent(in) :: nx,ny real(kind=4),dimension(:,:),intent(in) :: a,b real(kind=4),dimension(:,:),intent(out) :: c integer :: i,j!$acc kernels do j = 1,ny do i = 1,nx c(i,j) = a(i,j) + b(i,j) end do end do!$acc end kernels end subroutine calc OpenACC コンパイラは配列 (a, b, c) を shared 変数として自動で転送してくれるはずだが 44 C F

45 C F OpenACC: kernels n データサイズがわからずコンパイルエラー C 言語では配列サイズの指定がほぼ必須! $ make pgcc -O3 -acc -Minfo=accel -ta=tesla,cc60 -c main.c PGC-S-0155-Compiler failed to translate accelerator region (see -Minfo messages): Could not find allocated-variable index for symbol (main.c: 13) calc: 14, Complex loop carried dependence of a->,c->,b-> prevents parallelization Accelerator serial kernel generated Accelerator kernel generated Generating Tesla code 14, #pragma acc loop seq 15, #pragma acc loop seq 15, Accelerator restriction: size of the GPU copy of c,b,a is unknown Complex loop carried dependence of a->,c->,b-> prevents parallelization PGC-F-0704-Compilation aborted due to previous errors. (main.c) PGC/x86-64 Linux : compilation aborted make: *** [main.o] Error 2 $ make pgfortran -O3 -mp -acc -ta=tesla,cc60 -Minfo=accel -c main.f90 calc: 13, Generating implicit copyin(b(:nx,:ny)) Generating implicit copyout(c(:nx,:ny)) Generating implicit copyin(a(:nx,:ny)) 14, Loop is parallelizable 15, Loop is parallelizable Accelerator kernel generated Generating Tesla code 14,!$acc loop gang, vector(4)! blockidx%y threadidx%y 15,!$acc loop gang, vector(32)! blockidx%x threadidx%x 45 C F データサイズを検知して自動転送 Fotran ではサイズ情報が配列に付随するため

46 OpenACC: kernels C n 03_kernels_copy : calc ü void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ const unsigned int n = nx * ny; #pragma acc kernels copy(a[0:n], b[0:n], c[0:n]) for (unsigned int j=0; j<ny; j++) { allocate, H -> D for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; D->H, deallocate ü kernels data ü copy openacc_basic/03_kernels_copy ü GPU CPU 46

47 OpenACC: kernels n 03_kernels_copy : ü CPU kernels openacc_basic/03_kernels_copy C F C int main(int argc, char *argv[]) {... #pragma acc kernels copyout(b[0:n], c[0:n]) { for (unsigned int i=0; i<n; i++) { b[i] = b0; for (unsigned int i=0; i<n; i++) { c[i] = 0.0;... F program main...!$acc kernels copyout(b,c) do j = 1,ny do i = 1,nx b(i,j) = b0 end do end do c(:,:) = 0.0!$acc end kernels... end program 47

48 OpenACC: kernels C n ü $ make pgcc -O3 -acc -Minfo=accel -ta=tesla,cc60 -c main.c calc: 13, Generating copy(a[:n],c[:n],b[:n]) 14, Complex loop carried dependence of a-> prevents parallelization Loop carried dependence due to exposed use of c[:n] prevents parallelization Complex loop carried dependence of c->,b-> prevents parallelization Accelerator scalar kernel generated Accelerator kernel generated Generating Tesla code 14, #pragma acc loop seq 15, #pragma acc loop seq 15, Complex loop carried dependence of a->,c->,b-> prevents parallelization Loop carried dependence due to exposed use of c[:i1+n] prevents parallelization main: 43, Generating copyout(c[:n],b[:n]) 45, Loop is parallelizable Accelerator kernel generated Generating Tesla code 45, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ 48, Loop is parallelizable Accelerator kernel generated Generating Tesla code 48, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ pgcc -O3 -acc -Minfo=accel -ta=tesla,cc60 main.o -o run 48

49 OpenACC: kernels n ü pgfortran -O3 -mp -acc -ta=tesla,cc60 -Minfo=accel -c main.f90 calc: 13, Generating copyin(a(:,:)) Generating copyout(c(:,:)) Generating copyin(b(:,:)) 14, Loop is parallelizable 15, Loop is parallelizable Accelerator kernel generated Generating Tesla code 14,!$acc loop gang, vector(4)! blockidx%y threadidx%y 15,!$acc loop gang, vector(32)! blockidx%x threadidx%x main: 61, Generating copyout(c(:,:),b(:,:)) 62, Loop is parallelizable 63, Loop is parallelizable Accelerator kernel generated Generating Tesla code 62,!$acc loop gang, vector(4)! blockidx%y threadidx%y 63,!$acc loop gang, vector(32)! blockidx%x threadidx%x 68, Loop is parallelizable Accelerator kernel generated Generating Tesla code 68,!$acc loop gang, vector(4)! blockidx%y threadidx%y!$acc loop gang, vector(32)! blockidx%x threadidx%x pgfortran -O3 -mp -acc -ta=tesla,cc60 -Minfo=accel main.o -o run 49 F

50 Tips: n n n n foo(&a[0],&a[1]) n n n n Fortran n C n これってデータ独立? void foo(float *a, float *b){ for (int i=0; i<n; i++) b[i] = a[i]; インデックス計算 for (int i=0; i<n; i++){ j = i % 10; b[j] = a[i]; 間接参照 for (int i=0; i<n; i++){ b[idx[i]] = a[i]; 50

51 OpenACC: loop C n 04_loop ü 03_kernels loop independent openacc_basic/04_loop void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ const unsigned int n = nx * ny; #pragma acc kernels copy(a[0:n], b[0:n], c[0:n]) #pragma acc loop independent for (unsigned int j=0; j<ny; j++) { #pragma acc loop independent for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; // main 関数内 #pragma acc kernels copyout(b[0:n], c[0:n]) { #pragma acc loop independent for (unsigned int i=0; i<n; i++) { b[i] = b0; #pragma acc loop independent for (unsigned int i=0; i<n; i++) { c[i] = 0.0; 51

52 OpenACC: loop F n 04_loop ü 03_kernels loop independent subroutine calc(nx, ny, a, b, c)...!$acc kernels copyin(a,b) copyout(c)!$acc loop independent do j = 1,ny!$acc loop independent do i = 1,nx c(i,j) = a(i,j) + b(i,j) end do end do!$acc end kernels end subroutine openacc_basic/04_loop! main 関数内!$acc kernels copyout(b,c)!$acc loop independent do j = 1,ny!$acc loop independent do i = 1,nx b(i,j) = b0 end do end do c(:,:) = 0.0!$acc end kernels 各次元について loop 指示文を指定する ( 並列サイズなどを指定したいなど ) 場合 do 文で書き下す必要がある 52

53 OpenACC: loop C n ü openacc_basic/04_loop $ make pgcc -O3 -acc -Minfo=accel -ta=tesla,cc60 -c main.c calc: 13, Generating copy(a[:n],c[:n],b[:n]) 15, Loop is parallelizable 17, Loop is parallelizable Accelerator kernel generated Generating Tesla code 15, #pragma acc loop gang, vector(4) /* blockidx.y threadidx.y */ 17, #pragma acc loop gang, vector(32) /* blockidx.x threadidx.x */ main: 45, Generating copyout(c[:n],b[:n]) 48, Loop is parallelizable Accelerator kernel generated Generating Tesla code 48, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ 52, Loop is parallelizable Accelerator kernel generated Generating Tesla code 52, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ Fortran 版は既に並列化されていたため省略 loop independent をつける事による挙動の変化はない ( 少なくとも PGI compiler ver では ) 53

54 OpenACC: loop n 04_loop ü $ qsub./run.sh $ cat run.sh.o?????? mean = Time = [sec] openacc_basic/04_loop ü calc GPU CPU void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ const unsigned int n = nx * ny; #pragma acc kernels copy(a[0:n], b[0:n], c[0:n]) #pragma acc loop independent allocate, H -> D for (unsigned int j=0; j<ny; j++) { #pragma acc loop independent for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; D->H, deallocate 54

55 OpenACC: data C n 05_data ü 04_loop data // main 関数内 #pragma acc data copyin(a[0:n]) create(b[0:n]) copyout(c[0:n]) { #pragma acc kernels copyout(b[0:n], c[0:n]) { #pragma acc loop independent for (unsigned int i=0; i<n; i++) { b[i] = b0; #pragma acc loop independent for (unsigned int i=0; i<n; i++) { c[i] = 0.0; for (unsigned int icnt=0; icnt<nt; icnt++) { calc(nx, ny, a, b, c); openacc_basic/05_data a: allocate, H -> D b: allocate c: allocate present a: deallocate b: deallocate c: D->H, deallocate ü copy/copyin/copyout/create present OpenACC2.5 ü a, b, c 55

56 OpenACC: data F n 05_data ü 04_loop data! main 関数内!$acc data copyin(a) create(b) copyout(c)!$acc kernels copyout(b,c)!$acc loop independent do j = 1,ny!$acc loop independent do i = 1,nx b(i,j) = b0 end do end do openacc_basic/05_data a: allocate, H -> D b: allocate c: allocate present c(:,:) = 0.0!$acc end kernels do icnt = 1,nt call calc(nx, ny, a, b, c) end do!$acc end data a: deallocate b: deallocate c: D->H, deallocate ü copy/copyin/copyout/create present OpenACC2.5 ü a, b, c 56

57 OpenACC: data n 05_data ü $ qsub./run.sh $ cat run.sh.o?????? mean = Time = [sec] openacc_basic/05_data 57

$OpenACC: present C n 06_present ü 05_data present openacc_basic/06_present void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ const unsigned int n = nx * ny;$

58 OpenACC: present C n 06_present ü 05_data present openacc_basic/06_present void calc(unsigned int nx, unsigned int ny, const float *a, const float *b, float *c){ const unsigned int n = nx * ny; #pragma acc kernels present(a, b, c) #pragma acc loop independent present for (unsigned int j=0; j<ny; j++) { #pragma acc loop independent for (unsigned int i=0; i<nx; i++) { const int ix = i + j*nx; c[ix] += a[ix] + b[ix]; F subroutine calc(nx, ny, a, b, c)...!$acc kernels present(a, b, c)!$acc loop independent do j = 1,ny!$acc loop independent do i = 1,nx c(i,j) = a(i,j) + b(i,j) end do end do!$acc end kernels end subroutine u データ転送の振る舞いは変化しないため性能変化はなし u present ではメモリ確保データ転送をしないため配列サイズの指定は不要 u コードとしては見通しがよい 58

59 n ü ü ü CUDA 1 double sum = 0.0; for (unsigned int i=0; i<n; i++) { sum += array[i];

60 ( n loop reduction ü reduction double sum = 0.0; #pragma acc kernels #pragma acc loop reduction(+:sum) for (unsigned int i=0; i<n; i++) { sum += array[i]; n Reduction ü acc loop reduction(+:sum) ü n ü : +, : 0 ü : *, : 1 ü : max, : least ü : min, : largest 60

61 OpenACC: reduction C n 07_reduction ü 06_present reduction // main 関数内 for (unsigned int icnt=0; icnt<nt; icnt++) { calc(nx, ny, a, b, c); openacc_basic/07_reduction #pragma acc kernels #pragma acc loop reduction(+:sum) for (unsigned int i=0; i<n; i++) { sum += c[i]; ü data c create n 07_reduction ü $ make pgcc -O3 -acc -Minfo=accel -ta=tesla,cc60 -c main.c ( 省略 ) main: ( 省略 ) 67, Loop is parallelizable Accelerator kernel generated Generating Tesla code 67, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ Generating reduction(+:sum) 61

62 OpenACC: reduction n 07_reduction ü 06_present reduction sum = 0!$acc kernels present(c)!$acc loop reduction(+:sum) do j = 1,ny!$acc loop reduction(+:sum) do i = 1,nx sum = sum + c(i,j) end do end do!$acc end kernels ü data c create n 07_reduction ü $ make pgfortran -O3 -mp -acc -ta=tesla,cc60 -Minfo=accel -c main.f90 ( 省略 ) main: ( 省略 ) 86, Loop is parallelizable Accelerator kernel generated Generating Tesla code 84,!$acc loop gang, vector(4)! blockidx%y threadidx%y 86,!$acc loop gang, vector(32)! blockidx%x threadidx%x Generating reduction(+:sum) openacc_basic/07_reduction reduction 62 F

63 OpenACC: reduction n 07_reduction ü ü c GPU $ qsub./run.sh $ cat run.sh.o?????? mean = Time = [sec] openacc_basic/07_reduction 63

64 OpenACC C n OpenACC 3 ü kernels GPU ü data - ü loop #pragma acc data copyin(a[0:n]) create(b[0:n], c[0:n]) { #pragma acc kernels { #pragma acc loop independent for (unsigned int i=0; i<n; i++) { b[i] = b0; #pragma acc loop independent for (unsigned int i=0; i<n; i++) { c[i] = 0.0; for (unsigned int icnt=0; icnt<nt; icnt++) { calc(nx, ny, a, b, c); #pragma acc kernels #pragma acc loop reduction(+:sum) for (unsigned int i=0; i<n; i++) { sum += c[i]; openacc_basic/07_reduction 64

65 OpenACC F n OpenACC 3 ü kernels GPU ü data - ü loop!$acc data copyin(a) create(b,c)!$acc kernels present(b,c)!$acc loop independent do j = 1,ny!$acc loop independent do i = 1,nx b(i,j) = b0 end do end do c(:,:) = 0.0!$acc end kernels do icnt = 1,nt call calc(nx, ny, a, b, c) end do! 続く! 続き sum = 0!$acc kernels present(c)!$acc loop reduction(+:sum) do j = 1,ny!$acc loop reduction(+:sum) do i = 1,nx sum = sum + c(i,j) end do end do!$acc end kernels!$acc end data openacc_basic/07_reduction 65

66 OPENACC 66

67 C n 3 OpenACC ü openacc_diffusion/01_original n 3 CPU OpenACC kernels, data, loop GPU for(int k = 0; k < nz; k++) { for (int j = 0; j < ny; j++) { for (int i = 0; i < nx; i++) { const int ix = nx*ny*k + nx*j + i; const int ip = i == nx - 1? ix : ix + 1; const int im = i == 0? ix : ix - 1; const int jp = j == ny - 1? ix : ix + nx; const int jm = j == 0? ix : ix - nx; const int kp = k == nz - 1? ix : ix + nx*ny; const int km = k == 0? ix : ix - nx*ny; diffusion.c, diffusion3d fn[ix] = cc*f[ix] + ce*f[ip] + cw*f[im] + cn*f[jp] + cs*f[jm] + ct*f[kp] + cb*f[km]; openacc_diffusion/01_original 67

68 F n 3 OpenACC ü openacc_diffusion/01_original n 3 CPU OpenACC kernels, data, loop GPU do k = 1, nz do j = 1, ny do i = 1, nx diffusion.f90, diffusion3d w = -1; e = 1; n = -1; s = 1; b = -1; t = 1; if(i == 1) w = 0 if(i == nx) e = 0 if(j == 1) n = 0 if(j == ny) s = 0 if(k == 1) b = 0 if(k == nz) t = 0 fn(i,j,k) = cc * f(i,j,k) + cw * f(i+w,j,k) & + ce * f(i+e,j,k) + cs * f(i,j+s,k) + cn * f(i,j+n,k) & + cb * f(i,j,k+b) + ct * f(i,j,k+t) end do end do end do openacc_diffusion/01_original 68

69 n ü ü n ü 69

70 n ü ü 3 C 1 ü 2 n ü : nx * ny * nz 3 ü : nt 70

71 n 2 平均後の自分自身の値上下左右の値自分自身の値の 4 倍 j 最初の状態 12 回目の平均後 i 繰り返し平均化を行うとインクが拡散します 71

72 n 2 72

73 CPU n CPU $ cd openacc_diffusion/01_original $ make $ qsub./run.sh # cat run.sh.o?????? time( 0) = time( 100) = time( 200) = time(1000) = time(1100) = time(1200) = time(1300) = time(1400) = time(1500) = time(1600) = Time = [sec] Performance= 2.17 [GFlops] Error[128][128][128] = e-06 n OpenACC 73

74 OpenACC (0): Makefile n Makefile OpenACC acc C CC = pgcc CXX = pgc++ GCC = gcc RM = rm -f MAKEDEPEND = makedepend CFLAGS = -O3 -acc -Minfo=accel -ta=tesla,cc60 GFLAGS = -Wall -O3 -std=c99 CXXFLAGS = $(CFLAGS) LDFLAGS =... F F90 = pgfortran RM = rm -f FFLAGS... = -O3 -mp -acc -ta=tesla,cc60 -Minfo=accel 74

75 OpenACC (1): kernels C n diffusion3d kernels #pragma acc kernels copyin(f[0:nx*ny*nz]) copyout(fn[0:nx*ny*nz]) for(int k = 0; k < nz; k++) { for (int j = 0; j < ny; j++) { for (int i = 0; i < nx; i++) { const int ix = nx*ny*k + nx*j + i; const int ip = i == nx - 1? ix : ix + 1; const int im = i == 0? ix : ix - 1; const int jp = j == ny - 1? ix : ix + nx; const int jm = j == 0? ix : ix - nx; const int kp = k == nz - 1? ix : ix + nx*ny; const int km = k == 0? ix : ix - nx*ny; fn[ix] = cc*f[ix] + ce*f[ip] + cw*f[im] + cn*f[jp] + cs*f[jm] + ct*f[kp] + cb*f[km]; return (double)(nx*ny*nz)*13.0; diffusion.c, diffusion3d make 75

76 OpenACC (1): kernels F n diffusion3d kernels!$acc kernels copyin(f) copyout(fn) do k = 1, nz do j = 1, ny do i = 1, nx w = -1; e = 1; n = -1; s = 1; b = -1; t = 1; if(i == 1) w = 0 if(i == nx) e = 0 if(j == 1) n = 0 if(j == ny) s = 0 if(k == 1) b = 0 if(k == nz) t = 0 fn(i,j,k) = cc * f(i,j,k) + cw * f(i+w,j,k) & + ce * f(i+e,j,k) + cs * f(i,j+s,k) + cn * f(i,j+n,k) & + cb * f(i,j,k+b) + ct * f(i,j,k+t) end do end do end do!$acc end kernels diffusion.f90, diffusion3d make 76

77 OpenACC (2): loop C n diffusion3d loop #pragma acc kernels copyin(f[0:nx*ny*nz]) copyout(fn[0:nx*ny*nz]) #pragma acc loop independent for(int k = 0; k < nz; k++) { #pragma acc loop independent for (int j = 0; j < ny; j++) { #pragma acc loop independent for (int i = 0; i < nx; i++) { const int ix = nx*ny*k + nx*j + i; const int ip = i == nx - 1? ix : ix + 1; const int im = i == 0? ix : ix - 1; const int jp = j == ny - 1? ix : ix + nx; const int jm = j == 0? ix : ix - nx; const int kp = k == nz - 1? ix : ix + nx*ny; const int km = k == 0? ix : ix - nx*ny; fn[ix] = cc*f[ix] + ce*f[ip] + cw*f[im] + cn*f[jp] + cs*f[jm] + ct*f[kp] + cb*f[km]; return (double)(nx*ny*nz)*13.0; diffusion.c, diffusion3d make qsub./run.sh 77

78 OpenACC (2): loop F n diffusion3d loop!$acc kernels copyin(f) copyout(fn)!$acc loop independent do k = 1, nz!$acc loop independent do j = 1, ny!$acc loop independent do i = 1, nx w = -1; e = 1; n = -1; s = 1; b = -1; t = 1; if(i == 1) w = 0 if(i == nx) e = 0 if(j == 1) n = 0 if(j == ny) s = 0 if(k == 1) b = 0 if(k == nz) t = 0 fn(i,j,k) = cc * f(i,j,k) + cw * f(i+w,j,k) & + ce * f(i+e,j,k) + cs * f(i,j+s,k) + cn * f(i,j+n,k) & + cb * f(i,j,k+b) + ct * f(i,j,k+t) end do end do end do!$acc end kernels diffusion.f90, diffusion3d make qsub./run.sh 78

79 OpenACC (3): (1) C n diffusion3d present main data #pragma acc kernels present(f, fn) #pragma acc loop independent for(int k = 0; k < nz; k++) { #pragma acc loop independent for (int j = 0; j < ny; j++) { #pragma acc loop independent for (int i = 0; i < nx; i++) { const int ix = nx*ny*k + nx*j + i; const int ip = i == nx - 1? ix : ix + 1; const int im = i == 0? ix : ix - 1; const int jp = j == ny - 1? ix : ix + nx; const int jm = j == 0? ix : ix - nx; const int kp = k == nz - 1? ix : ix + nx*ny; const int km = k == 0? ix : ix - nx*ny; diffusion.c, diffusion3d fn[ix] = cc*f[ix] + ce*f[ip] + cw*f[im] + cn*f[jp] + cs*f[jm] + ct*f[kp] + cb*f[km]; return (double)(nx*ny*nz)*13.0; present 79

80 OpenACC (3): (1) F n diffusion3d present main data!$acc kernels copyin(f) copyout(fn)!$acc loop independent do k = 1, nz!$acc loop independent do j = 1, ny!$acc loop independent do i = 1, nx diffusion.f90, diffusion3d end do end do end do!$acc end kernels w = -1; e = 1; n = -1; s = 1; b = -1; t = 1; if(i == 1) w = 0 if(i == nx) e = 0 if(j == 1) n = 0 if(j == ny) s = 0 if(k == 1) b = 0 if(k == nz) t = 0 fn(i,j,k) = cc * f(i,j,k) + cw * f(i+w,j,k) & + ce * f(i+e,j,k) + cs * f(i,j+s,k) + cn * f(i,j+n,k) & + cb * f(i,j,k+b) + ct * f(i,j,k+t) present 80

81 OpenACC (4): (2) C n diffusion3d present main data #pragma acc data copy(f[0:n]) create(fn[0:n]) { start_timer(); main.c, main for (; icnt<nt && time + 0.5*dt < 0.1; icnt++) { if (icnt % 100 == 0) fprintf(stdout, "time(%4d) = %7.5f n", icnt, time); flop += diffusion3d(nx, ny, nz, dx, dy, dz, dt, kappa, f, fn); swap(&f, &fn); time += dt; elapsed_time = get_elapsed_time(); copy/create make OpenACC openacc_diffusion/02_openacc 81

82 OpenACC (4): (2) F n diffusion3d present main data!$acc data copy(f) create(fn) call start_timer() main.f90, main do icnt = 0, nt-1 if(mod(icnt,100) == 0) write (*,"(A5,I4,A4,F7.5)"), "time(",icnt,") = ",time flop = flop + diffusion3d(nx, ny, nz, dx, dy, dz, dt, kappa, f, fn) call swap(f, fn) time = time + dt if(time + 0.5*dt >= 0.1) exit end do elapsed_time = get_elapsed_time()!$acc end data copy/create make OpenACC openacc_diffusion/02_openacc 82

83 PGI_ACC_TIME OpenACC n n n PGI OpenACC PGI_ACC_TIME Linux PGI_ACC_TIME 1 $ export PGI_ACC_TIME=1 $./run Reedbush PGI_ACC_TIME $ cat run.sh.... /etc/profile.d/modules.sh module load pgi/18.7 export PGI_ACC_TIME=1./run openacc_diffusion/03_openacc _pgi_acc_time 83

84 PGI_ACC_TIME OpenACC n $ cat run.sh.e?????? Accelerator Kernel Timing data /lustre/pz0115/z30115/lecture/lecture_samples/openacc_diffusion/03_open acc_pgi_acc_time/main.c main NVIDIA devicenum=0 time(us): 6,359 38: data region reached 2 times 38: data copyin transfers: 1 device time(us): total=3,327 max=3,327 min=3,327 avg=3,327 55: data copyout transfers: 1 device time(us): total=3,032 max=3,032 min=3,032 avg=3,032 /lustre/pz0115/z30115/lecture/lecture_samples/openacc_diffusion/03_open acc_pgi_acc_time/diffusion.c diffusion3d NVIDIA devicenum=0 time(us): 101,731 19: compute region reached 1638 times 25: kernel launched 1638 times grid: [4x128x32] block: [32x4] device time(us): total=101,731 max=64 min=62 avg=62 elapsed time(us): total=136,255 max=540 min=81 avg=83 19: data region reached 3276 times 84

85 OpenACC ICCG 85

86 ICCG OpenACC n : openacc_iccg/ ü ü ü ü openacc_iccg/01_original openacc_iccg/02_setup OpenMP コードこれを改変します OpenACC 化を行うための準備をした OpenMP コード openacc_iccg/03_unified Unified memory 機能を用いた OpenACC コード ( 後述 ) openacc_iccg/04_data_present openacc_iccg/05_exclude_data cpy_time Unified memory 機能を用いない OpenACC コード CPU-GPU 間のデータ転送時間を時間計測から除いた OpenACC コード openacc_iccg/06_optimized 上を OpenACC で出来る範囲内で最適化 OpenACC コード 86

87 CPU n $ cat./run.sh #! /bin/sh #PBS -q h-lecture #PBS -l select=1:mpiprocs=1:ompthreads=18 #PBS -W group_list=gt00 #PBS -l walltime=00:05:00 openacc_iccg/01_original cd $PBS_O_WORKDIR. /etc/profile.d/modules.sh module load pgi/18.7 numactl./run -n 128 -c -20 -nt 18 n ü INPUT.DAT -nx 128 NX = 128 -n 128 NX = NY = NZ = 128 -c -20 NCOLORtot = -20 -e 1.0e-8 EPSICCG = 1.0e 08 -nt 18 PEsmpTOT = 18 87

88 CPU n $ qsub./run.sh $ cat./run.sh.o????? ### THREAD number= 18 You have elements. How many colors do you need? #COLOR must be more than 2 and #COLOR must not be more than CM if #COLOR.eq. 0 RCM if #COLOR.eq.-1 CMRCM if #COLOR.le.-2 => color number: 20 openacc_iccg/01_original C run.sh.e???? ### CM-RCM ### FINAL COLOR NUMBER E-01 sec. (assemble) E E E E E-09 N= E+01 sec. (solver) OpenACC 88

89 OpenACC Makefile n C CC = pgcc OPTFLAGS= -O3 -mp TARGET = run F F90 = pgfortran F90OPTFLAGS= -O3 mp TARGET = run openacc_iccg/02_setup CC = pgcc OPTFLAGS= -O3 -acc -Minfo=accel - ta=tesla:cc60 TARGET = run F90 = pgfortran F90OPTFLAGS= -O3 -acc -Minfo=accel -ta=tesla:cc60 TARGET = run 89

90 OpenACC n #! /bin/sh #PBS -q u-lecture #PBS -l select=1:mpiprocs=1:ompthreads=18 #PBS -W group_list=gt00 #PBS -l walltime=00:05:00 cd $PBS_O_WORKDIR. /etc/profile.d/modules.sh module load pgi/18.7./run -n 128 -c -20 -nt 18 #! /bin/sh #PBS -q h-lecture #PBS -l select=1:mpiprocs=1:ompthreads=1 #PBS -W group_list=gt00 #PBS -l walltime=00:05:00 cd $PBS_O_WORKDIR openacc_iccg/02_setup. /etc/profile.d/modules.sh module load pgi/18.7 numactl./run -n 128 -c -20 -nt 1 90

91 OpenACC C n ü solver_iccg_mc.c openacc_iccg/02_setup extern int 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) main.c if(solve_iccg_mc(iceltot, NL, NU, NPL, NPU, indexl, iteml, indexu, itemu, D, BFORCE, PHI, AL, AU, NCOLORtot, PEsmpTOT, SMPindex, SMPindexG, EPSICCG, &ITR, &IER)) goto error; solver_iccg_mc.h extern int 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); 91

92 OpenACC n OpenACC 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 でのターゲットループ 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では同期を取るためにカーネルを閉じるしかない! 92

93 OpenACC 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; 配列のサイズを一々書くのめんどくさい! 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 配列のサイズ情報を書く必要はない 93

94 Unified Memory n Unified memory ü CPU GPU ucpu/gpu GPU/CPU ü NVIDIA GPU OpenACC n OpenACC ü PGI compiler -ta=tesla,cc60,managed ü Data Unified memory u OpenACC ü CPU 94

95 Unified Memory OpenACC n openacc_iccg/03_unified C CC = pgcc OPTFLAGS= -O3 mp -acc -Minfo=accel -ta=tesla:cc60 TARGET = run F F90 = pgfortran F90OPTFLAGS= -O3 -acc -Minfo=accel -ta=tesla:cc60 TARGET = run CC = pgcc OPTFLAGS= -O3 -acc -Minfo=accel - ta=tesla:cc60,managed TARGET = run F90 = pgfortran F90OPTFLAGS= -O3 -acc -Minfo=accel -ta=tesla:cc60,managed TARGET = run 95

96 Unified Memory OpenACC 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 openacc_iccg/03_unified 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 #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;!$omp parallel do private(i)!$acc kernels!$acc loop independent 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 96

97 Unified Memory run.sh.exxxxx ### CMRCM 標準エラー出力 run.sh.oxxxxx ### CM-RCM ### FINAL COLOR NUMBER e-01 sec. (assemble) e e e e e-09 N= e+00 sec. (solver) ### FINAL COLOR NUMBER E-01 sec. (assemble) E E E E E-09 N= E+00 sec. (solver) n n n ü OpenACC ü Unified memory GPU solver CPU 97

98 present solver_iccg_mc.c solver_iccg_mc.f openacc_iccg/04_data_ present #pragma acc enter data copyin(d[:n],indexl[:n+1]) copyin(al[:npl],iteml[:npl],au[:npu]) copyin(indexu[:n+1],itemu[:npu]) copyin(smpindex[0:ncolortot*pesmptot]) create(x[:n],w[:4][:n]) #pragma omp parallel for private (i) #pragma acc kernels present(x,w) #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 acc exit data delete(d[:n],indexl[:n+1],w[:4][:n]) delete(al[:npl],iteml[:npl],au[:npu]) delete(indexu[:n+1],itemu[:npu]) delete(smpindex[0:ncolortot*pesmptot]) copyout(x[:n]) Makefile から maneged を外すのを忘れずに!$acc enter data!$acc& copyin(b,d,indexl,iteml)!$acc& copyin(indexu,itemu,al,au)!$acc& create(x,w)!$omp parallel do private(i)!$acc kernels present(x,w,b)!$acc loop independent 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 exit data!$acc& delete(b,d,indexl,iteml)!$acc& delete(indexu,itemu,al,au,w)!$acc& copyout(x) Data kernels 98

99 n EX1 Makefile INPUT.dat CPU n EX2 (OMP DO ) OpenACC n n CPU n n n EX3 n EX4 n PGI_ACC_TIME export PGI_ACC_TIME=0 ( ) 99

100 n GPU OpenACC n ICCG OpenACC n ICCG OpenACC n kernels n n PGI_ACC_TIME n 100

101 Q & A n 1 n n 101

OpenACCによる並列化

OpenACCによる並列化実習 OpenACC による ICCG ソルバーの並列化 1 ログイン Reedbush へのログイン $ ssh reedbush.cc.u-tokyo.ac.jp l txxxxx Module のロード $ module load pgi/17.3 cuda ログインするたびに必要です! ワークディレクトリに移動 $ cdw ターゲットプログラム /srcx OpenACC 用のディレクトリの作成