CUDA by Example積の集約の構成

作者: Jason Sanders,Edward Kandrot,株式会社クイープ
出版社/メーカー: インプレスジャパン
発売日: 2011/02/14
メディア: 単行本（ソフトカバー）
購入: 1人クリック: 36回
この商品を含むブログ (11件) を見る

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];
    }
}