GPUコンピューティング入門 2015.08.26 エヌビディア合同会社 CUDAエンジニア 村上真奈
エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 2
エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 3
18 1993 年創立 共同創業者 社長兼 CEO : ジェンスン フアン 1999 年 銘柄コード NVDA で NASDAQ に株式上場 1999 年に GPU を発明し 現在までに10 億点以上を出荷 2015 年会計年度の収益 : 46.8 億ドル世界中に 9,300 名の従業員 7,300 件の特許取得済み資産本社 : カリフォルニア州サンタクララ 4
3 ゲーミング オートモーティブ エンタープライズ HPC クラウド 当社は ビジュアル コンピューティングが必要不可欠で重要な価値を持つ大規模な市場に特化し プロセッサのプラット フォーム ソフトウェア システム サービスを提供しています 当社はPCテクノロジ データセンター テクノロジ モバイル テ クノロジの革新に取り組んでいます そして 当社の発明は さまざまな業界のOEM製品の原動力となっています 5
6
東京工業大学 TSUBAME 2.5 4,224 枚の Tesla K20X 単精度理論性能値で日本 No.1 スパコン 17PFLOPS SP 7
NVIDIA GPU OFFERS TOP LEVEL COMPUTATIONAL PERFORMANCE WITH HIGH ENERGY EFFICIENVY From SC TOP500 Nov., 2014 Rank Country Site System Cores Rmax 1 China National Super Computer Center in Guangzhou Tianhe-2 (MilkyWay-2) - TH-IVB-FEP Cluster, Intel Xeon E5-2692 12C 2.200GHz, TH Express-2, Intel Xeon Phi 31S1P (TFlop/s) Rpeak (TFlop/s) Power (kw) 3,120,000 33,862.70 54,902.40 17,808 2 US DOE/SC/Oak Ridge National laboratory Titan-Cray XK7, Opt. 6274 16C 2.2GHz, NVIDIA K20x 560,640 17,590.00 27,112.50 8,209 3 US DOE/NNSA/LLNL Sequoia - BlueGene/Q, Power BQC 16C 1.60 GHz, Custom 1,572,864 17,173.20 20,132.70 7,890 4 Japan RIKEN Advanced Institute for Computational Science (AICS) 5 US DOE/SC/Argonne National Laboratory K computer, SPARC64 VIIIfx 2.0GHz, Tofu interconnect Mira - BlueGene/Q, Power BQC 16C 1.60GHz, Custom 705,024 10,510.00 11,280.40 12,660 786,432 8,586.60 10,066.30 3,945 In GREEN500 the most energy efficient super computers, NVIDIA GPU drives 8 systems out of TOP 10. 8
REAL WORLD EXAMPLE Rendering 30-second Animation at Renault 9
Deep Learning における GPU の活用 Deep Learning に GPU を活用 Input Result 110 28% 26% 60 16% 12% 7% 0 0 4 2010 2011 2012 2013 2014 person dog chair GPU 対応した Deep Learning 用ツール Caffe Torch Theano Cuda-convnet cudnn cublas 10
SGEMM / W GPU ロードマップ 72 60 Volta 48 36 Pascal 24 Maxwell 12 0 Tesla Fermi Kepler 2008 2010 2012 2014 2016 2018 11
TESLA KEPLER FAMILY WORLD S FASTEST AND MOST EFFICIENT HPC ACCELERATORS GPUs Single Precision Peak (SGEMM) Double Precision Peak (DGEMM) Memory Size Memory Bandwidth (ECC off) PCIe Gen System Solution CFD, BioChemistry, Neural Networks, High Energy Physiscs, Graph analytics, Material Science, BioInformatics, M&E K80 K40 8.74 TF (5.6TF) 4.29 TF (3.22TF) 2.91TF (1.87TF) 1.43 TF (1.33 TF) 24 GB 480GB/s (240GB/s x2) 12 GB 288 GB/s Gen 3 Gen 3 Server + Workstation Server + Workstation Weather & Climate, Physics, BioChemistry, CAE, Material Science K20X K20 3.95 TF (2.90 TF) 3.52 TF (2.61 TF) 1.32 TF (1.22 TF) 1.17 TF (1.10 TF) 6 GB 250 GB/s Gen 2 Server only 5 GB 208 GB/s Gen 2 Server + Workstation Image, Signal, Video, Seismic K10 4.58 TF 0.19 TF 8 GB 320 GB/s Gen 3 Server only 12
M6000 K6000 K5200 K4200 K2200 K620 K420 # CUDA Cores 3072 2880 2304 1344 640 384 192 Single Precision 5.2 TFLOPs 3.1 TFLOPs 2.1 TFLOPs 1.3 TFLOPs 0.8 TFLOPs 0.3 TFLOPs PCIe Gen 3.0 2.0 Memory Size 12GB 12 GB 8 GB 4 GB 4 GB 2 GB 1 GB Memory BW 317 GB/s 288 GB/s 192 GB/s 173 GB/s 80 GB/s 29 GB/s 29 GB/s Slots + Display Connectors THE NEW QUADRO FAMILY 2x DP * + 2x DVI 2x DP * + 2x DVI 2x DP * + 2x DVI 2x DP * + DVI * * 2x DP + DVI DP + DVI * DP + DVI Max Resolution 4096 x 2160 3840 x 2160 Max Displays 4 4 4 4 4 4 4 Pro Features SDI, SYNC, STEREO, MOSAIC, NVIEW MOSAIC, NVIEW Board Power 250W 225 W 150 W 108 W 68 W 45 W 41 W * DisplayPort 1.2 multi-streaming can be used to drive multiple displays from a single DP connector 13
TFLOPS TFLOPS 前世代比 3 倍の性能 1.4 1.2 1 0.8 0.6 0.4 0.2 0 3.5 3 2.5 2 1.5 1 0.5 0 Double Precision FLOPS (DGEMM) 1.33 TFLOPS 0.40 TFLOPS Tesla M2090 Tesla K40 Single Precision FLOPS (SGEMM) 3.22 TFLOPS 0.89 TFLOPS Tesla M2090 Tesla K40 Tesla M2090 Tesla K40 CUDA コア数 512 2880 倍精度演算性能 DGEMM 単精度演算性能 SGEMM 665 G 400 GF 1.33 TF 0.89 TF 1.43 TF 1.33 TF 4.29 TF 3.22 TF メモリバンド幅 178 GB/s 288 GB/s メモリサイズ 6 GB 12 GB 消費電力 225W 235W 14
NVIDIA GPU SCALABLE ARCHITECTURE FROM SUPER COMPUTER TO MOBILE Tegra Tesla In Super Computers Quadro In Work Stations GeForce In PCs Mobile GPU In Tegra 17
2015 TEGRA X1 MOBILE SUPERCHIP 256-core Maxwell GPU 8-core 64-bit CPU 4Kp60 10-bit H.265/VP9 19
CPU: Quad ARM Cortex A57/A53 64/32b CPU that delivers Performance and Power Efficiency GPU: Next Generation 256- Core Maxwell GPU that deliver Class-Leading Performance and Power Efficiency End-to-End 4k 60fps Pipeline that delivers Premium 4K Experience Built on 20nm Process Technology TEGRA X1 OVERVIEW 20
Advancements BRIDGING THE GAP Maxwell Tesla Fermi Kepler Tegra K1 Tegra X1 GEFORCE ARCHITECTURE Tegra 4 Tegra 3 MOBILE ARCHITECTURE 21
GFLOPS WORLD S 1 ST TERAFLOPS MOBILE PROCESSOR 1200 Tegra X1 (FP16) Tegra X1 1000 Core i7 GPU GPU CPU 800 CPU FP16/INT16 600 400 Tegra K1 200 Tegra 2 Tegra 3 Tegra 4 0 TIME Note: 4790K Core i7, CPU @ 4GHz, GPU 22 @ 350 MHz
エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 24
NVIDIA GPU の歴史 2010 Fermi 3 Billion Transistors 2012 Kepler 7 Billion Transistors GPU 統合シェーダ + CUDA 25
PCI Express GPU の構造 GPU Giga Thread Engine SM SM SM SM L2 Cache DRAM 26
GPU アーキテクチャ概要 PCI I/F SM SM SM SM SM SM SM SM ホスト接続インタフェース Giga Thread Engine SM に処理を割り振るスケジューラ DRAM (384-bit, GDDR5) SM SM SM SM SM SM SM 全 SM PCI I/F からアクセス可能なメモリ ( デバイスメモリ, フレームバッファ ) Kepler GK110 L2 cache (1.5MB) 全 SM からアクセス可能な R/W キャッシュ SM (Streaming Multiprocessor) 並列 プロセッサ 27
SM (STREAMING MULTIPROCESSOR) CUDA core GPU スレッドはこの上で動作 Kepler: 192 個 Other units DP, LD/ST, SFU Register File (65,536 x 32bit) Shared Memory/L1 Cache (64KB) Kepler GK110 Read-Only Cache(48KB) 28
COMPUTE CAPABILITY GPU コアアーキテクチャのバージョン CUDA GPUs : https://developer.nvidia.com/cuda-gpus アーキテクチャは進化する 高効率の命令実行 省消費電力 29
SM ARCHITECTURE VS COMPUTE CAPABILITY Instruction Cache Scheduler Scheduler Dispatch Dispatch Register File Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Load/Store Units x 16 Special Func Units x 4 Interconnect Network 64K Configurable Cache/Shared Mem Uniform Cache Fermi CC 2.0 : 32 cores / SM Kepler CC 3.5 : 192 cores / SMX Maxwell CC 5.0 : 128 cores / SMM 30
GPU コンピューティングとは? GPUは何の略? Graphics Processing Unit 3DCG 等 画像データ処理の為のデバイス GPUによる汎用コンピューティングのこと 計算科学など様々な用途でGPUを利用する 31
ヘテロジニアス コンピューティング CPU 逐次処理に最適化 GPU Accelerator 並列処理に最適化 32
GPU アプリケーションの例 画像処理コンピュータビジョン医療画像防衛計算化学 気象金融工学バイオ数値解析 33
GPU アクセラレーションの実現方法 アプリケーション GPU ライブラリ OpenACC ディレクティブ CUDA ライブラリを呼び出すだけ簡単に高速化を実現 既存コードにディレクティブを挿入して高速化 重要なコードを CUDA で記述最も自由度が高い 34
簡単 GPU アクセラレーションの実現方法 ライブラリ ライブラリを呼び出すだけで 高速化が可能ライブラリとして提供されている機能のみ高速化が可能 OpenACC 既存の C 言語や Fortran のコードにディレクティブを挿入するだけで簡単に高速化 最適化はコンパイラが行う為 細かいチューニングを行う事は出来ない 高速化 CUDA 自由度が最も高く 細かいチューニングが可能 CUDA でのプログラミングを学ぶ必要がある 35
エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 36
GPU アクセラレーションの実現方法 アプリケーション GPU ライブラリ OpenACC ディレクティブ CUDA ライブラリを呼び出すだけ簡単に高速化を実現 既存コードにディレクティブを挿入して高速化 重要なコードを CUDA で記述最も自由度が高い 37
OPENACC 標準的な GPU ディレクティブ シンプル : ディレクティブを挿入するのみ コードを変更する事なく高速化 オープン : OpenACC はマルチコアプロセッサで並列化を行う為のオープン標準 柔軟 : GPU ディレクティブは 高い並列性を保ちつつ同一コードで複数のアーキテクチャに対応可能 38
OpenACC メンバーとパートナー 39
コンパイラとツール 2013 年 12 月 ~ 2014 年 1 月 ~ 2015 年 ( 予定 ) コンパイラ OpenACC 2.0 対応 デバッグツール 40
簡単に高速 自動車金融生命科学 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 倍 41
大学関係者の方は無償で使用可能に 簡単に始められる 下記のサイトから OpenACC toolkit をダウンロード https://developer.nvidia.com/openacc PGI コンパイラ /MPI/CUDA など一式が簡単にインストール可能 42
実行モデル アプリケーション コード $acc parallel GPU CPU 逐次部分は CPU コードを生成 $acc end parallel 計算の重い部分 並列部分は GPU コードを生成 43
OpenACC ディレクティブ CPU GPU コンパイラへシンプルなヒント Program myscience... serial code...!$acc kernels do k = 1,n1 do i = 1,n2... parallel code... enddo enddo!$acc end kernels... End Program myscience コンパイラがコードを並列化 コンパイラへの OpenACC ヒント 並列部はGPUで 逐次処理はCPUで動作 Fortran または C言語 のオリジナルコード 44
OpenMP と OpenACC の比較 OpenMP OpenACC CPU CPU GPU main() { double pi = 0.0; long i; main() { double pi = 0.0; long i; CPUコアに計算処理を分散 #pragma omp parallel for reduction(+:pi) for (i=0; i<n; i++) { double t = (double)((i+0.05)/n); pi += 4.0/(1.0+t*t); printf( pi = %f\n, pi/n); #pragma acc kernels for (i=0; i<n; i++) { double t = (double)((i+0.05)/n); pi += 4.0/(1.0+t*t); printf( pi = %f\n, pi/n); GPU コアに計算処理を分散 45
OpenACC ディレクティブ構文 C/C++ #pragma acc 指示行 [ 節 [,] 節 ] ] { structured block Fortran!$acc 指示行 [ 節 [,] 節 ] ] { structured block!$acc end directive 46
OpenACC構文: parallel 指示行 parallel : 並列に実行される領域を指示行で指定 #pragma acc parallel for(int i=0;i<n;i++){ a[i] = 0.0; b[i] = 1.0; c[i] = 2.0; kernel 1 Kernel(カーネル): GPU上で実行される 関数 47
OpenACC 構文 : kernels 指示行 kernels : 複数のカーネルを作成 #pragma acc kernels for(int i=0;i<n;i++){ a[i] = 0.0; b[i] = 1.0; c[i] = 2.0; #pragma acc kernels for(int i=0;i<n;i++){ a[i] = b[i] + c[i]; kernel 1 kernel 2 Kernel( カーネル ): GPU 上で実行される関数 48
[C tips]: restrict 修飾子 コンパイラに対して明示的に restrict 修飾子を指定 ポインタのエイリアシングを制限 例 ) float *restrict ptr OpenACC コンパイラに restrict 修飾子をつけ変数の独立性を伝える 独立性の保障がないとコンパイラは並列化を行う事が出来ない http://en.wikipedia.org/wiki/restrict 49
例 :SAXPY (Y=A*X+Y) Trivial first example Apply a loop directive Learn compiler commands int main(int argc, char **argv) { int N = 1<<20; // 1 million floats if (argc > 1) N = atoi(argv[1]); #include <stdlib.h> void saxpy(int n, float a, float *x, float *restrict y) { #pragma acc kernels for (int i = 0; i < n; ++i) y[i] = a * x[i] + y[i]; *restrict: y は x のエイリアスでない と明示的に指定 float *x = (float*)malloc(n * sizeof(float)); float *y = (float*)malloc(n * sizeof(float)); for (int i = 0; i < N; ++i) { x[i] = 2.0f; y[i] = 1.0f; saxpy(n, 3.0f, x, y); return 0; 50
C 言語 :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 データの移動 51
Fortran: SAXPY (Y=A*X+Y) 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)... 52
コンパイルオプション C: pgcc acc -ta=nvidia -Minfo=accel o saxpy_acc saxpy.c Fortran: pgf90 acc -ta=nvidia -Minfo=accel o saxpy_acc saxpy.f90 ターゲットに nvidia を指定 コンパイラが GPU 用のコードを生成する際の情報を表示する 53
簡単にコンパイル OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) $ pgcc acc { ta=nvidia Minfo=accel 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);... 54
簡単に実行 OpenMP / OpenACC void saxpy(int n, float a, float *x, float *restrict y) $ pgcc -Minfo -acc { saxpy.c saxpy: $ nvprof./a.out #pragma acc kernels copy(y[:n]) copyin(x[:n]) ==10302== 16, Generating NVPROF #pragma is present_or_copy(y[:n]) profiling omp parallel process for 10302, command:./a.out ==10302== Generating Profiling for present_or_copyin(x[:n]) (int application: i = 0; i./a.out < n; ++i) ==10302== Generating Profiling y[i] Tesla result: += code a*x[i]; Time(%) 19, Loop Time is parallelizable Calls Avg Min Max Name 62.95% Accelerator 3.0358ms kernel 2 generated 1.5179ms 1.5172ms 1.5186ms [CUDA memcpy HtoD] 31.48% 19, 1.5181ms #pragma... acc loop 1 1.5181ms gang, vector(128) 1.5181ms /* 1.5181ms blockidx.x [CUDA threadidx.x memcpy DtoH] */ 5.56% 268.31us saxpy(n, 3.0, 1 x, 268.31us y); 268.31us 268.31us saxpy_19_gpu... 55
例 : ヤコビ反復法 正しい値になるように反復計算を行う 隣接点の平均値で値を更新 連立一次方程式を解く為のオーソドックスな手法 例 : 2 次元ラプラス方程式 : 2 f(x, y) = 0 A(i,j+1) A(i-1,j) A(i,j) A(i+1,j) A k+1 i, j = A k(i 1, j) + A k i + 1, j + A k i, j 1 + A k i, j + 1 4 A(i,j-1) 56
ヤコビ反復法 ( アルゴリズム ) 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]; 57
並列領域 (OpenMP) while ( error > tol ) { error = 0.0; #pragma omp parallel for shared(m, n, Anew, 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 omp parallel for shared(m, n, Anew, A) for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; 58
並列領域 (OpenACC) 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 複数のカーネル 59
[PGI tips] コンパイラメッセージ $ pgcc acc ta=nvidia 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 60
並列領域 (KERNELS CONSTRUCT) while ( error > tol ) { error = 0.0; Parallels と 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; 並列走行の開始 $ pgcc acc ta=nvidia error = max(error, -Minfo=accel abs(anew[j][i] -jacobi.c A[j][i]); jacobi: 59, Generating present_or_copyout(anew[1:4094][1:4094]) Parallels Kernels 複数のGPUカーネル Generating present_or_copyin(a[:][:]) #pragma acc kernels Generating code{ for (int j = 1; j <Tesla N-1; j++) for (int = 1; i < M-1; i++) { 61, Loop iis parallelizable A[j][i] = Anew[j][i]; 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 61
データ転送 (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] + $ pgcc acc ta=nvidia A[j-1][i] -Minfo=acc + A[j+1][i]) jacobi.c * 0.25; jacobi: error = max(error, abs(anew[j][i] - A[j][i]); 59, Generating present_or_copyout(anew[1:4094][1:4094]) Generating present_or_copyin(a[:][:]) #pragma Generating acc kernels Tesla code for 61, (int Loop j = is 1; parallelizable j < N-1; j++) { for (int i = 1; i < M-1; i++) { 63, Loop is parallelizable A[j][i] = Anew[j][i]; 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
コードの解析 実行状況を確認 ボトルネックはどの部分か? 実行時間の内訳を調べる 63
-ta=nvidia,time コンパイルオプションに -ta=nvidia,time を追加して コンパイル 実行 /home/openacc/c/jacobi.c jacobi NVIDIA devicenum=0 Kernel 実行 :196ms time(us): 4,595,922 44: compute region reached 200 times 46: kernel launched 200 times grid: [32x4094] block: [128] device time(us): total=196,036 max=1,053 min=931 avg=980 データコピー (H->D):1087ms elapsed time(us): total=201,618 max=1,084 min=958 avg=1,008 46: reduction kernel launched 200 times grid: [1] block: [256] device time(us): total=39,356 max=206 min=187 avg=196 elapsed time(us): total=42,155 max=227 min=200 avg=210 44: data region reached 200 times 44: data copyin transfers: 800 device time(us): total=1,087,027 max=1,374 min=1,354 データコピーがボトルネック avg=1,358 53: compute region reached 200 times 55: kernel launched 200 times 64 grid: [32x4094] block: [128]
NVIDIA Visual Profiler (NVVP) を使用 65
NVVP による解析 : データ転送がボトルネック 1 cycle 利用率 : 低い GPU kernel GPU kernel 66
計算処理とデータ転送 CPU Memory データ転送 GPU Memory PCI 計算オフロード 計算オフロード データ転送 両方を考慮する必要がある 67
OpenACC 構文 : データ指示行 copy ( X ) copyin(list) + copyout(list) copyin ( X ) アクセラレータ領域に入る際に GPU 上に X 用のメモリを確保し ホストから GPU( デバイス ) へ X を転送する copyout ( X ) アクセラレータ領域に入る際に GPU 上に X 用のメモリを確保し アクセラレータ領域から出る時に GPU( デバイス ) からホストへ X を転送する create ( X ) アクセラレータ領域に入る時に GPU 上に X 用のメモリが確保される ( 転送はされない ) present ( X ) アクセラレータ領域に入る時に X が既にデバイス上に存在することを示す 68
OpenACC 構文 : データ指示行 pcopy ( X ) present (X) + copy(x) pcopyin ( X ) present (X) + copyin(x) pcopyout ( X ) present (X) + copyout(x) pcreate ( X ) present (X) + create(x) 69
データ転送 (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]) 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: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 70
データ転送 (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]; copyin (Host GPU) copyout (Host GPU) copy create present pcopyin pcopyout pcopy pcreate 71
過剰なデータ転送 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]; 72
Host while ( error > tol ) { error = 0.0; 過剰なデータ転送 GPU #pragma acc kernels \ pcopy(anew[:][:]) \ pcopyin(a[:][:]) { #pragma acc kernels \ pcopy(a[:][:]) \ pcopyin(anew[:][:]) { copyin copyout 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]); for (int j = 1; j < N-1; j++) { for (int i = 1; i < M-1; i++) { A[j][i] = Anew[j][i]; 73
データ領域 (data construct) #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]); #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 74
データ領域 (data CONSTRUCT) #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]); #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 75
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]; 76
データ転送の削減 (NVVP) 1 cycle 稼働率 : 高い 77
GPU アクセラレーションの実現方法 アプリケーション GPU ライブラリ OpenACC ディレクティブ CUDA ライブラリを呼び出すだけ簡単に高速化を実現 既存コードにディレクティブを挿入して高速化 重要なコードを CUDA で記述最も自由度が高い 78
CUDA とは? Compute Unified Device Architectureの略 NVIDIA GPU 上の汎用並列計算プラットフォーム Linux Windows MacOS X(+Android) で動作 現在 7.0が最新 7.5RCも公開中 79
CUDA 開発 実行環境 ライブラリ ミドルウェア 開発環境 cufft cublas cusparse cudnn cusolver curand Thrust NPP NVRTC MATLAB Mathematica etc.. NVCC (CUDA compiler) プログラミング言語 CUDA-GDB C C++ C++11 Fortran Java Python etc.. Profiler NVIDIA-GPUs( ハードウェア ) Nsight IDE HyperQ Dynamic Parallelism GPU Direct 82
CUDA 開発 実行環境 ライブラリ ミドルウェア 開発環境 cufft cublas cusparse cudnn cusolver curand Thrust NPP NVRTC MATLAB Mathematica etc.. NVCC (CUDA compiler) プログラミング言語 CUDA-GDB C C++ C++11 Fortran Java Python etc.. Profiler NVIDIA-GPUs( ハードウェア ) Nsight IDE HyperQ Dynamic Parallelism GPU Direct 83
進化するハードウェア NVIDIA-GPUS 84
進化するハードウェア NVIDIA-GPUS Hyper-Q Dynamic Parallelism GPU Direct 85
CUDA 開発 実行環境 ライブラリ ミドルウェア 開発環境 cufft cublas cusparse cudnn cusolver curand Thrust NPP NVRTC MATLAB Mathematica etc.. NVCC (CUDA compiler) プログラミング言語 CUDA-GDB C C++ C++11 Fortran Java Python etc.. Profiler NVIDIA-GPUs( ハードウェア ) Nsight IDE HyperQ Dynamic Parallelism GPU Direct 87
プログラミング言語 C C++ Python Fortran その他 CUDA C CUDA C++(C++11),Thrust PyCUDA CUDA Fortran F#, MATLAB, Mathematica, 88
CUDA 開発 実行環境 ライブラリ ミドルウェア 開発環境 cufft cublas cusparse cudnn cusolver curand Thrust NPP NVRTC MATLAB Mathematica etc.. NVCC (CUDA compiler) プログラミング言語 CUDA-GDB C C++ C++11 Fortran Java Python etc.. Profiler NVIDIA-GPUs( ハードウェア ) Nsight IDE HyperQ Dynamic Parallelism GPU Direct 90
CUDA ライブラリ cudnn ディープニューラルネットワーク計算用ライブラリ cusolver 線形代数演算 LAPACK 用ライブラリ curand 乱数生成ライブラリ cusparse 疎行列計算用ライブラリ cufft 高速フーリエ変換ライブラリ cublas 線形代数計算用ライブラリ NPP 動画像処理 信号処理用ライブラリ Thrust C++ テンプレートライブラリ 91
CUDA を使用したソフトウェア MATLAB Mathematica ArrayFire OpenCV etc Caffe torch theano 92
CUDA 開発 実行環境 ライブラリ ミドルウェア 開発環境 cufft cublas cusparse cudnn cusolver curand Thrust NPP NVRTC MATLAB Mathematica etc.. NVCC (CUDA compiler) プログラミング言語 CUDA-GDB C C++ C++11 Fortran Java Python etc.. Profiler NVIDIA-GPUs( ハードウェア ) Nsight IDE HyperQ Dynamic Parallelism GPU Direct 94
開発環境 DEBUG & ANALYSIS NVCC CUDA-GDB CUDA-MEMCHECK Nsight IDE Profiler CUDA 用コンパイラ CUDA 用デバッガ (Linux,Mac) GPUメモリエラーチェックツール CUDA 統合開発環境 (Linux,Windows) CUDA 解析ツール 95
NSIGHT VISUAL STUDIO EDITION 96
ここまでの復習 CUDAでは 様々なプログラミング言語やライブラリを使う事が可能 ケースによって最適なものを選択すれば良い 既存のライブラリやミドルウェアを有効活用する CUDAはロードマップが存在し 進化し続けている よりプログラミングしやすく パフォーマンスが出やすいように 97
CUDA C/C++ アプリケーション入門 今回は CUDA C/C++ で説明します 98
典型的な装置構成 PC GPU CPU につながった外部演算装置 CPU ( 数コア ) 制御 PCIe Giga Thread Engine SM SM SM SM L2 Cache ホスト側 DRAM 転送 DRAM 99
典型的な実行例 CPU プログラム開始 GPU は CPU からの制御で動作する データ転送 CUDA カーネル実行 完了待ち データ転送 入力データは CPU GPU へと転送 GPU 結果は GPU CPU と転送 GPU での演算 GPU 上に常駐するプログラムはない 100
CUDA C/C++ 用語 GPU で実行される関数をカーネル (kernel) と呼ぶ CPU で実行されるコードをホストコード GPU で実行されるコードをデバイスコードと呼ぶ データ並列を表現する為に以下の概念を用いる グリッド (grid) ブロック (block) スレッド (thread) 101
グリッド ブロック スレッド グリッド (grid) ブロックをまとめた物 ブロック (block) スレッドをまとめた物 1ブロックあたり最大 1024スレッド スレッド (thread) カーネルを動作させる最小単位 Block0 Thread Block1 Thread Grid Block2 Thread Block Thread n 103
グリッド ブロック スレッド CUDA GPU Block0 Thread SM GPU SM Block1 Thread core Grid Block2 Thread SM SM Block n Thread 105
カーネル実行の流れ Giga Thread Engine がブロックを SM に割り当てる Grid Block0 Giga Thread Engine Block1 Block2 Block3 Block4 BlockN Block4 107
カーネル実行の流れ SM の中のスケジューラがコアにスレッドを投入する Grid Block 0 Thread ワープを投入 32スレッド単位で投入 Thread Thread Thread BlockN 108
Block 1 SM BLOCK は SM 上で実行 複数の SM にまたがらない (SM 中では 複数 Block が実行される場合もある ) Block 内部では SMX のリソースを活用可能 各々の Block は 独立に 非同期に処理を実行する 実行順序の保証はない Block 間の通信 同期は行わない 109
例 : 一次元配列の加算 配列 A と配列 B の加算結果を配列 C に書き込む [0] [1] [2] [3] [4] [5] [6] [7] [8] [9] [10] [11] [12] [13] [14] [15] A 10 1 8 7 14 13 2 5 6 15 3 9 12 11 0 4 + + + + + + + + + + + + + + + + B 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 = = = = = = = = = = = = = = = = C 10 2 10 10 18 18 8 12 14 24 13 20 24 24 14 19 110
CPU 例 : 一次元配列の加算 (CPU) 配列の 0 番から逐次加算していく C[0] = A[0] + B[0]; C[1] = A[1] + B[1]; C[2] = A[2] + B[2]; for(int i=0 C[3] ; = i<nmatrixsize A[3] + B[3]; ; i++) { C[i] C[4] = = A[i] A[4] + B[i]; + B[4]; C[5] = A[5] + B[5]; 111
例 : 一次元配列の加算 (GPU) [0] [1] [2] [3] [4] [5] [6] [7] [8] [9] [10] [11] [12] [13] [14] [15] A 10 1 8 7 14 13 2 5 6 15 3 9 12 11 0 4 + + + + + + + + + + + + + + + + B 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 = = = = = = = = = = = = = = = = C 10 2 10 10 18 18 8 12 14 24 13 20 24 24 14 19 T0 T1 T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 T13 Block0 Block1 Block2 BlockN 112
ブロック ID とスレッド ID ブロック ID とスレッド ID から インデックス ( グローバル ID) を生成する インデックスを用いて各スレッドから グローバルメモリへアクセスする index = blockdim.x * blockidx.x + threadidx.x; 8 6 1 + 2 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 Thread 0 1 2 3 4 5 0 1 2 3 4 5 0 1 2 3 4 5 Block 0 1 2 113
例 : 一次元配列の加算 (GPU) ホスト側 (CPU) int main(int argc,char** argv){ MatrixAdd<<<N, M>>>(C,A,B); デバイス側 (GPU) global void MatrixAdd(float* C,const float* A,const float* B){ int i = blockdim.x * blockidx.x + threadidx.x; C[i] = A[i] + B[i]; 114
GPU 側メモリの確保 復習 : 典型的な実行例 CPU GPU は CPU からの制御で動作する データ転送 CUDA カーネル実行 完了待ち データ転送 入力データは CPU GPU へと転送 GPU 結果は GPU CPU と転送 GPU での演算 GPU 上に常駐するプログラムはない 115
ホスト側から呼び出す API cudamalloc GPU 上の DRAM( グローバルメモリ ) にメモリの確保を行う cudafree cudamalloc で取得したメモリの解放を行う cudamemcpy CPU->GPU GPU->GPU GPU->CPU のメモリ転送を行う cudadevicesynchronize CUDA カーネルが終了するまで待つ 116
cudamemcpy() メモリは ホスト デバイス の二種類 enum cudamemcpykind cudamemcpyhosttodevice cudamemcpydevicetohost cudamemcpydevicetodevice cudamemcpyhosttohost (cudamemcpydefault : UVA) 117
int main() { 略 サンプルコード ( ホスト ) int matrixsize= 256 * 100; float *A, *B, *C; cudamalloc(&a,sizeof(float)*matrixsize); cudamalloc(&b,sizeof(float)*matrixsize); cudamalloc(&c,sizeof(float)*matrixsize); cudamemcpy(a,ha, sizeof(float)*matrixsize, cudamemcpyhosttodevice); cudamemcpy(b,hb, sizeof(float)*matrixsize, cudamemcpyhosttodevice); MatrixAdd<<<matrixSize/256, 256>>>(C, A, B, matrixsize); cudadevicesynchronize(); cudamemcpy(hc, C, sizeof(float)*matrixsize, cudamemcpydevicetohost); cudafree(a); cudafree(b); cudafree(c); 略 118
サンプルコード ( デバイス ) global void MatrixAdd(float* C,const float* A,const float* B,const int size){ int i = blockdim.x * blockidx.x + threadidx.x; if( i < size){ C[i] = A[i] + B[i]; 119
例 :RGB->YUV 変換を考える Y U V = 0.299 0.587 0.114 0.169 0.331 0.500 0.500 0.419 0.081 R G B 121
1. GPU のメモリ構造 最適化の為に理解する事 2. スレッド (thread) 構成と占有率 (Occupancy) 122
1. GPU のメモリ構造 最適化の為に理解する事 2. スレッド (thread) 構成と占有率 (Occupancy) 123
GPU のメモリ階層 SM Threads アクセスが速い SMEM L1 Read TEX only L2 cache DRAM アクセスが遅い 124
Global Memory Local Memory GPU 上のメモリ キャッシュ レジスタ GPU 内部の記憶域 GPU 上の DRAM すべての SM からアクセス可能 Thread スコープのメモリ GPU 上の DRAM スレッド内部の配列 レジスタスピル時の記憶域 L2 Cache L1(Kepler のみ ) L2 Shared Memory SM 内部のメモリ Blockスコープでアクセス なし 手動管理のキャッ低レイテンシのRead/Write シュとして用いる場合あスレッド間のデータ共有り Texture Memory テクスチャユニット経由でアクセスするメモリ L1(Texture) L2 Read-only Data Cache Read Only でアクセスできる Global Memory L1(Texture) L2 Constant Memory 定数を収めるメモリ ブロードキャストアクセスに特化 Registers SM 内部のレジスタ 演算可能 なし SM 内部のキャッシュ 125
READ-ONLY(RO) CACHE SM Threads TEX Texture API SMEM L1 Read TEX only CUDA Arrays 一般的な Read-Only キャッシュとして使用可能 L2 cache Kepler 以降 コンパイラに指示 DRAM 126 12
RO DATA CACHE 使い方 型修飾子 : const restrict を付ける global kernel( int* output, const int* restrict input ) input ) {... output[idx] =... + input[idx + delta] +...;... 127
GLOBAL MEMORY SM SMEM Threads L1 Read TEX only GPU 上のメモリの中で最もポピュラーなメモリ メモリサイズは大きく アクセスコストは高い L2 cache Global DRAM Memory 128
コアレスアクセス 連続するスレッドは連続するメモリアクセスになるようにデバイスメモリは32byte,64byte,128byteの単位でロード ストア CUDA7.0 現在 thread0 thread1 thread2 thread3 thread4 thread5 128 160 192 Device Memory 129
コアレスアクセス 連続するスレッドは連続するメモリアクセスになるようにデバイスメモリは32byte,64byte,128byteの単位でロード ストア CUDA7.0 現在 thread0 thread1 thread2 thread3 thread4 thread5 128 160 192 Device Memory 130
コアレスアクセス 連続するスレッドは連続するメモリアクセスになるようにデバイスメモリは32byte,64byte,128byteの単位でロード ストア CUDA7.0 現在 thread0 thread1 thread2 thread3 thread4 thread5 128 160 192 Device Memory 131
height パディングを考慮したメモリの確保 x 方向の先頭アドレスが 32byte の倍数になるようにパディング 例 : RGB 24byte padding = 32 (3*width%32) width padding 132
2 次元メモリ確保 転送 API cudamallocpitch width バイトのメモリを height 行分 取得する 行は パディングを考慮した pitch バイトで整列する cudamemcpy2d cudamallocpitch で取得したパディングを考慮したメモリ (Dst) に Src のメモリ ( パディングなし ) をコピーする 133
サンプルコード uchar4 *src, *dimage; size_t spitch, dpitch, dpitchinpixel; // ピッチつきで メモリをアロケート cudamallocpitch(&dimage, &dpitch, width * sizeof(uchar4), height); dpitchinpixel = dpitch / sizeof(uchar4); // ピッチを変換しつつ ホスト デバイスへと メモリ転送 cudamemcpy2d(dimage, dpitch, src, spitch, width * sizeof(uchar4), height,cudamemcpyhosttodevice); 134
1. GPU のメモリ構造 最適化の為に理解する事 2. スレッド (thread) 構成と占有率 (Occupancy) 135
復習 : 一次元配列の加算 global void MatrixAdd(float *A, const float *B,const float *C) { グローバルID int i = threadidx.x + blodkdim.x * blockidx.x; if ( i >= N j >=N ) return; C[i][i] = A[i][j] + B[i][j]; 総スレッド数 1ブロックあたりのスレッド数... MatrixAdd<<< N/128, 128>>>(A, B, C);... 136
復習 : ブロック ID とスレッド ID ブロック ID とスレッド ID から インデックス ( グローバル ID) を生成する インデックスを用いて各スレッドから グローバルメモリへアクセスする index = blockdim.x * blockidx.x + threadidx.x; 8 6 1 + 2 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 Thread 0 1 2 3 4 5 0 1 2 3 4 5 0 1 2 3 4 5 Block 0 1 2 137
ブロック ID とスレッド ID( 二次元 ) BLOCK (M,0) BLOCK (M,N-2) BLOCK (M,N-1) BLOCK (M,N) index_x = blockdim.x * blockidx.x + threadidx.x; Index_y = blockdim.y * blockidx.y + threadidx.y; thread BLOCK (15,0) (M-1,N) thread (15,15) BLOCK (1,0) BLOCK (1,1) thread (1,0) thread (1,1) BLOCK (0,0) BLOCK (0,1) BLOCK (0,2) thread (0,0) BLOCK thread (0,1) (0,N) thread (0,15) 138
二次元配列の加算 global void MatrixAdd(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 ) return; C[i][i] = A[i][j] + B[i][j]; 1ブロックあたり16*16=256スレッド... dim3 sizeblock( 16, 16 ); dim3 numblocks( N/sizeBlock.x, N/sizeBlock.y ); MatrixAdd<<< numblocks, sizeblock >>>(A, B, C);... 139
例 :RGB->YUV 変換を考える 1スレッドで1pixelぶんの処理を行うピクセルの数だけスレッドを作成例 ) 1920*1080 = 2,073,600 スレッド 3840*2160 = 8,294,400 スレッド 140
例 :RGB->YUV 変換を考える thread7 thread6 thread5 thread4 thread3 thread2 thread1 int x = blockdim.x * blockidx.x + threadidx.x; int y = blockdim.y * blockidx.y + threadidx.y; if ((x < w) && (y < h)) { //Global Memory(Src) から 4byte ロード uchar4 urgb = gsrc[index]; //Global Memory(Dst) へ変換後の値を 4byte ストア gdst[idx] = RGB2YUV(uRGB.x, urgb.y, urgb.z); Height thread0 Width 141
ブロックサイズの決定 x = BlockDim.x * BlockIdx.x + threadidx.x (0<= x < width) y = BlockDim.x * BlockIdx.x + threadidx.x (0<= y < height) グリッド ブロックサイズの例 ) 960 threads / block 128 threads / block 32 threads / block? height width 142
ブロックサイズの決定 占有率を 100% にする ブロックサイズ ( ブロック辺りのスレッドの数 ) は少ない方が良い ブロックは横長の方が良い 143
占有率 (OCCUPANCY) とは? マルチプロセッサで同時に実行されるワープの数を同時に実行できるワープの最大数で除算したもの 144
BLOCKDIM の決定 ( 占有率から ) 項目 値 最大のBlock 数 / SMX 16 最大のThread 数 / SMX 2048 最大のThread 数 / Block 1024 SMX あたり 2048 Thread 走らせたい Occupancy ( 占有率 ) = 100 % Occupancy = 100 % を満たす Block あたりのスレッド数は 2048 Thread / 16 Block = 128 Thread / Block 2048 Thread / 8 Block = 256 Thread / Block 2048 Thread / 4 Block = 512 Thread / Block 2048 Thread / 2 Block = 1024 Thread / Block 145
BLOCKDIM の決め方 (BLOCK の粒度から ) Grid = 4096 Thread の実行例を考えてみる Block : 256 Thread 1024 Thread で比較 3 SMX / GPU 1 SMX あたり 1 Block が実行可能とする SMX 0 SMX 1 SMX 2 Block Block Block Block Block Block Block Block Block Block Block Block 256 Thread / Block Block Block Block Block t SMX 0 SMX 1 SMX 2 Block Block Block 1024 Thread / Block Block Block サイズは小さいほうが得 128 Threads / Block 146 t
復習 : カーネル実行の流れ Giga Thread Engine がブロックを SM に割り当てる Grid Block0 Giga Thread Engine Block1 Block2 Block3 Block4 BlockN Block4 148
復習 : カーネル実行の流れ SM の中のスケジューラがコアにスレッドを投入する Grid Block 0 Thread 32スレッド単位でワープを投入投入 Thread Thread Thread BlockN 149
Block Warp 32 GPU Thread CUDA cores Warp ワープ (WARP) : 並列実行の最少単位 - ワープ (Warp) : 32 GPU スレッド 1 命令を Warp (32 スレッド ) が 並列に処理 SIMT (Single Instruction Multiple Thread) SW SMX Warp Thread Thread Thread Thread Thread Thread Core Core Core Core Core 1 命令を 32 並列実行 150
BLOCKDIM の決め方 (SMX の構造から ) Warp Scheduler x 4 : 1 clock あたり 4 Warp に対する命令発行 Block のサイズは 128 Thread の倍数が望ましい (128 Thread = 32 Thread/Warp x 4 Warp) 152
タイルは横長がよい タイルの横幅は 32(Warp の幅 ) の倍数がよい 32 より小さい場合 16 もしくは 8 を使う Thread : 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 Memory : threadidx.x 153
blockdim.y RGB Y 変換時のバンド幅 : TESLA K20C blockdim.x 1 2 4 8 16 32 64 128 256 512 1024 1 1.4 2.8 5.6 11.2 22.1 43.9 78.5 119.8 119.3 115.4 87.7 2 2.6 5.2 10.4 Occupancy < 100 % 20.6 40.7 77.9 119.8 119.4 115.3 87.4-4 4.8 9.6 19.2 37.8 74.0 119.4 118.2 114.2 87.3 - - 8 8.4 16.7 33.3 69.6 115.0 117.9 111.9 87.1 - - - 16 13.4 26.3 60.6 106.7 115.0 114.3 87.2 - - - - 32 17.7 40.4 81.1 103.9 110.9 86.9 - - - - - 64 20.7 41.7 79.8 99.0 83.5 - - - - - - 128 20.7 41.6 75.6 75.3 - - - - - - - 256 20.7 41.0 60.3 - - - - - - - - 512 20.5 37.6 - - - - - - - - - 1024 19.1 - - - - - - - - - - blockdim.x < 8 値 : バンド幅 (GB/sec) Tesla K20c (ECC off) 154
RGB->YUV 変換 ( ホスト ) /* value radix で割って 切り上げる */ int divroundup(int value, int radix) { return (value + radix 1) / radix; /* griddim, blockdim を 2 次元 (x, y 方向 ) に初期化 */ dim3 blockdim(128, 1); /* divroundup() は 切り上げの割り算 */ dim3 griddim(divroundup(width, blockdim.x),divroundup(height, blockdim.y)); RGB2YUV<<<gridDim, blockdim>>>(ddst, dsrc, ); 155
RGB->YUV 変換 ( デバイス ) device inline uchar4 rgb_2_yuv(unsigned char R, unsigned char G, unsigned char B){ float fy,fu,fy; unsigned char uy,uu,uv; fy = 0.299f * value.x + 0.587f * value.y + 0.114f * value.z; uy = (unsigned char)min(255, (int)y); U と Y の処理は省略 make_uchar4(uy, uu, uv, 0); global void RGB2YUV (uchar4 *gdst, const uchar4 *gsrc, int w, int h){ int x = blockdim.x * blockidx.x + threadidx.x; int y = blockdim.y * blockidx.y + threadidx.y; if ((x < w) && (y < h)) { int index = y * width + x; //Global Memory(Src) から 4byte ロード uchar4 urgb = gsrc[index]; //Global Memory(Dst) へ変換後の値を 4byte ストア gdst[idx] = rgb_2_yuv(urgb.x, urgb.y, urgb.z); 156
まとめ グローバルメモリはコアレスアクセスする 二次元の場合は cudamallocpitch を使う事でメモリアライメントを考慮したメモリ確保が可能 メモリの Load のみの場合は Read Only Data Cache を活用 占有率 (Occupancy) と Block 内のスレッド構成を意識 Block サイズは 128 が適当 ( 単純なカーネルの場合 ) Block の横幅は 32 の倍数 無理な場合 16, 8 を選択 (4 byte / pixel の場合 ) 157
158
159
160
Appendix. CUDA ダウンロードサイト https://developer.nvidia.com/cuda-toolkit OpenACC toolkit https://developer.nvidia.com/openacc OpenACC オンライン講座 http://info.nvidianews.com/gettingstartedwithpgiopenacccompiler_reg- Landing-Page.html GPU コンピューティング Facebook ページ https://www.facebook.com/nvidiagpucomputing 161
Thankyou 162
Thank you 173