在Harris‘s Parallel Reduction教程中，翘曲循环展开是如何工作的？

Question

我正在关注Mark Harris的reduction in CUDA演讲。我已经到了优化步骤#5，我被warpReduce()函数的主要逻辑搞糊涂了：

__device__ void warpReduce(volatile int* sdata, int tid) {
  sdata[tid] += sdata[tid + 32]; // line A
  sdata[tid] += sdata[tid + 16]; // line B
  sdata[tid] += sdata[tid + 8];
  sdata[tid] += sdata[tid + 4];
  sdata[tid] += sdata[tid + 2];
  sdata[tid] += sdata[tid + 1];
}

我的问题是关于A行:为什么我们需要sdata[tid] += sdata[tid + 32]？如果是tid < 32，那么它应该从sdata[tid] += sdata[tid + 16]开始？否则会超出范围吗？

Answer 1

解释是，在对warpReduce()函数的调用中，每个warp处理两个输入元素，因此每个warp 32*2 = 64个元素。

看看你链接的幻灯片中的幻灯片14 -你会看到线程的数量是他们正在处理的元素数量的一半。

但我同意这有点令人惊讶/困惑，因为在之前的幻灯片中，are s的条件是s < blockDim.x，所以只处理bit Dim.x元素。