第 4 章 · GPU 硬件架构

学习目标

• 理解 SM 内部结构：CUDA core / Tensor Core / 寄存器堆 / shared mem
• 掌握 SIMT 执行模型与 warp divergence
• 明白什么是 占用率、什么因素决定占用率
• 看完知道 Volta → Ampere → Hopper 在 Tensor Core 上的演进

前置知识

已完成 Ch03，能计算 thread 的全局索引，并理解 block、warp 与 lane 的基本关系。

核心概念

4.1 SM 解剖图

A100 (Ampere) 上有 108 个 SM，每个 SM 内部长这样：

关键数字（A100 / sm_80）：

4.2 SIMT 执行模型

SM 调度的最小单位不是 thread 而是 warp = 32 个 thread。warp 内 32 个 thread同步执行同一条指令但操作不同的数据（寄存器/内存地址）。这就是 SIMT (Single Instruction Multiple Threads)。

graph LR
    PC["Program Counter"] --> Decode["Decode"]
    Decode -->|"1 cycle"| Disp["Dispatch"]
    Disp --> T0(("T0"))
    Disp --> T1(("T1"))
    Disp --> T2(("T2"))
    Disp --> Tn(("... T31"))
    style PC fill:#f3f1e8,stroke:#8b1538
    style Disp fill:#f3f1e8,stroke:#2f5d3a

SM 同时持有多个 warp（最多 64 个）。每周期硬件挑一个 ready 的 warp 发射指令。 当某个 warp 在等 global memory（几百周期），SM 切到别的 warp 干活——这就是"GPU 用并行隐藏延迟"。

关键代码

4.3 Warp Divergence — 同 warp 走两条路

如果 warp 内 32 个 thread 走了 if/else 两支：

if (threadIdx.x & 1) {        // 偶数 lane 走 else，奇数 lane 走 if
    x = x * 2;
} else {
    x = x + 1;
}

硬件没办法同时执行两条不同的指令，所以它串行跑两遍：先让走 if 的 lane 干活（else 的 lane 掩码屏蔽），再切到 else 的 lane。代价 = 两支的耗时之和。

运行 warp_divergence.cu，先核对两种布局输出一致，再记录目标 GPU 的时间：

uniform   (warp-aligned branch) : TODO(on GPU)
divergent (lane-aligned branch) : TODO(on GPU)
slowdown                        : TODO(on GPU)

两支工作量接近时，divergence 会让 warp 分别执行两条路径；实际代价还受指令、访存与编译器 predication 影响，必须实测。

4.4 Occupancy

定义：实际驻留 warp 数 / SM 最大支持 warp 数。高 occupancy 让 SM 有更多 warp 可以切换以隐藏延迟。

占用率受三个资源约束（取最严格的那个）：

1. 线程数：block 大小 × block/SM 数 ≤ 2048（A100）
2. 寄存器：每 thread 用 N 个寄存器 → SM 上 64K/N 个 thread 上限
3. Shared memory：每 block 用 M KB → SM 上 192/M 个 block 上限

运行 occupancy_probe.cu，CUDA Runtime 会替你算：

--- block = 256 ---
  light_kernel    block= 256  active blocks/SM= 8  warps/SM=64/64  occ=100.0%
  heavy_kernel    block= 256  active blocks/SM= 4  warps/SM=32/64  occ= 50.0%

heavy_kernel 因为本地数组占了很多寄存器（编译器把 float acc[64] 放到了寄存器），SM 上的驻留 block 数被腰斩。

运行结果

4.5 架构演进 (5 分钟扫盲)

对 LLM 推理重要程度（粗略）：

• Tensor Core（Volta 起）→ 为支持的数据类型与 shape 提供专用 MMA 路径
• bf16（Ampere）→ 训练 / 推理动态范围大
• cp.async（Ampere）→ 让 shared memory 加载与计算重叠，FlashAttention 的基础
• FP8（Hopper/Ada）→ 降低输入字节并启用对应 Tensor Core 路径；端到端吞吐需按模型与硬件实测
• TMA（Hopper）→ kernel 不再手写 load 循环，硬件批量搬运 tile

4.6 工业实战：MIG、ECC、功耗管理

把 GPU 上线到生产，光会写 kernel 不够，还要会调"GPU 设置"。这一节是数据中心运维必备。

4.6.1 MIG (Multi-Instance GPU) — 一卡当七卡用

A100 / H100 / H200 支持把一张 GPU 切成最多 7 个独立的"小 GPU"（叫 MIG instance），每个有独立 SM、独立显存、独立 cache。用途：多租户推理服务，让小模型共享大卡。

# 1) 启用 MIG 模式 (需要 root + 没人在用 GPU)
sudo nvidia-smi -mig 1

# 2) 创建 MIG instances (A100-40G 可切 1g.5gb × 7)
sudo nvidia-smi mig -cgi 19,19,19,19,19,19,19 -C
#                      ^^^ profile id, 19 = 1g.5gb

# 3) 看现在的切片
nvidia-smi -L
# GPU 0: NVIDIA A100-SXM4-40GB
#   MIG 1g.5gb     Device 0: (UUID: MIG-...)
#   MIG 1g.5gb     Device 1: (UUID: MIG-...)
#   ...

# 4) 关 MIG 模式
sudo nvidia-smi -mig 0

陷阱：MIG 开启后，cudaGetDeviceCount 返回的是 MIG instance 数而不是物理 GPU 数；代码看到的算力是切片大小，跟物理 GPU 不一样。新手调试半天找不到原因，记得 nvidia-smi -L 看现在到底切了没。

4.6.2 ECC — 单 bit 翻转保护

数据中心 GPU 默认开 ECC（错误检测与纠正）。它给每 64 bit 数据加 8 bit 校验，能纠 1 bit 错误、检测 2 bit 错误。代价：

• 显存少 ~6%（A100-80G 实际可用 ~76 GB）
• HBM 带宽降 ~5%（要传校验位）
• 极小幅算力影响（< 2%）

nvidia-smi -e 0           # 关 ECC (需重启 GPU)
nvidia-smi -e 1           # 开
nvidia-smi -q -d ECC      # 查累计错误数
# 关注:
#   Volatile Single Bit ECC Errors  (本次启动后)
#   Aggregate Single Bit ECC Errors (从出厂)

是否关 ECC：

• 训练：必开。一次 bit 翻转可能让 loss NaN 浪费几小时。
• 推理：可关。bit 翻转影响一个 token 不致命。生产推理服务 80% 都关，多 ~5GB 显存。
• 看到 Volatile Single Bit Errors > 0：观察增长率，> 100/天 该送修。

4.6.3 功耗与时钟管理

GPU 默认动态调频：温度高就降频、负载低就降频。生产服务希望延迟稳定，常用做法是锁定时钟：

# 查支持的时钟
nvidia-smi -q -d SUPPORTED_CLOCKS | head -40

# 锁定 GPU 时钟到 1410 MHz, mem 1215 MHz (A100 典型最大)
sudo nvidia-smi -ac 1215,1410

# 锁定功耗墙到 300W (默认 400W, 省电用)
sudo nvidia-smi -pl 300

# 解锁
sudo nvidia-smi -rac
sudo nvidia-smi -rgc

锁频后 benchmark 数字更稳定（消除 thermal throttling 的抖动），适合 CI 性能回归测试。

4.6.4 nvidia-smi 读图：理解每列

+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 535.86.05    Driver Version: 535.86.05    CUDA Version: 12.2     |   ← 驱动 + 驱动支持的最高 CUDA
|-----------------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M | Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap |         Memory-Usage | GPU-Util  Compute M. |
|                                |                      |               MIG M. |
|=========================================+======================+======================|
|   0  NVIDIA A100-SXM...  On    | 00000000:00:04.0 Off |                    0 |
| N/A   42C    P0    72W / 400W  |   1234MiB / 81920MiB |     45%      Default |
|                                |                      |             Disabled |
+-----------------------------------------+----------------------+----------------------+

要关注的字段:
- Persistence-M = On    避免重复初始化开销（是否启用按部署策略决定）
- Pwr 72W/400W         功耗远低于 cap → mem-bound 或 launch 等
- 1234MiB/81920MiB     显存占用 (有时是其他进程, nvidia-smi 看不到 driver 占用)
- GPU-Util 45%         注意陷阱: 高 GPU-Util 不代表算力跑满
- MIG M. Disabled      MIG 未开

看进程占了多少：

nvidia-smi pmon -i 0 -d 1     # 每秒打印每进程的 sm/mem/enc/dec util
nvidia-smi --query-compute-apps=pid,used_memory --format=csv

4.6.5 各代 GPU 在 LLM 工业部署中的定位

对 LLM 推理来说最大的几次跳跃：

1. Volta → Turing：消费端有了 TC，inference 落地变现实
2. Turing → Ampere：cp.async 让 FlashAttention 成为可能
3. Ampere → Hopper：TMA + Transformer Engine + fp8 让 100B+ 模型可服务
4. Hopper → Blackwell：NVFP4、TMEM 与更大的内存系统改变推理优化空间；成本要按质量、SLO 与真实 workload 测量

选 GPU 的时候，对应到你的模型规模和 latency 要求查上表。

4.7 研究前沿（2025-2026）：Blackwell 解剖

4.7.1 Blackwell B200 关键特性

Blackwell（sm_100）是 Hopper 之后的下一代，2024 末发布，2025 大规模量产。结构变化比 Ampere→Hopper 更激进：

4.7.2 Tensor Memory (TMEM) — Blackwell 最大改动

Hopper 及之前所有 GPU 用一个统一的 register file 存 fragment / accumulator。Blackwell 把matrix accumulator 抽出来到一块独立的 SRAM（叫 TMEM）：

SM 内部 (Blackwell):
  ┌─ Register file       ~64 K × 32-bit  (跟 Hopper 一样)
  ├─ Shared / L1         ~228 KB         (跟 Hopper 一样)
  ├─ Tensor Memory (TMEM)   256 KB          ← 新增, 专给 MMA accumulator
  └─ Tensor Cores (5th gen) + Transformer Engine 2

动机：fp4 / fp8 MMA 的 accumulator 是 fp32，每个 MMA 要消耗几十个寄存器持有 acc。把 acc 放 TMEM 释放了寄存器给其他用途，让 ILP 大幅提升。

用 TMEM 的新指令: TCGEN05

// PTX-level (CUTLASS 4.6.1 / CuTe DSL（2026-07 快照） 自动用)
tcgen05.mma.async        // 异步 MMA, accumulator 写 TMEM
tcgen05.ld               // 从 TMEM 读出 accumulator
tcgen05.cp               // shared memory ↔ TMEM

CUTLASS/CuTe 封装了许多 TCGEN05 与 TMEM 细节；研究 kernel 也可能直接使用更底层原语。是否接近峰值必须按 dtype、dense/sparse、shape 与完整 benchmark 记录判断。

4.7.3 2nd-gen Transformer Engine：动态精度

Hopper 的 1st-gen TE 只能在 fp16 / fp8 之间切。Blackwell 的 2nd-gen 可以：

• 每个张量自动选 fp4 / fp6 / fp8 基于范围分析
• 支持 microscaling (block size 16/32) — 每 16/32 元素一个 scale，比 per-tensor scale 精度高
• 训练 / 推理都能用，DeepSeek-V3 部分用 fp8 训练就是同思路（虽然 V3 用的还是 Hopper）

4.7.4 GB200 与 NVL72：rack 即 GPU

单 B200 已经很强，但 NVIDIA 的真正杀手锏是 NVL72：

1 个 GB200 Grace Blackwell Superchip = 1 Grace CPU + 2 Blackwell GPU
每 GPU 186 GB HBM3e；Superchip 共 372 GB、16 TB/s
1 个 GB200 NVL72 = 36 个 Superchip = 72 GPU
整柜 13.4 TB HBM3e、576 TB/s；峰值按 NVIDIA 表中的 sparse | dense 口径阅读

能干什么：

• 1.8T 参数 fp4 模型一柜推理（之前需要跨节点）
• Llama 4 Behemoth 2T 训练，跨 rack 通信压力小
• DeepSeek-V3 671B MoE TP=16 + EP=72，单柜服务

4.7.5 Rubin（2026 中后期）— 下一代

NVIDIA 已公布 Rubin 路线图：

• Rubin (R100, 2026)：3nm，HBM4，再一轮 ~2× 算力
• Rubin Ultra (R200, 2027)：4 die / GPU, ~8 PF fp4
• NVL576 整柜：训练顶级模型

核心趋势：单卡再难有 2× 跳跃，整柜带宽与 cluster 调度是新战场。

4.7.6 AMD / Google TPU 现状（2026）

2026 工业现实：推理硬件与软件栈已有多种选择；本教程不保留无一手来源的市场份额或竞品等价数字。CUDA 仍是本课程的研发起点，但选型应以可移植性、容量、SLO、成本和 workload benchmark 为依据。

自检清单

练习题

1. 跑 warp_divergence.cu，改 iters 看耗时如何线性变化。
2. 给 occupancy_probe.cu 加一个 shared_kernel，用 __shared__ float buf[16384]（= 64 KB），看 SM 上能驻留几个 block。
3. 查阅你 GPU 的 Compute Capability 表，记下：每 SM 最大寄存器、shared mem、warp 数。

常见坑

• 把 occupancy 当成唯一目标，而忽略寄存器复用、ILP 和内存流量。
• 只按产品名推断 compute capability；部署前应查询实际设备属性和目标 fatbin。
• 用跨 GPU、跨精度的峰值数字直接预测真实 kernel 延迟。

4.10 CUDA 官方手册精讲（CUDA Programming Guide 13.2（核验：2026-07-20））

SIMT 深读、Warp Scheduler、Thread Scope、寄存器堆

深入 SIMT — Independent Thread Scheduling

4.2 节给出的"32 lane 同步执行同一条指令"是 Pascal (sm_60) 及之前的描述。 Volta (sm_70) 起 NVIDIA 给每个 thread 都加了独立的 program counter 和 call stack，称为 Independent Thread Scheduling (ITS)。

ITS 带来的能力：

• warp 内一个 thread 在等锁，另一个 thread 可以继续干活，不再 100% lock-step
• 过去会死锁的"warp 内细粒度 mutex"现在合法
• 编译器/驱动里有一个 schedule optimizer，会动态把可一起执行的 active thread 重新打包成 SIMT 单元

// ✅ 正确（sm_70+ 必须）
    unsigned mask = __activemask();         // 当前 warp 内活跃的 lane bitmask
    int v = __shfl_sync(mask, val, 0);       // 显式同步 + broadcast lane 0 的值
    __syncwarp(mask);                        // 任何依赖 warp-内一致性的点都补一次

    // ❌ 旧写法 (deprecated, 在 sm_70+ 上行为未定义)
    int v = __shfl(val, 0);
    

所以本教程后续章节里你会看到大量 __shfl_xor_sync(0xffffffff, ...)——0xffffffff 表示"全 warp 参与"，这是约定俗成的安全 mask。

warp scheduler 与零开销切换

"GPU 用并行隐藏延迟"是口号。具体怎么做？关键在两点：

1. 每个 warp 的 context（PC + 寄存器 + call stack）常驻片上。SM 上有最多 64 个 warp 的 context 同时存在，切换不需要 push/pop（不像 CPU 上下文切换）。
2. 每个 SM 配备多个 warp scheduler。A100 上每 SM 有 4 个 scheduler，每周期每个 scheduler 各挑一个 ready warp 发射 1 条指令——所以 A100 SM 每周期最多同时执行 4 个 warp。

所以"驻留 64 warp / 实际同时执行 4 warp"的真正含义是：scheduler 在 64 个候选里挑 4 个最 ready 的。这就是延迟隐藏：

    sequenceDiagram
        participant W0 as Warp 0
        participant W1 as Warp 1
        participant Wn as ...Warp 63
        participant SM as Scheduler
        SM->>W0: 发射 LDG (~500 cyc)
        Note over W0: 等 HBM
        SM->>W1: 发射 FMA
        SM->>Wn: 发射 FMA
        Note over W0: HBM 数据到达
        SM->>W0: 发射 FMA (用 LDG 结果)
    

结论：

• kernel 想跑满 SM 吞吐，需要 scheduler 永远有 ready warp 可挑。warp 太少（occupancy 低）→ 全部在等 HBM → ALU 空闲
• 但 warp 不是越多越好。每 thread 寄存器太挤会被迫 spill；FlashAttention v2 故意保持 25-50% occupancy 换取每 thread 大寄存器，让 ILP 主导

用 cuOccupancyMaxActiveBlocksPerMultiprocessor（runtime API：cudaOccupancyMaxActiveBlocksPerMultiprocessor）在编译期之外、launch 之前就能让 driver 替你算出当前 kernel 在当前 GPU 上能跑几个 active block：

int max_blocks_per_sm;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &max_blocks_per_sm,
        my_kernel,        // __global__ 函数指针
        /*blockSize=*/256,
        /*dynamicSMemSize=*/0);
    printf("active blocks / SM = %d\n", max_blocks_per_sm);

    // 进阶: 让 CUDA 帮你挑最优 block size
    int min_grid, best_block;
    cudaOccupancyMaxPotentialBlockSize(
        &min_grid, &best_block, my_kernel, 0, 0);
    

分歧、谓词执行与 ILP（指令级并行）

4.5.1 编译器何时把分支编成谓词（predication）

不是所有 if/else 都会变成"两支都跑"的串行。如果两支都很短（典型阈值：每支 <= 7 条指令），nvcc 会用 PTX 的 predicate 寄存器把两支都编译成无跳转的指令流：

// 源码
    if (cond) x = a * b;
    else      x = a + b;

    // nvcc 编译后 (伪 PTX, 一条 setp + 两条带谓词的指令)
    setp.ne.f32   %p1, %cond, 0f00000000;       // p1 = (cond != 0)
    @%p1   fma.f32     %x, %a, %b, 0f00000000;   // p1=true 时执行
    @!%p1  add.f32     %x, %a, %b;               // p1=false 时执行
    

谓词执行的代价：两支的指令都流过流水线，但只有匹配 lane 写回。所以哪怕没分歧，"两支都跑"的代价也存在。Nsight Compute 的 smsp__sass_branch_targets_threads_uniform.pct 接近 100% 说明几乎全部走了谓词路径。

4.5.2 当两支不等长，divergence 代价不是 2×

4.3 节实验比较等长 if/else 分支；一般情况下：

实际耗时 ≈ max(T_if, T_else) × 1   if 全 warp 走同一支
             ≈ T_if + T_else           if warp 内分歧
    

所以最坏情况：if 路径 95%、else 路径 5%，但只要 warp 内有一个 lane 走了 else，整 warp 等到两支都跑完。这就是为什么大语言模型 attention mask 实现里"用 -∞ 填充"比"用 if 跳过"快——x + (-inf) 永远是无分歧的算术。

4.5.3 ILP — 单 thread 内的并行

SIMT 之外还有 thread 内的指令级并行。同一 thread 内独立指令可以并行流水：

// ILP 友好: 4 个独立 FMA, 编译器可重排
    float a0 = x[0]*y[0] + b[0];
    float a1 = x[1]*y[1] + b[1];
    float a2 = x[2]*y[2] + b[2];
    float a3 = x[3]*y[3] + b[3];

    // ILP 不友好: 串行依赖
    float a = 0;
    a = a + x[0]*y[0];
    a = a + x[1]*y[1];
    a = a + x[2]*y[2];
    a = a + x[3]*y[3];
    

NVIDIA SM 是 in-order 发射，不做 branch prediction 也不做 speculative execution（不像 CPU）。靠的就是 ILP + warp 切换两条路一起隐藏延迟。这也是为什么 FlashAttention 故意展开内层循环——让编译器排出更宽的 ILP 窗口，每 thread 持有更多寄存器中的 accumulator。

Thread Scope 与 scoped atomic

"原子操作"四个字隐藏着一个问题：对谁原子？ 同一个 block 内的 lane？整个 GPU？跨 GPU？scope 越大、硬件需要走的 cache coherence path 越长、代价越高。

CUDA 12+ 的 libcu++ 提供四种 scope，对应内存层级中的不同 "coherency point"：

4.6.1 block-scope 计数器示例

#include <cuda/atomic>

    __global__ void block_scoped_counter() {
        // Block-scoped atomic avoids unnecessary wider-scope ordering.
        __shared__ cuda::atomic<int, cuda::thread_scope_block> counter;

        if (threadIdx.x == 0) {
            counter.store(0, cuda::memory_order_relaxed);
        }
        __syncthreads();

        // 每 thread 加 1, 取旧值
        int old = counter.fetch_add(1, cuda::memory_order_relaxed);
        // ... 用 old ...
    }
    

4.6.2 producer-consumer：acquire/release 排序

当一个 thread 要把"数据准备好了"通知给另一个 thread，relaxed 不够，要用 release（写端）+ acquire（读端）保证数据写在 ready 信号之前对消费者可见：

__global__ void producer_consumer() {
        __shared__ int data;
        __shared__ cuda::atomic<bool, cuda::thread_scope_block> ready;

        if (threadIdx.x == 0) {
            data = 42;
            ready.store(true, cuda::memory_order_release);   // 写端
        } else {
            while (!ready.load(cuda::memory_order_acquire)) { // 读端
                // spin
            }
            int v = data;   // 保证看见 42 (acquire 同步了 release 之前的写)
        }
    }
    

这套 API 在 LLM 实现里到处出现：FlashAttention 的 row-max 协调（block scope）、PagedAttention 的 ref count（device scope）、跨 GPU NCCL all-reduce 的 barrier（system scope）。

寄存器堆：占用率与 spill 的真正瓶颈

4.7.1 每代 GPU 寄存器堆容量

关键点：每 SM 一共 64K 个 32-bit 寄存器，所有驻留 thread 共享。所以如果你每 thread 用 64 个寄存器，SM 上最多驻留 1024 thread = 32 warp（不到上限 64）。寄存器用得越多，occupancy 越低，但 ILP 越强。

4.7.2 spill 到底贵多少？

当编译器发现一个 thread 用的寄存器超过 -maxrregcount 或 hard limit (255)，多出的局部变量会被放到 local memory——名字叫"local"，其实是每 thread 独占的一段 global memory 分区。被 L1 cache 兜底，但延迟从 1 cycle 暴增到 30+ (L1 命中) 或 500+ (L1 miss)。

# 编译时显示每 kernel 的寄存器/local memory/shared mem 用量
    nvcc -Xptxas -v -arch=sm_80 my_kernel.cu

    # 输出节选:
    # ptxas info    : Used 96 registers, 0 stack, 16 bytes smem, 332 bytes cmem[0]
    #                                    ^^^^^^^^^ stack > 0 = spill 到 local !
    

4.7.3 三种控制寄存器用量的方式

// 1) 编译期硬限制 (整个编译单元)
    //    nvcc -maxrregcount=64 -arch=sm_80 my.cu

    // 2) per-kernel 软提示: __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM)
    //    告诉编译器"我准备每 SM 跑 minBlocksPerSM 个 block"
    //    编译器据此反推每 thread 寄存器预算 = 64K / (maxThreadsPerBlock * minBlocksPerSM)
    __global__ __launch_bounds__(256, 4)
    void my_kernel(/*...*/) {
        // 256 thread * 4 block/SM = 1024 thread → 每 thread 最多 64 reg
    }

    // 3) 用 const + 编译期 unroll 让编译器复用寄存器
    #pragma unroll
    for (int i = 0; i < 8; ++i) acc[i] += x * y[i];
    

记住这条经验法则：看 Nsight Compute 的 Stack Frame 字段，> 0 就该 debug。

异步执行：硬件 barrier 与 async proxy

Ampere 起 NVIDIA 在 SM 内增加了三种真正并行的硬件单元，把"等数据"和"算"在物理上拆开：

4.8.1 Async Thread / Async Proxy 概念

这是 PDF §3.2.2.3.1 的核心定义，看 CUTLASS 源码必备：

• Async thread：CUDA thread 发起一个异步操作（如 cp.async），硬件想象成"由另一个虚拟 thread 执行"，原 CUDA thread 不阻塞
• Generic proxy：普通 load/store 走的通道。LDGSTS 模型上仍属于这个 proxy
• Async proxy：TMA、wgmma.mma_async、tcgen05.* 走的独立通道。同一地址的 generic 写和 async 读之间必须显式 fence_proxy_async，否则顺序未定义

    graph LR
        T["CUDA thread"]
        T --"普通 LDG/STS"--> GP[Generic Proxy]
        T --"cp.async (LDGSTS)"--> AT1["Async thread\n@ generic proxy"]
        T --"cp.async.bulk (TMA)"--> AT2["Async thread\n@ async proxy"]
        GP --> L1[L1/Shared]
        AT1 --> L1
        AT2 --> L1
        style AT2 fill:#f3f1e8,stroke:#8b1538
    

4.8.2 Asynchronous barrier（split-arrive-wait）

普通 __syncthreads() 把"到达"和"等待"合二为一。Ampere 起硬件支持 split barrier——arrive 不阻塞，可以先做别的事，最后再 wait：

#include <cuda/barrier>
    #include <cooperative_groups.h>

    __global__ void split_arrive_wait(int n_iter, float* data) {
        using barrier_t = cuda::barrier<cuda::thread_scope_block>;
        __shared__ barrier_t bar;
        auto block = cooperative_groups::this_thread_block();

        if (block.thread_rank() == 0)
            init(&bar, block.size());
        block.sync();

        for (int i = 0; i < n_iter; ++i) {
            // 1) arrive — 不阻塞, 返回 token
            auto token = bar.arrive();

            // 2) 此时可以做与 barrier 无关的工作 (隐藏 barrier 等待开销!)
            compute_unrelated(data, i);

            // 3) wait — 真正阻塞直到全 block 都 arrive
            bar.wait(std::move(token));
        }
    }
    

用途：在 producer-consumer 流水线里，producer 提交完异步 copy 立刻 arrive，然后干别的；consumer 真正需要数据时才 wait。FlashAttention v3 整套 pipeline 就建立在这个机制上。

4.8.3 Pipeline — 多 stage producer-consumer 抽象

libcu++ 提供 cuda::pipeline 把 split barrier + async copy 封装成 producer/consumer FIFO：

#include <cuda/pipeline>
    #include <cooperative_groups/memcpy_async.h>

    template <size_t STAGES = 2>
    __global__ void double_buffer_kernel(int* g_out, const int* g_in, int n) {
        extern __shared__ int smem[];     // STAGES * block.size() * sizeof(int)
        auto block = cooperative_groups::this_thread_block();
        auto pipe  = cuda::make_pipeline();

        // 预先填满 STAGES-1 个 stage
        for (size_t s = 0; s < STAGES - 1; ++s) {
            pipe.producer_acquire();
            cuda::memcpy_async(block, smem + s * blockDim.x,
                               g_in + s * blockDim.x,
                               sizeof(int) * blockDim.x, pipe);
            pipe.producer_commit();
        }

        // 主循环: 一边 consume, 一边异步拉下一批
        for (int i = 0; i < n; ++i) {
            pipe.consumer_wait();
            compute(smem + (i % STAGES) * blockDim.x);
            pipe.consumer_release();

            // 立刻发起下一次 copy
            pipe.producer_acquire();
            cuda::memcpy_async(/*...*/);
            pipe.producer_commit();
        }
    }
    

下一章导览

第 5 章进入性能优化的核心：内存层级与合并访问。高性能 kernel 往往需要大量精力减少和重排访存。