当对共享内存执行原子操作时，为什么不能对共享内存使用“ volatile”关键字？

Question

我有一段CUDA代码，其中线程在共享内存上执行原子操作。我一直在思考，因为原子操作的结果无论如何对于该块的其他线程都是立即可见的，因此最好指示编译器具有共享内存volatile。
所以我改变了

__global__ void CoalescedAtomicOnSharedMem(int* data, uint nElem)
{
    __shared__ int smem_data[BLOCK_SIZE];
    uint tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for ( uint i = tid; i < nElem; i += blockDim.x*gridDim.x){
        atomicAdd( smem_data+threadIdx.x, 6);
    }
}

至

__global__ void volShared_CoalescedAtomicOnSharedMem(int* data, uint nElem)
{
    volatile __shared__ int smem_data[BLOCK_SIZE];
    uint tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for ( uint i = tid; i < nElem; i += blockDim.x*gridDim.x){
        atomicAdd( smem_data+threadIdx.x, 6);
    }
}

发生以下更改时发生以下编译时错误：

error: no instance of overloaded function "atomicAdd" matches the argument list
        argument types are: (volatile int *, int)

为什么不volatile支持将地址用作原子操作的参数？是因为编译器一经识别到将要进行原子操作，便已将共享内存视为易失性？

Answer 1

编程指南volatile中给出了限定符的定义。它指示编译器始终为该访问生成读取或写入，并且永远不会将其“优化”到寄存器或其他优化中。

由于原子操作保证作用于实际的内存位置（共享的或全局的），因此两者的组合是不必要的。因此，volatile未提供针对限定符原型化的原子函数的版本。

如果您有一个内存位置已声明为，只需在将地址传递给原子函数时volatile将其转换为相应的非类型即可。volatile行为将如预期。（示例）

因此，原子操作可以在volatile按此附带条件指定的位置上进行操作。

您在代码中的某处使用原子访问了特定位置这一简单事实并不意味着编译器会将其他地方的每个访问视为隐式volatile。如果您需要volatile其他地方的行为，请明确声明。