为什么在共享内存上执行原子操作时无法使用"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行为，请显式声明它。

Answer 2

前面的海报已经正确地识别了问题：没有定义接受volatile参数的atomicAdd函数。

你的问题是为什么会出现这种情况，我猜你的库开发人员只是忽略了那个接口。想象一下volatile、const的所有组合，以及可能的参数和潜在接口的数量开始爆炸。

为什么不支持将volatile地址作为原子操作的参数？

原子操作不是C/C++的一部分。在您的案例中，它们是在一个可能是用汇编语言实现的库中实现的。

是因为编译器在识别出将要对共享内存进行原子操作时就已经将其视为易失性内存了吗？

不，这是库编写器定义函数接口的方式。