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


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

}


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


CUDA - C

Applications

Libraries

Easy to use

Most Performance

Programming Languages

Most Performance

Most Flexibility

Easy to use

Portable code

Compiler Directives


CPUとGPU の違い


• 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


並列処理

並列計算は、

タスク並列とデータ並列

に分けられる。

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


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


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上のメモリを解放 }


CPU

Host Memory

GPU

Device Memory

Part 1

Part 3

Part 2

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

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


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


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


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

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


-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）


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


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 メモリに限定される。



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);


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] にコピーする。


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 に対して絶えず非同期。


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; と変更可能である。


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

…

0 1 2 254 255

…

Threadの管理(Block)

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

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


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


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

… 1 2 254 255

Thread Block N-1

0


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)


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

• >nvcc HelloWorld.cu

• >./a.out


#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


C言語の拡張

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

return 値は void 限定

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

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


host host __host__ float HostFunc()‏

host device __global__ void KernelFunc()‏

device device __device__ float DeviceFunc()‏

Only callable from the: Executed on the:

__host__ と __device__


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


ベクトル加法(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];

}


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;


Warp スケジューリング

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

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

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


…

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


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


同時実行可能な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 の制限）


Warp内のthread実行仕組み

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

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

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

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

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

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


Threadの中の条件分岐

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

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

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

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


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 の倍数


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


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％以下


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


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


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

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


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


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


所要時間計測

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)

計測範囲


エラー処理（API）

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

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


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 が表示される。


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


Developer Tools - Profilers

NSIGHT NVVP NVPROF

3rd Party

NVIDIA Provided

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

VampirTrace TAU


NVIDIA’s Visual Profiler (NVVP) Timeline

Guided

System Analysis


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


Documents

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

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