CUDA

CUDA-再看规约-一段规约

一阶段规约

使用原子操作和shared memory的组合可以避免第二个kernel的调用。使用一个（global还是shared）变量记录哪个block已经做完了自己的工作。一旦所有的block都完成后，使用一个block执行最有的规约。

【CUDA专家手册 269页】

更简洁的一段规约

在之前的二段归约中，每一阶段的第一步是每个block的每个thread得到部分和，由于block间是不能同步的，所以才有了第二阶段。

现在是使用一个寄存器partialSum来存储每个thread对应的自己的部分和，最后使用 原子操作对这些部分和再求和：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
__global__ void reductionKernel(int* out, int* in, size_t N){
    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    int stride = blockDim.x * gridDim.x;

    // 每个thread得到自己的部分和，
    int partialSum = 0;
    for (size_t i = tid; i<N; i+=stride){
        partialSum += in[i];
    }

    // 使用原子操作，求部分和的和
    atomicAdd(out, partialSum);
}

void func(int* res, int* in, size_t N){
    // 初始化
    cudaMemset(res, 0, sizeof(int));
    reductionKernel<<<numBlock, threadsPerBlock>>>(answer, in ,N);
}

注意一点，要写入out数组，那么out数组的初始值必须设为0；