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

Question

我正在学习 Mark Harris 的优化实施 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;?

Answer 1

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

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

网格迈出的每一步。因此，如果我们根据每步处理的元素数量来衡量网格步幅，则该数字是线程中网格维度的两倍。

线程中的网格大小：

gridSize = blockSize*gridDim.x;

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

gridSize = 2*blockSize*gridDim.x;

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

How to understand the stride size for parallel reduction when doing multiple add?

cuda

reduction