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 );