GPU 内存模型 一致性：弱序、可见性层级与 atomic 内存序
在 CUDA 线程层次：grid、thread block 与 thread 的三级组织 里，几万线程同时往同一块全局内存写数据；在 GPU 内存层级：寄存器、shared memory、L2 到 HBM 的容量带宽延迟金字塔 里，这些写要穿过 store buffer、L1、L2 才到 HBM。本篇回答一个容易被「单机单线程直觉」带偏的问题：线程 A 写下去的值，线程 B 什么时候、以什么顺序才能看见？GPU 为了吞吐，默认给出的答案比 CPU 宽松得多——这就是 GPU 内存模型 一致性的核心。相关硬件调度背景见 GPU 架构深入：SIMT、warp 调度，CPU 侧缓存协议背景见 缓存一致性与 MESI。
本文要回答的问题
- 为什么同一个线程「先写 data 再写 flag」，别的线程可能先看到 flag 后看到 data？
__syncthreads()、__threadfence()、volatile、atomicAdd分别保证什么、不保证什么？- thread/block/device/system 四个内存作用域，和可见性层级是什么关系？
- relaxed/acquire/release 这些内存序在 CUDA 里怎么配对使用？
- GPU 弱内存模型和 x86 的 TSO 强模型差在哪，什么样的代码构成数据竞争？
一、为什么 GPU 默认是弱内存模型
CPU 为了让单线程跑得快，也允许乱序执行，但对程序员承诺「看起来像按程序顺序执行」；x86 更进一步用 TSO（Total Store Order，全存储定序）保证：一个核心的写按顺序对其他核心可见，只是读可以提前。GPU 连这个承诺都不给默认值。
原因是吞吐优先。一个 SM 上几十个 warp 抢发射槽，访存请求由内存子系统批量合并、重排后发出（合并访存见 GPU 架构深入）。如果强制每笔写都按线程的程序顺序、立刻对全部线程可见，store buffer 不能攒批、不能跨过慢请求先发快请求，warp 切换隐藏延迟的设计就被废了一半。于是 GPU 的取舍是：
- 单线程内：编译器和硬件保证该线程自己看到的是程序顺序的结果（它读自己刚写的地址能读回来），这是最基本的程序因果。
- 跨线程：默认没有任何可见性顺序保证。线程 A 的两次写，B 看到的顺序可能颠倒；A 写了的值，B 可能在一段时间内还从自己的 L1 旧副本里读；同 warp 内的线程因为锁步执行，现象上「碰巧」有序，但这不是模型承诺，换架构、换编译器版本都可能变。
- kernel 的边界除外：kernel 启动前的写对 kernel 内全部线程可见，kernel 结束后的同步点（如
cudaDeviceSynchronize）保证全部写对 CPU 可见。这是少数免费的全局顺序。
用一个最小的「消息传递」例子体会这件事。线程 A 执行 data = 42; flag = 1;，线程 B 执行 if (flag == 1) assert(data == 42);。在 TSO 的 x86 上这个断言基本不会失败（A 的两笔写按序出去）；在 GPU 默认模型下断言可能失败：flag 的写恰好赶上一次空闲的访存批次先到了 L2，data 的写还堵在本 SM 的 store buffer 里；或者 B 的读从自己 L1 的旧行里取到了 data 的初值。编译器层面还可能叠加一种重排：它发现 flag 先被循环读取，干脆把指令顺序对调以减少寄存器占用。硬件重排和编译器重排是两类独立的威胁，fence 类原语的价值正在于同时压住两者。
一句话：弱内存模型不是硬件缺陷，而是把「要不要强制顺序」的选择权和开销交还给程序员。你不需要顺序时（绝大多数独立元素的数组运算），一分钱不花；需要时，显式插入 fence、barrier 或带内存序的 atomic。
二、内存可见性层级：四个作用域
CUDA 把「写对谁可见」划成嵌套的四个作用域，自小到大：
| 作用域 | 存储位置 | 可见范围 | 对应硬件层级 |
|---|---|---|---|
| thread scope | 寄存器、局部内存 | 仅本线程 | 每线程私有 |
| block scope | shared memory | 同一 thread block 内所有线程 | 片上 SRAM，SM 私有 |
| device scope | global memory | 同一 GPU 上所有 block、所有线程 | L2 + HBM |
| system scope | 可映射/统一内存 | GPU 与 CPU、其他 GPU | PCIe/NVLink 对端 |
作用域越大，让写「真正可见」的代价越高：flush 的缓存层级更多、要通知的 agent 更多。下图画出可见性向外扩散的层级结构，每个层级默认只保证自己范围内的事，跨层必须靠显式原语。
图中有两个要点。其一，作用域是嵌套的，能保证 system 级可见的原语自然覆盖 device 和 block，反过来不成立——__threadfence_block() 出了 block 什么都不保证。其二，作用域与内存类型是两件事：global memory 上同样可以只做 block 作用域的 atomic（cuda::atomic_ref<int, cuda::thread_scope_block>），因为你只关心本 block 的协作线程；作用域说的是「对谁建立顺序」，不是「数据放哪」。
system 作用域在多 GPU 和统一内存场景下尤其绕不开。用统一内存（managed memory）让 CPU 和 GPU 共享同一指针时，页面可以在两端之间迁移，迁移点本身提供了顺序，但你若在 GPU kernel 运行期间让 CPU 并发访问同一块缓冲，就必须用 system scope 的 atomic 或 fence 协调，否则两端各自看到迁移前的快照。多 GPU 之间通过 NVLink 做 P2P 读写也是同理：本 GPU 上 device 级 fence 只保证到本卡 L2，对兄弟卡的缓存毫无约束力。这也是为什么同一套无锁算法从 CPU 移植到 GPU、再扩展到多卡时，往往要把每个原语的作用域逐档放大，每放大一档都对应真实的硬件跳转。
三、block 内一致性：barrier 一举两得
同一个 block 的线程共享一块片上 shared memory，它的一致性由 __syncthreads() 统一兜底。这个 barrier 同时承担两件事：
- 等待：所有线程都到达这一行之后，才允许任何线程继续。没有到达屏障的线程所做的写，不算数。
- 可见性：屏障处强制每个线程此前对 shared memory（以及经 global memory 传递的数据）的写，对过屏障后的所有同 block 线程可见；同时屏障前的读写不会被编译器/硬件重排到屏障后，反之亦然。
经典模式是「一批线程写分片数据 → __syncthreads() → 另一批线程读」。矩阵分块乘法就是这个形状：block 里每个线程先把自己负责的 global memory 元素搬进 shared memory 的小矩阵，第一个屏障保证瓦片完整，所有线程再从瓦片取数做乘加；进入下一轮迭代前再放一个屏障，保证没有线程还在读旧瓦片时它就被新一轮覆盖。这里两个屏障缺一不可：前者管「写完才能读」，后者管「读完才能覆写」，漏掉任何一个，不同线程会在瓦片的新旧版本之间各取各的，算出没有任何规律的错误结果。
典型陷阱有三个：漏掉屏障导致读到半新半旧的中间态；屏障放在条件分支里且分支不被全 block 走，结果部分线程永远等不到（未定义行为，可能直接挂死）；把 barrier 当成跨 block 的工具——它只管 block 内，不同 block 之间不存在这种轻量屏障，只能走 atomic + global memory。还要区分 __syncthreads() 与 __syncwarp()：后者只管同一个 warp 内 32 个线程，Volta 之后 GPU 采用独立线程调度，warp 内线程不再保证锁步，warp 级别的协作（如 shuffle、共享寄存器交换）需要 __syncwarp() 明确同步；但它同样不是 block 级屏障，也不要因为「32 个线程本来一起跑」就省略它。
shared memory 本身不经 L1、按 bank 组织，bank 冲突影响的是性能不是正确性（多体存储原理见 多体存储、交叉访问与冲突）。反过来，bank 不冲突也不代表可见性正确：两个线程写同一个 bank 的不同地址，顺序问题照样存在。
四、fence 家族：强制可见性顺序
fence 不等待任何线程，它只是给当前线程的访存操作立顺序边界，并把本线程的写按作用域刷出去。三个变体按作用域排列：
| 原语 | 保证的作用域 | 大致开销 | 典型用途 |
|---|---|---|---|
__threadfence_block() | block | 低（编译器屏障 + 片上可见） | block 内经 global memory 的协作 |
__threadfence() | device | 中（刷到 L2 可见点） | 跨 block 生产者-消费者 |
__threadfence_system() | system | 高（涉及对端缓存点） | GPU 与 CPU 共享内存通信 |
fence 是单向、单边的：它约束的是执行它的这个线程自己的访存顺序，不会让另一个线程「停下来等」。具体说，线程执行 __threadfence() 时，它此前所有写被保证在 fence 之后的写之前变得对全设备可见，此前的读也有对应的顺序约束；但它既不等待这些写真正到达 HBM，也不对其他线程的访存下达任何命令。所以跨 block 的消息传递必须配合一个 atomic 标志位——fence 保证「data 先于 flag 可见」，对方通过 atomic 看到 flag 后再读 data。只插 fence 而对方用普通读去轮询 flag，编译器完全可能把普通读优化成只读一次寄存器，程序直接死循环。
fence 的开销差异值得量化感知：block 级多数情况下只是编译器屏障加片上写缓冲的排空，几乎不产生额外访存；device 级要确保写到达全设备共享的 L2 并使各 SM 的私有副本失效，代价与一次全局访存同量级；system 级还要让对端（CPU 或其他 GPU）的缓存副本失效，涉及 PCIe/NVLink 上的顺序事务，微秒级延迟都不稀奇。因此把 __threadfence_system() 写在循环热路径里，往往比不加同步的版本慢几十倍，应当只在真正跨设备发布结果时使用。
五、atomic 与内存序：relaxed/acquire/release
volatile 在 CUDA 里的语义比在 CPU 上窄：它只强制每次真的访存、不缓存进寄存器，并附带一些旧架构上的可见性效果；它不保证原子性，也不建立跨线程顺序。两个线程对同一地址做读-改-写，volatile 照样丢更新。把丢更新拆开看就很清楚：counter += 1 是三步——读旧值、加一、写新值。线程 A 和 B 同时读到旧值 7，各自写回 8，最终只加了一次。atomic 操作由硬件把这三步合成对该地址不可分割的一次事务（L2 上有专门的原子处理单元负责仲裁），同一时刻只有一个线程能改成功。需要读-改-写或跨线程顺序，唯一的正规工具是 atomic。
另一个常见误区是把 64 位访问想成「天然安全」。普通的 64 位写在硬件上不保证整体可见，另一个线程可能读到高低两半拼起来的撕裂值；跨线程共享的 64 位标志同样要用 atomicAdd（加 0）或 atomicExch 读。浮点 atomic 过去只支持加法，较新硬件才支持 min/max 与按位比较交换，需要 CAS 实现复合原子操作时要先核对目标架构的能力。
atomic 操作本身（atomicAdd、atomicCAS、atomicExch）保证对同一地址的互斥修改；内存序参数保证它与周围普通访存之间的顺序。自 C++ 语义引入 CUDA（cuda::atomic）后，常用四档：
| 内存序 | 保证 | 适用场景 |
|---|---|---|
| relaxed | 仅原子性，无顺序 | 只做计数、统计，结果与顺序无关 |
| release | 本线程此前的读写不允许重排到其后 | 生产者发布数据 |
| acquire | 本线程此后的读写不允许重排到其前 | 消费者获取数据 |
| acq_rel / seq_cst | 双向 / 全局全序 | 双方都改标志、需要最强顺序 |
release 与 acquire 必须配对在同一个变量上才建立同步关系。下图是标准的单标志消息传递：生产者先写 data，再以 release 发布 flag；消费者以 acquire 读到 flag 为 1 后，必然看到 data 的新值。
老代码里常见的 __threadfence() + atomicAdd(flag, 1, relaxed) 组合，等价于 release；新代码直接用语义明确的 release/acquire 即可。开销上，device 作用域的 global atomic 主要花在 L2 往返（几十到上百周期），system 作用域要走 PCIe/NVLink，能高一个数量级，能用 block scope（shared memory 上的 atomic）就别用 device scope。
除了单次延迟，更要警惕争用放大。如果几千个线程同时 atomicAdd 到同一个全局计数器，所有请求在 L2 同一地址上排队串行化，并行度直接塌缩，这种热点常被称为 atomic 串行点。处理思路按代价排序：先让线程在 shared memory 内按 block 归约（block 内只有几十路争用，且在片上），再由每 block 一个代表向全局计数器加一次，把几万次全局原子降到几百次；能换算法消除共享计数器（如分段计数最后归并）则更彻底。内存序在这里同样要按需选择：纯计数器用 relaxed，结果只取决于最终值；但「计数 + 发布数据」的组合必须 release，否则消费者可能看到计数已增加而数据还没就位。
六、与 CPU TSO 的对比及数据竞争
把两个模型并排看，差异才具体：
| 维度 | x86 CPU（TSO） | GPU（弱内存模型） |
|---|---|---|
| 写-写顺序 | 默认按序对其他核可见 | 默认不保证 |
| 读-读/读-写重排 | 读可提前于写，其余基本禁止 | 各类重排均可能 |
| 跨核可见性 | 硬件 MESI 自动传播 | 不做全局自动一致性，需 fence/atomic |
| 同步原语 | 带 lock 前缀指令 + 内存屏障 | atomic + 内存序 + 分层 fence |
| 默认取向 | 保护程序员直觉 | 最大化吞吐，顺序按需购买 |
TSO 的「强」来自硬件上的两道保险：每个核心有私有的 store buffer，写先进 buffer 再按 FIFO 顺序提交，所以同一核心的写绝不会乱序；同时缓存之间跑 MESI 协议，一个核心要写某行先拿到独占权，其他核心的副本立即失效。GPU 的 SM 数量是 CPU 核心的几十上百倍，若让每次写都在全芯片范围内做这种点对点协商，NoC 会被一致性流量占满，这正是它不做全局自动一致性的物理原因。还有一个叫 store buffering 的经典测试：两个线程各写自己的变量再读对方的变量，TSO 允许两人都读到旧值（写还在各自 buffer 里），而 GPU 上由于读也允许重排，能观察到的反直觉结果更多，只能靠 acq_rel 以上的顺序堵住。
正因为 GPU 没有 MESI 式的自动兜底（协议本身见 缓存一致性与 MESI），数据竞争的判据要严格记：两个线程访问同一内存位置，其中至少一个是写，且这些访问之间没有 happens-before 关系（barrier、release-acquire、kernel 边界等建立的顺序），就是数据竞争，结果是未定义行为——不是「随机值」这么温和，编译器可以据此假设竞争不发生而删掉整段代码。典型的三种错法：用普通写发布标志、用普通读轮询；有 fence 没有 atomic 配对；atomic 用了 relaxed 却指望它建立数据顺序。反过来，happens-before 是可以传递的：A 经 release-acquire 与 B 同步、B 再经屏障与 C 同步，则 A 之前的写对 C 也可见，复杂流水线可以据此逐环建立顺序，而不必在两端之间硬插一个全局 fence。
七、典型规则速查
落地时按这张表对号入座，可以覆盖绝大多数一致性问题：
| 场景 | 正确做法 | 错误做法 |
|---|---|---|
| block 内写后交换读 | __syncthreads() 分隔 | 期望 warp 锁步天然有序 |
| block 内经 global memory 协作 | __threadfence_block() + block atomic | 只用 __syncthreads() 跨 SM |
| 跨 block 生产-消费 | release 写标志，acquire 读标志 | fence 配普通轮询 |
| 多线程累加同一计数器 | relaxed atomicAdd | volatile 变量直接 += |
| GPU-CPU 共享缓冲 | system 作用域 atomic / __threadfence_system() | device fence 后直接通知 CPU |
| 纯独立数组元素运算 | 什么都不加 | 无差别加 fence 拖垮吞吐 |
最后一条最容易被忽视：一致性原语都是有代价的串行点。先判断访问之间是否真的存在「必须被对方按特定顺序看到」的因果，没有因果就不加同步——这正是弱模型把选择权交给你的意义。
这一篇在知识体系里的位置
本篇向上承接 GPU 内存层级：寄存器、shared memory、L2 到 HBM 的容量带宽延迟金字塔（数据流经哪些层级）和 CUDA 线程层次：grid、thread block 与 thread 的三级组织（谁在同时访问），解释这些线程间的可见性规则；与 GPU 架构深入：SIMT、warp 调度 是「性能视角」与「正确性视角」的互补。CPU 侧的自动一致性机制对照见 缓存一致性与 MESI，物理带宽与 HBM 接口实现见 HBM 高带宽内存与物理接口。
总结
GPU 内存模型 一致性的本质是「为吞吐放松默认顺序」：单线程内自洽，跨线程的可见性与定序一律要显式购买——block 内用 __syncthreads() 兼顾等待与可见，跨 block 用 release/acquire 配对的 atomic 发布消息，__threadfence 按 block/device/system 作用域刷写，volatile 只防优化不保证原子与顺序；凡是无 happens-before 关系的并发写冲突都是数据竞争，而没有因果关系的访问则一分钱同步开销都不该花。