CUDA関数修飾子

CUDAでは修飾子によりその関数がGPU側，CPU側どちらで実行されるかを指定する． 例えば，

__device__
inline void func(int i)
{
    // 処理
}

のように__device__と指定することでGPU側(デバイス)で実行される関数となる．

修飾子の種類†

修飾子の種類を実行場所と呼び出し元でまとめると，

その他注意事項は以下．

• 修飾子を何も指定しなければ __host__ となる．
• __global__関数は返値が void でなければならない．
• __global__関数，<<< >>> オペレータによりスレッド数などを指定した上で実行されなければならない．
• __global__関数は非同期(asynchronous)なので，デバイスで実行が完了する前にホストに処理が返ってくる．
• __global__と__host__は同時指定できない．
• __device__と__host__は同時に指定することができる．

  __device__ __host__
  void func(int i)
  {
  #if __CUDA_ARCH__ == 100
      // compute capability 1.0用の処理
  #elif __CUDA_ARCH__ == 200
      // compute capability 2.0用の処理
  #elif !defined(__CUDA_ARCH__)
      // ホストでの処理
  #endif
  }

  この場合，ホスト関数，デバイス関数両方がコンパイル時に生成される． __CUDA_ARCH__はホストとデバイス，および，compute capabilityの違いにより異なる処理を行いたい場合に用いることができる．

インライン関数†

compute capability 1.* では__device__関数は常にインライン化されるが， compute capability 2.* ではコンパイラが適切と判断したもののみインライン化される．

ユーザが明示的にインライン化を指定したい場合は，__forceinline__修飾子を用いる． 逆にインライン化したくない場合は__noinline__修飾子を用いる． compute capability 1.* ではポインタ引数を持つ関数と引数が多い関数に関しては__noinline__は無視される可能性がある． compute capability 2.* では__noinline__はほとんどの場合有効である．