CUDA 缩减优化

Question

我正在尝试进行 Nvidia Reduction 中看到的所有优化。我已经实现了前四个部分，但我仍然坚持第 22 号幻灯片的第 5 部分。

我无法理解为什么提供的代码可以在没有任何 syncthreads() 的情况下工作。线程可以访问输出中的相同内存位置。

此外，幻灯片表明，如果变量未设置为易失性，代码将无法运行。波动在这方面有何帮助？如果我不想调用内核，最好的编程方式是什么？

我也把那个代码放在这里供参考。

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

for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
if (tid < s)
sdata[tid] += sdata[tid + s];
__syncthreads();
}

if (tid < 32) warpReduce(sdata, tid);

在此先感谢您的帮助。如果需要更多信息，请发表评论。

Answer 1

代码依赖于所谓的warp 同步编程。在 warp 中避免 __syncthreads() 是一种常见的做法。然而，这种行为没有记录在案，实际上现在 NVIDIA 强烈反对编写依赖于这种行为的代码。

来自Kepler tuning guide：

The absence of an explicit synchronization in a program where different threads communicate via memory constitutes a data race condition or synchronization error. Warp-synchronous programs are unsafe and easily broken by evolutionary improvements to the optimization strategies used by the CUDA compiler toolchain

您提到的示例包含在 CUDA 工具包附带的示例中。如果你查看最近的版本，你会发现这部分减少现在是通过 warp shuffle 操作实现的计算能力 >= 3.0 并使用 __syncthreads() 作为旧设备你会期望。在较旧的示例中（例如在 CUDA 工具包 6.0 中），它仍然使用 warp 同步技术实现。

如果你还想学习 warp 同步编程，我推荐 this answer。

CUDA 缩减优化

CUDA reduction optimizations

cuda

nvidia

reduction