CUDA3.1をざっとまとめてみる

簡単にprogramming guide 3.1を読んでみたので備忘録も兼ねて軽くまとめる。コード量とかあるので基本Runtime APIで。あと、結構落としてる箇所があると思うけど、まぁキニシナイ。

CC1.x-2.0共通

device emulation

消されたようです。もうデバッガとかが揃ってきたからいらないと判断したんでしょう。実際あれがスレッド作るときなんかひどかったからな・・・。

Surface Reference

確保したCUDA ArrayからSurface Referenceが作れる。Textureと似た操作感。

surface< void, 2 > surRef;

// host code:
...
cudaArray cuArray;
  // malloc (ry
cudaBindSurfaceToArray( surRef, cuArray );
...

// device code:
...
uchar4 data;
surf2Dread( &data, surRef, threadIdx.x, threadIdx.y ); // read
surf2Dwrite( data, surRef, threadIdx.x, threadIdx.y ); // write
...

簡単に言えばread/writeできるTexture Referenceと同じ。ただ、実態はCUDA Array。GlobalとTexture、ArrayとSurfaceのペアということらしい。Textureがread onlyなのに対し、Surfaceはwriteにも使えるのが大きい。L1/L2のキャッシュが絡んでくるので__threadfence系が重要になってくる。
テンプレート第2引数で次元を指定できる。これはTextureと同じ感じ。ただし1または2次元まで。第1引数がvoid固定で説明されてるけどなんなんだろう。これに関しては何も書かれてないっぽい。あとでヘッダ読む。

Surface Referenceの最大サイズはTexture Referenceと同じ。
ちなみにArrayのmalloc時にcudaArraySurfaceLoadStoreフラグを指定する必要あり。

__forceinline__

いままでの__noinline__の類ですね。強制的にinline展開させる指示子です。CC2.0だとrecursive callとかfunction pointerがある関係で必要になったんでしょう。

CC2.0以降のみ

formatted output

はい。キモイのキター。__global__/__device__関数でprintf()が使えますー。ナニコレキモーイ。関数の仕様はC89とほとんど同じっぽい。
デバッグで効力を発揮しますね。

// device code:
...
printf( "threadIdx.x = %d\n", thraedIdx.x );
...

数万スレッドで実行すると愚直に数万の出力が出るので、if文で適当に間引かないと逆に死にます。

hostに転送するバッファのサイズの上限を設けることができるらしい。ただしprintf()しか使わないバッファ。デフォルトだと1MByteだって。

cudaThreadSetLimit( cudaLimitPrintfFifoSize, size );

recursive call/function pointer

__device__の再帰ができます。そして関数ポインタも取れます。いろいろ遊べそうですね。ただ、あんまりやりすぎると相当遅くなるはずです。アーキテクチャ的に。