1. No category

WINDOWSで始めるCUDA
エヌビディアジャパン CUDAエンジニア 森野慎也
GTC Japan 2014
内容
 GPUのハードウエア構造を理解する
— GPU Diagram
— Compute Capability
 CUDAのプログラミングモデルを理解する
— Grid, Block, Warp, Thread
 Visual Studio + Nsightで、例題を実行する
— 配列の和
GPUの構造
 Giga Thread Engine
— 処理を、SMに割り振る
GPU
Giga Thread Engine
PCI Express
 SM
— 「並列」プロセッサ
 L2 Cache
— R/W可能な二次キャッシュ
SM
SM
SM
L2 Cache
 DRAM
— すべてのSMとPCI Expressから
アクセス可能なメモリ
 PCI Express
— PC(ホスト)との接続インターフェース
DRAM
SM
…
KEPLER GK110 ブロックダイアグラム
アーキテクチャ
 最大 15 SMX ユニット
SMX = KeplerのSM
 71億トランジスタ
 1 TFLOP以上の
倍精度演算性能
 1.5 MB L2 Cache
 384-bit GDDR5
STREAMING MULTIPROCESSOR EXTREME
 Kepler
192 CUDA cores / SMX
 GPU内部の並列プロセッサ
— SMXの単位で処理を実行。
— CUDA coreは、単体で動作しない。
COMPUTE CAPABILITY
 GPUコアアーキテクチャのバージョン
— CUDA GPUs : https://developer.nvidia.com/cuda-gpus
 アーキテクチャは進化する
— 低消費電力
— 高効率の命令実行
SM 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 (1st gen)
CC 5.0 : 128 cores / SMM
CUDAプログラミングモデル
 並列処理のみ実行可能
— 100万スレッド以上での並列動作
— Massively Parallel
 SMアーキテクチャによらず、同じプログラムが動作する
— GPUハードウエアの共通概念
— プログラミングモデルとの対応
装置構成
GPU
PC
CPU
(数コア)
GPU: CPUにつながった
外部演算装置
Giga Thread Engine
制御
PCIe
ホスト側DRAM
SMX
SMX SMX SMX
L2 Cache
DRAM
転送
…
典型的な実行例
プログラム
開始
 GPUはCPUからの制御
で動作する。
CPU
データ
転送
CUDA
カーネル
実行
完了
待ち
GPU
GPUでの演算
データ
転送
 入力データは
CPU→GPUへと転送。
 結果は、
GPU→CPUと転送
 GPU上に常駐する
プログラムはない。
2.1 カーネル実行の階層
GPU
CPU
Grid
データ
転送
Thread
Thread
Thread
Thread
Block1
Thread
Thread
Thread
Thread
Block2
Thread
Thread
Thread
Thread
Block
n
CPUから呼び出される。
Blockにより構成される。
Block
Threadにより構成される。
Thread数は、Grid内部で一定
…
…
CUDA
カーネル
実行依頼
Grid
Block0
Thread
Thread
Thread
Thread
Thread
最小の実行単位
BLOCK は、SMX上で実行
プログラミングモデル
Grid
Block0
GPU
Block1
Block2
SMX
Block3
Block4
SMX
※ Blockの実行順序は保証されない。
※ 特定のSMXへのBlock割り当てはできない。
SMX
Block5
Block6
SMX
…
Block N
BLOCK は SMX上で実行
 Block ⇒ 1 SMX
— 複数のSMXにまたがらない。
（SMX中では、複数Blockが実行される場合もある。）
— Block内部では、SMXのリソースを活用可能
 各々のBlockは、独立に処理を実行する。
— 実行順序の保証はない。
— Block間の通信・同期は行わない。
STREAMING MULTIPROCESSOR EXTREME (略図)
 192 Cores/SMX
Compute
Capability 3.5
SMX (簡略化しています)
レジスタ 64 K個 (256 KB)
Core
Core
Core
0
Core
1
2
3
Core
Core
Core
0
Core
1
2
3
Core
SFU LD/ST DP
Core Core Core Core
Core
Core
Core
Core
0
0
0
0
Core
Core
Core
Core
1
1
1
1
2
2
2
2
3
3
3
3
共有メモリ
L1 Cache
64 KB
テクスチャ
キャッシュ
48 KB
Core Core
15
15
Core
15
SFU
15
LD/ST
DP
15
15
 SFU
Special Function Unit
 LD/ST
Load/Store
 DP
倍精度演算ユニット
WARP : 命令発行の単位
 1命令を Warp が、32並列で処理
— SIMT (Single Instruction Multiple Thread)
Thread
Core
Core
Core
Core
Core
CUDA cores
Thread
Thread
Thread
Thread
Thread
SMX
…
32 GPU Thread
Warp
…
Warp
Block
Warp
SW
1命令を
32並列実行
CUDAプログラム実行の概要
SM(X)
CPU
Grid
Block
Block
Block
Warp (32 Thread)
Warp (32 Thread)
Warp (32 Thread)
Grid
カーネル全体、全てのBlockを含む
Block
SMX内部で、実行される。
Blockサイズは、Grid内で、一定。
Warp
命令発行の単位。32並列で実行。SIMT
Thread
GPUスレッド。ソースコードに対応。
SM
GPUの並列プロセッサ
プログラミングの基礎
 ホストプログラミング
— メモリ転送、カーネルの実行
 カーネルプログラミング
— GPU上の関数の実装
CUDA ホストAPI (抜粋)
機能
CUDA Runtime API
メモリ確保・解放 cudaError_t cudaMalloc(void ∗∗ devPtr, size_t size)
cudaError_t cudaFree(void *)
メモリ転送
cudaError_t
cudaMemcpy (void ∗ dst, const void ∗ src, size_t count,
enum cudaMemcpyKind kind)
同期
cudaError_t cudaDeviceSynchronize(void)
エラーチェック
cudaError_t
const char∗ cudaGetErrorString (cudaError_t error)
CUDA C/C++ カーネル
__global__
void myKernel(int a, float *pb, …) {
/* device code */
}
 ホストから呼び出し可能なデバイス側の関数
— __global__を修飾子として持つ
— 戻り値は、voidでなければならない。
 並列度に対応する回数、カーネルが呼び出されるイメージ。
2.3 プログラム例
 配列の和
c[i] = a[i] + b[i]
 メモリの取り扱い
 基本的なカーネルの実装
装置構成
GPU
PC
CPU
(数コア)
GPU: CPUにつながった
外部演算装置
Giga Thread Engine
制御
PCIe
SMX
SMX SMX SMX
L2 Cache
DRAM
(Global Memory)
ホスト側DRAM
転送
…
配列の和：メモリの扱い
GPU
ホスト
float *a, *b, *c をアロケート
float *da, *db, *dc をアロケート (デバイスメモリ)
*a, *bに値を設定
ホスト->デバイス転送 a-> da, b->db
カーネル実行依頼
カーネル
dc[i] = da[i] + db[i]
ホスト <- デバイス転送 c <- dc
結果表示・検証
float *da, *db, *dc を開放 (デバイスメモリ)
float *a, *b, *c を開放
並列化(カーネル設計)
 複数のブロックに配分して、和をとる。
— 図は、1 ブロックあたり、4スレッドとした場合
Block[0]
a[i]
b[i]
c[i]
0
1
2
Block[1]
3
4
5
6
Block[2]
7
8
9
Block[3]
10 11
12 13
14 15
+ + + +
+ + + +
+ + + +
+ + + +
15 14
11 10
7
3
13 12
9
8
6
5
4
2
1
0
GLOBAL ID
— Global ID
 Grid内部でのスレッド番号
 Grid内で一意
blockDim.x * blockIdx.x + threadIdx.x
— blockIdx


Block番号
Grid内で一意
— threadIdx


Block内のスレッド番号
Block内で一意
Global ID
0
1
2
3
blockIdx
0
blockIdx
4
5
6
7
1
threadIdx
0
1
2
3
Thread
Thread
Thread
Thread
threadIdx
0
1
2
3
Thread
Thread
Thread
Thread
カーネル実装例
__global__
void addArrayKernel(float *dc, const float *da, const float *db, int size) {
/* Global IDを算出 */
int globalID = blockDim.x * blockIdx.x + threadIdx.x;
if (globalID < size) { /* 範囲チェック */
/* 自スレッド担当の要素のみ、処理 */
dc[globalID] = da[globalID] + db[globalID];
}
}
ブロック数の指定
 カーネルはブロック数でスケールする
— ブロックごとのスレッド数は一定
/* gridDim * blockDim個のスレッドを起動する */
int blockDim = 256;
int gridDim = (size + blockDim – 1) / blockDim;
addArrayKernel<<<gridDim, blockDim>>>(dc, da, db, size);
動かしてみる
まとめ
 GPUのハードウエア
— Giga Thread Engine, SMX, SIMT
 CUDAプログラミングモデル
— Grid, Block, Warp, Thread
 基本的なプログラミング
— ホストプログラミング : メモリ転送、カーネル実行
— カーネルプログラミング : Global ID