• 本文
• リロード
• 差分
• バックアップ
• 名前変更
• リファラ
• 新規ページ
• 編集

CUDA/用語

CUDA/カーネル

Top / 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];
  }

組み込み変数

• 名称	型	変数名
  ブロックID(グリッド内のブロック番号)	uint3	blockIdx
  スレッドID(ブロック内のスレッド番号)	uint3	threadIdx
  スレッドブロック数	dim3?	gridDim
  スレッドブロック内スレッド数	dim3?	blockDim

例

足し算(１次元配列)

• __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を持ち、メモリアドレスの計算と制御の判定に使用

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”)は不可
• 再帰処理はできない
• 静的変数は使えない

修飾子

• __global__? - カーネルを示す関数識別子
• __device__? - カーネル関数内で使用する関数
• __host__? - ホストで実行する関数

関連

• CUDA/スレッドブロック

ソフトウェア開発ツール

• http://www.nvidia.co.jp/object/tesla_software_jp.html

メモ

• ホスト側からカーネル関数を呼び出す際、スレッドブロック数へ1を指定してしまうと、1つのマルチプロセッサ(SM)しか動作してくれない

参考

• http://gpu.fixstars.com/index.php/CUDA_文法_2