Harris 的 Parallel Reduction 教程中的 warp 循环展开是如何工作的？

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

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

查看您链接到的幻灯片中的第 14 张幻灯片 - 您会看到线程数是它们正在处理的元素数的一半。

但我同意这有点 surprising/confusing，因为在之前的幻灯片中，加法偏移量 s 具有条件 s < blockDim.x，因此只有 blockDim.x处理元素。

Harris 的 Parallel Reduction 教程中的 warp 循环展开是如何工作的？

How does the warp loop unrolling work in Harris' Parallel Reduction tutorial?

cuda

loop-unrolling