cuda - 并行归约示例

Question

我从斯坦福找到了这个使用共享内存的并行缩减代码。

该代码是 1<<18 个元素的示例，它等于 262144 并得到正确的结果。

为什么对于某些数量的元素我得到正确的结果，而对于其他数量的元素，如 200000 或 25000，我得到的结果与预期不同？在我看来，它总是在指定所需的线程块

score 2 · Accepted Answer

// launch a single block to compute the sum of the partial sums
block_sum<<<1,num_blocks,num_blocks * sizeof(float)>>>

此代码导致错误。

假设 numblocks 是 13，

然后在内核 blockDim.x / 2 中将是 6，并且

if(threadIdx.x < offset)
{
    // add a partial sum upstream to our own
    sdata[threadIdx.x] += sdata[threadIdx.x + offset]; 
}

只会添加导致错误的前 12 个元素。

当元素计数为 200000 或 250000 时，num_blocks 将是奇数并导致错误，对于偶数 num_blocks 它将正常工作

score 0 · Accepted Answer

该内核对内核的阻塞参数（网格和线程块大小）敏感。您是否使用足够的线程来调用它来覆盖输入大小？

使用 for 循环来制定这样的内核更加健壮 - 而不是：

unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

就像是：

for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
      i < N;
      i += blockDim.x*gridDim.x ) {
    sum += in[i];
}

CUDA 手册中的源代码有很多“阻塞不可知”代码的示例。减少代码在这里：

2 回答 2