Runguo Li

Runguo Li

M.S. student in Information Science @ UIUC, advised by Prof. Minjia Zhang. Working on ML systems for large-scale models: MoE inference and expert streaming, distributed training and rollout, GPU compiler correctness.

H100 · 异步与屏障（Asynchronicity and barriers）

目录 · ← l2 · l4 →

H100 · 异步与屏障（Asynchronicity and barriers）

本文档基于课程讲义《3. Asynchronicity and barriers.pdf》（Lesson 3，50 页）整理， 系统介绍 H100 的异步执行模型与同步原语（mbarrier）。 前置阅读：H100-架构介绍.md（TMA、异步拷贝、Tensor Core）与 H100-Clusters-数据类型-内联PTX-状态空间.md（PTX、状态空间、mapa）。

目录

  1. 同步世界与异步的动机
  2. 阻塞 vs 非阻塞操作
  3. GPU 中的异步：H100 的真正天赋
  4. 异步操作的三阶段模式
  5. 两大问题与 mbarrier 的引入
  6. mbarrier：硬件加速的同步原语
  7. Proxy（代理）模型
  8. 冒险（Hazard）：RAW 与 WAR
  9. fence（栅栏）
  10. mbarrier 完整流水线
  11. phase 与 expected transaction count
  12. mbarrier 初始化与关键注意点
  13. 同步：try_wait / test_wait / acquire
  14. producer / consumer warp 分工
  15. mbarrier 复用模式
  16. mbarrier 不止用于异步拷贝（WAR 防护）
  17. arrive 的变体与 release 语义
  18. barrier.cluster（跨块屏障）
  19. 异步组（bulk_group / commit_group / wait_group）
  20. namedbarriers（命名屏障）
  21. 总结与学习衔接

1. 同步世界与异步的动机

任何系统都有两类主要操作：

  • Doing（做）：处理信息、解方程、做饭……
  • Fetching（取）：取数据、读题、备料……

同步世界里二者串行：你被”阻塞”直到当前任务完成——做饭时不做别的，做完才做下一件。

异步（Asynchronicity）把”请求（Request）”与”结果（Result）”解耦： 做饭时去做别的事，过一会儿收到信号再回来用做好的饭。→ 在”做一件事的时间里完成了两件事”， 这就是延迟隐藏（Latency Hiding）。

本质：如果”取数据”的代价高于”处理数据”的时间，就必须让取下一项与处理当前项重叠。 延迟隐藏的核心是：不被某条指令阻塞，能把当前任务放到后台、去干别的活。

2. 阻塞 vs 非阻塞操作

 阻塞 / 同步非阻塞 / 异步
行为暂停发起线程直到操作完全完成立即把控制权还给发起线程（操作尚未完成）
控制权不返回，无法执行下一行立即返回，继续执行
同步点隐式同步点，保证任务完成后才继续无隐式同步，需后续显式等待

3. GPU 中的异步：H100 的真正天赋

  • 现代计算里，”做”（算力）极快，”取”（访存）慢得痛苦——取数据的时间远高于计算它的时间。
  • 因此高性能架构的目标不只是让数学更快，而是确保数学永不停止。
  • H100 常被夸”算力强（FLOPS）”，但它真正的天才在于异步架构：确保巨大的 Tensor Core 永不等数据。
  • 实现方式：为 Tensor Core 和 TMA 使用非阻塞指令。

4. 异步操作的三阶段模式

阶段内容
Stage 1 初始化一个线程发起异步操作，立刻去执行下一条指令
Stage 2 跟踪与并行执行系统某部分跟踪该操作（TMA 用 mbarrier，wgmma 用内部硬件 scoreboard）；同时其它单元并行工作
Stage 3 同步后台操作完成后，进行同步

4.1 H100 里的同步步骤

  • warp 发指令极快（纳秒级）。
  • 异步指令到达执行单元后，立刻尝试读内存。
  • 瓶颈是内存带宽：从 HBM 搬到 shared/registers 是几百纳秒的高延迟操作（相对逻辑核心）。

4.2 由此产生的两大问题

  1. 发令者（~ns）比搬数据者（~几百 ns）快 → 若拷贝与计算同时发指令，Tensor Core 要等几百 ns 等数据到达 → 存在延迟隐藏的空间。
  2. 大 kernel 里指令队列塞满”数据还没到”的命令时，若无同步，Tensor Core 可能在数据到达前就执行了操作。

所以我们需要：① 保证对的数据上做对的操作；② 一个能实现延迟隐藏的系统。 答案就是 mbarrier。

5. 两大问题与 mbarrier 的引入

mbarrier 解决了上面的两个问题：

  1. 让慢操作独立于快操作发起——等快单元启动时数据已就位。
  2. 防止快执行单元在错误的数据上操作。

6. mbarrier：硬件加速的同步原语

mbarrier 是驻留在 Shared Memory 中的、硬件加速的同步原语，用于跟踪异步内存事务的完成。

  • 与传统 barrier（等”线程”到达）不同，mbarrier 是等”数据”到达。
  • 它把 Producer（发起拷贝的一方） 与 Consumer（等数据的一方） 解耦， 实现分阶段（split-phase）、”发完即忘（fire-and-forget）”的内存流水线。
  • 用法：Producer 在 barrier 里设定期望的传输大小，硬件在后台干活，事务完成后 barrier 就打开。

一句话：mbarrier 让”快的发令者”尊重”慢的物理硬件”——用硬件加速的计数来协调”数据到了没”。

7. Proxy（代理）模型

在 H100 / CUDA PTX 内存模型里，Proxy 用于区分”谁在执行内存操作”：

  • 两个内存操作发生在同一个 proxy（如 Generic Proxy）→ 硬件保证它们（大体）安全且有序。
  • 操作 A 在 Proxy 1、操作 B 在 Proxy 2 → 硬件停止检查，认为二者完全无关，让它们乱序、并行地跑。

7.1 Generic Proxy vs Async Proxy

 Generic Proxy（通用代理）Async Proxy（异步代理）
谁写 kernel 的线程执行顺序 load/store硬件机制做 bulk 异步拷贝 / Tensor Core 运算
举例普通线程读写cp.async.bulk、wgmma
关系线程只是”踢一脚”命令就继续完全独立运行，不知道线程在干嘛，线程也不知道它何时完成

7.2 两条不同的内存通路（关键！）

  • Generic Proxy 走 SM 的 L1 cache。
  • Async Proxy 绕过 L1，直接与 L2 / HBM 交互。
  • Generic Proxy 的 store 先停留在本地 store buffer 或 L1，不会立即对系统其余部分（含 TMA）可见。

8. 冒险（Hazard）：RAW 与 WAR

由于两条 proxy 走不同通路，产生两类经典数据冒险：

8.1 RAW（Read After Write，读后写冒险）

  • 线程（Generic Proxy）写了地址 X → 立刻让 Async Proxy 去读同一地址 X。
  • 但新数据还卡在 Generic Proxy 的 L1 里 → Async Proxy 会从 L2/DRAM 读到旧数据（stale）。

8.2 WAR（Write After Read，写后读冒险）

  • 线程（或 CTA）用 Generic Proxy 读了 X，把旧版本”锚”在 L1。
  • 之后 Async Proxy 通过 L2/HBM 更新了 X，但没有更新/失效 Generic Proxy 的 L1 视图。
  • 两种失败模式：
    1. 之后的 Generic store/writeback/eviction 把 Async Proxy 写的新值覆盖（旧行胜出）；
    2. 之后的 Generic load 一直返回旧值。

结论：跨 proxy 访问同一地址时，必须用 proxy fence 同步（见 §9）。

9. fence（栅栏）

fence 是显式的”排序/可见性”点，约束内存效果何时变得可观察——尤其当异步操作不会像你假设的那样 与普通 load/store 自动排序时。

  • NVIDIA 明确指出：跨 proxy 需要 proxy fence 才能正确排序。
  • fence 是有作用域的（.cta / .cluster / .gpu / .sys），作用域决定”谁必须看到该排序”， 对应层级中的一致性点（如 L1 vs L2）。
  • 两种类型：普通 fence 与 跨 proxy fence。

9.1 Release fence（生产者侧）

单向规则：所有在 release 之前（程序顺序）的内存操作（尤其是写），对”与之同步的其它线程”而言， 都可见于 release 之后出现的任何操作之前。 它防止先前的写被延迟/重排到 release 点之后。

9.2 Acquire fence（消费者侧）

相反的单向规则：acquire 之后（程序顺序）的内存操作（尤其是读）不允许被观察到发生在它之前。 acquire 之后，线程保证能观察到”由匹配的 release 变为可见”的那些写。

9.3 Cross-proxy fence

  • 跨多个 proxy 访问同一地址时，需要跨 proxy fence。
  • 对 async proxy 用 fence.proxy.async 同步 generic 与 async proxy 之间的内存。
  • 作用：排空/排序 generic-proxy 的 shared-memory 写可见性，让 async proxy 不读到旧视图。
  • 它不是”集体 flush”，而是每线程排序——所以仍需一次 block 同步（如 __syncthreads()）， 确保所有写者都做完，再由被选中的线程启动 TMA。

10. mbarrier 完整流水线

步骤操作
Step 1在 shared memory 中初始化 mbarrier，用 mbarrier.init 创建带期望线程到达数的 barrier
Step 2用 mbarrier.arrive 为即将发起的异步操作设定期望事务数（expect_tx）；每次调用都记录一次指令、递减到达数
Step 3发起异步操作（如 cp.async.bulk）
Step 4发起其它操作并等待（线程休眠，直到”线程到达数”和”expected_tx”都归零）
Step 5翻转 phase、重置到达数，进行下一次操作

11. phase 与 expected transaction count

  • phase：barrier 当前可复用的状态/周期。它是单个 bit，每完成一个周期就翻转一次。
  • transaction count：你正在执行的异步操作的大小。
  • expected transaction count：异步操作”还差多少工作”。
    • 因为是异步操作，硬件会在异步操作推进时自动递减事务计数（把 barrier 挂到异步操作上就为这个行为）。
  • 复用 barrier：在操作末尾翻转 phase、并增大事务计数；线程到达数会按初始化时设定的值自动重置。

12. mbarrier 初始化与关键注意点

mbarrier.init 是”几乎所有 bug 的源头”——这里错了，其它都没意义。

mbarrier.init.shared::cta.b64 [addr], count;
  • addr：mbarrier 在状态空间中的内存地址（shared memory）。
  • count：期望的线程到达数——也是复用 barrier（翻转 phase）时 barrier 重置到的值。

12.1 必须牢记的点

  • 作用域通常用 shared::cta：只有同一线程块内的线程能”看到”该 barrier。
  • 地址是 shared memory 指针，用 __cvta 创建。
  • 用 shared::cluster 可让 barrier 对整个 cluster 的所有线程可见；此时地址需用 mapa PTX 指令获取。
  • mbarrier 对象必须 64 位对齐；非对齐访问可能造成静默数据损坏。

12.2 设置期望事务数（mbarrier.arrive.expect_tx）

  • 把到达数减 1。
  • 把 tx_count 字节加到 barrier 的待处理事务数上。
  • 累加语义：若调用两次 arrive.expect_tx 各带 4096，则 barrier 总共期望 8192 字节。

12.3 占位符 _

  • 指令里的 _ 是 phase_out 的占位符——本可拿到一个含”当前 phase 与事务数”的 64 位编码 token。
  • 它用于某些复杂同步场景；我们丢弃 phase_out，因为把信息写进 producer 线程的寄存器里更省 （寄存器压力昂贵，翻转 1 位整数比写 64 位 token 便宜）。

13. 同步：try_wait / test_wait / acquire

发起异步拷贝后，怎么知道操作完成了？

  • 老方法 barrier.sync：阻塞线程直到操作完成——破坏异步性。
  • Hopper 的方法：用一个能返回”完成 true / 未完成 false”的指令 → 实现延迟隐藏： 线程检查 barrier、发现没好就重复，直到操作结束、phase 翻转。这正是 mbarrier.try_wait.parity 的工作。

13.1 mbarrier.try_wait.parity

mbarrier.try_wait.parity.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint;
  • phaseParity：传入 0 或 1，代表”必须确认已完成的那个 phase”。
  • 硬件把你的 parity 与 barrier 当前内部 parity 比较：
    • 相等 → 该 phase 仍在处理 → 返回 false（继续等）；
    • 不等 → barrier 已进入下一 phase → 返回 true（继续）。
  • 成功返回意味着 barrier parity 已”翻转”，即你请求的 parity 现在指的是刚完成的前一个 phase。
  • suspendTimeHint：可选立即数，告诉调度器”条件不满足时让线程让出多久”。

13.2 mbarrier.test_wait.parity

  • 与 try_wait 类似，用 phase parity 位（0/1）而非原始整数计数跟踪进度。
  • waitComplete（目标）：1 位谓词寄存器——barrier 完成 → 1，仍在忙 → 0。

13.3 mbarrier.wait 带 .acquire

  • 做 try_wait 的一切，但有一个关键区别：若当前同步轮次已完成，.acquire 语义会建立严格的 memory fence， 保证 async proxy 或其它线程的所有数据写在继续前一定可见。
  • 例：从 shared memory 加载数据到寄存器（手动喂 Tensor Core）时，.acquire 确保这些 LD 指令 不会在数据有效前发射——否则寄存器里是垃圾。
  • 即使 wgmma 直接读 shared memory，指令本身也需要描述符与内存状态一致；.acquire 保证依赖链被遵守。

14. producer / consumer warp 分工

  • Producer warp：只有单个线程在执行 mbarrier.init、mbarrier.arrive.expect_tx、cp.async.bulk 等指令； warp 里其它线程只是闲着。
  • Consumer warp：执行 wait 操作，等 TMA 拷贝完成后才能处理数据。

15. mbarrier 复用模式

  • phase 声明在寄存器里，不是 shared memory！ 每个线程各自维护一个 int phase。
  • 循环的最后一步做 phase ^= 1（翻转）。
  • 通过翻转 phase，同一个 mbarrier 对象可被反复复用（下一轮 tile）。

16. mbarrier 不止用于异步拷贝（WAR 防护）

mbarrier 不只是为 cp.async.bulk 服务的。考虑：

Producer 写 shared memory → Consumer 用 WGMMA 读 shared memory
→ Producer 想用下一个 tile 覆盖同一块 shared memory
  • 若 Producer 在 WGMMA 还在读 sA 时就覆盖它 → WAR 冒险，数学结果就是垃圾。

16.1 如何防 WAR

  1. 创建带”期望线程到达数”的 barrier。
  2. Consumer 忙着读数据，Producer 在 barrier 上自旋（spin）。
  3. Consumer 执行完最后一条读指令。
  4. Consumer 执行递减线程到达数的指令。
  5. barrier 状态翻转 → Producer 现在可以写新数据。

递减到达数的指令是 mbarrier.arrive（见 §17）。

17. arrive 的变体与 release 语义

17.1 mbarrier.arrive

  • 只是把 barrier 的待处理线程到达数减 count。

17.2 编译器重排序（为什么需要 release）

编译器常为效率重排操作：

a = 1; b = 2; c = a + b;   // 编译器可能先做 b=2 再做 a=1，认为不影响 c
  • 多线程代码里重排会破坏正确性：不同线程会观察到不同顺序。
    // 线程 a           // 线程 b
d[0][0] = 42.0f;    while (flag == 0);
flag = 1;           float x = d[0][0];

    若 flag=1 被重排到 d[0][0]=42 之前，线程 b 会读到未初始化的 d[0][0]。

17.3 mbarrier.arrive.release

  • .release 后缀提供 release 语义：单向硬件 fence——它上方的内存写不能重排到它下方。
  • 用法：线程完成”产出将被其它线程消费的数据”之后，调用 mbarrier.arrive.release。

17.4 mbarrier.arrive_drop（及其后缀）

  • 行为类似标准 arrive（递减当前 phase 的待处理到达数）。
  • 但更重要：它永久递减 barrier 的期望到达数——若之后翻转 phase 复用 barrier， “期望到达数”将是”原值 − 已发生的 drop 数”。
  • 后缀组合：
    • arrive_drop.expect_tx：设定期望事务数，并永久 + 临时地把到达数减 1。
    • .sem（sem 可为 .release 或 .relaxed）：
      • .relaxed：不强制任何内存排序；
      • .release：保证该线程在到达前做的所有写，对任何等待该 barrier 的线程可见。
    • .noComplete：让硬件执行到达（递减 pending/expected 计数）但不触发 phase 完成， 即使”pending 计数归零”的完成条件已满足。

17.5 销毁 barrier（mbarrier.inval）

  • 正式使 mbarrier 对象（64 位同步原语）失效，抹掉硬件对该 barrier 状态的跟踪。
  • 释放该 shared memory 地址，使其可被安全覆盖或改作他用。
  • 它把通用指针转换成 32 位 shared memory 偏移，确保 GPU 硬件寻址到正确的本地内存 bank。

17.6 完整流水线回顾

1. mbarrier.init          ：在 shared memory 创建 barrier，设期望线程数，起始 phase=0
2. mbarrier.arrive.expect_tx：producer 告诉 barrier 还要期待即将到来的异步操作的字节数
3. cp.async.bulk          ：发起拷贝，然后 mbarrier.try_wait.parity 0 自旋，直到 phase 0 完成
4. Flip                   ：所有期望字节（和线程）到达 → barrier 自动翻到 phase 1
5. 开始计算
6. 计算结束 → 手动翻转 phase → 开始下一个 k tile（复用 barrier）

18. barrier.cluster（跨块屏障）

  • 标准 barrier（如 bar.sync / __syncthreads）只能同步单 block 内的线程，无法跨 cluster 协调 producer/consumer 块。
  • barrier.cluster 用于同步同一 cluster 内共调度的不同线程块，并保证 barrier 之前对 DSMEM 的写在 barrier 之后对所有块可见。
  • 关键用例：TMA Multicast。初始时 block 0 可能在跑而 block 1 还没跑； 若 block 0 直接写 block 1 会崩溃。启动 barrier.cluster 就能确保每个 block 物理上都已就位。

18.1 两个重要指令

  • barrier.cluster.arrive：线程发出”到达”信号，不停止，继续执行不依赖其它块数据的独立指令（数学、本地寄存器操作）。
  • barrier.cluster.wait：阻塞指令，线程停在此处，直到 cluster 内其它所有线程/块都发出了 arrive。 一旦解除阻塞，就保证其它块写的所有数据现在可安全读取。

H100 上，因为 Block A 能写 Block B 的内存，就需要 barrier 防止 Block B 在 A 写完之前读。

19. 异步组（bulk_group / commit_group / wait_group）

19.1 bulk_group

  • 把 bulk_group 挂到 cp.async.bulk 指令上，就能用 cp.async.bulk.commit_group 和 cp.async.bulk.wait_group。
  • 用途：批量发起一堆指令，再统一等待。

19.2 cp.async.bulk.commit_group

  • 用 cp.async.bulk 发起一批操作（配 bulk_group）。
  • 之前没有 commit_group → 这些拷贝是”未提交”的；commit 即”把它们打包成一个组”。
  • 打包成组的目的：之后能让其它单元等到这 N 个单元执行完。
  • 该指令之后发起的指令属于下一组（或只是普通指令）。
  • 可以有多个组，每组多个 cp.async.bulk 操作。

19.3 cp.async.bulk.wait_group<N>

  • 一批操作已发起并提交后，让要用其输出的其它单元等待。
  • wait_group<N> 等到”最多还剩 N 个已提交的组处于 pending”：
    • wait_group<0>：等所有组完成；
    • wait_group<2>：等到你发起的所有操作里只剩最近 2 个仍 pending。
  • 计数指”仍 pending 的组”，不是”已完成的组”。

19.4 cp.async.bulk.wait_group.read

  • 既停顿执行，又强制一个 Acquire Fence。
  • 对 read-after-write 场景至关重要（否则可能读到旧数据）。

20. namedbarriers（命名屏障）

  • __syncthreads() 是全 block 的汇合点。
  • 命名屏障让你在一个 block 内创建多个独立的同步点，让不同 warp 子集各自协调，不必拖住无关 warp。
  • PTX 明确允许不同 warp 用不同的操作使用同一个命名屏障，例如混合 .arrive 与 .sync 来构建 producer/consumer 流水线。
bar.sync a, b;   // a = 屏障 ID，b = 参与线程数

21. 总结与学习衔接

21.1 核心脉络速记

概念一句话
异步动机取数据远慢于算数据 → 让”取下一块”与”算当前块”重叠
ProxyGeneric（走 L1）vs Async（绕过 L1 直连 L2/HBM），二者互不自动有序
HazardRAW（async 读到 L1 里的旧数据）、WAR（L1 旧值覆盖/读旧）
fencerelease（生产者）/ acquire（消费者）/ fence.proxy.async（跨 proxy）
mbarrier硬件原语，等数据到达而非等线程到达
phase1 bit，每完成一轮翻转；复用即翻转 phase
try_wait.parity非阻塞轮询，返回 true/false，保留延迟隐藏
wait.acquire完成后强制 acquire fence，保证数据可见
arrive_drop永久递减期望到达数
barrier.cluster跨块同步，TMA Multicast 的前置
commit/wait_group批量异步拷贝的分组与等待（wait_group<0> 全等）
namedbarriersblock 内多个独立同步点

21.2 与本课程其他内容的衔接

本文概念对应后续专题
TMA 异步拷贝 + mbarrier《4. cuTensorMap》《5. cp.async.bulk》
wgmma scoreboard / 异步组《6. WGMMA-1》《7. Wgmma part 2》
cluster / DSMEM / barrier.cluster《2. Clusters…》（已整理）、《9. Multi GPU》
异步流水线（double buffering）《8. Kernel Design》（软件流水 / GEMM 优化）

21.3 一句话记忆

H100 的异步 = “发令者（快）”与”搬数者（慢）”解耦； mbarrier 是协调二者的硬件原语（等数据、翻 phase、可复用、跨 proxy 需 fence）； 把拷贝（cp.async.bulk）与计算（wgmma）都变成后台流水线，让 Tensor Core 永不等数据。

参考来源：3. Asynchronicity and barriers.pdf（Lesson 3，50 页）。