最終更新:2014-02-10 (月) 11:42:06 (3745d)
CUDA/カーネル
http://gpu.fixstars.com/index.php/CUDA_文法_2
ホストから呼び出され、デバイス上で実行される関数。
文法
kernel<<<スレッドブロック数, 各スレッドブロック内のスレッド数>>>(カーネル引数1,カーネル引数2, ...);
呼び出す側
kernel_func<<<1, 512>>>(カーネル引数1,カーネル引数2, ...);
- の時は512個のスレッドブロックx1つのSMで処理 = 512個のスレッド。
- 512個のkernel_funcスレッドがGPU上で並列動作
kernel_func<<<2, 256>>>(カーネル引数1,カーネル引数2, ...);
- とすると2つのSMで256個ずつ処理。
呼び出される側
- GPUで動作するカーネル関数コードとして記述する内容は、1スレッド分の処理のみ
__global__ void kernel_func(float *A, float *B, float *C) { C[0] = A[0] + B[0]; }
組み込み変数
例
足し算(1次元配列)
__global__ void kernel_func(float *A, float *B, float *C) { int index = blockIdx.x * blockDim.x + threadIdx.x; C[index] = A[index] + B[index]; }
2次元配列へのアクセス
__global__ void kernel_mat_func(引数) { int indexCol = blockIdx.x * blockDim.x + threadIdx.x; int indexRow = blockIdx.y * blockDim.y + threadIdx.y; // do something... }
スレッドブロックに分割するのは次の理由から
- GPU内部は複数のマルチプロセッサ(SM)に分割されている。
- 1つのマルチプロセッサ(SM)で同時に扱えるスレッド数には制限がある。
- CUDAはスレッドブロックをマルチプロセッサ(SM)へ割り当てる単位にしている。
CUDAのカーネルとスレッド
- アプリケーションの並列部分をデバイス上でカーネルとして実行
- 1個のデバイスで一度に実行できるカーネルは1個
- Fermiからは、同じコンテクストの中の依存性のないカーネルプログラムをGPUの中で同時に走らせることができるようになった。
- たくさんのスレッドで1個のカーネルを処理
- 各スレッドがIDを持ち、メモリアドレスの計算と制御の判定に使用
- 1個のデバイスで一度に実行できるカーネルは1個
CUDA/スレッドブロック
- 単純に並列動作させたいスレッドの総数を指定するのではなく、わざわざスレッドブロックに分割するのは次の理由からです。
- GPU内部は複数のマルチプロセッサ(SM)に分割されている。
- 1つのマルチプロセッサ(SM)で同時に扱えるスレッド数には制限がある。
- CUDAはスレッドブロックをマルチプロセッサ(SM)へ割り当てる単位にしている。
- ホスト側からカーネル関数を呼び出す際、スレッドブロック数へ1を指定してしまうと、1つのマルチプロセッサ(SM)しか動作してくれない
ソフトとハードの対応
カーネルの呼び出し方
kernel<<<dim3 dG, dim3 dB, shared-size, stream>>>
- dG - グリッドあたりのブロックの次元とサイズ
- グリッドで起動されるブロック数:dG.x * dG.y
- dB - ブロックあたりのスレッドの次元とサイズ
- ブロックあたりのスレッド数:dB.x * dB.y * dB.z
メモ
- カーネルはC関数+多少の制約
- ホストメモリはアクセスできない
- 戻り値 型はvoid
- 可変引数(“vaargs”)は不可
- 再帰処理はできない
- 静的変数は使えない
修飾子
関連
ソフトウェア開発ツール
メモ
- ホスト側からカーネル関数を呼び出す際、スレッドブロック数へ1を指定してしまうと、1つのマルチプロセッサ(SM)しか動作してくれない