CUDA缩减优化
CUDA缩减优化
提问于 2016-05-10 22:49:46
我正在尝试执行在Nvidia还原中看到的所有优化。我已经实现了前四部分，但我坚持第5部分在幻灯片22。
我无法理解为什么所提供的代码可以在没有任何sync线程()的情况下工作。线程对输出中相同的内存位置具有访问权限。
此外，幻灯片还表明，如果变量没有设置为易失性，代码将无法工作。反复无常在这方面有什么帮助呢？如果我不想调用内核，那么最好的编程方法是什么呢？
我也把这段代码放在这里作为参考。
__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);提前谢谢你的帮助。如果需要进一步的信息，请发表评论。
回答 1
Stack Overflow用户
回答已采纳
发布于 2016-05-10 23:13:51
代码依赖于所谓的翘曲同步编程。在翘曲中避免__syncthreads()是一种常见的做法。然而，这种行为是没有记录的，实际上现在NVIDIA强烈反对编写依赖于这种行为的代码。
来自开普勒调谐导轨
在不同线程通过内存进行通信的程序中，缺少显式同步构成了数据竞争条件或同步错误。曲时同步程序是不安全的，很容易被CUDA编译器工具链使用的优化策略的进化改进所破坏。
您提到的示例包含在CUDA工具包附带的示例中。如果您查看最近的版本，您会发现，这部分简化现在是用翘曲洗牌操作来实现的，用于计算能力>= 3.0，并像您预期的那样对旧设备使用__syncthreads()。在旧的示例中(例如，在CUDA工具包6.0中)，它仍然是用warp同步技术实现的。
如果您还想学习有关曲同步编程的知识，我建议使用这个答案。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/37155072
复制相关文章
相似问题
腾讯云开发者
