CUDA：在扭曲减少和 volatile 关键字中答案

【问题标题】：CUDA: In warp reduction and volatile keywordCUDA：在扭曲减少和 volatile 关键字中
【发布时间】：2014-01-18 14:33:09
【问题描述】：

阅读以下问题及其答案后
LINK

我的脑海里还有一个问题。来自我的 C/C++ 背景；我知道使用volatile 有它的缺点。并且在答案中还指出，在 CUDA 的情况下，如果不使用 volatile 关键字，优化可以用寄存器替换共享数组以保存数据。

我想知道在计算（总和）减少时会遇到哪些性能问题。例如

__device__ void sum(volatile int *s_data, int tid)
{
    if (tid < 16)
    {
        s_data[tid] += s_data[tid + 16];
        s_data[tid] += s_data[tid +  8];
        s_data[tid] += s_data[tid +  4];
        s_data[tid] += s_data[tid +  2];
        s_data[tid] += s_data[tid +  1];
    }
}

我正在使用减少翘曲。由于所有带有 in warp 的线程都是同步的，因此我相信没有必要使用syncthreads() 构造。

我想知道删除关键字 volatile 是否会弄乱我的总和（由于 cuda 优化）？我可以在没有volatile 关键字的情况下使用这样的缩减吗？

由于我多次使用此缩减功能，volatile 关键字会导致性能下降吗？

【问题讨论】：

标签： c++ cuda reduction

【解决方案1】：

从该代码中删除 volatile 关键字可能在 Fermi 和 Kepler GPUS 上破坏该代码。这些 GPU 缺乏直接在共享内存上操作的指令。相反，编译器必须向寄存器发出加载/存储对。

在这种情况下，volatile 关键字的作用是让编译器尊重加载-操作-存储循环，而不是执行将s_data[tid] 的值保留在寄存器中的优化。保持寄存器中累积的总和会破坏使扭曲级别共享内存总和正常工作所需的隐式内存同步。

【讨论】：

你能给我一些代码，我可以看到这个故障。我正在使用基于 Fermi Arch 的 GPU。正如有人建议的那样，我在没有volatile 的情况下执行了这个减少，结果总和仍然正确。通过代码，我可以看到您提到的负载和将如何兑现。
在“Fermi CUDA 应用程序兼容性指南”pdf 中有完整的讨论。请注意，我说可以中断。是否会取决于编译器、优化设置和内核中的寄存器压力。仅仅因为您的代码与您正在使用的编译器一起工作并不意味着它可以以相同的方式与每个版本的 CUDA 工具包一起工作，或者它对您可能编写的每个内核都是安全的。
因此，最好使用volatile 来确保求和始终按预期工作。感谢我正在寻找的确切答案。