交换运算符的单块并行约简

在 CUDA 中并行缩减的最简单方法是分配一个块来执行任务：

static const int arraySize = 10000;
static const int blockSize = 1024;

__global__ void sumCommSingleBlock(const int *a, int *out) {
    int idx = threadIdx.x;
    int sum = 0;
    for (int i = idx; i < arraySize; i += blockSize)
        sum += a[i];
    __shared__ int r[blockSize];
    r[idx] = sum;
    __syncthreads();
    for (int size = blockSize/2; size>0; size/=2) { //uniform
        if (idx<size)
            r[idx] += r[idx+size];
        __syncthreads();
    }
    if (idx == 0)
        *out = r[0];
}

...

sumCommSingleBlock<<<1, blockSize>>>(dev_a, dev_out);

当数据大小不是很大（几个 thousants 元素）时，这是最可行的。当减少是一些更大的 CUDA 计划的一部分时，通常会发生这种情况。如果输入从一开始就匹配 blockSize，则可以完全删除第一个 for 循环。

请注意，在第一步中，当元素多于线程时，我们完全独立地添加内容。只有当问题减少到 blockSize 时，实际的并行减少才会触发。相同的代码可以应用于任何其他可交换的关联运算符，例如乘法，最小值，最大值等。

注意，可以使算法更快，例如通过使用扭曲级并行缩减。