Slide 1

Transcription

1 OPENACC の現状 Akira Naruse NVIDAI Developer Technologies

2 アプリを GPU で加速する方法 Application CUDA OpenACC Library 主要処理を CUDA で記述高い自由度 既存コードにディレクティブを挿入簡単に加速 GPU 対応ライブラリにチェンジ簡単に開始

3 OPENACC CPU GPU Program myscience... serial code...!$acc kernels do k = 1,n1 do i = 1,n2... parallel code... enddo enddo!$acc end kernels... serial code End Program myscience ヒントの追加 簡単 : 既存のコードにコンパイラへのヒントを追加 強力 : そこそこの労力で コンパイラがコードを自動で並列化 オープン : 複数コンパイラベンダが 複数アクセラレータをサポート NVIDIA, AMD, Intel( 予定 ) 既存の C/Fortran コード

4 実行モデル アプリケーション コード $acc parallel GPU CPU 逐次部分は CPU コードを生成 $acc end parallel 計算の重い部分 並列部分は GPU コードを生成

5 SAXPY (Y=A*X+Y, C/C++) OpenMP void saxpy(int n, float a, float *x, float *restrict y) { #pragma omp parallel for for (int i = 0; i < n; ++i) y[i] += a*x[i];... saxpy(n, 3.0, x, y);... void saxpy(int n, float a, float *x, float *restrict y) { #pragma acc parallel copy(y[:n]) copyin(x[:n]) for (int i = 0; i < n; ++i) y[i] += a*x[i];... saxpy(n, 3.0, x, y);... OpenACC

6 SAXPY (Y=A*X+Y, FORTRAN) OpenMP subroutine saxpy(n, a, X, Y) real :: a, X(:), Y(:) integer :: n, i!$omp parallel do do i=1,n Y(i) = a*x(i)+y(i) enddo!$omp end parallel do end subroutine saxpy OpenACC subroutine saxpy(n, a, X, Y) real :: a, Y(:), Y(:) integer :: n, i!$acc parallel copy(y(:)) copyin(x(:)) do i=1,n Y(i) = a*x(i)+y(i) enddo!$acc end parallel end subroutine saxpy... call saxpy(n, 3.0, x, y) call saxpy(n, 3.0, x, y)...

7 OPENMP との併用 OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) { #pragma acc parallel copy(y[:n]) copyin(x[:n]) #pragma omp parallel for for (int i = 0; i < n; ++i) y[i] += a*x[i];... saxpy(n, 3.0, x, y);...

8 簡単にコンパイル OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) $ pgcc -Minfo -acc { saxpy.c saxpy: #pragma acc parallel copy(y[:n]) copyin(x[:n]) 16, Generating #pragma present_or_copy(y[:n]) omp parallel for Generating for present_or_copyin(x[:n]) (int i = 0; i < n; ++i) Generating y[i] Tesla += code a*x[i]; 19, Loop is parallelizable Accelerator kernel generated 19, #pragma... acc loop gang, vector(128) /* blockidx.x threadidx.x */ saxpy(n, 3.0, x, y);...

9 簡単に実行 OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) $ pgcc -Minfo -acc { saxpy.c $ nvprof./a.out saxpy: #pragma acc kernels copy(y[:n]) copyin(x[:n]) ==10302== NVPROF is profiling process 10302, command:./a.out 16, Generating #pragma present_or_copy(y[:n]) omp parallel for ==10302== Profiling application:./a.out Generating for present_or_copyin(x[:n]) (int i = 0; i < n; ++i) ==10302== Profiling result: Generating y[i] Tesla += code a*x[i]; Time(%) Time Calls Avg Min Max Name 19, Loop is parallelizable 62.95% ms ms ms ms [CUDA memcpy HtoD] Accelerator kernel generated 31.48% ms ms ms ms [CUDA memcpy DtoH] 19, #pragma... acc loop gang, vector(128) /* blockidx.x threadidx.x */ 5.56% us us us us saxpy_19_gpu saxpy(n, 3.0, x, y);...

10 簡単に高速 Automotive Financial Life Science Real-Time Object Detection Global Manufacturer of Navigation Systems Valuation of Stock Portfolios using Monte Carlo Global Technology Consulting Company Interaction of Solvents and Biomolecules University of Texas at San Antonio 40 時間で 5 倍 4 時間で 2 倍 8 時間で 5 倍

11 コンパイラとツール 2013 年 10 月 ~ 2013 年 12 月 ~ 2014 年 1 月 ~ 2015 年 ( 予定 ) コンパイラ OpenACC 2.0 対応 デバッグツール

12 SPEC ACCEL 15 本の OpenACC ベンチマーク

13 NCAR-CISL, ORNL / CESM CAM-SE (HOMME) LANL / POP MPAS-O NASA / GEOS-5 NOAA-GFDL / CFSv2 NOAA-GFDL / MOM6 UKMO / HadGEM3 UM NEMO GungHo MPI-M / MPI-ESM ECHAM6 ICON-ATM MPIOM ICON-OCE RIKEN, UniTokyo / NICAM IPSL / DYNAMICO NCAR-M3 / WRF MPAS-A or NIM DWD, MCH / COSMO UniMiami / OLAM 気象 天候 海洋モデル 気候 (C) 天候 (W) 海洋 (O) UKMO / UM GungHo ECMWF / IFS PantaRhei DWD / GME ICON NOAA-NCEP / GFS NIM? EC, CMC / GEM USNRL / NAVGEM NOAA-ESRL / FIM NIM DWD, MPI-M / ICON NOAA-ESRL / NIM NCAR / MPAS-A NCAR-M3 / WRF USNRL / COAMPS DWD, MCH / COSMO MFR / AROME MFR, ICHEC / HARMONIE HIRLAM + ALADIN JAMSTEC-JMA / ASUCA CAS-CMA / GRAPES UniMiami / OLAM MPAS-A or NIM LANL / POP MPAS-O NOAA-GFDL / MOM6 CNRS, STFC/ NEMO USNRL / HYCOM MIT / MITgcm LANL / MPAS-O MPI-M / ICON-OCE Rutgers-UCLA / ROMS UNC-ND / ADCIRC GPU Development (8) CAM-SE, GEOS-5, NEMO, WRF, COSMO, NIM, FIM, GRAPES GPU Evaluation (15) POP, ICON, NICAM, OLAM, GungHo, PantaRhei, ASUCA, HARMONIE, COAMPS, HYCOM, MITgcm, ROMS, ADCIRC, DYNAMICO, MOM6 GPU Not Started (7) MPAS-A, MPAS-O, GFS, GEM, NAVGEM, AROME, ICON-OCE Indicates Next-Gen Model

14 OPENACC への移行 Model Focus GPU Approach Collaboration NCAR / WRF NWP/Climate-R (1) OpenACC, (2) CUDA (1) NCAR-MMM, (2) SSEC UW-M DWD / COSMO NWP/Climate-R CUDA+OpenACC CSCS, MeteoSwiss (MCH) ORNL / CAM-SE Climate-G CUDA-F OpenACC ORNL, Cray NCAR / CAM-SE Climate- G CUDA,CUDA-F,OpenACC NCAR-CISL NOAA / NIM&FIM NWP/Climate-G F2C-ACC,OpenACC NOAA-ESRL, PGI NASA / GEOS-5 Climate-G CUDA-F OpenACC NASA, PGI CNRS / NEMO Ocean GCM OpenACC STFC UKMO / GungHo NWP/Climate-G OpenACC STFC, UKMO in future? USNRL / HYCOM Ocean GCM OpenACC US Naval Research Lab RIKEN / NICAM Climate-G OpenACC RIKEN, UniTokyo UNC / ADCIRC Storm Surge OpenACC (AmgX?) LSU LONI NOAA / MOM6 Ocean GCM OpenACC NOAA-GFDL NASA / FV-Core Atmospheric GCM OpenACC NASA, NOAA-GFDL Other Evaluations: US COAMPS, MPAS, ROMS, OLAM; Europe ICON, IFS, HARMONIE; DYNAMICO Asia-Pacific ASUCA (JP), GRAPES (CN)

15 OPENACC でどこまで出来るの?

16 例 : JACOBI ITERATION while ( error > tol ) { error = 0.0; for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i])); A(i-1,j) A(i,j+1) A(i,j) A(i+1,j) A(i,j-1) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i];

17 並列領域 (KERNELS CONSTRUCT) while ( error > tol ) { error = 0.0; #pragma acc kernels for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; Parallels と Kernels 並列領域を指示 Parallels 並列実行スタート Kernels 複数のカーネル

18 並列領域 (KERNELS CONSTRUCT) while ( error > tol ) { error = 0.0; #pragma acc kernels for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); $ pgcc -Minfo=acc -acc jacobi.c jacobi: #pragma acc kernels for (int j = 1; j < N-1; j++) { 60, Loop carried scalar dependence for 'error' at line 64 for (int i = 1; i < M-1; i++) { A[j][i]... = Anew[j][i]; Accelerator scalar kernel generated 61, Loop carried scalar dependence for 'error' at line Accelerator scalar kernel generated Parallels と Kernels 並列領域を指示 Parallels 並列走行の開始 Kernels 複数の GPU カーネル

19 リダクション (REDUCTION CLAUSE) while ( error > tol ) { error = 0.0; #pragma acc kernels #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; 演算の種類 + 和 * 積 Max 最大 Min 最小 ビット和 & ビット積 ^ XOR 論理和 && 論理積

20 リダクション (REDUCTION CLAUSE) while ( error > tol ) { error = 0.0; 演算の種類 + 和 #pragma acc kernels #pragma acc loop reduction(max:error) * 積 for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) Max 最大 $ pgcc for -Minfo=acc (int i = 1; -acc i < M-1; jacobi.c i++) { jacobi: Anew[j][i] = (A[j][i+1] + A[j][i-1] + Min 最小 A[j-1][i] + A[j+1][i]) * 0.25; 59, Generating present_or_copyout(anew[1:4094][1:4094]) error = max(error, abs(anew[j][i] - A[j][i]); ビット和 Generating present_or_copyin(a[:][:]) Generating Tesla code & ビット積 61, Loop is parallelizable ^ XOR #pragma acc kernels 63, Loop is parallelizable for (int j = 1; j < N-1; j++) { 論理和 for (int Accelerator i = 1; i kernel < M-1; i++) generated { A[j][i] 61, #pragma = Anew[j][i]; acc loop gang /* blockidx.y */ && 論理積 63, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ Max reduction generated for error

21 データ転送方法 (DATA CLAUSE) while ( error > tol ) { error = 0.0; #pragma acc kernels #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) $ pgcc for -Minfo=acc (int i = 1; -acc i < M-1; jacobi.c i++) { jacobi: Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; 59, Generating present_or_copyout(anew[1:4094][1:4094]) error = max(error, abs(anew[j][i] - A[j][i]); Generating present_or_copyin(a[:][:]) Generating Tesla code 61, Loop is parallelizable #pragma acc kernels 63, Loop is parallelizable for (int j = 1; j < N-1; j++) { for (int Accelerator i = 1; i kernel < M-1; i++) generated { A[j][i] 61, #pragma = Anew[j][i]; acc loop gang /* blockidx.y */ 63, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ Max reduction generated for error

22 データ転送方法 (DATA CLAUSE) while ( error > tol ) { error = 0.0; #pragma acc kernels \ pcopyout(anew[1:n-2][1:m-2]) pcopyin(a[0:n][0:m]) #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels \ pcopyout(a[1:n-2][1:m-2]) pcopyin(anew[1:n-2][1:m-2]) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; copyin (Host GPU) copyout (Host GPU) copy create present pcopyin pcopyout pcopy pcreate

23 データ転送方法 (DATA CLAUSE) while ( error > tol ) { error = 0.0; #pragma acc kernels \ pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels \ pcopy(a[:][:]) pcopyin(anew[:][:]) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; copyin (Host GPU) copyout (Host GPU) copy create present pcopyin pcopyout pcopy pcreate

24 データ転送がボトルネック (NVVP) 1 cycle 稼働率 : 低い GPU kernel GPU kernel

25 Host while ( error > tol ) { error = 0.0; 過剰なデータ転送 GPU #pragma acc kernels \ pcopy(anew[:][:]) \ pcopyin(a[:][:]) { copyin copyout #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels \ pcopy(a[:][:]) \ pcopyin(anew[:][:]) { copyin copyout for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i];

26 データ領域 (DATA CONSTRUCT) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels pcopy(a[:][:]) pcopyin(anew[:][:]) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; copyin (CPU GPU) copyout (CPU GPU) copy create present pcopyin pcopyout pcopy pcreate

27 Host #pragma acc data \ pcopy(a) create(anew) while ( error > tol ) { error = 0.0; 適正なデータ転送 copyin GPU #pragma acc kernels \ pcopy(anew[:][:]) \ pcopyin(a[:][:]) { #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]); #pragma acc kernels \ pcopy(a[:][:]) \ pcopyin(anew[:][:]) { copyout for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i];

28 データ転送の削減 (NVVP) 1 cycle 稼働率 : 高い

29 2 つの処理 CPU Memory データ転送 GPU Memory PCI 計算オフロード 計算オフロード データ転送 両方を考慮する必要がある

30 カーネルチューニング

31 カーネルチューニング (LOOP CONSTRUCT) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]);...

32 カーネルチューニング (LOOP CONSTRUCT) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; Gang Worker Vector SIMD 幅 #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) for (int j = 1; j < N-1; j++) { $ pgcc -Minfo=acc -acc jacobi.c Independent #pragma acc loop reduction(max:error) jacobi: for (int i = 1; i < M-1; i++) { Collapse 59, Anew[j][i] Generating = (A[j][i+1] present_or_copyout(anew[1:4094][1:4094]) + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; Seq Generating present_or_copyin(a[:][:]) error = max(error, abs(anew[j][i] - A[j][i]); Generating Tesla code... 61, Loop is parallelizable... 63, Loop is parallelizable Accelerator kernel generated 61, #pragma acc loop gang /* blockidx.y */ 63, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ Max reduction generated for error

33 カーネルチューニング (LOOP CONSTRUCT) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) gang vector(1) for (int j = 1; j < N-1; j++) { #pragma acc loop reduction(max:error) gang vector(128) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]);... Gang Worker Vector SIMD 幅 Collapse Independent Seq Cache Tile

34 実行条件設定 (VECTOR CLAUSE) #pragma acc loop gang vector(4) for (j = 0; j < 16; j++) { #pragma accloop gang vector(16) for (i = 0; i < 16; i++) {... #pragma acc loop gang vector(8) for (j = 1; j < 16; j++) { #pragma accloop gang vector(8) for (i = 0; i < 16; i++) {... i i 4 x 16 4 x 16 8 x 8 8 x 8 j 4 x 16 j 8 x 8 8 x 8

35 カーネルチューニング (LOOP CONSTRUCT) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) \ collapse(2) gang vector(128) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]);... Gang Worker Vector SIMD 幅 Collapse Independent Seq Cache Tile...

36 カーネルチューニング (LOOP CONSTRUCT) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop reduction(max:error) independent for (int jj = 1; jj < NN-1; jj++) { int j = list_j[jj]; #pragma acc loop reduction(max:error) for (int i = 1; i < M-1; i++) { Anew[j][i] = (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]) * 0.25; error = max(error, abs(anew[j][i] - A[j][i]);... Gang Worker Vector SIMD 幅 Collapse Independent Seq Cache Tile...

37 カーネルチューニング (LOOP CONSTRUCT) #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop seq for (int k = 3; k < NK-3; k++) { #pragma acc loop for (int j = 0; j < NJ; j++) { #pragma acc loop for (int i = 0; i < NI; i++) { Anew[k][j][i] = func( A[k-1][j][i], A[k-2][j][i], A[k-3][j][i], A[k+1][j][i], A[k+2][j][i], A[k+3][j][i],... ); Gang Worker Vector SIMD 幅 Collapse Independent Seq Cache Tile...

38 MPI とは簡単に併用できるの?

39 MPI 並列 (HALO EXCHANGE) A(i,j+1) A(i-1,j) A(i,j) A(i+1,j) A(i,j-1) ブロック分割 各プロセスは 1 ブロック担当 境界部 (halo) のデータ交換

40 MPI JACOBI ITERATION #pragma acc data pcopy(a) create(anew) while ( error > tol ) { #pragma acc kernels pcopy(anew) pcopyin(a) calc_new_a( Anew, A,... ); #pragma acc kernels pcopy(a) pcopyin(anew) update_a( A, Anew );

41 MPI JACOBI ITERATION #pragma acc data pcopy(a) create(anew) while ( error > tol ) { pack_data_at_boundary( send_buf, A,... ); exchange_data_by_mpi( recv_buf, send_buf,... ); 1. 送信データの梱包 GPU MPI 2. データの交換 unpack_data_to_halo( A, recv_buf,... ); #pragma acc kernels pcopy(anew) pcopyin(a) calc_new_a( Anew, A,... ); #pragma acc kernels pcopy(a) pcopyin(anew) update_a( A, Anew ); 3. 受信データの開梱 GPU

42 MPI JACOBI ITERATION #pragma acc data pcopy(a) create(anew) while ( error > tol ) { #pragma acc kernels pcopyin(a) pcopyout(send_buf) pack_data_at_boundary( send_buf, A,... ); exchange_data_by_mpi( recv_buf, send_buf,... ); #pragma acc kernels pcopy(a) pcopyin(recv_buf) unpack_data_to_halo( A, recv_buf,... ); 1. GPU 上でデータを送信バッファに梱包し Host に転送 2. 隣接プロセスとデータ交換 GPU MPI #pragma acc kernels pcopy(anew) pcopyin(a) calc_new_a( Anew, A,... ); #pragma acc kernels pcopy(a) pcopyin(anew) update_a( A, Anew ); 3. GPU に転送 GPU 上で受信バッファのデータを開梱 GPU

43 MPI JACOBI ITERATION (NVVP) 1 cycle Pack MPI Upck データ梱包 MPI データ開梱

44 オーバーラップ (ASYNC/WAIT CLAUSE) while ( error > tol ) { #pragma acc kernels pcopyin(a) pcopyout(send_buf) pack_data_at_boundary( send_buf, A,... ); exchange_data_by_mpi( recv_buf, send_buf,... ); #pragma acc kernels pcopy(a) pcopyin(recv_buf) unpack_data_to_halo( A, recv_buf,... ); #pragma acc kernels pcopy(anew) pcopyin(a) calc_new_a( Anew, A,... ); #pragma acc kernels pcopy(a) pcopyin(anew) update_a( A, Anew );

45 オーバーラップ (ASYNC/WAIT CLAUSE) while ( error > tol ) { #pragma acc kernels pcopyin(a) pcopyout(send_buf) pack_data_at_boundary( send_buf, A,... ); #pragma acc kernels pcopy(anew) pcopyin(a) calc_new_a_inside( Anew, A,... ); exchange_data_by_mpi( recv_buf, send_buf,... ); #pragma acc kernels pcopy(a) pcopyin(recv_buf) unpack_data_to_halo( A, recv_buf,... ); #pragma acc kernels pcopy(anew) pcopyin(a) calc_new_a_at_boundary( Anew, A,... ); 内部 境界部 #pragma acc kernels pcopy(a) pcopyin(anew) update_a( A, Anew );

46 オーバーラップ (ASYNC/WAIT CLAUSE) while ( error > tol ) { #pragma acc kernels pcopyin(a) pcopyout(send_buf) async(2) pack_data_at_boundary( send_buf, A,... ); #pragma acc kernels pcopy(anew) pcopyin(a) async(1) calc_new_a_inside( Anew, A,... ); #pragma acc wait(2) exchange_data_by_mpi( recv_buf, send_buf,... ); #pragma acc kernels pcopy(a) pcopyin(recv_buf) async(2) unpack_data_to_halo( A, recv_buf,... ); #pragma acc kernels pcopy(anew) pcopyin(a) async(2) calc_new_a_at_boundary( Anew, A,... ); #pragma acc kernels pcopy(a) pcopyin(anew) wait(1,2) update_a( A, Anew );

47 オーバーラップ (NVVP) 1 cycle Pack MPI Upck

48 OPENACC って 実際に使われているの?

49 NICAM 気象 気候モデル by 理研 AICS/ 東大 膨大なコード ( 数十万行 ) ホットスポットがない ( パレートの法則 ) 特性の異なる 2 種類の処理 力学系 メモリバンド幅ネック 物理系 演算ネック

50 Performance (GFLOPS) NICAM: 力学系 (NICAM-DC) OpenACC による GPU 化 主要サブルーチンは 全て GPU 上で動作 (50 以上 ) MPI 対応済み 1.E+05 1.E+04 Tsubame 2.5 (GPU:K20X) K computer Tsubame 2.5 (CPU:WSM) (*) weak scaling 2 週間 1.E+03 良好なスケーラビリティ Tsubame 2.5, 最大 2560 GPUs 1.E+02 Scaling factor: E+01 1.E+00 1.E+01 1.E+02 1.E+03 1.E+04 Number of CPUs or GPUs Weak scaling

51 Measured Performance (GFLOPS) NICAM: 力学系 (NICAM-DC) 1.E+05 1.E+04 Tsubame 2.5 (GPU:K20X) K computer Tsubame 2.5 (CPU:WSM) 1.E+03 1.E+02 1.E+01 1.E+02 1.E+03 1.E+04 1.E+05 1.E+06 Aggregate Peak Memory Bandwidth (GB/s)

52 Speedup vs. CPU 1-core NICAM: 物理系 (SCALE-LES) Atmospheric radiation transfer 物理系の中で 最も重い計算 OpenACCによるGPU 対応 完了 core 2 core 4 core 10 core 1 GPU 2 GPUs 4 GPUs Xeon E5-2690v2(3.0GHz,10-core) Tesla K40 (*) PCI データ転送時間込み, グリッドサイズ :1256x32x32

53 Time (sec) SEISM3D 地震シミュレーション by 東大地震研 ( 古村教授 ) 主要サブルーチンのGPU 対応が完了 メモリバンド幅ネック 3 次元モデル (2 次元分割 ) 隣接プロセス間通信 SEISM3D (480x480x1024, 1K steps) GPU の実行時間内訳 x speedup ( アプリ全体 ) Others (CPU, MPI and so on) [CUDA memcpy DtoH] [CUDA memcpy HtoD] (other subroutines) update_vel_pml update_vel K: 8x SPARC64 VIIIfx CPU: 8x Xeon E5-2690v2 134 GPU: 8x Tesla K GPU: 8x Tesla K40 update_stress_pml update_stress diff3d_*

54 性能 (M grids/sec) SEISM3D 10,000 1,000 Tesla K40 SX9 FX10 K Xeon E5-2* v2 (IVB) Xeon E5-4* (SDB) Xeon X7* (NHL EX) ,000 10,000 トータルピークメモリバンド幅 (GB/s)

55 Speedup vs. 1 CPU core FFR/BCM 次世代 CFD by 理研 AICS/ 北大 ( 坪倉准教授 ) MUSCL_bench: MUSCL スキームに基づく Flux 計算 ( とても複雑な計算 ) CFD 計算の主要部分 (60-70%) OpenACC による GPU 対応 完了 core 2 core 5 core 10 core 1 GPU Xeon E5-2690v2(3.0GHz,10-core) Tesla K40 (*) PCI データ転送時間込み サイズ :80x32x32x32

56 まとめ OpenACC の現状を紹介 簡単 : 既存コードへのディレクティブ追加 強力 : 少ない労力で GPU 利用可能 オープン : 採用事例の増加

57 CUDA 6 の強化ポイント Akira Naruse NVIDAI Developer Technologies

58 CUDA 6 ユニファイド メモリ XT ライブラリドロップイン ライブラリ GPUDirect RDMA 開発ツール

59 ユニファイドメモリ Now 開発者から見えるメモリモデル ユニファイドメモリ ホストメモリ GPU メモリ

60 煩雑なメモリマネジメント void sortfile(file *fp, int N) { char *data; data = (char *)malloc(n); fread(data, 1, N, fp); qsort(data, N, 1, compare); use_data(data); free(data); CPU コード GPU コード void sortfile(file *fp, int N) { char *data char *d_data; data = (char *)malloc(n); cudamalloc(&d_data, N); fread(data, 1, N, fp); cudamemcpy(d_data, data, N,..); qsort<<<...>>>(d_data,n,1,compare); cudadevicesynchronize(); cudamemcpy(data, d_data, N,..); use_data(data); cudafree(d_data); free(data);

61 メモリマネジメントを簡素化 void sortfile(file *fp, int N) { char *data; data = (char *)malloc(n); fread(data, 1, N, fp); qsort(data, N, 1, compare); use_data(data); CPU コード ユニファイドメモリ (CUDA6) void sortfile(file *fp, int N) { char *data cudamallocmanaged(&d_data, N); fread(data, 1, N, fp); qsort<<<...>>>(d_data,n,1,compare); cudadevicesynchronize(); use_data(data); free(data); cudafree(data);

62 メモリマネジメントの統合 ( 将来 ) CPU コード将来? void sortfile(file *fp, int N) { char *data; data = (char *)malloc(n); fread(data, 1, N, fp); qsort(data, N, 1, compare); use_data(data); void sortfile(file *fp, int N) { char *data data = (char *)malloc(n); fread(data, 1, N, fp); qsort<<<...>>>(d_data,n,1,compare); cudadevicesynchronize(); use_data(data); free(data); free(data);

63 DEEP COPY struct dataelem { int prop1; int prop2; char *text; ; CPU Memory dataelem prop1 prop2 Hello World *text GPU Memory

64 DEEP COPY コピーが 2 回必要 struct dataelem { int prop1; int prop2; char *text; ; CPU Memory dataelem prop1 prop2 Hello World *text dataelem prop1 prop2 Hello World *text GPU Memory

65 DEEP COPY void launch(dataelem *elem) { dataelem *g_elem; char *g_text; int textlen = strlen(elem->text); cudamalloc(&g_elem, sizeof(dataelem)); cudamalloc(&g_text, textlen); cudamemcpy(g_elem, elem, sizeof(dataelem)); cudamemcpy(g_text, elem->text, textlen); cudamemcpy(&(g_elem->text), &g_text, sizeof(g_text)); dataelem prop1 prop2 *text CPU Memory 実際は 3 回必要 Hello World kernel<<<... >>>(g_elem); dataelem prop1 prop2 Hello World *text GPU Memory

66 DEEP COPY ( ユニファイドメモリ ) void launch(dataelem *elem) { kernel<<<... >>>(elem); CPU Memory Unified Memory dataelem prop1 prop2 *text Hello World GPU Memory

67 連結リスト CPU Memory key key key key data data data data next next next next GPU Memory

68 連結リスト 毎回 全部転送 PCI のバンド幅ネック CPU Memory 最初は全部転送 以降は更新箇所だけ転送 key data next key data next key data next key data next とても複雑な処理 CPU メモリにデータを配置 GPU は PCI 経由のアクセス 全部を転送? PCI 経由 遅い GPU Memory

69 連結リスト ( ユニファイドメモリ ) CPU からも GPU からもリスト操作が可能 挿入 削除 リスト更新後に CPU メモリと GPU メモリ間の明示的な同期は不要 CPU Memory 通常のメモリアクセス Unified Memory key key key key CPU と GPU から同時アクセスは NG 排他制御必要 data next data next data next data next 通常のメモリアクセス GPU Memory

70 ロードマップ CUDA 6: 簡単に利用 Next: 最適化 単一のポインタ Memcpy 記述不要 ホスト側プログラムとデータ構造を共有 プリフェッチ データ移動ヒント OS サポートの追加 Pascal システムアロケータの統合 スタックメモリの統合 メモリコヒーレンシを HW でアクセラレート

71 XT ライブラリ

72 cublas-xt and cufft-xt XT ライブラリ 明示的なデータ転送の指示は不要 必要な GPU メモリはライブラリが確保 マルチ GPU に自動対応 マルチ GPU 向けのコード記述は不要 GPU メモリ容量を超えるサイズに対応 (out-of-core) カーネル実行とデータ転送をオーバーラップ (BLAS level 3)

73 CUBLAS 行列積コード cublas cublashandle_t handle; cublascreate(&handle); cudamalloc(&d_a,..); cudamalloc(&d_b,..); cudamalloc(&d_c,..); cudasetmatrix(.., d_a,.., A,..); cudasetmatrix(.., d_b,.., B,..); cublasdgemm(handle,.., d_a,.., d_b,.., d_c,..); cudagetmatrix(.., d_c,.., C,..); cudafree(d_a); cudafree(d_b); cudafree(d_c); cublasdestroy(handle);

74 CUBLAS CUBLAS-XT cublas 行列積コード cublas-xt cublashandle_t handle; cublascreate(&handle); cudamalloc(&d_a,..); cudamalloc(&d_b,..); cudamalloc(&d_c,..); cudasetmatrix(.., d_a,.., A,..); cudasetmatrix(.., d_b,.., B,..); cublasdgemm(handle,.., d_a,.., d_b,.., d_c,..); cudagetmatrix(.., d_c,.., C,..); cudafree(d_a); cudafree(d_b); cudafree(d_c); cublasdestroy(handle); cublasxthandle_t handle; cublasxtcreate(&handle); cublasxtdgemm(handle,.., A,.., B,.., C,..); cublasxtdestroy(handle);

75 CUBLAS-XT API 使用 GPU cublasxtdeviceselect() GPU 数 使用 GPU IDs ブロッキングサイズ cublasxtsetblockdim() ブロッキングサイズの設定 cublasxtgetbloskdim() ( 現設定の取得 ) CPU GPU ハイブリッド実行 cublasxtsetcpuroutine() CPU 版 BLAS の設定 cublasxtsetcpuratio() CPU 比率の設定 Pinned メモリ cublasxtsetpinningmemmode() Pinned メモリの設定 cublasxtgetpinningmemmode() ( 現設定の取得 )

76 GFLOPS CUBLAS-XT 全てのBLAS level 3 ルーチンをサポート 行列サイズがGPUメモリ容量超でもOK (out-of-core) cublas ZGEMM Performance on 2 GPUs K20c 2 K20c In-core Out-of-core Matrix Size (NxN)

77 CUBLAS-XT (NVVP)

78 ドロップイン ライブラリ

79 ドロップイン ライブラリ 標準ライブラリ API での GPU 利用を可能に NVBLAS BLAS level 3 関数呼び出しを 自動的に cublas に置き換え cublas 利用のためのソース変更は不要 CPU dgemm(.., A,.., B,.., C,..); 使い方 NVBLAS を入れて再コンパイル Linux は LD_PRELOAD 設定で使用可能 ( 最コンパイル不要 )

80 ドロップイン ライブラリ 標準ライブラリ API での GPU 利用を可能に NVBLAS BLAS level 3 関数呼び出しを 自動的に cublas に置き換え cublas 利用のためのソース変更は不要 CPU NVBLAS dgemm(.., A,.., B,.., C,..); dgemm(.., A,.., B,.., C,..); 使い方 NVBLAS を入れて再コンパイル Linux は LD_PRELOAD 設定で使用可能 ( 最コンパイル不要 )

81 NVBLAS_LOGFILE NVBLAS (LINUX) nvblas.log 設定ファイル (nvblas.conf) NVBLAS_CPU_BLAS_LIB libmkl_intel_lp64.so \ libmkl_core.so \ libmkl_intel_thread.so NVBLAS_GPU_LIST 0 # ALL, ALL0 NVBLAS_TILE_DIM 2048 NVBLAS_AUTOPIN_MEM_ENABLED $ LD_PRELOAD=/usr/local/cuda-6.0/lib64/libnvblas.so./a.out

82 fp64 GFlops/s BLAS level 3 使用のアプリに適用可能 Octave, Scilab, など NVBLAS R 言語での行列乗算 nvblas, 4x K20X GPUs MKL, 6-core Xeon E CPU matrix dimension

83 NVBLAS デモ

84 CUDA 6 ユニファイド メモリ XT ライブラリドロップイン ライブラリ GPUDirect RDMA 開発ツール

85 CUDA 6 並列コンピューティングを簡単に CUDA Registered Developer Program developer.nvidia.com/cuda-toolkit

86 CUDA 6.5 RC 64-bit ARM マシン Microsoft Visual Studio 2013 (VC12) cufft callbacks cusparse (BSR 格納形式 ) CUDA 占有率計算 API CUDA FORTRAN デバッグ機能 アプリケーションリプレイモード (Visual Profile and nvprof) Nvprune ユーティリティ (object サイズ削減 )