CUDA by Example積の集約の構成

CUDA by Example 汎用GPUプログラミング入門

5.2章のDot Productの積について、プログラムを読んだだけでは理解できなかったので図を書いて確かめてみた。

__global__ void dot (float *a, float *b, float *c)
{
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float temp = 0;
    while (tid < N) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }

    // Set Cache
    cache[cacheIndex] = temp;
    // synchronize threads in the block
    __syncthreads ();

    // This code, threadsPerBlock should be power of 2
    int i = blockDim.x / 2;
    while (i != 0) {
        if (cacheIndex < i) {
            cache[cacheIndex] += cache[cacheIndex+i];
        }
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0) {
        c[blockIdx.x] = cache[0];
    }
}

1. まずは、全体の配列の内、半分のスレッドが、残り半分の配列の値を加算する。
2. 次に、4分の1の配列の中で、4/2分の配列の値を加算する。
3. これを最後の配列になるまで繰替えす。