CUDA系统拆解-07-同步、原子操作与内存一致性：并发正确性怎么保证

本文是「CUDA系统拆解」系列第 07 篇。
系列导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
上一篇：CUDA系统拆解-06-共享内存与Coalescing：访存优化先抓哪几个抓手
下一篇：CUDA系统拆解-08-Occupancy、寄存器压力与Launch参数：调优不是把占用率拉满

1. 这篇解决什么问题

  • 为什么并行程序比串行程序更容易出错。
  • 同步在 CUDA 里到底在保护什么。
  • __syncthreads() 真正的作用是什么，为什么它只在 block 内有效。
  • 原子操作解决了什么问题，为什么它又常常带来性能代价。
  • 什么叫可见性和顺序性，为什么这些问题不能只靠“感觉”判断。

2. 先记住的核心结论

  • 并行程序更难，不是因为线程多，而是因为多个线程会同时读写共享状态，导致协调、冲突、可见性问题。
  • __syncthreads() 是 block 级 barrier，它的意义不只是“等一下”，而是建立 block 内的阶段边界。
  • __syncthreads() 只在 block 内有效，不能拿它做跨 block 的同步。
  • 原子操作解决的是共享地址上的并发写冲突，但通常会把热点更新串行化，所以正确不等于高性能。
  • block 独立是 CUDA 的重要设计，这让调度和扩展更灵活，但也意味着很多全局协作要拆成多阶段。
  • 内存一致性和顺序性要单独理解，因为“某线程已经写了”不等于“其他线程现在一定能按你期望的顺序看到”。

3. 正文讲解

3.1 为什么并行程序更难

串行程序里，你通常有一种很强的默认直觉：

  • 先执行 A
  • 再执行 B
  • 再执行 C

所以只要代码顺序写出来了，你就会自然相信后面的语句能看到前面的结果。

并行程序不是这样。到了 CUDA 里：

  • 很多线程同时执行
  • 不同线程推进速度未必一样
  • 不同 warp 之间没有天然顺序
  • 不同 block 之间更没有默认同步

于是问题就变成了：

当多个线程同时读写同一份数据时，谁先写、谁后读、谁能看到什么，不再是天然确定的。

所以并行程序真正难的地方，不是把工作分出去，而是把共享状态管对。

3.2 并行正确性里最核心的三类问题

这一篇最重要的总图可以压成三类。

第一类，协调问题。
也就是线程之间要不要在某个阶段对齐。
典型情况是：

  • 有的线程还没写完
  • 另一些线程已经准备开始读

第二类，冲突问题。
也就是多个线程同时更新同一个地址，导致结果丢失或不确定。

第三类，可见性和顺序问题。
也就是一个线程已经写了数据，但另外的线程什么时候、以什么顺序能可靠看到这次写入。

同步、atomic 和一致性，本质上都在处理这三类问题。

3.3 什么是 race condition

race condition 不是说“大家都在竞争 GPU 资源”，而是：

程序结果依赖于多个线程的执行先后顺序，而这个顺序本身又不受控制。

最经典的例子就是很多线程一起做：

  • 读一个共享变量
  • 改它
  • 再写回去

如果两个线程都先读到旧值，再分别写回，就会发生“丢更新”。

所以你要建立一个很稳的判断：

只要多个线程可能同时写同一个地址，就要立刻警惕 race condition。

3.4 同步到底在解决什么问题

很多人把同步理解成“让线程停一下，等别人”。
这只是表面。

同步更本质的作用有两层：

  • 让一组线程在时间上对齐，形成阶段边界
  • 让这个阶段边界之前的写入，在边界之后对相关线程可可靠使用

也就是说，同步不只是控制时间，更是在建立数据依赖。

这就是为什么同步通常出现在：

  • shared memory 协作
  • 分阶段归约
  • tile 搬运之后
  • 多轮迭代中间

3.5 __syncthreads() 真正是什么意思

__syncthreads() 是 CUDA 里最重要的 block 级 barrier。

它的直觉可以理解成：

同一个 block 内的线程，都必须先走到这里，之后才能继续往后执行。

但更重要的是它的语义含义：

  • 它建立了 block 内的阶段边界
  • 让 barrier 前的 shared memory 写入，可以在 barrier 后被同 block 线程可靠消费

所以 __syncthreads() 最典型的使用场景，不是“为了好看加一个同步”，而是：

  • 先把一块数据搬进 shared memory
  • 同步
  • 再让别的线程读取和复用

3.6 为什么 __syncthreads() 只在 block 内有效

这是这篇必须吃透的点。

__syncthreads() 只能同步：

  • 同一个 block 内的线程

它不能同步：

  • 不同 block
  • 整个 grid
  • 跨 kernel 的阶段

原因不是 API 做得不够强，而是 CUDA 的设计就是这样：

block 是独立调度单位。

不同 block：

  • 可能在不同 SM 上
  • 可能不是同时运行
  • 甚至可能前后分批被调度

如果 CUDA 默认支持轻松的跨 block barrier，调度成本和硬件复杂度都会高很多。

所以 block 独立不是“限制”，而是 CUDA 可扩展和可调度的基础。

3.7 为什么 block 独立是重要设计

block 独立至少带来三件好事：

  • 调度更灵活，SM 可以自由接活
  • 同一 kernel 更容易适配不同规模 GPU
  • 编程模型更清晰，天然鼓励“局部协作、全局分阶段”

这也是为什么很多 CUDA 算法的结构都长这样：

  1. block 内先做局部处理
  2. block 内同步
  3. 产出局部结果
  4. 用下一阶段 kernel 再做更高层合并

所以你以后看到 reduction、prefix sum、softmax 分阶段实现时，要知道这不是“写得麻烦”，而是在适配 CUDA 的协作边界。

3.8 __syncthreads() 的一个典型陷阱

__syncthreads() 不是随便塞进分支里就行。

最危险的情况是：

  • block 内只有一部分线程会执行到 barrier
  • 另一部分线程永远到不了

这样就会产生死锁或未定义行为。

所以非常重要的一条规则是：

对同一个 block 来说，凡是要到达 __syncthreads() 的控制流，必须保证相关线程一致地到达。

这一点在写条件分支、边界处理和循环时尤其容易踩坑。

3.9 原子操作到底解决什么问题

同步解决的是“何时协作”。
原子操作解决的是“同时更新同一个地址怎么办”。

例如一个共享计数器、一个全局和、一个直方图桶，如果很多线程都对同一位置做读改写，那么普通写法会丢更新。

原子操作的含义就是：

对某个共享地址上的读改写过程，提供不可分割的并发保护。

所以 atomicAdd、atomicMax 这类原语的本质，不是“更高级的加法”，而是：

  • 正确处理共享热点更新
  • 防止多个线程互相覆盖结果

3.10 为什么原子操作对，但常常慢

原子操作能保证正确，但它常常会引入一个新的现实：

热点地址上的更新会变得高度串行化。

如果几千个线程都在抢同一个地址做 atomicAdd，那就意味着：

  • 逻辑上结果是对的
  • 但硬件上大家都在争同一个共享点

于是吞吐会明显下降。

所以原子操作的成熟理解应该是：

  • 它是正确性工具
  • 不是默认高性能方案

高性能实现通常会想办法减少 atomic 的次数和冲突范围。

3.11 为什么常见优化是“先局部，再全局”

这条思路你以后会反复看到。

最差的做法通常是：

  • 每个线程直接对同一个全局地址做 atomic

更好的做法通常是：

  1. warp 内先聚合
  2. block 内再聚合
  3. 最后只让少量线程做全局 atomic

这样做的意义非常直接：

  • 降低全局热点
  • 减少 atomic 次数
  • 让大部分协作先在更近、更便宜的层级完成

所以很多高性能 reduction、softmax、统计类 kernel 的主线，本质上都是：

先局部合并，再少量全局提交。

3.12 内存一致性和顺序性到底在说什么

这个词听起来抽象，但你可以先把它翻译成一个非常具体的问题：

一个线程写了数据，别的线程什么时候能看到？看到时顺序是不是你想要的那个顺序？

这里要特别警惕两个误解：

  • 写了，不等于别的线程立刻可靠可见
  • 看到了，不等于顺序就是你以为的顺序

所以一致性和顺序性要单独理解。
它们不是“补充知识”，而是并行正确性的基础。

对当前阶段来说，你至少要建立下面这个直觉：

  • shared memory 协作通常要靠 barrier 建立“先写后读”的阶段边界
  • 跨 block 的复杂协作，通常不该在一个 kernel 内硬做
  • kernel 边界本身就是很强的全局阶段边界

3.13 为什么很多跨 block 协作要拆成多个 kernel

这是 CUDA 很重要的工程哲学。

因为 block 间没有天然 barrier，所以很多算法会故意拆成多阶段：

  1. 第一个 kernel 做局部工作
  2. kernel 结束，形成全局阶段边界
  3. 第二个 kernel 再读取前一阶段结果

这种写法的好处是：

  • 正确性更清晰
  • 调度更稳定
  • 更符合 CUDA 的 block 独立模型

所以以后看到“多 kernel 分阶段实现”，不要第一反应觉得它笨。很多时候，这恰恰是更对的设计。

4. 和 AI 推理的关系

这篇和 AI 推理的关系非常直接，因为推理里的很多 kernel 都有局部协作和共享状态问题。

典型场景包括：

  • softmax、LayerNorm、RMSNorm：要做 block 或 warp 内归约
  • attention tile 计算：先搬数据、同步，再进行阶段性计算
  • 统计、计数、buffer slot 分配：经常涉及 atomic
  • 推理框架里的共享元数据更新：本质上也是并发读写协调问题

所以做 AI infra / 推理时，你不能只把这些当成“CUDA 基础题”，而要知道：

  • barrier 放错了会直接错
  • atomic 用多了会直接拖性能
  • block 边界理解错了，算法结构就会从根上出问题

5. 常见误区

  • 加了 __syncthreads() 就一定安全。不是，它只解决 block 内 barrier 问题。
  • __syncthreads() 可以拿来同步整个 kernel。不是，它不跨 block。
  • atomic 能解决所有并发问题。不是，它只解决特定共享更新的正确性问题。
  • 为了保险，多加同步总没错。不是，过多 barrier 会拖慢执行，而且不一定真解决逻辑问题。
  • 结果偶尔对，说明写法大体没问题。不是，并发 bug 最危险的就是不稳定和偶发。
  • 多 kernel 分阶段实现说明实现不高级。不是，这往往正是在适配 CUDA 的同步边界。

6. 复习自测

  • 为什么并行程序比串行程序更难？
  • 同步在 CUDA 里到底在保护什么？
  • __syncthreads() 的真实作用是什么，为什么它只在 block 内有效？
  • 为什么 block 独立是 CUDA 的重要设计？
  • 原子操作到底解决什么问题，为什么它常常会带来性能代价？
  • 为什么高性能实现通常更喜欢“先局部聚合，再少量全局提交”？
  • 内存一致性和顺序性为什么要单独理解？
  • 为什么很多跨 block 协作更适合拆成多个 kernel 分阶段完成？

系列导航