1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin

Windows で始める CUDA 入門 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 L1 キャッシュなど 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.4

1. Nsight Visual Studio Edition Visual Studio での CUDA 開発 ビルド デバッグ プロファイル CUDA Toolkit に含まれる (CUDA 5.5 から ) 開発者登録が不要になりました

1. 装置構成 マザーボード CPU ( 数コア ) DRAM チップセット PCIe GPU: CPU につながった GPU プロセッサ (~ 数千コア ) DRAM 外部演算装置

1. 典型的な実行例 CPU プログラム開始 GPU は CPU からの制御で動作する GPU データ転送 GPU プログラム実行依頼 完了待ち GPU での演算 データ転送 入力データは CPU GPU へと転送 結果は GPU CPU と転送

1. CUDA カーネル カーネル = GPU 上のプログラム GPU 向け言語 (C++ ベース ) にて記述される 特別なコンパイラ (NVCC) でコンパイル 100 万スレッドオーダーでの並列動作 Massively parallel 並列度の階層構造 GPU のアーキテクチャに密接に関係

1. カーネル実行の階層 CPU GPU Grid データ転送 Block0 Block1 Thread Thread CPU からの呼び出し単位 Block に分解される Block GPU プログラム実行依頼 Grid Block2 Thread 一定数の Thread を持つ GPU 上の並列プロセッサ (SMX) 内部でで実行される Block n Thread Thread 最小の実行単位

1. Warp Block Warp0 Thread Warp 32 GPU-threads Warp1 Thread HW 上の実行単位 Warp2 Thread Warp n Thread

1. Streaming Multiprocessor extreme SMX ( 簡略化しています ) レジスタ 64 K 個 (256 KB) 192 Cores/SMX Compute Capability 3.5 Core Core Core Core 0 1 3 2 Core Core Core Core 0 1 3 2 Core Core Core Core 0 1 3 2 SFU Core Core Core 0 1 3 2 LD/ST Core Core Core 0 1 3 2 DP Core Core Core 0 1 3 2 共有メモリ L1 Cache 64 KB テクスチャキャッシュ 48 KB SFU Special Function Unit LD/ST Load/Store Core 15 Core 15 Core 15 SFU 15 LD/ST 15 DP 15 DP 倍精度演算ユニット

1. 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スレッド カーネルは スレッド単位の視点で書く

1. CUDA プログラム実行の概要 SM(X) CPU Grid Block Block Block Warp Warp Warp Grid Block CPU からの呼び出し単位 Block に分解 SM 上の実行単位 Warp に分解 SM 共有メモリのスコープ Warp CUDA 固有の並列単位 32 GPU threads 条件分岐の粒度 SM(X) ハードウエア上の並列プロセッサ

2. プログラミングの基礎 ホストプログラミング メモリ転送 カーネルの実行 カーネルプログラミング GPU 上の関数の実装

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 デバイス メモリ構成 ホスト GPU SM(X) CPU ホストメモリ PCIe デバイスメモリ ( グローバルメモリ ) CPU 側ホストホストメモリ GPU 側デバイスグローバルメモリ

2.2 配列の和 : メモリの扱い ホスト GPU float *a, *b, *c をアロケート float *da, *db, *dc をアロケート ( デバイスメモリ ) *a, *bに値を設定ホスト-> デバイス転送 a-> da, b->db カーネル実行依頼 カーネル dc[i] = da[i] + db[i] ホスト <- デバイス転送 c <- dc 結果表示 検証 float *a, *b, *c を開放 float *da, *db, *dc を開放 ( デバイスメモリ )

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 スレッドとした場合 Block[0] Block[1] Block[2] Block[3] a[i] b[i] 0 1 2 3 + + + + 15 14 13 12 4 5 6 7 + + + + 11 10 9 8 8 9 10 11 + + + + 7 6 5 4 12 13 14 15 + + + + 3 2 1 0 c[i]

2.2 Global ID / Local ID / Block ID Global ID Grid 内で一意 blockdim.x * blockidx.x + threadidx.x Local ID Block 内で一意 threadidx.x Global ID 0 1 2 3 Block ID (blockidx) 0 Local ID (threadidx) 0 Thread 1 Thread 2 Thread 3 Thread Block ID blockidx.x (OpenCL から概念を拝借 ) 4 5 6 7 Block ID 1 Local ID 0 Thread 1 Thread 2 Thread 3 Thread

2.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; addarraykernel<<<griddim, blockdim>>>(dc, da, db, size);

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

3. Visual Studio 2010 によるビルド ビルドルールの追加

3. Visual Studio 2010 によるビルド Compute Capability の設定

3. Visual Studio 2010 によるビルド ライブラリ指定

NVIDIA Japan CUDA Monthly Seminar NVIDIA Japan では 毎月 CUDA の無償セミナーを実施しています 是非 ご参加ください 申し込み : http://www.nvidia.co.jp/object/event-calendar-jp.html 場所 : NVIDIA Japan 赤坂オフィス