cuda-tutorial

2.2 デバイスコード

2.2.1 カーネル

デバイス上でプログラムを走らせる関数のうち、ホスト側で呼び出されるものをカーネルと言う。 先程の

 __global__ void addVec(float *vec1_dev, float *vec2_dev){
  unsigned int i_global = blockIdx.x * blockDim.x + threadIdx.x;             
  vec1_dev[i_global] += vec2_dev[i_global];
} 
...
int main(){
  ...
   //ホスト関数内でカーネルは呼ばれる
   addVec<<<NB,NT>>>(vec1_dev, vec2_dev);
  ...
}

はまさにカーネルだ。カーネルには

• __global__ 修飾語をつけることでホスト関数と区別する
• ホスト関数内で呼ばれる
• デバイスメモリにアクセスし、ホストメモリにアクセス出来ない　←注意
• 返り値がない
• スレッドが持つID等の情報に参照できる

といった特徴がある。

上のカーネルを例に取ると、引数の float *vec1_dev 等はデバイスメモリを前提にしている。また、 blockIdx.x 、 threadIdx.x が実際に動くスレッドの情報を関数に渡してくれる。

2.2.2 カーネル上で動く関数

複数のカーネル上で似たような処理を書くことになった場合、それを関数という形でまとめたくなる。ただし、カーネル内でホスト側のの関数を呼び出すことは出来ない。そこでデバイス上で呼び出されてデバイス上で動作する関数が必要で、CUDAはそれをちゃんとサポートしている。

引数を2乗する関数を例に、その使用例を以下に記す。

__device__ float square_dev(float x){
  return x * x;
}
__global__ void kernel(...){
  ...
  x = square_dev(x);
  ...
}

• __device__ 修飾語をつける
• 「カーネル内でのみ」呼び出し可能　←注意

のような特徴がある。__device__で修飾された関数は通常カーネル内でインライン展開される。

Links

• 次へ
• 前へ
• 目次