CUDA减少优化

Question

我正在努力做到在Nvidia Reduction上看到的所有优化.我已经实现了前四个部分,但我在第22个幻灯片中遇到了第5部分.

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

此外,幻灯片表明,如果变量未设置为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-synchronous编程.__syncthreads()在经线中避免这种情况是常见的做法.但是,这种行为没有记录,实际上现在NVIDIA强烈反对编写依赖于该行为的代码.

从开普勒调整指南:

在不同线程通过存储器通信的程序中缺少显式同步构成数据争用条件或同步错误.Warp同步程序是不安全的,并且很容易通过CUDA编译器工具链使用的优化策略的进化改进来打破

您提到的示例包含在CUDA工具包附带的示例中.如果你查看最近的版本,你会发现减少的这部分现在用于计算能力> = 3.0的warp shuffle操作实现,并且__syncthreads()可以像你期望的那样用于旧设备.在较旧的样本中(例如在CUDA工具包6.0中),它仍然使用warp同步技术实现.

如果您仍想了解warp-synchronous编程,我推荐这个答案.