thrustのdevice_vectorとdevice_ptrと"raw"デバイスポインタ

cudaを使ってnvidiaのGPUで演算するときに、慣れるまでdevice_ptrや"raw"デバイスポインタの必要性をなかなか理解できなかったので、まとめてみました。

メモリ管理方法が違う

• クラス自身が管理
  • device_vector
• 自分でnew/delete, malloc/freeを管理
  • device_ptr (device_newやdevice_mallocで作る。)
  • "raw"デバイスポインタ (cudaMallocで作る)

利用できる位置が違う

• ホスト関数でのみ利用可能
  • device_vector
• ホスト関数でもデバイス関数でも利用可能
  • device_ptr
• デバイス関数でのみ値を参照可能
  • "raw"デバイスポインタ

例

h_arrays.cu

__host__
void h_func(thrust::device_vector<int>& ioVec,
            thrust::device_ptr<float>& ioPtr,
            double* ioRawPtr,
            const int inSize)
{
  // (1) device_vector
  ioVec.resize(inSize); // device_vectorの恩恵は、__host__ならでは。
  ioVec[inSize-1] = 3; // こんなに簡単にdeviceの値を書き換えることができる
  thrust::transform(ioVec.begin(), ioVec.end(),
            ioVec.begin(),
            [] __device__ (int v) -> int {
              return v*v;
            }); // transformやfor_eachでdevice_vectorが大活躍するのも__host__ならでは。

  // (2) device_ptr
  ioPtr[inSize-1] = .0f; // __host__でこそ、device_ptrの利便性が高い
  const float sum = thrust::reduce(ioPtr, ioPtr + inSize); // "raw"デバイスポインタだと無理。

  // (3) rawデバイスポインタ
  device_ptr<double> d_ptr = thrust::device_pointer_cast(ioRawPtr);
  d_ptr[inSize-1] = 3.14; // __host__でrawデバイスポインタの値を参照する唯一の方法
}

d_arrays.cu

__device__
void d_func(thrust::device_ptr<float>& ioPtr,
            double* ioRawPtr, const int inSize) // device_vectorは使えない
{
  // (1) device_vector
  // device_vectorは使えない

  // (2) device_ptr
  ioPtr[inSize-1] = .0f; // __device__だと何となく普通。
  const float sum = thrust::reduce(thrust::device, ioPtr, ioPtr + inSize); // execution policy

  // (3) rawデバイスポインタ
  d_ptr[inSize-1] = 3.14; // __device__なのでアクセスできる。
}

言い忘れましたが、device_ptrと"raw"デバイスポインタは、device_pointer_cast()とraw_pointer_cast()を使って、相互変換することができます。

余談

もう"raw"デバイスポインタは必要性がないんじゃないかな。cudaMallocを捨てて、device_mallocを利用することでホスト側からもデバイス側からもアクセスできてバグが減りそう。