OpenACC CUDA による GPU コンピューティング Akira Naruse, 19 th Jul. 2018
成瀬彰 (Naruse, Akira) 自己紹介 2013 年 ~: NVIDIA シニア デベローパーテクノロジー エンジニア 1996~2013 年 : 富士通研究所 研究員など 専門 興味 : 並列処理 性能最適化 スパコン HPC GPU コンピューティング DeepLearning 詳しくは github.com/anaruse 2
AGENDA GPU Computing Volta GPUs OpenACC CUDA 3
GPU Computing 4
GPU コンピューティング Low latency + High throughput CPU GPU 5
アプリケーション実行 アプリケーション コード CPU Do i=1,n 計算の重い部分 GPU 逐次部分は CPU 上で実行 End do 並列部分は GPU で実行 6
GPU の構造 (TESLA P100) 大量の CUDA コア並列性が鍵 64 CUDA core/sm 3584 CUDA core/chip 56 SM/chip 7
GPU APPLICATIONS 数百のアプリケーションが GPU に対応 www.nvidia.com/object/gpu-applications.html 8
LEADING APPLICATIONS 9
DL FRAMEWORKS 10
アプリを GPU 対応する方法 Application Library OpenACC CUDA GPU 対応ライブラリにチェンジ簡単に開始 既存コードにディレクティブを挿入簡単に加速 主要処理を CUDA で記述高い自由度 11
NVIDIA LIBRARIES cublas cusparse cusolver AmgX cufft curand nvgraph Thrust Performance Primitives cudnn TensorRT NCCL 12
PARTNER LIBRARIES Computer Vision Audio and Video Matrix, Signal and Image Linear Algebra Math, Signal and Image Graph Sparse direct solvers Linear Algebra Linear Algebra Computational Geometry Sparse Iterative Methods Real-time visual simulation 13
アプリを GPU 対応する方法 Application Library OpenACC CUDA GPU 対応ライブラリにチェンジ簡単に開始 既存コードにディレクティブを挿入簡単に加速 主要処理を CUDA で記述高い自由度 14
SAXPY (Y=A*X+Y) 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 15
アプリを GPU 対応する方法 Application Library OpenACC CUDA GPU 対応ライブラリにチェンジ簡単に開始 既存コードにディレクティブを挿入簡単に加速 主要処理を CUDA で記述高い自由度 16
CPU SAXPY (Y=A*X+Y) CUDA void saxpy(int n, float a, float *x, float *y) { for (int i = 0; i < n; ++i) y[i] += a*x[i]; global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx; if (i < n) y[i] += a*x[i]; saxpy(n, 3.0, x, y); size_t size = sizeof(float) * N; cudamemcpy(d_x, x, size, cudamemcpyhosttodevice); cudamemcpy(d_y, y, size, cudamemcpyhosttodevice); saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y); cudamemcpy(y, d_y, size, cudamemcpydevicetohost); 17
Volta GPUs 18
Performance / W GPU ロードマップ Volta Pascal Maxwell Tesla Fermi Kepler 2008 2010 2012 2014 2016 2018 19
TESLA V100 の概要 Volta Architecture Improved NVLink & HBM2 Volta MPS Improved SIMT Model Tensor Core Most Productive GPU Efficient Bandwidth Inference Utilization New Algorithms 125 Programmable TFLOPS Deep Learning Deep Learning と HPC 両方に最適な GPU 20
P100 に対する相対性能 VOLTA HPC 性能を大きく向上 HPC アプリケーション性能 System Config Info: 2X Xeon E5-2690 v4, 2.6GHz, w/ 1X Tesla P100 or V100. V100 measured on pre-production hardware. Summit Supercomputer 200+ PetaFlops ~3,400 Nodes 10 Megawatts 21
VOLTA 米国トップスパコンのエンジン SUMMIT AI Exascale Today Performance Leadership 200 PF Accelerated Science ACME DIRAC FLASH GTC HACC LSDALTON NAMD 20 PF NUCCOR NWCHEM QMCPACK RAPTOR SPECFEM XGC 3+EFLOPS Tensor Ops 10X Perf Over Titan 5-10X Application Perf Over Titan 22
TESLA V100 トランジスタ数 :21B 815 mm 2 80 SM 5120 CUDA コア 640 Tensor コア HBM2 32 GB, 900 GB/s NVLink 300 GB/s *full GV100 chip contains 84 SMs 23
GPU ピーク性能比較 : P100 vs v100 P100 V100 性能 UP トレーニング性能 10 TOPS 125 TOPS 12x インファレンス性能 21 TFLOPS 125 TOPS 6x FP64/FP32 5/10 TFLOPS 7.8/15.6 TFLOPS 1.5x HBM2 バンド幅 720 GB/s 900 GB/s 1.2x NVLink バンド幅 160 GB/s 300 GB/s 1.9x L2 キャッシュ 4 MB 6 MB 1.5x L1 キャッシュ 1.3 MB 10 MB 7.7x 24
STREAM: Triad- Delivered GB/s HBM2 メモリ 使用効率 UP 実効バンド幅 1.5 倍 V100 measured on pre-production hardware. HBM2 stack P100 V100 76% 95% 使用効率 25
VOLTA NVLINK P100 V100 リンク数 4 6 バンド幅 / リンク トータルバンド幅 (*) バンド幅は双方向 40 GB/s 50 GB/s 160 GB/s 300 GB/s DGX1V 26
VOLTA GV100 SM GV100 FP32 ユニット 64 FP64 ユニット 32 INT32 ユニット 64 Tensor コア 8 レジスタファイル 統合 L1 共有メモリ 256 KB 128 KB Active スレッド 2048 (*) SM あたり 27
VOLTA GV100 SM 生産性の向上 命令セットを一新スケジューラを2 倍命令発行機構をシンプルに L1キャッシュの大容量 高速化 SIMTモデルの改善テンソル計算の加速 最もプログラミングの簡単な SM 28
OpenACC 29
OPENACC プログラミング 概要紹介 プログラムの OpenACC 化 OpenACC 化事例 30
CPU GPU OPENACC 簡単 : 既存のコードにコンパイラへのヒントを追加 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 既存の C/Fortran コード ヒントの追加 強力 : 相応の労力で コンパイラがコードを自動で並列化 オープン : 複数コンパイラベンダが 様々なプロセッサをサポート NVIDIA GPU, AMD GPU, x86 CPU, 31
GPU コンピューティング アプリケーション コード CPU Do i=1,n 計算の重い部分 OpenACC GPU 逐次部分は CPU 上で実行 End do 並列部分は GPU で実行 32
CPU void saxpy(int n, float a, float *x, float *y) { for (int i = 0; i < n; ++i) y[i] += a*x[i]; saxpy(n, 3.0, x, y); SAXPY (Y=A*X+Y) CUDA global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx; if (i < n) y[i] += a*x[i]; size_t size = sizeof(float) * N; cudamemcpy(d_x, x, size, cudamemcpyhosttodevice); cudamemcpy(d_y, y, size, cudamemcpyhosttodevice); saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y); cudamemcpy(y, d_y, size, cudamemcpydevicetohost); 33
SAXPY (y=a*x+y) 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 omp acc データの移動 34
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) FORTRAN も同様 35
簡単にコンパイル OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) $ pgcc -acc { 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); 36
簡単に実行 OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) $ pgcc -Minfo { -acc saxpy.c $ nvprof./a.out saxpy: #pragma acc parallel copy(y[:n]) copyin(x[:n]) ==10302== NVPROF #pragma is profiling omp parallel process for 10302, command:./a.out 16, Generating present_or_copy(y[:n]) ==10302== Profiling for (int application: Generating present_or_copyin(x[:n]) i = 0; i./a.out < n; ++i) ==10302== Profiling Generating y[i] result: Tesla += code a*x[i]; Time(%) Time 19, Loop is Calls Avg Min Max Name parallelizable 62.95% 3.0358ms 2 1.5179ms 1.5172ms 1.5186ms [CUDA memcpy HtoD] Accelerator kernel generated 31.48% 1.5181ms 19, #pragma 1 1.5181ms 1.5181ms 1.5181ms [CUDA memcpy DtoH] acc loop gang, vector(128) /* blockidx.x threadidx.x */ 5.56% 268.31us saxpy(n, 3.0, 1 x, 268.31us y); 268.31us 268.31us saxpy_19_gpu 37
OPENACC プログラミング 概要紹介 プログラムの OpenACC 化 OpenACC 化事例 38
事例 : ヤコビ反復法 while ( error > tol ) { error = 0.0; A(i,j+1) 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) A(i,j-1) A(i+1,j) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; 39
並列領域の指定 (parallel/kernels ディレクティブ ) while ( error > tol ) { error = 0.0; Parallel と Kernels #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]; Parallel OpenMP と親和性 開発者主体 Kernels 複数 kernel の生成 コンパイラ主体 40
[PGI tips] コンパイラメッセージ $ pgcc acc Minfo=accel jacobi.c jacobi: 44, Generating copyout(anew[1:4094][1:4094]) Generating copyin(a[:][:]) Generating Tesla code 45, Loop is parallelizable 46, Loop is parallelizable Accelerator kernel generated 45, #pragma acc loop gang /* blockidx.y */ 46, #pragma acc loop gang, vector(128) /* blockidx.x threadidx.x */ 49, Max reduction generated for error 41
並列領域の指定 (kernels) 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; $ pgcc error -Minfo=acc = max(error, abs(anew[j][i] -acc jacobi.c - A[j][i]); jacobi: 並列領域の指定 Parallels と Kernels Parallels OpenMP と親和性 59, Generating present_or_copyout(anew[1:4094][1:4094]) #pragma acc Generating kernels present_or_copyin(a[:][:]) 開発者主体 for (int j Generating = 1; j < N-1; Tesla j++) code { for (int i = 1; i < M-1; i++) { Kernels 61, Loop is parallelizable A[j][i] = Anew[j][i]; 63, Loop is parallelizable 複数 kernelの生成 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 42
データの転送 (data clause) 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; $ pgcc error -Minfo=acc = max(error, abs(anew[j][i] -acc jacobi.c - A[j][i]); jacobi: 並列領域の指定 Parallels と Kernels Parallels OpenMP と親和性 59, Generating present_or_copyout(anew[1:4094][1:4094]) #pragma acc Generating kernels present_or_copyin(a[:][:]) 開発者主体 for (int j Generating = 1; j < N-1; Tesla j++) code { for (int i = 1; i < M-1; i++) { Kernels 61, Loop is parallelizable A[j][i] = Anew[j][i]; 63, Loop is parallelizable 複数 kernelの生成 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 43
データの転送 (data clause) while ( error > tol ) { error = 0.0; #pragma acc kernels \ pcopyout(anew[1:4094][1:4094]) pcopyin(a[:][:]) 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 \ pcopyout(a[1:4094][1:4094]) pcopyin(anew[1:4094][1:4094]) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; N=M=4096 copyin (Host GPU) copyout (Host GPU) copy create present 44
配列形状の指定 配列は 全要素でなく 一部だけ指定して転送することも可能 注意 : C/C++ と Fortran では指定方法が異なる C/C++: array[ start : size ] float Anew[4096][4096] pcopyout( Anew[1:4094][1:4094]) pcopyin( A[:][:]) ) Fortran: array( start : end ) real Anew(4096,4096) pcopyout( Anew(2:4095, 2:4095) ) pcopyin( A(:,:) ) 45
データの転送 (data clause) while ( error > tol ) { error = 0.0; #pragma acc kernels \ pcopy(anew[:][:]) pcopyin(a[:][:]) 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 \ 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]; N=M=4096 copyin (Host GPU) copyout (Host GPU) copy create present 46
[PGI tips] PGI_ACC_TIME $ PGI_ACC_TIME=1./a.out Accelerator Kernel Timing data /home/anaruse/src/openacc/jacobi/c/task1-solution/jacobi.c jacobi NVIDIA devicenum=0 time(us): 649,886 44: data region reached 200 times 44: data copyin transfers: 800 device time(us): total=14,048 max=41 min=15 avg=17 53: data copyout transfers: 800 device time(us): total=11,731 max=43 min=6 avg=14 44: compute region reached 200 times 46: kernel launched 200 times grid: [32x4094] block: [128] device time(us): total=382,798 max=1,918 min=1,911 avg=1,913 elapsed time(us): total=391,408 max=1,972 min=1,953 avg=1,957 46: reduction kernel launched 200 times grid: [1] block: [256] device time(us): total=48,235 max=242 min=241 avg=241 elapsed time(us): total=53,510 max=280 min=266 avg=267 統計データ 47
[PGI tips] PGI_ACC_NOTIFY $ PGI_ACC_NOTIFY=3./a.out upload CUDA data file=/home/anaruse/src/openacc/jacobi/c/task1- solution/jacobi.c function=jacobi line=44 device=0 variable=a bytes=16777216 launch CUDA kernel file=/home/anaruse/src/openacc/jacobi/c/task1- solution/jacobi.c function=jacobi line=46 device=0 num_gangs=131008 num_workers=1 vector_length=128 grid=32x4094 block=128 shared memory=1024 download CUDA data file=/home/anaruse/src/openacc/jacobi/c/task1- solution/jacobi.c function=jacobi line=53 device=0 variable=anew bytes=16736272 トレースデータ 48
NVIDIA VISUAL PROFILER (NVVP) 49
NVVP による解析 : データ転送がボトルネック 1 cycle 利用率 : 低い GPU kernel GPU kernel 50
過剰なデータ転送 while ( error > tol ) { error = 0.0; #pragma acc kernels \ pcopy(anew[:][:]) pcopyin(a[:][:]) 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 \ 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]; 51
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++) { 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]; 52
データ領域の指定 (data ディレクティブ ) #pragma acc data pcopy(a, Anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) 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]); copyin (CPU GPU) copyout (CPU GPU) copy create present #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]; 53
データ領域の指定 (data ディレクティブ ) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) 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]); copyin (CPU GPU) copyout (CPU GPU) copy create present #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]; 54
Host #pragma acc data \ pcopy(a) create(anew) while ( error > tol ) { error = 0.0; 適正なデータ転送 copyin GPU #pragma acc kernels \ pcopy(anew[:][:]) \ pcopyin(a[:][:]) { 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 \ 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]; 55
データ転送が減少 (NVVP) 1 cycle 利用率 : 高い 56
2 つの処理 CPU Memory データ転送 GPU Memory PCI 計算オフロード 計算オフロード データ転送 両方を考慮する必要がある 57
その他のデータ管理方法 float *array; Init( ) { array = (float*)malloc( ); input_array( array ); #pragma enter data copyin(array) Fin( ) { #pragma exit data copyout(array) output_array( array ); free( array ); Enter data Copyin Create Exit data Copyout Delete 58
その他のデータ管理方法 #pragma acc data pcopy(a,b) for (k=0; k<loop; k++) { #pragma acc kernels present(a,b) for (i=0; i<n; i++) { A[i] = suba(i,a,b); #pragma acc update self(a[0:1]) output[k] = A[0]; Update self CPU GPU Update device CPU GPU A[N-1] = input[k]; #pragma acc update device(a[n-1:1]) #pragma acc kernels present(a,b) for (i=0; i<n; i++) { B[i] = subb(i,a,b); 59
Unified Memory の利用 その他のデータ管理方法 PGI コンパイラオプション : -acc ta=,managed, プログラム実行中に動的に確保される配列は (C: malloc, Fortran: allocate) Unified Memory で管理される OpenACC のデータディレクティブで移動を指示する必要がない 60
リダクション ( 縮約計算 ) 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]); 61
リダクション ( 縮約計算 ) 演算の種類 while ( error > tol ) { error = 0.0; + 和 * 積 Max 最大 Min 最小 jacobi: ビット和 & ビット積 ^ XOR 論理和 && 論理積 #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; $ pgcc -Minfo=acc error = max(error, -acc jacobi.c abs(anew[j][i] - A[j][i]); 59, Generating present_or_copyout(anew[1:4094][1:4094]) Generating present_or_copyin(a[:][:]) #pragma acc kernels for (int Generating j = 1; j < Tesla N-1; j++) code{ 61, for (int Loop i is = 1; parallelizable i < M-1; i++) { 63, A[j][i] Loop is = Anew[j][i]; 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 62
リダクション (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++) { 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]); 演算種類 (C/C++) + 和 * 積 max 最大 min 最小 ビット和 & ビット積 ^ XOR 論理和 && 論理積 63
並列方法の指示 #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++) { 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]); 64
並列方法の指示 #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++) { $ pgcc -Minfo=acc -acc jacobi.c #pragma acc loop reduction(max:error) jacobi: for (int i = 1; i < M-1; i++) { 59, Anew[j][i] Generating = (A[j][i+1] present_or_copyout(anew[1:4094][1:4094]) + A[j][i-1] + Generating A[j-1][i] present_or_copyin(a[:][:]) + A[j+1][i]) * 0.25; 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 65
並列方法の指示 (loop ディレクティブ ) #pragma acc data pcopy(a) create(anew) while ( error > tol ) { error = 0.0; #pragma acc kernels pcopy(anew[:][:]) pcopyin(a[:][:]) #pragma acc loop gang vector(1) reduction(max:error) for (int j = 1; j < N-1; j++) { #pragma acc loop 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 66
実行条件設定 (gang, vector) #pragma acc loop gang vector(4) for (j = 0; j < 16; j++) { #pragma acc loop gang vector(16) for (i = 0; i < 16; i++) { #pragma acc loop gang vector(8) for (j = 1; j < 16; j++) { #pragma acc loop 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 67
ループを融合 (collapse) #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 68
ループを融合 (collapse) #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(128) for (int ji = 0; ji < (N-2)*(M-2); ji++) { j = (ji / (M-2)) + 1; i = (ji % (M-2)) + 1; 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 69
並列実行可能 (independent) #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]; 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 70
逐次に実行 (seq) #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 71
OPENACC プログラミング 概要紹介 プログラムの OpenACC 化 OpenACC 化事例 72
OPENACC ACCELERATES COMPUTATIONAL SCIENCE LSDalton PowerGrid CloverLeaf INCOMP3D Quantum Chemistry Aarhus University 12X speedup 1 week Medical Imaging University of Illinois 40 days to 2 hours Comp Hydrodynamics AWE 4X speedup Single CPU/GPU code CFD NC State University 4X speedup NekCEM MAESTRO CASTRO FINE/Turbo Comp Electromagnetics Argonne National Lab 2.5X speedup 60% less energy Astrophysics Stony Brook University 4.4X speedup 4 weeks effort CFD NUMECA International 10X faster routines 2X faster app 73
Speedup vs CPU LSDALTON Large-scale application for calculating highaccuracy molecular energies Lines of Code Modified Minimal Effort # of Weeks Required # of Codes to Maintain <100 Lines 1 Week 1 Source Big Performance LS-DALTON CCSD(T) Module Benchmarked on Titan Supercomputer (AMD CPU vs Tesla K20X) 11.7x OpenACC makes GPU computing approachable for domain scientists. Initial OpenACC implementation required only minor effort, and more importantly, no modifications of our existing CPU implementation. Janus Juul Eriksen, PhD Fellow qleap Center for Theoretical Chemistry, Aarhus University 7.9x ALANINE-1 13 ATOMS 8.9x ALANINE-2 23 ATOMS ALANINE-3 33 ATOMS https://developer.nvidia.com/openacc/success-stories 74
NUMECA FINE/TURBO Commercial CFD Application CHALLENGE Accelerate 20 year old highly optimized code on GPUs OpenACC enabled us to target routines for GPU acceleration without rewriting code, allowing us to maintain portability on a code that is 20-years old David Gutzwiller, Head of HPC NUMECA International SOLUTION Accelerated computationally intensive routines with OpenACC RESULTS Achieved 10x or higher speed-up on key routines Full app speedup of up to 2x on the Oak Ridge Titan supercomputer Total time spent on optimizing various routines with OpenACC was just five person-months https://developer.nvidia.com/openacc/success-stories 75
POWERGRID Advanced MRI Reconstruction Model CHALLENGE Produce detailed and accurate brain images by applying computationally intensive algorithms to MRI data Reduce reconstruction time to make diagnostic use possible Now that we ve seen how easy it is to program the GPU using OpenACC and the PGI compiler, we re looking forward to translating more of our projects Brad Sutton, Associate Professor of Bioengineering and Technical Director of the Biomedical Imaging Center University of Illinois at Urbana-Champaign SOLUTION Accelerated MRI reconstruction application with OpenACC using NVIDIA GPUs RESULTS Reduced reconstruction time for a single high-resolution MRI scan from 40 days to a couple of hours Scaled on Blue Waters at NCSA to reconstruct 3000 images in under 24 hours https://developer.nvidia.com/openacc/success-stories 76
INCOMP3D 3D Fully Implicit CFD Solver OpenACC is a highly effective tool for programming fully implicit CFD solvers on GPU to achieve true 4X speedup Lixiang Luo, Researcher Aerospace Engineering Computational Fluid Dynamics Laboratory North Carolina State University (NCSU) CHALLENGE Accelerate a complex implicit CFD solver on GPU SOLUTION Used OpenACC to run solver on GPU with minimal code changes. RESULTS Achieved up to 4X speedups on Tesla GPU over parallel MPI implementation on x86 multicore processor Structured approach to parallelism provided by OpenACC allowed better algorithm design without worrying about GPU architecture * CPU Speedup on 6 cores of Xeon E5645, with additional cores performance reduces due to partitioning and MPI overheads https://developer.nvidia.com/openacc/success-stories 77
NEKCEM Computational Electromagnetics Application CHALLENGE Enable NekCEM to deliver strong scaling on next-generation architectures while maintaining portability The most significant result from our performance studies is faster computation with less energy consumption compared with our CPU-only runs. Dr. Misun Min, Computation Scientist Argonne National Laboratory SOLUTION Use OpenACC to port the entire program to GPUs RESULTS 2.5x speedup over a highly tuned CPU-only version GPU used only 39 percent of the energy needed by 16 CPUs to do the same computation in the same amount of time https://developer.nvidia.com/openacc/success-stories 78
MAESTRO & CASTRO Astrophysics 3D Simulation On the reactions side, accelerated calculations allow us to model larger networks of nuclear reactions for similar computational costs as the simple networks we model now 2 weeks to learn OpenACC 2 weeks to modify code Adam Jacobs, PhD candidate in the Department of Physics and Astronomy at Stony Brook University CHALLENGE Pursuing strong scaling and portability to run code on GPU-powered supercomputers SOLUTION OpenACC compiler on OLCF s Titan supercomputer RESULTS 4.4x faster reactions than on a multi-core computer with 16 cores Accelerated calculations allow modeling larger networks of nuclear reactions for similar computational costs as simple networks https://developer.nvidia.com/openacc/success-stories 79
Speedup vs 1 CPU Core CLOVERLEAF Performance Portability for a Hydrodynamics Application CHALLENGE Application code that runs across architectures without compromising on performance We were extremely impressed that we can run OpenACC on a CPU with no code change and get equivalent performance to our OpenMP/MPI implementation. Wayne Gaudin and Oliver Perks Atomic Weapons Establishment, UK https://developer.nvidia.com/openacc/success-stories SOLUTION Use OpenACC to port the CloverLeaf mini app to GPUs and then recompile and run the same source code on multi-core CPUs RESULTS Same performance as optimized OpenMP version on x86 CPU 4x faster performance using the same code on a GPU Benchmarked Intel(R) Xeon(R) CPU E5-2690 v2 @ 3.00GHz, Accelerator: Tesla K80 80
Earthquake Simulation ( 理研 東大地震研 ) WACCPD 2016(SC16 併催 WS): Best Paper Award http://waccpd.org/wp-content/uploads/2016/04/sc16_waccpd_fujita.pdf より引用 81
NICAM 気象 気候モデル by 理研 AICS/ 東大 膨大なコード ( 数十万行 ) ホットスポットがない ( パレートの法則 ) 特性の異なる 2 種類の処理 力学系 メモリバンド幅ネック 物理系 演算ネック 82
Performance (GFLOPS) NICAM: 力学系 (NICAM-DC) OpenACC による GPU 化 主要サブルーチンは 全て GPU 上で動作 (50 以上 ) MPI 対応済み 2 週間 良好なスケーラビリティ Tsubame 2.5, 最大 2560 GPUs Scaling factor: 0.8 1.E+05 1.E+04 1.E+03 1.E+02 1.E+01 1.E+00 1.E+01 1.E+02 1.E+03 1.E+04 Number of CPUs or GPUs Weak scaling Tsubame 2.5 (GPU:K20X) K computer Tsubame 2.5 (CPU:WSM) (*) weak scaling 83 Courtesy of Dr. Yashiro from RIKEN AICS
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) 84 Courtesy of Dr. Yashiro from RIKEN AICS
Speedup vs. CPU 1-core NICAM: 物理系 (SCALE-LES) Atmospheric radiation transfer 物理系の中で 最も重い計算 OpenACC による GPU 対応済 Better 160 140 120 100 80 60 40 20 0 151 76.0 37.8 1.00 1.99 3.88 8.51 1 core 2 core 4 core 10 core 1 GPU 2 GPUs 4 GPUs Xeon E5-2690v2(3.0GHz,10-core) Tesla K40 85 (*) PCIデータ転送時間込み, グリッドサイズ :1256x32x32
Time (sec) SEISM3D 地震シミュレーション by 古村教授 ( 東大地震研 ) 主要サブルーチンの GPU 対応が完了 メモリバンド幅ネック 3 次元モデル (2 次元分割 ) 隣接プロセス間通信 600 500 400 SEISM3D (480x480x1024, 1K steps) 605 459 3.4x speedup ( アプリ全体 ) 140 120 100 GPU の実行時間内訳 Others (CPU, MPI and so on) [CUDA memcpy DtoH] [CUDA memcpy HtoD] Better 300 200 100 0 K: 8x SPARC64 VIIIfx CPU: 8x Xeon E5-2690v2 134 GPU: 8x Tesla K40 80 60 40 20 0 GPU: 8x Tesla K40 (other subroutines) update_vel_pml update_vel update_stress_pml update_stress diff3d_* 86
Speedup vs. 1 CPU core FFR/BCM ( 仮称 ) 次世代 CFD コード by 坪倉准教授 ( 理研 AICS/ 北大 ) MUSCL_bench: MUSCL スキームに基づく Flux 計算 ( とても複雑な計算 ) CFD 計算の主要部分 (60-70%) OpenACC による GPU 対応 完了 Better 35 30 25 20 15 10 5 0 33.21 8.30 4.55 1.00 1.93 1 core 2 core 5 core 10 core 1 GPU Xeon E5-2690v2(3.0GHz,10-core) Tesla K40 87 (*) PCIデータ転送時間込み サイズ :80x32x32x32
Time (second) CM-RCM IC-CG IC-CG 法のベンチマークコード by 中島教授 ( 東大 ) CM-RCM 法 (Cyclic Multi-coloring Reverse Cuthill-Mckee) を使用 メインループ内のサブルーチンを全て OpenACC で GPU 化 4 3 3.36 CM-RCM ICCG (100x100x100) 2 1.58 1.53 1.39 1 0.655 Better 0 OpenMP OpenMP OpenMP OpenMP OpenACC Opteron 6386SE (2.8GHz,16core) SPARC64 Ixfx (1.85GHz,16core) Xeon E5-2680v2 (2.8GHz,10core) Xeon-Phi 5110P Tesla K40 88 Courtesy of Dr. Ohshima from U-Tokyo
GFLOPS CCS-QCD QCD コード by 石川准教授 ( 広島大 ) BiCGStab 計算を全て OpenACC で GPU 化 データレイアウトを変更 : AoS SoA Better 90 80 70 60 50 40 30 20 10 0 CCS-QCD: BiCGStab Total FLOPS 42.5 32.3 Xeon E5-2690v2(3.0GHz,10core,OpenMP) 53.3 Tesla K40(OpenACC) 57.9 60.8 63.4 24.3 22.1 20.9 19.7 8x8x8x32 8x8x8x64 8x8x16x64 8x16x16x64 16x16x16x64 Problem Size 89
CUDA 90
プログラミングモデル アーキテクチャ 性能 Tips CUDA プログラミング 91
GPU コンピューティング CPU CPU Memory GPU Memory GPU PCI 高スループット指向のプロセッサ 分離されたメモリ空間 92
CPU GPU プログラム GPU void saxpy(int n, float a, float *x, float *y) { for (int i = 0; i < n; ++i) y[i] += a*x[i]; global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx; if (i < n) y[i] += a*x[i]; saxpy(n, 3.0, x, y); size_t size = sizeof(float) * N; cudamemcpy(d_x, x, size, cudamemcpyhosttodevice); cudamemcpy(d_y, y, size, cudamemcpyhosttodevice); saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y); cudadevicesynchronize(); cudamemcpy(y, d_y, size, cudamemcpydevicetohost); 93
GPU 実行の基本的な流れ CPU 入力データ転送 GPU GPU は CPU からの制御で動作 GPU カーネル投入 同期 GPU 上で演算 入力データ : CPU から GPU に転送 (H2D) GPU カーネル : CPU から投入 出力データ : GPU から CPU に転送 (D2H) 出力データ転送 94
CPU GPU プログラム GPU void saxpy(int n, float a, float *x, float *y) { for (int i = 0; i < n; ++i) y[i] += a*x[i]; 入力データ転送 global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx; if (i < n) y[i] += a*x[i]; saxpy(n, 3.0, x, y); カーネル起動 同期 出力データ転送 size_t size = sizeof(float) * N; cudamemcpy(d_x, x, size, cudamemcpyhosttodevice); cudamemcpy(d_y, y, size, cudamemcpyhosttodevice); saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y); cudadevicesynchronize(); cudamemcpy(y, d_y, size, cudamemcpydevicetohost); 95
GPU プログラム (Unified Memory) CPU GPU void saxpy(int n, float a, float *x, float *y) { for (int i = 0; i < n; ++i) y[i] += a*x[i]; global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx; if (i < n) y[i] += a*x[i]; カーネル起動 saxpy(n, 3.0, x, y); 同期 saxpy<<< N/128, 128 >>>(N, 3.0, x, y); cudadevicesynchronize(); 96
void saxpy(int n, float a, float *x, float *y) { CPU for (int i = 0; i < n; ++i) y[i] += a*x[i]; saxpy(n, 3.0, x, y); GPU カーネル GPU global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx.x; if (i < n) y[i] += a*x[i]; saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y); Global スレッド ID GPU カーネル : 1 つの GPU スレッドの処理内容を記述 基本 : 1 つの GPU スレッドが 1 つの配列要素を担当 97
Execution Configuration ( ブロック数とブロックサイズ ) スレッド ID global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx.x; if (i < n) y[i] += a*x[i]; ブロックサイズ saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y); ブロック ID ブロック数 ブロックサイズ ブロック数 x ブロックサイズ = 配列要素数 98
y[i] = a*x[i] + y[i] スレッド階層 ( スレッド ブロック グリッド ) x[] 0 127 128 255 256 383 384 y[] 0 127 128 255 256 383 384 スレッド (global) (local) 0 127 128 0 127 255 256 0 127 383 384 0 ブロック0 ブロック1 ブロック2 グリッド ブロックサイズ ( スレッド数 / ブロック ) は カーネル毎に設定可能 推奨 : 128 or 256 スレッド 99
Execution Configuration ( ブロック数とブロックサイズ ) global void saxpy(int n, float a, float *x, float *y) { int i = threadidx.x + blodkdim.x * blockidx.x; if (i < n) y[i] += a*x[i]; saxpy<<< N/256, N/128, 64, 256 128 64 >>>(N, 3.0, d_x, d_y); ブロック数 ブロックサイズ ブロック数 x ブロックサイズ = 配列要素数 100
2D 配列の GPU カーネル例 global void MatAdd(float A[N][N], float B[N][N], float C[N][N]) { int i = threadidx.x + blodkdim.x * blockidx.x; int j = threadidx.y + blodkdim.y * blockidy.y; if ( i < N && j < N ) C[i][i] = A[i][j] + B[i][j]; Global スレッド ID (x) Global スレッド ID (y) ブロックサイズ (x,y) dim3 sizeblock( 32, 64, 16, 16 84 ); dim3 numblocks( N/sizeBlock.x, N/sizeBlock.y ); MatAdd<<< numblocks, sizeblock >>>(A, B, C); ブロック数 (x,y) ブロックサイズ ( ブロック形状 ) は 1D~3D で表現可能 101
ブロック マッピング スレッド マッピング (0,0) dim3 sizeblock(16,16) (0,0) dim3 sizeblock(32,8) (0,0) (1,0) (0,0) (1,0) (2,0) (31,7) (0,1) (1,1) (15,15) (0,2) (1,2) (0,1) (1,1) (2,1) (0,3) (1,3) (0,2) (1,2) (2,2) (0,4) (1,4) ブロック ID(blockIdx) スレッド ID(threadIdx) 102
プログラミングモデル アーキテクチャ 性能 Tips CUDA プログラミング 103
GPU アーキテクチャ概要 Pascal GP100 PCI I/F ホスト接続インタフェース Giga Thread Engine SM に処理を割り振るスケジューラ DRAM I/F (HBM2) 全 SM PCI I/F からアクセス可能なメモリ ( デバイスメモリ, フレームバッファ ) L2 cache (4MB) 全 SM からアクセス可能な R/W キャッシュ SM (Streaming Multiprocessor) 並列 プロセッサ GP100: 最多 60 104
SM (Stream Multi-Processor) CUDA コア GPU スレッドはこの上で動作 GP100: 64 個 /SM Other units LD/ST, SFU, etc Pascal GP100 レジスタ (32bit): 64K 個 共有メモリ : 64KB Tex/L1 キャッシュ 105
GPU カーネル実行の流れ CPU が GPU に グリッドを投入 具体的な投入先は Giga Thread Engine グリッド ブロック スレッド スレッド ブロック 106
GPU カーネル実行の流れ Giga Thread Engine(GTE) が SM に ブロックを投入 GTE は ブロックスケジューラ グリッドをブロックに分解して ブロックを 空いている SM に割当てる グリッド ブロック スレッド スレッド ブロック 107
ブロックを SM に割り当て 各ブロックは 互いに独立に実行 ブロック間では同期しない 実行順序の保証なし 1 つのブロックは複数 SM にまたがらない 1 つの SM に 複数ブロックが割当てられることはある グリッド ブロック ブロック ブロック ブロック 108
GPU カーネル実行の流れ SM 内のスケジューラが スレッドを CUDA コアに投入 グリッド ブロック スレッド スレッド ブロック 109
GPU カーネル実行の流れ SM 内のスケジューラが ワープを CUDA コアに投入 ワープ : 32 スレッドの塊 ブロックをワープに分割 実行可能なワープを 空 CUDA コアに割当てる グリッド ブロック ワープ ワープ ブロック 110
ワープの CUDA コアへの割り当て ワープ内の32スレッドは 同じ命令を同期して実行 各ワープは 互いに独立して実行 同じブロック内のワープは 明示的に同期可能 ( syncthreads()) グリッド SIMT (Single Instruction Multiple Threads) ブロック ワープ ワープ ブロック ワープ 111
GPU アーキの変化を問題としないプログラミングモデル Pascal, CC6.0 64 cores /SM Kepler, CC3.5 192 cores /SM Maxwell, CC5.0 128 cores /SM 112
プログラミングモデル アーキテクチャ 性能 Tips CUDA プログラミング 113
リソース使用率 (Occupancy) SM の利用効率を上げる SM に割当て可能なスレッド数を 上限に近づける レジスタ使用量 (/ スレッド ) できる限り減らす DP(64bit) は 2 レジスタ消費 レジスタ割り当て単位は 8 個 レジスタ使用量と 割当て可能なスレッド数の関係 32 レジスタ : 2048(100%), 64 レジスタ : 1024(50%) 128 レジスタ :512(25%), 256 レジスタ : 256(12.5%) CUDA コア数 : 64 最大スレッド数 : 2048 最大ブロック数 : 32 共有メモリ : 64KB レジスタ数 (32-bit): 64K 個 リソース量 /SM (GP100) 114
リソース使用率 (Occupancy) SM の利用効率を上げる SM に割当て可能なスレッド数を 上限に近づける スレッド数 (/ ブロック ) 64 以上にする 64 未満だと最大ブロック数がネックになる CUDA コア数 : 64 最大スレッド数 : 2048 最大ブロック数 : 32 共有メモリ : 64KB レジスタ数 (32-bit): 64K 個 リソース量 /SM (GP100) 共有メモリ使用量 (/ ブロック ) できる限り減らす 共有メモリ使用量と 割当て可能なブロック数の関係 32KB:2 ブロック, 16KB:4 ブロック, 8KB:8 ブロック 115
空き時間を埋める リソース使用率 (Occupancy) (*) 操作 : GPUカーネル, データ転送 CUDAストリーム ( キュー) 同じCUDAストリームに投入した操作 : 投入順に実行 別のCUDAストリームに投入した操作 : 非同期に実行 ( オーバラップ実行 ) [CUDA ストリームの効果例 ] GPU カーネルとデータ転送がオーバラップして同時に実行されている 116
デバイスメモリへのアクセスは まとめて コアレス アクセス 32 スレッド ( ワープ ) のロード / ストアをまとめて メモリトランザクションを発行 トランザクションサイズ : 32B, 64B or 128B トランザクション数は 少ないほど良い 配列 128B 境界 128B 境界 Best 128B x1 0 1 2 3 28 29 30 31 Good 0 1 14 15 16 17 30 31 64B x2 Bad 32B x32 0 1 2 117
分岐を減らす ワープ内のスレッドが別パスを選択すると遅くなる ワープ内のスレッドは 命令を共有 (SIMT) ワープ内のスレッドが選んだ全パスの命令を実行 あるパスの命令を実行中 そのパスにいないスレッドは inactive 状態 Path divergence を減らす できる限り 同ワープ内のスレッドは同じパスを選択させる true 0 処理 X 1 A[id] > 0 2 3 処理 Y false 118
まとめ GPU Computing Volta GPUs OpenACC CUDA 119