c++ - CUDA:减少变形和 volatile 关键字

Question

阅读以下问题及其答案后
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 关键字会导致性能下降吗？

Answer 1

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

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