Slide 1

Transcription

1 GPUコンピューティング入門 エヌビディア合同会社 CUDAエンジニア 村上真奈

2 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 2

3 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 3

4 年創立 共同創業者 社長兼 CEO : ジェンスン フアン 1999 年 銘柄コード NVDA で NASDAQ に株式上場 1999 年に GPU を発明し 現在までに10 億点以上を出荷 2015 年会計年度の収益 : 46.8 億ドル世界中に 9,300 名の従業員 7,300 件の特許取得済み資産本社 : カリフォルニア州サンタクララ 4

5 3 ゲーミング オートモーティブ エンタープライズ HPC クラウド 当社は ビジュアル コンピューティングが必要不可欠で重要な価値を持つ大規模な市場に特化し プロセッサのプラット フォーム ソフトウェア システム サービスを提供しています 当社はPCテクノロジ データセンター テクノロジ モバイル テ クノロジの革新に取り組んでいます そして 当社の発明は さまざまな業界のOEM製品の原動力となっています 5

6 6

7 東京工業大学 TSUBAME 2.5 4,224 枚の Tesla K20X 単精度理論性能値で日本 No.1 スパコン 17PFLOPS SP 7

8 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 E C 2.200GHz, TH Express-2, Intel Xeon Phi 31S1P (TFlop/s) Rpeak (TFlop/s) Power (kw) 3,120,000 33, , ,808 2 US DOE/SC/Oak Ridge National laboratory Titan-Cray XK7, Opt C 2.2GHz, NVIDIA K20x 560,640 17, , ,209 3 US DOE/NNSA/LLNL Sequoia - BlueGene/Q, Power BQC 16C 1.60 GHz, Custom 1,572,864 17, , ,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, , , ,432 8, , ,945 In GREEN500 the most energy efficient super computers, NVIDIA GPU drives 8 systems out of TOP 10. 8

9 REAL WORLD EXAMPLE Rendering 30-second Animation at Renault 9

10 Deep Learning における GPU の活用 Deep Learning に GPU を活用 Input Result % 26% 60 16% 12% 7% person dog chair GPU 対応した Deep Learning 用ツール Caffe Torch Theano Cuda-convnet cudnn cublas 10

11 SGEMM / W GPU ロードマップ Volta Pascal 24 Maxwell 12 0 Tesla Fermi Kepler

12 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 K 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 K 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 K TF 0.19 TF 8 GB 320 GB/s Gen 3 Server only 12

13 M6000 K6000 K5200 K4200 K2200 K620 K420 # CUDA Cores Single Precision 5.2 TFLOPs 3.1 TFLOPs 2.1 TFLOPs 1.3 TFLOPs 0.8 TFLOPs 0.3 TFLOPs PCIe Gen 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 x 2160 Max Displays 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

14 TFLOPS TFLOPS 前世代比 3 倍の性能 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 コア数 倍精度演算性能 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

15 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

16 2015 TEGRA X1 MOBILE SUPERCHIP 256-core Maxwell GPU 8-core 64-bit CPU 4Kp60 10-bit H.265/VP9 19

17 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

18 Advancements BRIDGING THE GAP Maxwell Tesla Fermi Kepler Tegra K1 Tegra X1 GEFORCE ARCHITECTURE Tegra 4 Tegra 3 MOBILE ARCHITECTURE 21

19 GFLOPS WORLD S 1 ST TERAFLOPS MOBILE PROCESSOR 1200 Tegra X1 (FP16) Tegra X Core i7 GPU GPU CPU 800 CPU FP16/INT Tegra K1 200 Tegra 2 Tegra 3 Tegra 4 0 TIME Note: 4790K Core i7, 4GHz, GPU 350 MHz

20 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 24

21 NVIDIA GPU の歴史 2010 Fermi 3 Billion Transistors 2012 Kepler 7 Billion Transistors GPU 統合シェーダ + CUDA 25

22 PCI Express GPU の構造 GPU Giga Thread Engine SM SM SM SM L2 Cache DRAM 26

23 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

24 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

25 COMPUTE CAPABILITY GPU コアアーキテクチャのバージョン CUDA GPUs : アーキテクチャは進化する 高効率の命令実行 省消費電力 29

26 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

27 GPU コンピューティングとは? GPUは何の略? Graphics Processing Unit 3DCG 等 画像データ処理の為のデバイス GPUによる汎用コンピューティングのこと 計算科学など様々な用途でGPUを利用する 31

28 ヘテロジニアス コンピューティング CPU 逐次処理に最適化 GPU Accelerator 並列処理に最適化 32

29 GPU アプリケーションの例 画像処理コンピュータビジョン医療画像防衛計算化学 気象金融工学バイオ数値解析 33

30 GPU アクセラレーションの実現方法 アプリケーション GPU ライブラリ OpenACC ディレクティブ CUDA ライブラリを呼び出すだけ簡単に高速化を実現 既存コードにディレクティブを挿入して高速化 重要なコードを CUDA で記述最も自由度が高い 34

31 簡単 GPU アクセラレーションの実現方法 ライブラリ ライブラリを呼び出すだけで 高速化が可能ライブラリとして提供されている機能のみ高速化が可能 OpenACC 既存の C 言語や Fortran のコードにディレクティブを挿入するだけで簡単に高速化 最適化はコンパイラが行う為 細かいチューニングを行う事は出来ない 高速化 CUDA 自由度が最も高く 細かいチューニングが可能 CUDA でのプログラミングを学ぶ必要がある 35

32 エヌビディアのGPUについて (20 分 ) GPUコンピューティングとは?(10 分 ) OpenACC 入門 (90 分 ) CUDA 入門 (90 分 ) Agenda 36

33 GPU アクセラレーションの実現方法 アプリケーション GPU ライブラリ OpenACC ディレクティブ CUDA ライブラリを呼び出すだけ簡単に高速化を実現 既存コードにディレクティブを挿入して高速化 重要なコードを CUDA で記述最も自由度が高い 37

34 OPENACC 標準的な GPU ディレクティブ シンプル : ディレクティブを挿入するのみ コードを変更する事なく高速化 オープン : OpenACC はマルチコアプロセッサで並列化を行う為のオープン標準 柔軟 : GPU ディレクティブは 高い並列性を保ちつつ同一コードで複数のアーキテクチャに対応可能 38

35 OpenACC メンバーとパートナー 39

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

37 簡単に高速 自動車金融生命科学 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

38 大学関係者の方は無償で使用可能に 簡単に始められる 下記のサイトから OpenACC toolkit をダウンロード PGI コンパイラ /MPI/CUDA など一式が簡単にインストール可能 42

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

40 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

41 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

42 OpenACC ディレクティブ構文 C/C++ #pragma acc 指示行 [ 節 [,] 節 ] ] { structured block Fortran!$acc 指示行 [ 節 [,] 節 ] ] { structured block!$acc end directive 46

43 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

44 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

45 [C tips]: restrict 修飾子 コンパイラに対して明示的に restrict 修飾子を指定 ポインタのエイリアシングを制限 例 ) float *restrict ptr OpenACC コンパイラに restrict 修飾子をつけ変数の独立性を伝える 独立性の保障がないとコンパイラは並列化を行う事が出来ない 49

46 例 :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

47 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

48 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

49 コンパイルオプション 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

50 簡単にコンパイル 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

51 簡単に実行 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 ms kernel 2 generated ms ms ms [CUDA memcpy HtoD] 31.48% 19, ms #pragma... acc loop ms gang, vector(128) ms /* ms blockidx.x [CUDA threadidx.x memcpy DtoH] */ 5.56% us saxpy(n, 3.0, 1 x, us y); us us saxpy_19_gpu... 55

52 例 : ヤコビ反復法 正しい値になるように反復計算を行う 隣接点の平均値で値を更新 連立一次方程式を解く為のオーソドックスな手法 例 : 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 A(i,j-1) 56

53 ヤコビ反復法 ( アルゴリズム ) 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

54 並列領域 (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

55 並列領域 (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

56 [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

57 並列領域 (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

58 データ転送 (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

59 コードの解析 実行状況を確認 ボトルネックはどの部分か? 実行時間の内訳を調べる 63

60 -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]

61 NVIDIA Visual Profiler (NVVP) を使用 65

62 NVVP による解析 : データ転送がボトルネック 1 cycle 利用率 : 低い GPU kernel GPU kernel 66

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

64 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

65 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

66 データ転送 (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

67 データ転送 (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

68 過剰なデータ転送 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

69 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

70 データ領域 (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

71 データ領域 (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

72 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

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

74 GPU アクセラレーションの実現方法 アプリケーション GPU ライブラリ OpenACC ディレクティブ CUDA ライブラリを呼び出すだけ簡単に高速化を実現 既存コードにディレクティブを挿入して高速化 重要なコードを CUDA で記述最も自由度が高い 78

75 CUDA とは? Compute Unified Device Architectureの略 NVIDIA GPU 上の汎用並列計算プラットフォーム Linux Windows MacOS X(+Android) で動作 現在 7.0が最新 7.5RCも公開中 79

76 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

77 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

78 進化するハードウェア NVIDIA-GPUS 84

79 進化するハードウェア NVIDIA-GPUS Hyper-Q Dynamic Parallelism GPU Direct 85

80 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

81 プログラミング言語 C C++ Python Fortran その他 CUDA C CUDA C++(C++11),Thrust PyCUDA CUDA Fortran F#, MATLAB, Mathematica, 88

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 90

83 CUDA ライブラリ cudnn ディープニューラルネットワーク計算用ライブラリ cusolver 線形代数演算 LAPACK 用ライブラリ curand 乱数生成ライブラリ cusparse 疎行列計算用ライブラリ cufft 高速フーリエ変換ライブラリ cublas 線形代数計算用ライブラリ NPP 動画像処理 信号処理用ライブラリ Thrust C++ テンプレートライブラリ 91

84 CUDA を使用したソフトウェア MATLAB Mathematica ArrayFire OpenCV etc Caffe torch theano 92

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 94

86 開発環境 DEBUG & ANALYSIS NVCC CUDA-GDB CUDA-MEMCHECK Nsight IDE Profiler CUDA 用コンパイラ CUDA 用デバッガ (Linux,Mac) GPUメモリエラーチェックツール CUDA 統合開発環境 (Linux,Windows) CUDA 解析ツール 95

87 NSIGHT VISUAL STUDIO EDITION 96

88 ここまでの復習 CUDAでは 様々なプログラミング言語やライブラリを使う事が可能 ケースによって最適なものを選択すれば良い 既存のライブラリやミドルウェアを有効活用する CUDAはロードマップが存在し 進化し続けている よりプログラミングしやすく パフォーマンスが出やすいように 97

89 CUDA C/C++ アプリケーション入門 今回は CUDA C/C++ で説明します 98

90 典型的な装置構成 PC GPU CPU につながった外部演算装置 CPU ( 数コア ) 制御 PCIe Giga Thread Engine SM SM SM SM L2 Cache ホスト側 DRAM 転送 DRAM 99

91 典型的な実行例 CPU プログラム開始 GPU は CPU からの制御で動作する データ転送 CUDA カーネル実行 完了待ち データ転送 入力データは CPU GPU へと転送 GPU 結果は GPU CPU と転送 GPU での演算 GPU 上に常駐するプログラムはない 100

92 CUDA C/C++ 用語 GPU で実行される関数をカーネル (kernel) と呼ぶ CPU で実行されるコードをホストコード GPU で実行されるコードをデバイスコードと呼ぶ データ並列を表現する為に以下の概念を用いる グリッド (grid) ブロック (block) スレッド (thread) 101

93 グリッド ブロック スレッド グリッド (grid) ブロックをまとめた物 ブロック (block) スレッドをまとめた物 1ブロックあたり最大 1024スレッド スレッド (thread) カーネルを動作させる最小単位 Block0 Thread Block1 Thread Grid Block2 Thread Block Thread n 103

94 グリッド ブロック スレッド CUDA GPU Block0 Thread SM GPU SM Block1 Thread core Grid Block2 Thread SM SM Block n Thread 105

95 カーネル実行の流れ Giga Thread Engine がブロックを SM に割り当てる Grid Block0 Giga Thread Engine Block1 Block2 Block3 Block4 BlockN Block4 107

96 カーネル実行の流れ SM の中のスケジューラがコアにスレッドを投入する Grid Block 0 Thread ワープを投入 32スレッド単位で投入 Thread Thread Thread BlockN 108

97 Block 1 SM BLOCK は SM 上で実行 複数の SM にまたがらない (SM 中では 複数 Block が実行される場合もある ) Block 内部では SMX のリソースを活用可能 各々の Block は 独立に 非同期に処理を実行する 実行順序の保証はない Block 間の通信 同期は行わない 109

98 例 : 一次元配列の加算 配列 A と配列 B の加算結果を配列 C に書き込む [0] [1] [2] [3] [4] [5] [6] [7] [8] [9] [10] [11] [12] [13] [14] [15] A B = = = = = = = = = = = = = = = = C

99 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

100 例 : 一次元配列の加算 (GPU) [0] [1] [2] [3] [4] [5] [6] [7] [8] [9] [10] [11] [12] [13] [14] [15] A B = = = = = = = = = = = = = = = = C T0 T1 T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 T13 Block0 Block1 Block2 BlockN 112

101 ブロック ID とスレッド ID ブロック ID とスレッド ID から インデックス ( グローバル ID) を生成する インデックスを用いて各スレッドから グローバルメモリへアクセスする index = blockdim.x * blockidx.x + threadidx.x; Thread Block

102 例 : 一次元配列の加算 (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

103 GPU 側メモリの確保 復習 : 典型的な実行例 CPU GPU は CPU からの制御で動作する データ転送 CUDA カーネル実行 完了待ち データ転送 入力データは CPU GPU へと転送 GPU 結果は GPU CPU と転送 GPU での演算 GPU 上に常駐するプログラムはない 115

104 ホスト側から呼び出す API cudamalloc GPU 上の DRAM( グローバルメモリ ) にメモリの確保を行う cudafree cudamalloc で取得したメモリの解放を行う cudamemcpy CPU->GPU GPU->GPU GPU->CPU のメモリ転送を行う cudadevicesynchronize CUDA カーネルが終了するまで待つ 116

105 cudamemcpy() メモリは ホスト デバイス の二種類 enum cudamemcpykind cudamemcpyhosttodevice cudamemcpydevicetohost cudamemcpydevicetodevice cudamemcpyhosttohost (cudamemcpydefault : UVA) 117

106 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

107 サンプルコード ( デバイス ) 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

108 例 :RGB->YUV 変換を考える Y U V = R G B 121

109 1. GPU のメモリ構造 最適化の為に理解する事 2. スレッド (thread) 構成と占有率 (Occupancy) 122

110 1. GPU のメモリ構造 最適化の為に理解する事 2. スレッド (thread) 構成と占有率 (Occupancy) 123

111 GPU のメモリ階層 SM Threads アクセスが速い SMEM L1 Read TEX only L2 cache DRAM アクセスが遅い 124

112 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

113 READ-ONLY(RO) CACHE SM Threads TEX Texture API SMEM L1 Read TEX only CUDA Arrays 一般的な Read-Only キャッシュとして使用可能 L2 cache Kepler 以降 コンパイラに指示 DRAM

114 RO DATA CACHE 使い方 型修飾子 : const restrict を付ける global kernel( int* output, const int* restrict input ) input ) {... output[idx] =... + input[idx + delta] +...;

115 GLOBAL MEMORY SM SMEM Threads L1 Read TEX only GPU 上のメモリの中で最もポピュラーなメモリ メモリサイズは大きく アクセスコストは高い L2 cache Global DRAM Memory 128

116 コアレスアクセス 連続するスレッドは連続するメモリアクセスになるようにデバイスメモリは32byte,64byte,128byteの単位でロード ストア CUDA7.0 現在 thread0 thread1 thread2 thread3 thread4 thread Device Memory 129

117 コアレスアクセス 連続するスレッドは連続するメモリアクセスになるようにデバイスメモリは32byte,64byte,128byteの単位でロード ストア CUDA7.0 現在 thread0 thread1 thread2 thread3 thread4 thread Device Memory 130

118 コアレスアクセス 連続するスレッドは連続するメモリアクセスになるようにデバイスメモリは32byte,64byte,128byteの単位でロード ストア CUDA7.0 現在 thread0 thread1 thread2 thread3 thread4 thread Device Memory 131

119 height パディングを考慮したメモリの確保 x 方向の先頭アドレスが 32byte の倍数になるようにパディング 例 : RGB 24byte padding = 32 (3*width%32) width padding 132

120 2 次元メモリ確保 転送 API cudamallocpitch width バイトのメモリを height 行分 取得する 行は パディングを考慮した pitch バイトで整列する cudamemcpy2d cudamallocpitch で取得したパディングを考慮したメモリ (Dst) に Src のメモリ ( パディングなし ) をコピーする 133

121 サンプルコード 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

122 1. GPU のメモリ構造 最適化の為に理解する事 2. スレッド (thread) 構成と占有率 (Occupancy) 135

123 復習 : 一次元配列の加算 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);

124 復習 : ブロック ID とスレッド ID ブロック ID とスレッド ID から インデックス ( グローバル ID) を生成する インデックスを用いて各スレッドから グローバルメモリへアクセスする index = blockdim.x * blockidx.x + threadidx.x; Thread Block

125 ブロック 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

126 二次元配列の加算 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);

127 例 :RGB->YUV 変換を考える 1スレッドで1pixelぶんの処理を行うピクセルの数だけスレッドを作成例 ) 1920*1080 = 2,073,600 スレッド 3840*2160 = 8,294,400 スレッド 140

128 例 :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

129 ブロックサイズの決定 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

130 ブロックサイズの決定 占有率を 100% にする ブロックサイズ ( ブロック辺りのスレッドの数 ) は少ない方が良い ブロックは横長の方が良い 143

131 占有率 (OCCUPANCY) とは? マルチプロセッサで同時に実行されるワープの数を同時に実行できるワープの最大数で除算したもの 144

132 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

133 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

134 復習 : カーネル実行の流れ Giga Thread Engine がブロックを SM に割り当てる Grid Block0 Giga Thread Engine Block1 Block2 Block3 Block4 BlockN Block4 148

135 復習 : カーネル実行の流れ SM の中のスケジューラがコアにスレッドを投入する Grid Block 0 Thread 32スレッド単位でワープを投入投入 Thread Thread Thread BlockN 149

136 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

137 BLOCKDIM の決め方 (SMX の構造から ) Warp Scheduler x 4 : 1 clock あたり 4 Warp に対する命令発行 Block のサイズは 128 Thread の倍数が望ましい (128 Thread = 32 Thread/Warp x 4 Warp) 152

138 タイルは横長がよい タイルの横幅は 32(Warp の幅 ) の倍数がよい 32 より小さい場合 16 もしくは 8 を使う Thread : Memory : threadidx.x 153

139 blockdim.y RGB Y 変換時のバンド幅 : TESLA K20C blockdim.x Occupancy < 100 % blockdim.x < 8 値 : バンド幅 (GB/sec) Tesla K20c (ECC off) 154

140 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

141 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 f * value.y f * 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

142 まとめ グローバルメモリはコアレスアクセスする 二次元の場合は cudamallocpitch を使う事でメモリアライメントを考慮したメモリ確保が可能 メモリの Load のみの場合は Read Only Data Cache を活用 占有率 (Occupancy) と Block 内のスレッド構成を意識 Block サイズは 128 が適当 ( 単純なカーネルの場合 ) Block の横幅は 32 の倍数 無理な場合 16, 8 を選択 (4 byte / pixel の場合 ) 157

143 158

144 159

145 160

146 Appendix. CUDA ダウンロードサイト OpenACC toolkit OpenACC オンライン講座 Landing-Page.html GPU コンピューティング Facebook ページ 161

147 Thankyou 162

148 Thank you 173