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

在 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 时,实际的并行减少才会触发。相同的代码可以应用于任何其他可交换的关联运算符,例如乘法,最小值,最大值等。

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