３次元デバイスメモリ配列から3次元cudaArrayへのコピー方法

そりゃ知っている人には当たり前なんだろうけど、知らない人には分かるわけない。
まあプログラムなんてそんなもの。


■
CUDAのデバイスの動的メモリ確保の手段は、入門書にはcudaMallocしか書いてないが、サンプルコードではむしろcudaMallocはあまり見かけない。とにかくコピー関数が色々あって、最初に見た時にそれの使い方がわからない。
幸い、最近はCUDAのドキュメントも充実してきたので大体は解決するが、今回のようにドキュメントをよく読まないと見落とすような落とし穴があったりする。
愚痴はここまでにして、今回の落とし穴を記録に残す。


■
事件はプログラムを3次元に拡張しようと、cudaMemcpy2DToArrayをcudaMemcpy3Dにしようとした時に起こった(cudaMemcpy3DToArrayを用意してればこんなことにはならなかったんじゃないのかNVIDIA)。
元の２次元のコードは次のとおりである。

typedef float2 DataType; 
DataType *data = NULL;
cudaArray *array = NULL;
size_t pitch = 0;

cudaMallocPitch((void **)&data, &pitch, sizeof(DataType)*SIZE_X, SIZE_Y);
cudaChannelFormatDesc desc = cudaCreateChannelDesc<DataType>();
cudaMallocArray(&array, &desc, SIZE_X, SIZE_Y);

cudaMemcpy2DToArray(array, 0, 0, data,
    pitch, SIZE_X*sizeof(DataType), SIZE_Y, cudaMemcpyDeviceToDevice);

このコードの2を3にすればいいんじゃないかという私のような甘い考えだと、このコードを修正するのに一日かかる。

今回の私の場合、arrayはテクスチャメモリにバインドされていて、テクスチャはtex2D関数で参照される。これを3次元にした時、テクスチャを参照するにはtex3D関数を用いるが、tex3D関数の戻り値にfloat3は用意されていない。float4しかない。ないものはしょうがないので、今回は3次元データとしてfloat4を使うことにした。

それぞれの関数に3次元用の関数が用意されている。

2D	3D
cudaMallocPitch	cudaMalloc3D
cudaMallocArray	cudaMalloc3DArray
cudaMemcpy2DToArray	cudaMemcpy3D

cudaMallocArrayにcudaMalloc3DArrayが対応するのはとてもよく分かる。
cudaMalloc3Dは一見どこにも関数名にPitch要素が含まれていないが、確保したメモリのポインタはしっかりcudaPitchedPtrで返してくる。
cudaMemcpy3Dは、もう2次元メモリコピー関数全てを受け入れたというような関数になっている。
cudaMemcpy3Dは、各情報を保持するcudaMemcpy3DParmsという構造体を渡して実行する。

以上を踏まえて、コードを書き換えてみる。

typedef float4 DataType; 
DataType *data = NULL;
cudaArray *array = NULL;
size_t pitch = 0;

cudaMalloc3D((void **)&data,
    make_cudaExtent(SIZE_X*sizeof(DataType), SIZE_Y, SIZE_Z));
pitch = ((cudaPitchedPtr *)&dvfield)->pitch;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<DataType>();
cudaMalloc3DArray(&array, &desc, make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z));

cudaMemcpy3DParms parms = { 0 };
parms.dstArray = array;
parms.srcPtr = *(cudaPitchedPtr *)&data;
parms.extent = make_cudaExtent(SIZE_X*sizeof(DataType), SIZE_Y, SIZE_Z);
parms.kind = cudaMemcpyDeviceToDevice;

cudaMemcpy3D(parms);

ピッチ数の取得法やcudaPitchedPtr*へのキャストが大変カオスになっているが、型を合わせるにはこうするしかない。綺麗に見せるならmake_cudaPitchedPtrとかするんだろうが、ポインタのキャストのほうが低コストだし、美しさより早さを求めなければCUDAを使う意味が無い。

が、問題はキャストではない。このコードはコンパイルこそ出来るが実行するとcudaMemcpy3Dでエラーを吐く。(cudaErrorInvalidValueとかcudaErrorInvalidPitchValueとか)
ドキュメントを読んでもエラーの原因がさっぱりわからず、CUDAの3Dデータを扱う情報が少ないことで心が折れかけたが、次のサイトが助けになった。
http://stackoverflow.com/questions/9399451/invalid-argument-in-cudamemcpy3d-using-width-in-bytes

CUDA ArrayではX次元のサイズをバイト数ではなく要素数で指定するらしい。確かに、改めてcudaMemcpy3Dの公式ドキュメントを改めて確認すると、If a CUDA array is participating in the copy, the extent is defined in terms of that array's elements.とある。
というわけで、正解はこうなる。

typedef float4 DataType; 
DataType *data = NULL;
cudaArray *array = NULL;
size_t pitch = 0;

cudaMalloc3D((cudaPitchedPtr *)&data,
    make_cudaExtent(SIZE_X*sizeof(DataType), SIZE_Y, SIZE_Z));
pitch = ((cudaPitchedPtr *)&dvfield)->pitch;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<DataType>();
cudaMalloc3DArray(&array, &desc, make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z));

cudaMemcpy3DParms parms = { 0 };
parms.dstArray = array;
parms.srcPtr = *(cudaPitchedPtr *)&data;
parms.extent = make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z);
parms.kind = cudaMemcpyDeviceToDevice;

cudaMemcpy3D(parms);

結局、ドキュメントは一行一行しっかり読みましょうという話。