CUDA基礎１

Transcription

1 CUDA 基礎 1 東京工業大学 学術国際情報センター 黄遠雄 2016/6/27 第 20 回 GPU コンピューティング講習会 1

2 ヘテロジニアス コンピューティング ヘテロジニアス コンピューティング (CPU + GPU) は広く使われている Financial Analysis Scientific Simulation Engineering Simulation Data Intensive Analytics Medical Imaging Digital Audio Processing Digital Video Processing Computer Vision Biomedical Informatics Electronic Design Automation Statistical Modeling Numerical Methods Ray Tracing Rendering Interactive Physics 2016/6/27 第 20 回 GPU コンピューティング講習会 2

3 GPU を用いてアプリケーションを高速化する 3 つの方法 Applications Libraries Compiler Directives Programming Languages Easy to use Most Performance Easy to use Portable code Most Performance Most Flexibility 2016/6/27 第 20 回 GPU コンピューティング講習会 3

4 GPU Accelerated Libraries Linear Algebra FFT, BLAS, SPARSE, Matrix NVIDIA cufft, cublas, cusparse Numerical & Math RAND, Statistics NVIDIA Math Lib NVIDIA curand Data Struct. & AI Sort, Scan, Zero Sum GPU AI Board Games GPU AI Path Finding Visual Processing Image & Video NVIDIA NPP NVIDIA Video Encode 2016/6/27 第 20 回 GPU コンピューティング講習会 4

5 Compiler directives:openacc Compiler directives for C, C++, and FORTRAN #pragma acc parallel loop copyin(input1[0:inputlength],input2[0:inputlength]), copyout(output[0:inputlength]) for(i = 0; i < inputlength; ++i) { } output[i] = input1[i] + input2[i]; 2016/6/27 第 20 回 GPU コンピューティング講習会 5

6 GPU Programming Languages Numerical analytics MATLAB, Mathematica, LabVIEW Fortran CUDA Fortran C CUDA C C++ CUDA C++ Python F# PyCUDA, Copperhead, Numba, NumbaPro Alea.cuBase 2016/6/27 第 20 回 GPU コンピューティング講習会 6

7 CUDA - C Applications Libraries Compiler Directives Programming Languages Easy to use Most Performance Easy to use Portable code Most Performance Most Flexibility 2016/6/27 第 20 回 GPU コンピューティング講習会 7

8 CPU と GPU の違い CPU DRAM Control Cache ALU ALU ALU ALU Powerful ALU Reduced operation latency Large caches Convert long latency memory accesses to short latency cache accesses Sophisticated control Branch prediction for reduced branch latency Data forwarding for reduced data latency GPU DRAM Small caches To boost memory throughput Simple control No branch prediction No data forwarding Energy efficient ALUs Many, long latency but heavily pipelined for high throughput Require massive number of threads to tolerate latencies Threading logic Thread state 2016/6/27 第 20 回 GPU コンピューティング講習会 8

9 DeviceQuery >./devicequery >sh rundevicequery.sh./devicequery Starting... CUDA Device Query (Runtime API) version (CUDART static linking) Detected 3 CUDA Capable device(s) Device 0: "Tesla K20Xm" CUDA Driver Version / Runtime Version 7.5 / 7.5 CUDA Capability Major/Minor version number: 3.5 Total amount of global memory: 5760 MBytes ( bytes) (14) Multiprocessors, (192) CUDA Cores/MP: 2688 CUDA Cores GPU Max Clock rate: 732 MHz (0.73 GHz) Memory Clock rate: 2600 Mhz Memory Bus Width: 384-bit L2 Cache Size: bytes Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096) 2016/6/27 第 20 回 GPU コンピューティング講習会 9

10 並列処理 並列計算は タスク並列とデータ並列 に分けられる データ並列の一例 : ベクトル加法 vector A A[0] A[1] A[2] A[N-1] vector B B[0] B[1] B[2] B[N-1] vector C C[0] C[1] C[2] C[N-1] 2016/6/27 第 20 回 GPU コンピューティング講習会 10

11 ベクトル加法 (C Code) // ベクトル加法関数 C = A + B void vecadd(float *A_h, float *B_h, float *C_h, int n) { int i; for (i = 0; i<n; i++) C_h[i] = A_h[i] + B_h[i]; } int main() { // A B C メモリ確保 // A B データ入力 / 初期化 (N 個 ) vecadd(a_h, B_h, C_h, N); } 2016/6/27 第 20 回 GPU コンピューティング講習会 11

12 GPU を用いたベクトル加法 GPU は単独では動かない host(cpu) で実行させ その中から CUDA API と GPU kernel 関数を call void vecadd(float *A_h, float *B_h, float *C_h, ( n int { int size = n* sizeof(float); float *A_d, *B_d, *C_d; // Part 1 // A B C GPU 上のメモリ確保 // A B のデータを host memory から device // memory にコピー Part 1 Part 2 // Part 2 // Kernel code を実行 ( ベクトルの加法 ) Host Memory CPU Device Memory GPU Part 3 } // Part 3 // 計算結果を回収 C のデータを device // memory から host memory にコピー // GPU 上のメモリを解放 2016/6/27 第 20 回 GPU コンピューティング講習会 12

13 CUDA プログラム実行の概念図 GPU device memory host (CPU) code メモリ ポインタ float *f_d, *f_h; CPU CUDA API cudamalloc(&f_d); host memory device code global func(f_d) { } cudamemcpy(f_h, f_d);.... kernel function func<<< N/256, 256>>>(f_d) 2016/6/27 第 20 回 GPU コンピューティング講習会 13

14 CUDA プログラムの実行モデル ヘテロジニアス (CPU + GPU) アプリケーション 逐次処理は Host C Code 並列計算は Device SPMD kernel Code ( host ) Serial Code ( device ) Parallel Kernel KernelA<<< Dg, Db >>>(args);... ( host ) Serial Code ( device ) Parallel Kernel KernelB<<< Dg, Db >>>(args); /6/27 第 20 回 GPU コンピューティング講習会 14

15 CUDA ソースコードのコンパイル CUDA のソースファイルは拡張子.cu を付ける CUDA Toolkit の nvcc でコンパイルする nvcc は CPU で実行するコードと GPU で実行する GPU kernel 関数のコード CUDA の API の部分を分離 CPU で実行するコードは gcc, g++ などにコンパイルを任せる GPU kernel 関数の部分を GPU 用にコンパイルする GPU 用の PTX コードも生成する Library をリンクして 実行ファイルを生成する CUDA core library (cuda) -lcuda CUDA runtime library (cudart) -lcudart Integrated C programs with CUDA extensions NVCC Compiler Host Code Device Code (PTX) Host C Compiler/ Linker Device Just-in-Time Compiler Heterogeneous Computing Platform with CPUs, GPUs, etc. 2016/6/27 第 20 回 GPU コンピューティング講習会 15

16 CUDA Compiler: nvcc 重要なコンパイル オプション -arch sm_52 --maxrregcount <N> Compute Capability に応じたコンパイルを行う DeviceQuery で確認し それ以下を指定す 可能なオプション :sm_20(default), sm_21, sm_30, sm_32, sm_35, sm_50 and sm_52 1 つの kernel 関数当たりに使用するレジスタ数を <N> に制限する このことにより 指定した並列数で thread が実行可能となるが 溢れた部分は local メモリ上に置かれ 実行速度は低下する -use_fast_math 高速な数学関数を利用する -G device コードに対して デバッグを可能にする --ptxas-options=-v レジスタやメモリの使用状況を表示する 2016/6/27 第 20 回 GPU コンピューティング講習会 16

17 CUDA Memory 確保 (1/3) メモリ ポインタ - は device (GPU) memory にも host (CPU) memory にも使える 例 ) 単精度実数 : float *f_d, *f_h; device 上にメモリを確保する runtime API cudamalloc(void **devptr, size_t count); devptr: count: デバイスメモリアドレスへのポインタ 確保したメモリのアドレスが書き込まれる 領域のサイズ 例 ) cudamalloc((void **)&f_d, sizeof(float)*n); f_d[n] の配列が GPU のメモリ上に確保される 2016/6/27 第 20 回 GPU コンピューティング講習会 17

18 CUDA Memory 確保 (2/3) host 側にメモリを確保する ( 通常 ) f_h = (float *) malloc(sizeof(float)*n); f_h = new float[n]; host 側に pinned メモリを確保する (C++) cudamallochost(void **ptr, size_t count); ptr: count: ホストメモリアドレスへのポインタ Page lock (pinned) された確保したメモリのアドレスが書き込まれる 領域のサイズ 例 ) cudamallochost((void **)&f_h, sizeof(float)*n); f_h[n] の配列が Host メモリ上に page lock (pinned) で確保される 通常の pageable メモリとして確保された場合より 転送速度が速い また 非同期通信の場合も page lock メモリに限定される 2016/6/27 第 20 回 GPU コンピューティング講習会 18

19 CUDA Memory 確保 (3/3) host 側にメモリを確保する ( 通常 ) f_h = (float *) malloc(sizeof(float)*n); f_h = new float[n]; (C++) host 側に確保したメモリを page-lock(pinned) にする cudahostregister(void *ptr, size_t count, unsigned int flags); ptr: count: flags: ホストメモリアドレスへのポインタ 領域のサイズ タイプを指定する定数 cudahostregisterdefault cudahostregisterportable cudahostregistermapped cudahostregisteriomemory 解除するには cudahostunregister(void *ptr); 2016/6/27 第 20 回 GPU コンピューティング講習会 19

20 CUDA データ転送 float *f_d, *f_h; cudamemcpy(void *dst, const void *src, size_t count, enum cudamemcpykind kind) dst: src: count: kind: 転送先メモリ アドレス転送元メモリ アドレス領域のサイズ転送タイプを指定する定数 cudamemcpyhosttodevice cudamemcpydevicetohost cudamemcpydevicetodevice cudamemcpydefault (Fermi GPU, CUDA 4.0 以降 ) 例 ) cudamemcpy (f_d, f_h, sizeof(float)*n, cudamemcpyhosttodevice); host 上の f_h[n] の配列のデータを device 上の f_d[n] にコピーする 2016/6/27 第 20 回 GPU コンピューティング講習会 20

21 CUDA 非同期データ転送 float *f_d, *f_h; cudamemcpyasync(void *dst, const void *src, size_t count, enum cudamemcpykind kind, cudastream_t stream) dst: src: count: kind: 転送先メモリ アドレス転送元メモリ アドレス領域のサイズ 転送タイプを指定する定数 cudamemcpyhosttohost cudamemcpyhosttodevice cudamemcpydevicetohost cudamemcpydevicetodevice 例 ) cudamemcpyasync (f_d, f_h, sizeof(float)*n, cudamemcpyhosttodevice, stream); Host 上の f_h[n] の配列のデータを Device 上の f_d[n] に Host や他の stream( 後述 ) に対して非同期でコピーする 2016/6/27 第 20 回 GPU コンピューティング講習会 21

22 GPU kernel-function call host code の中で次のように call する kernel_function<<< Dg, Db, Ns, S>>>(a, b, c,....); Dg: dim3 タイプの grid のサイズ指定 Db: dim3 タイプの block のサイズ指定 Ns: 実行時に指定する shared メモリのサイズ省略可 : 省略した場合は 0 が設定 S: 非同期実行の stream 番号省略可 : 省略した場合は 0 が設定され 同じ 0 に設定された GPU の kernel が同期実行となる Dg, Db で指定される数の thread が実行される kernel function の実行は CPU に対して絶えず非同期 2016/6/27 第 20 回 GPU コンピューティング講習会 22

23 dim3 宣言 kernel_function<<< Dg, Db, Ns, S>>>(a, b, c,....); の Dg, Db を dim3 で指定する dim3 a; dim3 a(n, m); dim3 a(n, m, k); 等価 等価 等価 dim3 a(1,1,1); dim3 a(n, m, 1); a.x = n; a.y = m; a.z = k; dim3 a(n0, m0, k0); は宣言と共に値の代入であり 随時 a.x = n1; a.y = m1; a.z = k1; と変更可能である 2016/6/27 第 20 回 GPU コンピューティング講習会 23

24 Thread の管理 CUDA Kernel は threads のまとまり (Grid) 単位で実行される Grid の中の全ての threads は同じ Kernel を実行する (Single Program Multiple Data) Thread は各自に ID の持っている i = blockidx.x * blockdim.x + threadidx.x; C[i] = A[i] + B[i]; 2016/6/27 第 20 回 GPU コンピューティング講習会 24

25 Thread の管理 (Block) Thread Block 0 Thread Block 1 Thread Block N i = blockidx.x * blockdim.x + threadidx.x; C[i] = A[i] + B[i]; i = blockidx.x * blockdim.x + threadidx.x; C[i] = A[i] + B[i]; i = blockidx.x * blockdim.x + threadidx.x; C[i] = A[i] + B[i]; いくつかの threads が一つの Block としてまとめられ 全て threads が複数の Block に分割されます 同一 Block 内の threads は shared memory を共有し atomic operations と barrier synchronization など同期が必要な演算を実行できます 異なる Block 間の threads 同期や shared memory の共有はできません 2016/6/27 第 20 回 GPU コンピューティング講習会 25

26 blockidx と threadidx kernel 関数 <<< 第 1 引数, 第 2 引数 >>> で指定 第 1 引数 :blockidx: 1D, 2D, or 3D (CUDA 4.0 以降 ) の範囲を指定 第 2 引数 :threadidx: 1D, 2D, or 3D の範囲を指定 多次元データを計算する場合にアドレス計算が簡単になる Image processing Solving PDEs on volumes device Grid Block (0, 0) Block (1, 0) Block (0, 1) Block (1, 1) Block (1,1) (1,0,0) (1,0,1) (1,0,2) (1,0,3) Thread (0,0,0) Thread (0,1,0) Thread (0,0,1) Thread (0,1,1) Thread (0,0,2) Thread (0,1,2) Thread (0,0,3) Thread Thread (0,0,0) (0,1,3) 2016/6/27 第 20 回 GPU コンピューティング講習会 26

27 実習 : 最初の CUDA プログラム #include <stdio.h> global void helloworld_kernel( void ) { printf("hello from GPU [thread %d of block %d] n", threadidx.x, blockidx.x); printf("good bye from GPU [thread %d of block %d] n", threadidx.x, blockidx.x); } int main(int argc, char *argv[]) { printf("hello from CPU n"); helloworld_kernel<<< 256, 1 >>>(); cudadevicesynchronize(); printf("good bye from CPU n"); } return 0; >nvcc HelloWorld.cu >./a.out Sample Code: simplecuda/helloworld 2016/6/27 第 20 回 GPU コンピューティング講習会 27

28 Built-in 変数 Device code の中で宣言せずに引用でき 書き換え不可 griddim blockidx blockdim threadidx griddim.x, griddim.y, griddim.z grid の各方向のサイズ blockidx.x, blockidx.y, blockidx.z block の各方向の index blockdim.x, blockdim.y, blockdim.z block の各方向のサイズ threadidx.x, threadidx.y, threadidx.z thread の各方向の index 2016/6/27 第 20 回 GPU コンピューティング講習会 28

29 C 言語の拡張 関数型の Qualifier () DeviceFunc device float () KernelFunc global void () HostFunc host float global device device 上でのみ実行される host 側からのみ call される return 値は void 限定 device 上でのみ実行される device からのみ call される Executed on the: device device host Only callable from the: device host host host host と device host 上でのみ実行される host 側からのみ call される ( 普通の CPU 上のプログラムの関数で 特に宣言する必要はない ) 両方同時に指定する事が可能 2016/6/27 第 20 回 GPU コンピューティング講習会 29

30 host と device host devcie void Cal_func(..) { C = A + B; } CPU for(..) { Cal_func( ); } GPU global kernel(.) { Cal_func( ); } 2016/6/27 第 20 回 GPU コンピューティング講習会 30

31 ベクトル加法 (Host Code) ( n void vecadd(float *A_h, float *B_h, float *C_h, int { int size = n* sizeof(float); float *A_d, *B_d, *C_d; // Part 1: A B C GPU 上のメモリ確保 cudamalloc( (void**) &A_d, size ); cudamalloc( (void**) &B_d, size ); cudamalloc( (void**) &C_d, size ); // Part 1: A B のデータを host memory から device memory にコピー cudamemcpy( A_d, A_h, size, cudamemcpyhosttodevice ); cudamemcpy( B_d, B_h, size, cudamemcpyhosttodevice ); // Part 2: Kernel code を実行 ( ベクトルの加法 ) vectoradd_kernel<<< n / BLOCK_SIZE, BLOCK_SIZE >>>( C_d, A_d, B_d ); } // Part 3: 計算結果を回収 C のデータを device memory から host memory にコピー cudamemcpy( c_h, c_d, size, cudamemcpydevicetohost ) // Part 3: GPU 上のメモリを解放 cudafree( A_d ); cudafree( B_d ); cudafree( C_d ); 2016/6/27 第 20 回 GPU コンピューティング講習会 31

32 ベクトル加法 (Device Kernel) vectoradd_kernel<<< n / BLOCK_SIZE, BLOCK_SIZE >>>( C_d, A_d, B_d ); blockdim の値と threadidx の値の範囲を決めます griddim の値と blockidx の値の範囲を決めます global void vectoradd_kernel ( float *C, // array pointer of the global memory float *A, // array pointer of the global memory float *B // array pointer of the global memory ) { unsigned int index = blockidx.x * blockdim.x + threadidx.x; } C[index] = A[index] + B[index]; Sample Code: simplecuda/simplevectoradd 2016/6/27 第 20 回 GPU コンピューティング講習会 32

33 2 次元データ アクセス NX*NY の 1 次元配列データであるが 2 次元的にアクセス dim3 grid(nx/16, NY/16), block(16, 16); iy blockdim.x = 16 grid サイズの最大値の制限から開放 NY ix ix = blockidx.x * blockdim.x + threadidx.x; iy = blockidx.x * blockdim.x + threadidx.x; index = iy * NX + ix; 2016/6/27 第 20 回 GPU コンピューティング講習会 NX

34 Warp スケジューリング Block 1 Warps Block 2 Warps Block 3 Warps t0 t1 t2 t31 t0 t1 t2 t31 t0 t1 t2 t31 block 内の thread は Streaming Multiprocessor によって Warp=32 thread 毎に並列実行される Warp: block 内の 32 thread のかたまり ( 例 :block 内に 256 thread = Warp 8 個 ) 1 Warp の thread の数は将来変わる可能性がある プログラム上には現れない 考慮しなくても正しいプログラムを書くことは可能 ただし 実行性能を引き出すためには考慮する必要がある 2016/6/27 第 20 回 GPU コンピューティング講習会 34

35 Warp 多次元分配 多次元 Thread blocks の場合 Warp は行優先で割り当てられる 最初 X それから Y 最後は Z 2016/6/27 第 20 回 GPU コンピューティング講習会 35

36 Compute Capability Compute Capability Technical Specifications 2.x Warp size 32 Maximum number of resident blocks per multiprocessor Maximum number of resident warps per multiprocessor Maximum number of resident threads per multiprocessor (= 64 x 32) Number of 32-bit registers per multiprocessor 32 K 64 K 128 K 64 K Maximum number of 32-bit registers per thread block 32 K 64 K 32 K Maximum number of 32-bit registers per thread Maximum amount of shared memory per multiprocessor Maximum amount of shared memory per thread block 48 KB 112 KB 64 KB 96 KB 64 KB 48 KB Number of shared memory banks 32 Amount of local memory per thread Constant memory size 512 KB 64 KB 2016/6/27 第 20 回 GPU コンピューティング講習会 36

37 同時実行可能な thread block 数 thread の使用する Register 数 : Nr Shared Memory 量 : Ns [byte] block 当たりの thread 数 : Db Compute Capability 5.2 (Maxwell 世代 GPU) の場合 : Warp per SM = Db/32 Active block = Min(32, 64/(Db/32), 98304/Ns, 65536/(Db*Nr) ) (block 当たりの最大 64 Warp) (Shared Memory の制限 ) (Register の制限 ) 2016/6/27 第 20 回 GPU コンピューティング講習会 37

38 Warp 内の thread 実行仕組み Warp 内の 32 thread は同一命令を実行 SIMD (Single Instruction Multiple Data) Warp 内の thread indices はインクリメント Warp 0 は thread 0 からスタート Warps 実行される順番は一定ではない 例えば Warp 8 が Warp 7 より先に実行される場合がある Threads の間の依存関係 ( 計算順序など ) は必ず syncthreads() を使う 2016/6/27 第 20 回 GPU コンピューティング講習会 38

39 Thread の中の条件分岐 プログラム上では任意の分岐を記述可能 ハードウェア上での分岐命令の処理 Warp 内全 thread が同一パスに分岐する場合は全 thread が分岐先 ( のみ ) を実行する Warp 内の thread が異なるパスに分岐する場合は全スレッドが両方の命令を実行 (diverged branch) し 最後に適合する方だけを採用する 性能低下の原因の一つ 2016/6/27 第 20 回 GPU コンピューティング講習会 39

40 Control Divergence Warp 内に条件分岐やループの違いがある場合 Control Divergence が発生する Control Divergence が発生しないようにするには Warp 内で 全ての if-then-else で同じ条件分岐とする 全てのループの反復回数を同じとする Control Divergence が発生する場合の例 : If ( threadidx.x > 2 ) { do_something(); } else{ do_other_thing(); } Thread ID によって違い操作 分岐単位 < warp size Control Divergence が発生しない場合の例 : If ( blockidx.x > 2 ) ) { do_something(); } else{ do_other_thing(); } block ID によって違い操作 分岐単位は block = warp size の倍数 2016/6/27 第 20 回 GPU コンピューティング講習会 40

41 例 : Vector Addition Kernel // Compute vector sum C = A + B Device Code // Each thread performs one pair-wise addition global void vectoradd_kernel(float* C, float* A, float* B, int n) { int i = threadidx.x + blockdim.x * blockidx.x; if(i<n) C[i] = A[i] + B[i]; } 2016/6/27 第 20 回 GPU コンピューティング講習会 41

42 Vector size (n=1,000) で計算すると BLOCK_SIZE = 256 に設定 8 warps per block Block 0, 1, 2 の全ての Threads は範囲内 i = 0 ~ 767 全て i < 1000 Block 3 の多くの Warps は Control Divergence が発生しない 6 warps は (i < 1000) の範囲内 Block 3 の中の一つの Warp は Control Divergence が発生する Thread ID 992 ~ 999 は範囲内 Thread ID 1000 ~ 1023 は範囲外 この例では Control Divergence の影響は少ない 1/32 warp は Control Divergence が発生する 性能に影響は 3% 以下 2016/6/27 第 20 回 GPU コンピューティング講習会 42

43 Memory Coalescing Coalesced Loads T 0 T 1 T 2 T 3 Coalesced Loads T 0 T 1 T 2 T Burst section Burst section Burst section Burst section 一つの warp の中の全ての threads が一斉にデータをロード 全てのデータアクセスが同じ burst section の場合 一回の DRAM request のみで終了する (fully coalesced) 2016/6/27 第 20 回 GPU コンピューティング講習会 43

44 Un-coalesced Accesses Un-coalesced Loads T 0 T 1 T 2 T 3 Un-coalesced Loads T 0 T 1 T 2 T Burst section Burst section Burst section Burst section データアクセスが別々の burst section になる場合 : Coalescing fails 何回かの DRAM requests が必要 アクセスは not fully coalesced 一部の転送されたデータは使われない 2016/6/27 第 20 回 GPU コンピューティング講習会 44

45 HEIGHT 例 : 行列の乗法のアクセスパターン M 0,0 M 0,1 M 0,2 M 0,3 M 1,0 M 1,1 M 1,2 M 1,3 M M 2,0 M 2,1 M 2,2 M 2,3 M 3,0 M 3,1 M 3,2 M 3,3 M 0,0 M 0,1 M 0,2 M 0,3 M 1,0 M 1,1 M 1,2 M 1,3 M 2,0 M 2,1 M 2,2 M 2,3 M 3,0 M 3,1 M 3,2 M 3,3 linearized order in increasing address A B T1 T2 (Thread) T1 (Thread) T2 i はループの変数 A は m n, B は n k のマトリックス Col = blockidx.x*blockdim.x + threadidx.x WIDTH A[Row*n+i] B[i*k+Col] 2016/6/27 第 20 回 GPU コンピューティング講習会 45

46 A Accesses are Not Coalesced Load iteration 1 T 0 T 1 T 2 T 3 Load iteration 0 T 0 T 1 T 2 T 3 A 0,0 A 0,1 A 0,2 A 0,3 A 1,1 A 1,0 A 1,2 A 1,3 A 2,0 A 2,1 A 2,2 A 2,3 A 3,0 A 3,1 A 3,2 A 3,3 Access direction in kernel code A 0,0 A 0,1 A 0,2 A 0,3 A 1,0 A 1,1 A 1,2 A 1,3 A 2,0 A 2,1 A 2,2 A 2,3 A 3,0 A 3,1 A 3,2 A 3,3 2016/6/27 第 20 回 GPU コンピューティング講習会 46

47 B accesses are coalesced N Load iteration 0 T 0 T 1 T 2 T 3 Load iteration 1 T 0 T 1 T 2 T 3 B 0,0 B 0,1 B 0,2 B 0,3 B 1,1 B 1,0 B 1,2 B 1,3 B 2,0 B 2,1 B 2,2 B 2,3 B 3,0 B 3,1 B 3,2 B 3,3 Access direction in kernel code B 0,0 B 0,1 B 0,2 B 0,3 B 1,0 B 1,1 B 1,2 B 1,3 B 2,0 B 2,1 B 2,2 B 2,3 B 3,0 B 3,1 B 3,2 B 3,3 2016/6/27 第 20 回 GPU コンピューティング講習会 47

48 所要時間計測 経過時間を計測することで GPU Computing のパフォーマンスを確認でき ハードウェア実行の様子を想像することができる また チューニングのためには必須 cudaevent_t start, stop; cudaeventcreate(&start); cudaeventcreate(&stop); float elapsedtime; cudaeventrecord(start,0); 計測範囲 kernel<<< grid, block>>>(a_d, b_d); elapsedtime に経過時間 (msec) cudaeventrecord(stop,0); cudaeventsynchronize(stop); cudaeventelapsedtime(&elapsedtime,start,stop); 2016/6/27 第 20 回 GPU コンピューティング講習会 48

49 エラー処理 (API) CUDA の API は全て return 値が cudaerror_t 型の error の status を返すようになっている cudaerror_t err = cudamemcpy( ); if (err!= cudasuccess) { fprintf(stderr, Memcopy failed: %s. n, cudageterrorstring(err)); } もし cudamalloc しないで cudamemcpy(); を実行してしまった場合などは invalid device pointer が返ってくる 2016/6/27 第 20 回 GPU コンピューティング講習会 49

50 エラー処理 ( kernel 関数 ) kernel 関数には return 値はない cudagetlasterror() で直前のエラーを拾い cudageterrorstring() でメッセージを表示させる vec_add<<<,,, >>>(... ); cudaerror_t err = cudagetlasterror(); if (err!= cudasuccess) { fprintf(stderr, kernel launch failed: %s n, cudageterrorstring(err)); exit(-1); } 例えば grid, block のサイズが最大値を超えていると invalid configuration argument が表示される 2016/6/27 第 20 回 GPU コンピューティング講習会 50

51 Device マネージメント API Device の情報を取得する API が準備されている cudagetdevicecount(int *count) システム ( ノード ) 上の CUDA の動作する GPU の個数を返す cudasetdevice(int device_no) それ以降の実行を device_no の GPU に向ける cudagetdevice(int *current_device) 現在指定されている GPU の device 番号を返す cudagetdeviceproperties(int *device, cudadeviceprop *prop) devicequery のような情報を prop のメンバーとして取得 Tips: これらは 同一ノード内に複数 GPU がある場合は必須 2016/6/27 第 20 回 GPU コンピューティング講習会 51

52 Developer Tools - Debuggers NSIGHT CUDA-GDB CUDA MEMCHECK NVIDIA Provided 3 rd Party /6/27 第 20 回 GPU コンピューティング講習会 52

53 Developer Tools - Profilers NSIGHT NVVP NVPROF NVIDIA Provided TAU VampirTrace 3 rd Party /6/27 第 20 回 GPU コンピューティング講習会 53

54 NVIDIA s Visual Profiler (NVVP) Timeline Guided System Analysis 2016/6/27 第 20 回 GPU コンピューティング講習会 54

55 IDE(NSIGHT) CUDA enabled Integrated Development Environment Source code editor: syntax highlighting, code refactoring, etc Build Manger Visual Debugger Visual Profiler Linux/Macintosh Editor = Eclipse Debugger = cuda-gdb with a visual wrapper Profiler = NVVP Windows Integrates directly into Visual Studio Profiler is NSIGHT VSE 2016/6/27 第 20 回 GPU コンピューティング講習会 55