GTC 2013 チュートリアル エヌビディアジャパン
CUDAエンジニア 森野慎也
1. GPUコンピューティング
GPUコンピューティング
— GPUによる、汎用コンピューティング — GPU = Graphics Processing Unit
CUDA
— Compute Unified Device Architecture
— NVIDIAのGPUコンピューティング環境 Linux・Windows・MacOS Xにて動作
Compute Capability
ハードウエアアーキテクチャのバージョン
— 「Fermi」(2.0以降) がターゲットです
機能 対応するデバイス 発表
1.0 初期バージョン 2006/11
1.1 globalメモリ上の32-bit atomic演算 GeForce 9XXX系 2007/10 1.2 メモリアクセスパターンの改善 実行スレッド数の増加、など。 GeForce GT240など 1.3 倍精度演算 GTX285, Tesla C1060など 2008/6 2.0 32 cores/SM、L１キャッシュなど GTX580、Tesla C2070など 2009/10 2.1 48 cores/SM GTX460、GTX560Ti など 3.0 192 cores/SMX GTX680, Tesla K10 2012/3 3.5 Dynamic Parallelism, 64 DP/SMX Tesla K20(X) 2012/11
最近のCUDAデバイス
デバイス名 コア数 ピーク演算性能 単精度/倍精度(FLOPS) メモリバンド幅 GB/sec Compute Capability 3.0 Quadro K5000 1536 2.1 T / 90 G 173 Tesla K10 1536 x 2 4.58 T / 0.19 T 320 GeForce GTX770 1536 3.21 T / 134 G 224.3 Compute Capability 3.5 Tesla K20X 2688 3.95 T / 1.31 T 250 Tesla K20 2496 3.52 T / 1.17 T 208 GeForce GTX Titan 2688 4.50 T / 1.31 T 288.41. Nsight Visual Studio Edition
Visual StudioでのCUDA開発
— ビルド・デバッグ・プロファイル — CUDA Toolkitに含まれる。 （CUDA 5.5から） 開発者登録が
不要になりました。
1. 装置構成
CPU (数コア) マザーボード DRAM チップ セットGPU
GPU: CPUにつながった
外部演算装置
プロセッサ(～数千コア) DRAMPCIe
1. 典型的な実行例
GPUはCPUから
の制御で動作す
る。
入力データは
CPU→GPUへと
転送。
結果は、
GPU→CPUと転
送
CPU
GPU
GPU
プログラム
実行依頼
GPUでの演算
データ
転送
完了
待ち
データ
転送
プログラム
開始
1. CUDAカーネル
カーネル = GPU上のプログラム
— GPU向け言語(C++ベース)にて記述される。 — 特別なコンパイラ(NVCC）でコンパイル 100万スレッドオーダーでの並列動作
— Massively parallel 並列度の階層構造
— GPUのアーキテクチャに密接に関係1. カーネル実行の階層
Grid CPUからの呼び出し単 位。 Blockに分解される。 Block 一定数のThread を持つ。 GPU上の並列プロセッ サ(SMX) 内部でで実行 される。 Thread 最小の実行単位CPU
GPU
GPU
プログラム
実行依頼
データ
転送
Grid Block0 Block1 Block2 Block n …Thread Thread Thread Thread Thread Thread Thread
Thread Thread Thread Thread
Thread
…
Thread Thread Thread Thread
1. Warp
Warp
— 32 GPU-threads — HW上の実行単位 Block …Thread Thread Thread Thread Thread Thread Thread
Thread Thread Thread Thread
Thread
Thread Thread Thread Thread Warp0
Warp1 Warp2
1. Streaming Multiprocessor eXtreme
192 Cores/SMX Compute Capability 3.5 SFU Special Function Unit LD/ST Load/Store DP 倍精度演算ユニ ット SMX (簡略化しています) レジスタ 64 K個 (256 KB) 共有メモリ L1 Cache 64 KB テクスチャ キャッシュ 48 KB Core 3 Core 2 Core 1 Core 0 Core 15 Core 3 Core 2 Core 1 Core 0 Core 15 Core 3 Core 2 Core 1 Core 0 Core 15 Core 3 Core 2 Core 1 SFU 0 SFU 15 Core 3 Core 2 Core 1 LD/ST 0 LD/ST 15 Core 3 Core 2 Core 1 DP 0 DP 151. Streaming Multiprocessor eXtreme
GPU内部の「並列プロセッサ」
— 本質的に並列 (しか実行できない） Blockは、SMX内部で動作
— GPUは、SMXの個数でスケールする。 高性能なGPU ⇒ 数多くのSMXを搭載 旧世代(Fermi以前):
Streaming Multiprocessor (SM)と呼びます
1. Grid・Block・Warp・Thread
Grid
— カーネル全体、全てのBlockを含む Block
— 「カーネル設計」時に、重要な粒度 — Blockのサイズはカーネル内で一定。実行個数は、変更可能 Warp
— 「高速なプログラムを書く」時に重要な粒度 — HWに密接に関連。分岐処理、メモリアクセスの粒度 Thread
— 個々のGPUスレッド — カーネルは、スレッド単位の視点で書くSM(X)
1. CUDAプログラム実行の概要
Grid CPUからの呼び出し単位 Blockに分解 Block SM上の実行単位 Warpに分解 SM・共有メモリのスコープWarp CUDA固有の並列単位 32 GPU threads・条件分岐の粒度
SM(X) ハードウエア上の並列プロセッサ
CPU
Grid
Block
Block
Block
Warp
Warp
Warp
2. プログラミングの基礎
ホストプログラミング
— メモリ転送、カーネルの実行
カーネルプログラミング
2.1 CUDA ホストプログラミング
メモリのアロケーション、解放
— cudaMalloc()/cudaFree() メモリコピー
— cudaMemcpy() カーネルの呼び出し
— 特殊な構文 同期
— cudaDeviceSynchronize()2.1 cudaMalloc() / cudaFree()
cudaError_t cudaMalloc(void ∗∗ devPtr, size_t size)
cudaError_t cudaFree(void *);
例：
float *devptr;
/* float型、1024個の要素分のデバイスメモリをアロケート */
cudaMalloc((void**)&devptr, sizeof(float) * 1024);
/* 解放 */
cudaFree(devptr);
2.1 cudaMemcpy()
cudaError_t
cudaMemcpy (void ∗ dst, const void ∗ src, size_t count,
enum cudaMemcpyKind kind)
例：
float src[1024] = {…..}
float *ddst;
cudaMalloc((void**)&ddst, sizeof(float) * 1024);
cudaMemcpy(ddst, src, sizeof(float) * 1024,
cudaMemcpyHostToDevice
);
srcから ddstに、float型、1024個の要素をコピーする。
2.1 cudaMemcpy()
メモリは、「ホスト」「デバイス」の二種類
enum cudaMemcpyKind cudaMemcpyHostToDevice cudaMemcpyDeviceToHost cudaMemcpyDeviceToDevice cudaMemcpyHostToHost (cudaMemcpyDefault : GPUdirect)2.1 カーネル呼び出し
カーネル呼び出しの構文
kernelName<<<GridDim, BlockDim>>>(引数);
GridDim : グリッド中のブロック数
BlockDim : ブロックあたりのスレッド数
引数は、複数個指定可能
例：
sample<<<1, 256>>>(x, y, z);
2.1 cudaDeviceSynchronize()
cudaError_t cudaDeviceSynchronize (void)
例：
someKernel<<<xxx, yyy>>>(a, b, c);
cudaDeviceSynchronize();
2.1 cudaError_t
エラーチェック
— 成功時は、cudaSuccessを返す。
エラーの場合、値を確認。
const char∗ cudaGetErrorString (cudaError_t error)
2.1 CUDAカーネル
__global__
void myKernel(int a, float *pb, …) { /* device code */ }
ホストから呼び出し可能なデバイス側の関数
— __global__を修飾子として持つ — 戻り値は、voidでなければならない。 通常のC/C++の構文が使用可能。
2.2 プログラム例
配列の和
c[i] = a[i] + b[i]
メモリの取り扱い
2.2 デバイス・メモリ構成
CPU側 ホスト ホストメモリ
GPU側 デバイス グローバルメモリ
CPUホスト
ホストメモリGPU
SM(X) デバイスメモリ (グローバルメモリ)PCIe
2.2 配列の和：メモリの扱い
float *a, *b, *c をアロケート
ホスト
GPU
カーネル
dc[i] = da[i] + db[i] *a, *bに値を設定 ホスト->デバイス転送 a-> da, b->db ホスト <- デバイス転送 c <- dc float *da, *db, *dc をアロケート (デバイスメモリ) 結果表示・検証 float *da, *db, *dc を開放 (デバイスメモリ) float *a, *b, *c を開放 カーネル実行依頼
2.2 配列の和 : ホストコード
int main() {
static const int size= 256 * 100; int memSize = sizeof(float) * size;
float *a, *b, *c, *da, *db, *dc; /* ホストもデバイスもメモリは同じポインタ型 */ /* ホスト側メモリの確保と値の初期化(略)*/
/* GPU側メモリをアロケート */
cudaMalloc(&da, memSize); cudaMalloc(&db,memSize); cudaMalloc(&dc, memSize);
cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice); /* メモリ転送(Host→Device) */ cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice);
/* カーネル(addArrayKernel)をここで呼ぶ */
cudaMemcpy(c, dc, memSize, cudaMemcpyDeviceToHost); /* メモリ転送(Host←Device) */ /* 表示などの処理 (略) */
cudaFree(da); cudaFree(db); cudaFree(dc); free(a); free(b); free(c);
2.2 並列化(カーネル設計)
複数のブロックに配分して、和をとる。
— 図は、1 ブロックあたり、4スレッドとした場合a[i]
b[i]
c[i]
Block[0]
0 1 2 3+ + + +
15 14 13 12Block[1]
4 5 6 7+ + + +
11 10 9 8Block[2]
8 9 10 11+ + + +
7 6 5 4Block[3]
12 13 14 15+ + + +
3 2 1 02.2 Global ID / Local ID / Block ID
Global ID、Grid内で一意
blockDim.x * blockIdx.x + threadIdx.x
Local ID、Block内で一意
threadIdx.x Block ID
blockIdx.x(OpenCLから概念を拝借)
0 1 2 3 Global ID 4 5 6 7 Local ID (threadIdx) 0 Thread 1 Thread 2 Thread 3 Thread 0 Block ID (blockIdx) Local ID 0 Thread 1 Thread 2 Thread 3 Thread 1 Block ID2.2 カーネル実装
__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]; }
2.2 ブロック数の指定
カーネルはブロック数でスケールする
— ブロックごとのスレッド数は一定
/* gridDim * blockDim個のスレッドを起動する */ int blockDim = 256;
int gridDim = (size + blockDim – 1) / blockDim;
2.2 配列の和 : ホストコード
int main() {
static const int size= 256 * 100; int memSize = sizeof(float) * size;
float *a, *b, *c, *da, *db, *dc; /* ホストもデバイスもメモリは同じポインタ型 */ /* ホスト側メモリの確保と値の初期化(略)*/
/* GPU側メモリをアロケート */
cudaMalloc(&da, memSize); cudaMalloc(&db,memSize); cudaMalloc(&dc, memSize);
cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice); /* メモリ転送(Host→Device) */ cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice);
int blockDim = 256; int gridDim = (size + blockDim – 1) / blockDim;
addArrayKernel<<<gridDim, blockDim>>>(dc, da, db, size);
// cudaDeviceSynchronize(); /* 同期。今回は、必須ではない。 */
cudaMemcpy(c, dc, memSize, cudaMemcpyDeviceToHost); /* メモリ転送(Host←Device) */ /* 表示などの処理 (略) */
cudaFree(da); cudaFree(db); cudaFree(dc); free(a); free(b); free(c);
