CUDA 缩减优化

Question

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

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

此外，幻灯片表明，如果变量未设置为 volatile，则代码将无法工作。不稳定在这方面有什么帮助？如果我不想调用内核，最好的编程方法是什么？

我也将这段代码放在这里以供参考。

__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 同步编程。在扭曲中避免__syncthreads() 是一种常见的做法。然而，这种行为是无证的，实际上现在 NVIDIA 强烈反对编写依赖于这种行为的代码。

来自Kepler tuning guide：

在不同线程通过内存进行通信的程序中没有显式同步构成数据竞争条件或同步错误。 Warp 同步程序是不安全的，并且很容易被 CUDA 编译器工具链使用的优化策略的进化改进破坏

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

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