CUDA:在减少翘曲和易失性关键字中

Question

阅读了以下
问题及其答案链接

我脑子里还剩下一个问题。从我在 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

从该代码中删除易失性关键字可能会破坏费米和开普勒 GPU 上的该代码。这些 GPU 缺少直接在共享内存上运行的指令。相反，编译器必须向寄存器发出加载/存储对，从寄存器发出加载/存储对。

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