CUDA为何无atomicRead？多线程内存访问行为的困惑与求证

两种说法的本质一致性 #

NVIDIA官方文档提到的「不同线程访问同一内存属于未定义行为」，特指非原子的并发读写场景（即一个线程写、另一个线程同时读/写）——这类操作可能导致读取到内存的中间撕裂态，结果完全不可控；而StackOverflow的回答「线程可读写同一内存但无法保证原子性」，只是表述角度不同：它承认操作可以执行，但明确结果不具备原子性（可能得到中间值），本质和官方文档的定义并不冲突。

示例代码的执行结果分析 #

你的示例代码中，线程A执行普通写X=2，线程B执行普通读循环while(X!=2)：

• 线程B不一定会终止，核心原因不是内存被「破坏」，而是两个关键问题：
  1. 缓存可见性：GPU的多级缓存架构可能导致线程B所在的SM核心无法及时看到线程A写入的最新值，即使A已经完成写入；
  2. 编译器优化：编译器可能将X的读操作优化为寄存器缓存，导致线程B永远读取旧值，进入死循环。
• 补充：对于32位对齐的int类型，硬件层面的单次读操作是原子的（不会读取到半写的撕裂值），但可见性问题依然会导致线程B无法感知更新。

为什么CUDA没有atomicRead？ #

CUDA中，对齐的32/64位内存读本身是原子的（无撕裂），但可见性无法保证。而atomicAdd(&X,0)或atomicLoad(&X)这类操作的核心作用，是利用原子指令的内存栅栏语义，强制刷新缓存，确保读取到全局内存的最新值，而非单纯为了避免读撕裂。因此不需要单独的atomicRead，用带内存栅栏语义的原子操作即可实现类似效果。

正确的同步实现方式 #

要保证线程B能感知到线程A的写入，需通过以下方式确保内存可见性：

• 方法一：使用原子操作替代普通读写

  __device__ int X = 1;
  
  __device__ void A() {
    atomicExch(&X, 2); // 原子写，附带内存栅栏
  }
  
  __device__ void B() {
    while (atomicLoad(&X) != 2) {} // 原子读，确保可见性
  }
  

• 方法二：使用内存栅栏指令

  __device__ int X = 1;
  
  __device__ void A() {
    X = 2;
    __threadfence(); // 确保写入全局内存，对其他线程可见
  }
  
  __device__ void B() {
    while (true) {
      __threadfence(); // 确保读取最新的全局内存值
      if (X == 2) break;
    }
  }
  

内容的提问来源于stack exchange，提问作者MaiaVictor