CUDAをはじめてみる(3)

コードは書けるようになったものの、あまりよくわかっていない。しかし、いちおう解説をしておこう。

GPUは、元々画像処理をするものなので、画像をGPUに送る。データは1次元だが、


const unsigned int size_x = 1024;
const unsigned int size_y = 1024;

というようにしておいて、上の例だと1024*1024の画像だと解釈する(シェーダプログラミングでは、4つのfloatをRGBとアルファチャンネルと解釈して、これを一括して処理していたが、CUDAではそのへんはどうするのか知らない。が、画像処理をするつもりはないので、あまり興味はない)。そして、


dim3 grid(size_x / BLOCK_DIM, size_y / BLOCK_DIM, 1);
dim3 threads(BLOCK_DIM, BLOCK_DIM, 1);

integral_sin<<< grid, threads >>>(d_data, size_x, size_y, nRep, dx);

とすると、GPUでは、BLOCKサイズごとに処理が行われる。ここでBLOCKサイズは、


#define BLOCK_DIM 16

としてあるので、16*16の画像ごとに処理される。そして、その1ピクセルごとにスレッドが同時に走ることになる。本当にそうなるのかはわからないが、そのように解釈しておく。
ちなみに、16というのはサンプルをそのまま使っているだけである。試してみたら、22まではちゃんと動いた。おそらく16のままでよいと思われる。
このように1次元のデータを画像として扱うのが前提となっているので、1次元のインデックスに戻すのに、


unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;
unsigned int index = yIndex * width + xIndex;

という手順を経るのである。