cuda – 为什么在共享内存上进行原子操作时,不能使用“volatile”关键字作为共享内存？

2023年11月6日 537次阅读

我有一段CUDA代码,其中线程在共享内存上执行原子操作.我在想,因为原子操作的结果会立即对块的其他线程可见,所以指示编译器使共享内存易失.

所以我改变了

__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函数的定义在
the programming guide中给出.它指示编译器始终为该访问生成读取或写入,并且永远不会将其“优化”为寄存器或其他优化.

由于原子操作是guaranteed to act on actual memory locations(共享或全局),因此两者的组合是不必要的.因此,未提供原型为volatile限定符的原子函数版本.

如果您的内存位置已声明为volatile,则只需将地址传递给原子函数时将其强制转换为相应的非易失性类型.行为将符合预期.(example)

因此,原子操作可以在指定为易失性的位置上操作.

您在代码中的某处使用atomics访问特定位置这一简单事实并不意味着编译器会将其他访问视为隐式volatile.如果您需要其他地方的易变行为,请明确声明它.

点赞