CUDA基礎1 - 東京工業大学...CUDA基礎1 東京工業大学 学術国際情報センター 黄...

CUDA基礎１

東京工業大学 学術国際情報センター

黄 遠雄

1 2016/6/27 第20回 GPU コンピューティング 講習会

ヘテロジニアス・コンピューティング

2

Financial Analysis

Scientific Simulation

Engineering Simulation

Data Intensive Analytics

Medical Imaging

Digital Audio Processing

Computer Vision

Digital Video Processing

Biomedical Informatics

Electronic Design

Automation

Statistical Modeling

Ray Tracing Rendering

Interactive Physics

Numerical Methods

ヘテロジニアス・コンピューティング（CPU + GPU）は広く使われている

2016/6/27 第20回 GPU コンピューティング 講習会

GPUを用いてアプリケーションを高速化する ３つの方法

Applications

Libraries

Easy to use

Most Performance

Programming Languages

Most Performance

Most Flexibility

Easy to use

Portable code

Compiler Directives

2016/6/27 第20回 GPU コンピューティング 講習会 3

GPU Accelerated Libraries

Linear Algebra FFT, BLAS,

SPARSE, Matrix

Numerical & Math RAND, Statistics

Data Struct. & AI Sort, Scan, Zero Sum

Visual Processing Image & Video

NVIDIA

cuFFT,

cuBLAS,

cuSPARSE

NVIDIA

Math

Lib

NVIDIA

cuRAND

NVIDIA

NPP

NVIDIA

Video

Encode

GPU AI –

Board

Games

GPU AI –

Path

Finding

2016/6/27 第20回 GPU コンピューティング 講習会 4

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

CUDA Fortran Fortran

CUDA C C

CUDA C++ C++

PyCUDA, Copperhead, Numba, NumbaPro Python

Alea.cuBase F#

MATLAB, Mathematica, LabVIEW Numerical analytics

GPU Programming Languages

2016/6/27 第20回 GPU コンピューティング 講習会 6

CUDA - C

Applications

Libraries

Easy to use

Most Performance

Programming Languages

Most Performance

Most Flexibility

Easy to use

Portable code

Compiler Directives

2016/6/27 第20回 GPU コンピューティング 講習会 7

CPUとGPU の違い

2016/6/27 第20回 GPU コンピューティング 講習会 8

• 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

DRAM

GPU

• 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

Cache

ALU Control

ALU

ALU

ALU

DRAM

CPU

DeviceQuery

9

./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 (6039339008 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: 1572864 bytes

Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536),

3D=(4096, 4096, 4096)

………

>./deviceQuery >sh runDeviceQuery.sh

2016/6/27 第20回 GPU コンピューティング 講習会

並列処理

並列計算は、

タスク並列 と データ並列

に分けられる。

データ並列の一例：ベクトル加法

10 2016/6/27 第20回 GPU コンピューティング 講習会

A[0] vector A

vector B

vector C

A[1] A[2] A[N-1]

B[0] B[1] B[2]

…

… B[N-1]

C[0] C[1] C[2] C[N-1] …

+ + + +

ベクトル加法（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

GPUを用いたベクトル加法

void vecAdd(float *A_h, float *B_h, float *C_h, int n)‏ { 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 2 // Kernel code を実行 （ベクトルの加法） // Part 3 // 計算結果を回収 – C のデータをdevice // memory からhost memory にコピー // GPU上のメモリを解放 }

2016/6/27 第20回 GPU コンピューティング 講習会 12

CPU

Host Memory

GPU

Device Memory

Part 1

Part 3

Part 2

• GPUは単独では動かない。host（CPU）で実行させ、その中から CUDA API と GPU kernel 関数を call

CUDAプログラム実行の概念図

2016/6/27 第20回 GPU コンピューティング 講習会 13

host (CPU) code

メモリ・ポインタ

float *f_d, *f_h;

CUDA API cudaMalloc(&f_d);

. . . .

kernel function func<<< N/256, 256>>>(f_d)

cudaMemcpy(f_h, f_d);

host memory device memory

device code __global__

func(f_d) { }

GPU CPU

CUDAプログラムの実行モデル

• ヘテロジニアス（CPU + GPU）アプリケーション

–逐次処理はHost C Code

–並列計算はDevice SPMD kernel Code

2016/6/27 第20回 GPU コンピューティング 講習会 14

Serial Code (host)‏

. . .

. . .

Parallel Kernel (device)‏

KernelA<<< Dg, Db >>>(args);

Serial Code (host)‏

Parallel Kernel (device)‏

KernelB<<< Dg, Db >>>(args);

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

2016/6/27 第20回 GPU コンピューティング 講習会 15

Integrated C programs with CUDA extensions

NVCC Compiler

Host C Compiler/ Linker

Host Code Device Code (PTX)

Device Just-in-Time Compiler

Heterogeneous Computing Platform with CPUs, GPUs, etc.

CUDA Compiler: nvcc

• 重要なコンパイル・オプション

2016/6/27 第20回 GPU コンピューティング 講習会 16

-arch sm_52 Compute Capability に応じたコンパイルを行う。DeviceQuery で確認し、それ以下を指定す。 可能なオプション：sm_20(default), sm_21, sm_30, sm_32, sm_35, sm_50 and sm_52

--maxrregcount <N> 1つのkernel 関数当たりに使用するレジスタ数を <N> に制限する。このことにより、指定した並列数でthreadが実行可能となるが、溢れた部分はlocal メモリ上に置かれ、実行速度は低下する。

-use_fast_math 高速な数学関数を利用する。

-G device コードに対して、デバッグを可能にする

--ptxas-options=-v レジスタやメモリの使用状況を表示する

CUDA Memory 確保（1/3）

2016/6/27 第20回 GPU コンピューティング 講習会 17

メモリ・ポインタ－は、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 のメモリ上に確保される

CUDA Memory 確保（2/3）

18

host 側にメモリを確保する（通常） f_h = (float *) malloc(sizeof(float)*n); f_h = new float[n]; (C++)

host 側にpinned メモリを確保する

cudaMallocHost(void **ptr, size_t count); ptr: ホストメモリアドレスへのポインタ。 Page lock （pinned）された確保したメモリのアドレスが書き込まれる count: 領域のサイズ

例） cudaMallocHost((void **)&f_h, sizeof(float)*n);

f_h[n] の配列が Host メモリ上にpage lock (pinned)で確保さ

れる。通常の pageable メモリとして確保された場合より、転送 速度が速い。また、非同期通信の場合も page lock メモリに限 定される。

2016/6/27 第20回 GPU コンピューティング 講習会

CUDA Memory 確保（3/3）

19

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 コンピューティング 講習会

CUDA データ転送

20

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 コンピューティング 講習会

CUDA 非同期データ転送

21

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 コンピューティング 講習会

GPU kernel-function call

22

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 コンピューティング 講習会

dim3 宣言

23

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 コンピューティング 講習会

Threadの管理

• CUDA Kernel はthreads の まとまり(Grid)単位で実行される – Gridの中の全てのthreadsは同じ Kernel を実行する（Single Program

Multiple Data）

– Threadは各自にIDの持っている

2016/6/27 第20回 GPU コンピューティング 講習会 24

i = blockIdx.x * blockDim.x + threadIdx.x; C[i] = A[i] + B[i];

…

0 1 2 254 255

…

Threadの管理(Block)

• いくつかのthreadsが一つのBlockとしてまとめられ、全てthreadsが複数のBlockに分割されます。 – 同一Block内のthreadsはshared memoryを共有し atomic operations と barrier synchronizationなど同期が必要な演算を実行できます

– 異なるBlock間のthreads同期やshared memoryの共有はできません

2016/6/27 第20回 GPU コンピューティング 講習会 25

i = blockIdx.x * blockDim.x + threadIdx.x;

C[i] = A[i] + B[i];

… 0 1 2 254 255

Thread Block 0

… 1 2 254 255

Thread Block 1

0

i = blockIdx.x * blockDim.x + threadIdx.x;

C[i] = A[i] + B[i];

… 1 2 254 255

Thread Block N-1

0

i = blockIdx.x * blockDim.x + threadIdx.x;

C[i] = A[i] + B[i];

…

… … …

blockIdx と threadIdx

• kernel 関数<<< 第１引数, 第2引数>>>で指定 – 第１引数：blockIdx: 1D, 2D, or 3D (CUDA 4.0以降)の範囲を指定

– 第２引数：threadIdx: 1D, 2D, or 3Dの範囲を指定

• 多次元データを計算する 場合にアドレス計算が 簡単になる – Image processing

– Solving PDEs on volumes

– …

26

device

Grid Block (0, 0)

Block (1, 1) Block (1, 0)

Block (0, 1)

Block (1,1)

Thread (0,0,0) Thread

(0,1,3) Thread (0,1,0)

Thread (0,1,1)

Thread (0,1,2)

Thread (0,0,0)

Thread (0,0,1)

Thread (0,0,2)

Thread (0,0,3)

(1,0,0) (1,0,1) (1,0,2) (1,0,3)

2016/6/27 第20回 GPU コンピューティング 講習会

実習：最初のCUDAプログラム

• >nvcc HelloWorld.cu

• >./a.out

2016/6/27 第20回 GPU コンピューティング 講習会 27

#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; }

Sample Code: simpleCUDA/HelloWorld

Built-in 変数

• Device code の中で宣言せずに引用でき、 書き換え不可

gridDim gridDim.x, gridDim.y, gridDim.z

grid の各方向のサイズ

blockIdx blockIdx.x, blockIdx.y, blockIdx.z

block の各方向のindex

blockDim blockDim.x, blockDim.y, blockDim.z

block の各方向のサイズ

threadIdx threadIdx.x, threadIdx.y, threadIdx.z

thread の各方向のindex

28 2016/6/27 第20回 GPU コンピューティング 講習会

C言語の拡張

• 関数型の Qualifier __global__ device 上でのみ実行される host 側からのみ call される

return 値は void 限定

__device__ device 上でのみ実行される device からのみ call される

__host__ host 上でのみ実行される host 側からのみ call される （普通の CPU 上のプログラムの関数 で、特に宣言する必要はない。） __host__ と __device__ 両方同時に指定する事が可能

29 2016/6/27 第20回 GPU コンピューティング 講習会

host host __host__ float HostFunc()‏

host device __global__ void KernelFunc()‏

device device __device__ float DeviceFunc()‏

Only callable from the: Executed on the:

__host__ と __device__

2016/6/27 第20回 GPU コンピューティング 講習会 30

CPU for(……..) { …… Cal_func(…); …… }

GPU __global__ kernel(….) { …… Cal_func(…); …… }

__host__ __devcie__ void Cal_func(……..) { C = A + B; }

ベクトル加法(Host Code)

void vecAdd(float *A_h, float *B_h, float *C_h, int n)‏ { 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

ベクトル加法(Device Kernel)

__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];

}

2016/6/27 第20回 GPU コンピューティング 講習会 32

vectorAdd_kernel<<< n / BLOCK_SIZE, BLOCK_SIZE >>>( C_d, A_d, B_d );

gridDimの値とblockIdx の値の範囲を決めます

blockDimの値とthreadIdx の値の範囲を決めます

Sample Code: simpleCUDA/simpleVectorAdd

2次元データ・アクセス NX*NYの1次元配列データであるが、2次元的にアクセス

NX

NY

blockDim.x = 16

dim3 grid(NX/16, NY/16), block(16, 16);

iy

ix gridサイズの最大値の制限から開放

ix = blockIdx.x * blockDim.x + threadIdx.x; iy = blockIdx.x * blockDim.x + threadIdx.x; index = iy * NX + ix;

2016/6/27 第20回 GPU コンピューティング 講習会

Warp スケジューリング

• block 内の thread は Streaming Multiprocessor によって

Warp＝32 thread毎に並列実行される – Warp: block 内の32 threadのかたまり （例：block 内に 256 thread = Warp 8 個）

– １ Warp の thread の数は将来変わる可能性がある。 – プログラム上には現れない – 考慮しなくても正しいプログラムを書くことは可能。ただし、実行性能を引き出すためには考慮する必要がある。

34 2016/6/27 第20回 GPU コンピューティング 講習会

…

t0 t1 t2 … t31

…

…

t0 t1 t2 … t31

… Block 1 Warps Block 2 Warps

…

t0 t1 t2 … t31

… Block 3 Warps

Warp 多次元分配

• 多次元Thread blocksの場合、Warpは行優先で割り当てられる – 最初X、それからY、最後はZ

35 2016/6/27 第20回 GPU コンピューティング 講習会

Compute Capability

36

Compute Capability

Technical Specifications 2.x 3.0 3.2 3.5 3.7 5.0 5.2 5.3

Warp size 32

Maximum number of resident blocks per multiprocessor

8 16 32

Maximum number of resident warps per multiprocessor

48 64

Maximum number of resident threads per multiprocessor

1536 2048 (= 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 63 255

Maximum amount of shared memory per multiprocessor

48 KB 112 KB 64 KB 96 KB 64 KB

Maximum amount of shared memory per thread block

48 KB

Number of shared memory banks 32

Amount of local memory per thread 512 KB

Constant memory size 64 KB

2016/6/27 第20回 GPU コンピューティング 講習会

同時実行可能なthread block数

37

thread の 使用する Register 数: Nr Shared Memory 量: Ns [byte] block 当たりの thread 数: Db Compute Capability 5.2 (Maxwell世代GPU)の場合：

Active block = Min(32, 64/(Db/32), 98304/Ns, 65536/(Db*Nr) )

Warp per SM = Db/32

（block当たりの最大64 Warp）

（Shared Memory の制限）

（Register の制限）

2016/6/27 第20回 GPU コンピューティング 講習会

Warp内のthread実行仕組み

• Warp 内の32 thread は同一命令を実行 SIMD (Single Instruction Multiple Data)

– Warp 内の thread indices はインクリメント

– Warp 0 は thread 0 からスタート

• Warps 実行される順番は一定ではない

–例えばWarp 8がWarp 7より先に実行される場合がある

– Threads の間の依存関係（計算順序など）は必ず__syncthreads()を使う

38 2016/6/27 第20回 GPU コンピューティング 講習会

Threadの中の条件分岐

• プログラム上では任意の分岐を記述可能。

• ハードウェア上での分岐命令の処理

– Warp内全threadが同一パスに分岐する場合 は 全threadが分岐先（のみ）を実行する。

– Warp内のthreadが異なるパスに分岐する場合は全スレッドが両方の命令を実行 (diverged branch)し、最後に適合する方だけを採用する。性能低下の原因の一つ。

2016/6/27 第20回 GPU コンピューティング 講習会 39

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

例: Vector Addition Kernel

// Compute vector sum C = A + B

// 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];

}

41

Device Code

2016/6/27 第20回 GPU コンピューティング 講習会

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

Memory Coalescing

• 一つのwarpの中の全ての threads が一斉に データをロード。全てのデータアクセスが同じburst sectionの場合、 一回のDRAM request のみで終了する。（fully coalesced）

43

2 1 0 3 5 4 6 7 9 8 10 11 13 12 14 15

Burst section Burst section Burst section Burst section

T0 T1 T2 T3

Coalesced Loads

T0 T1 T2 T3

Coalesced Loads

2016/6/27 第20回 GPU コンピューティング 講習会

Un-coalesced Accesses

• データアクセスが別々のburst sectionになる場合: – Coalescing fails – 何回かの DRAM requests が必要 – アクセスは not fully coalesced

• 一部の転送されたデータは使われない

44

2 1 0 3 5 4 6 7 9 8 10 11 13 12 14 15

Burst section Burst section Burst section Burst section

T0 T1 T2 T3

Un-coalesced Loads

T0 T1 T2 T3

Un-coalesced Loads

2016/6/27 第20回 GPU コンピューティング 講習会

例：行列の乗法のアクセスパターン

A B

WIDTH

(Thread) T1

(Thread) T2

A[Row*n+i] B[i*k+Col] i はループの変数 A は m × n, B は n × k のマトリックス

Col = blockIdx.x*blockDim.x + threadIdx.x

HEI

GH

T

T1 T2

M0,2

M1,1

M0,1 M0,0

M1,0

M0,3

M1,2 M1,3

M0,2 M0,1 M0,0 M0,3 M1,1 M1,0 M1,2 M1,3 M2,1 M2,0 M2,2 M2,3

M2,1 M2,0 M2,2 M2,3

M3,1 M3,0 M3,2 M3,3

M3,1 M3,0 M3,2 M3,3

M

linearized order in increasing address

2016/6/27 第20回 GPU コンピューティング 講習会 45

A Accesses are Not Coalesced

T0 T1 T2 T3

Load iteration 0

T0 T1 T2 T3

Load iteration 1

Access

direction in

kernel code

…

A0,2

A1,1

A0,1 A0,0

A1,0

A0,3

A1,2 A1,3

A2,1 A2,0 A2,2 A2,3

A3,1 A3,0 A3,2 A3,3

A0,2 A0,1 A0,0 A0,3 A1,1 A1,0 A1,2 A1,3 A2,1 A2,0 A2,2 A2,3 A3,1 A3,0 A3,2 A3,3

2016/6/27 第20回 GPU コンピューティング 講習会 46

B accesses are coalesced

N

T0 T1 T2 T3

Load iteration 0

T0 T1 T2 T3

Load iteration 1

Access

direction in

kernel code

B0,2

B1,1

B0,1 B0,0

B1,0

B0,3

B1,2 B1,3

B2,1 B2,0 B2,2 B2,3

B3,1 B3,0 B3,2 B3,3

B0,2 B0,1 B0,0 B0,3 B1,1 B1,0 B1,2 B1,3 B2,1 B2,0 B2,2 B2,3 B3,1 B3,0 B3,2 B3,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);

cudaEventRecord(stop,0); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsedTime,start,stop);

elapsedTime に 経過時間 （msec)

計測範囲

2016/6/27 第20回 GPU コンピューティング 講習会

エラー処理（API）

• CUDA の API は全て return 値が cudaError_t 型の error の status を返すようになっている。

• もし、cudaMalloc しないで、cudaMemcpy(); を実行してしまった場合などは、invalid device pointer が返ってくる。

2016/6/27 第20回 GPU コンピューティング 講習会 49

cudaError_t err = cudaMemcpy(…); if (err != cudaSuccess) { fprintf(stderr, “Memcopy failed: %s.¥n”, cudaGetErrorString(err)); }

エラー処理（ kernel関数）

50

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 コンピューティング 講習会

Device マネージメントAPI

51

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 コンピューティング 講習会

Developer Tools - Debuggers

NSIGHT CUDA-GDB CUDA MEMCHECK

3rd Party

NVIDIA Provided

https://developer.nvidia.com/debugging-solutions

2016/6/27 第20回 GPU コンピューティング 講習会 52

Developer Tools - Profilers

NSIGHT NVVP NVPROF

3rd Party

NVIDIA Provided

https://developer.nvidia.com/performance-analysis-tools

VampirTrace TAU

2016/6/27 第20回 GPU コンピューティング 講習会 53

NVIDIA’s Visual Profiler (NVVP) Timeline

Guided

System Analysis

2016/6/27 第20回 GPU コンピューティング 講習会 54

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