cuda - 进行多次添加时如何理解并行减少的步幅大小？

Question

我正在学习 Mark Harris 的 Optimizing Parallel Reduction的实现。

我对第 32 页上的这几行感到困惑：

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + threadIdx.x;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;
while (i < n){
sdata[tid] += g_idata[i] + g_idata[i+blockSize];
i += gridSize;}
__syncthreads();

为什么这里是网格步幅gridSize = blockSize*2*gridDim.x;？根据我的理解，我们应该一个一个地计算网格？所以应该是gridSize = blockSize*gridDim.x;？

标签： cudareduction

Answer 1

这里的答案与您上一个问题的答案非常相似。块（或网格）中的每个线程处理数据集中的两个元素：

sdata[tid] += g_idata[i] + g_idata[i+blockSize]
              element 1     element 2

对于网格的每一步。因此，如果我们根据每步处理的元素数量来衡量网格步长，则该数字是线程中网格尺寸的两倍。

线程中的网格大小：

gridSize = blockSize*gridDim.x;

每步处理的元素数量的网格大小：

gridSize = 2*blockSize*gridDim.x;