CUDA:在warp reduction和volatile关键字中

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];
    }
}

我正在使用减少经线.由于warp中的所有线程都是同步的,因此我认为不需要使用syncthreads()构造.

我想知道将删除关键字volatile弄乱我的总和(由于cuda优化)？没有volatile关键字可以使用这样的减少.

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

Answer 1

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

volatile关键字在此上下文中的作用是使编译器遵循加载 - 操作 - 存储周期,而不执行将保持s_data[tid]寄存器值的优化.保持寄存器中累加的总和将破坏使得该warp级共享存储器求和正确工作所需的隐式存储器同步.