最終更新:2014-02-10 (月) 11:42:06 (2371d)  

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(グリッド内のブロック番号)uint3blockIdx
    スレッドID(ブロック内のスレッド番号)uint3threadIdx
    スレッドブロック数dim3?gridDim
    スレッドブロック内スレッド数dim3?blockDim

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

CUDA/スレッドブロック

  • 単純に並列動作させたいスレッドの総数を指定するのではなく、わざわざスレッドブロックに分割するのは次の理由からです。
  • GPU内部は複数のマルチプロセッサ(SM)に分割されている。
  • 1つのマルチプロセッサ(SM)で同時に扱えるスレッド数には制限がある。
  • CUDAはスレッドブロックをマルチプロセッサ(SM)へ割り当てる単位にしている。
  • ホスト側からカーネル関数を呼び出す際、スレッドブロック数へ1を指定してしまうと、1つのマルチプロセッサ(SM)しか動作してくれない

ソフトとハードの対応

処理単位処理を行うハード
スレッドスレッドプロセッサ?(SP,CUDAコア)
スレッドブロック?ストリーミングマルチプロセッサ(SM)
グリッドデバイス

カーネルの呼び出し方

kernel<<<dim3 dG, dim3 dB, shared-size, stream>>>
  • dG - グリッドあたりのブロックの次元とサイズ
    • グリッドで起動されるブロック数:dG.x * dG.y
  • dB - ブロックあたりのスレッドの次元とサイズ
    • ブロックあたりのスレッド数:dB.x * dB.y * dB.z

メモ

  • カーネルはC関数+多少の制約
  • ホストメモリはアクセスできない
  • 戻り値 型はvoid
  • 可変引数(“vaargs”)は不可
  • 再帰処理はできない
  • 静的変数は使えない

このPDFへのリンク

このPDFへのリンク

修飾子

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

関連

ソフトウェア開発ツール

メモ

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

参考