Weile Luo
文章 标签 分类 关于我
Weile Luo

目录

CUDA Micro-benchmark 实战（下）：峰值算力、内存带宽与 Hopper 异步流水线

 收录于 博客
   约 4998 字   预计阅读 10 分钟 

在上一篇文章中，我们探讨了 CUDA 编译工具链、Warp 调度机制以及如何使用 Nsight Compute 进行性能瓶颈分析。

本篇不再罗列孤立的 CUDA 知识点，而是围绕一个问题展开：怎样设计一个可信的 GPU micro-benchmark，并让测量结果能够解释真实硬件行为？

全文依次讨论测量协议、计算吞吐、内存带宽、内存顺序、Inline PTX，最后用 Hopper 的 WGMMA 与 TMA 展示如何把计算和数据搬运组织成异步流水线。

1. 先建立可信的测量协议

Micro-benchmark 的代码通常很短，但短代码不代表结果可靠。GPU Boost、缓存命中、编译器消除、首次启动开销和计时范围都会改变结果。开始测量前，至少要回答下面几个问题。

1.1 测量对象和指标是什么？

同一条指令可以测两个不同指标：

  • Latency：一条依赖链从输入到输出需要多少周期。
  • Throughput：流水线稳定后，每周期最多完成多少条独立指令。

测 latency 时，要让下一次运算依赖上一次结果；测 throughput 时，则要使用多个互不依赖的 accumulator。把两者混在一起，最常见的结果就是“理论吞吐很高，但 benchmark 只能跑到一半”。

带宽测试也要先规定字节口径。例如 device-to-device copy 同时读取和写入 $N$ 字节，常用的有效带宽定义是：

$$ B_{\text{eff}}=\frac{N_{\text{read}}+N_{\text{write}}}{t} =\frac{2N}{t} $$

报告结果时必须说明使用的是单向字节数 $N/t$，还是读写合计 $2N/t$。

1.2 控制实验环境

建议为每次实验记录：

  • GPU 型号、驱动、CUDA Toolkit 和编译参数；
  • 实际 SM clock、memory clock、功耗上限、温度和 ECC 状态；
  • grid/block 配置、寄存器数、shared memory 和 occupancy；
  • 输入规模、数据类型、对齐方式和迭代次数。

运行前先 warm-up，避免把 context 初始化、JIT 或冷缓存开销算进结果。正式测量应重复多次，至少报告中位数；如果目标是逼近硬件上限，也可以同时报告最小值，但不能只挑一次最好结果。

对 DRAM 带宽测试，工作集应明显大于 L2 Cache；否则测到的可能是缓存带宽。对 cache benchmark 则相反，应主动扫描工作集大小并观察拐点。

1.3 防止测到“被优化掉的代码”

被测结果必须写回全局内存，且最好在 Host 端验证。否则编译器可能发现结果无人使用，删除整个循环。

此外，要同时检查 PTX 和 SASS：

1
2

nvcc -lineinfo -O3 benchmark.cu -o benchmark
cuobjdump --dump-sass benchmark

源代码里的 fmaf()、向量类型或 inline PTX，并不能保证最终机器码就是预期指令。Micro-benchmark 测的是实际执行的 SASS，不是 C++ 写法。

2. 计算吞吐：如何把 ALU 喂满？

2.1 理论峰值和实测值

某类运算的理论峰值可概括为：

$$ P_{\text{theory}}= N_{\text{unit}}\times \text{ops/cycle/unit}\times f_{\text{clock}} $$

$N_{\text{unit}}$、每周期吞吐和可用 pipeline 都依赖具体架构与数据类型。不要只拿营销页的 Boost Clock 代入；benchmark 期间的实际频率才是合理参照。

对于 FMA，a * b + c 按惯例计作 2 FLOPs。若每个线程有 $A$ 个独立 accumulator，循环执行 $I$ 次，共启动 $T$ 个线程，则实测吞吐为：

$$ P_{\text{measured}}=\frac{2AIT}{t} $$

这里的 $t$ 是只覆盖稳定计算区间的执行时间。

2.2 独立 accumulator 与依赖链

下面是 throughput benchmark 的核心形态：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

float x0 = input[tid];
float x1 = x0 + 1.0f;
float x2 = x0 + 2.0f;
float x3 = x0 + 3.0f;

#pragma unroll
for (int i = 0; i < ITERS; ++i) {
    x0 = fmaf(x0, a, b);
    x1 = fmaf(x1, a, b);
    x2 = fmaf(x2, a, b);
    x3 = fmaf(x3, a, b);
}

output[tid] = x0 + x1 + x2 + x3;

四条 FMA 链彼此独立，Warp Scheduler 可以在某条链等待结果时发射另一条链。独立链太少会暴露 RAW latency；太多又会增加寄存器压力、降低 occupancy。正确做法是扫描 accumulator 数量，而不是假设“展开越多越快”。

若要测 latency，则应只保留一条依赖链，并用总周期除以迭代次数。吞吐和延迟必须用不同 kernel。

FP16x2 指令能一次处理两个 half 元素，但“元素吞吐翻倍”不等于 latency 减半；打包、转换和目标架构的执行管线都会影响结果。

3. 内存带宽：指令更少不等于事务更少

3.1 合并、对齐和连续性

Warp 的内存请求会根据访问地址被拆分为一个或多个 memory sectors/transactions，具体粒度和缓存路径具有架构依赖性。因此不要把“连续 128 Bytes”简单等价为“永远只有一个事务”。

非合并访存示例
非合并访问：Warp lane 分散访问多个地址区域，产生更多内存事务
不对齐访存示例
不对齐访问：连续数据跨越对齐边界，可能需要额外 sector
不连续访存示例
不连续访问：固定 stride 使一个 Warp 覆盖更大的地址范围

三张图对应不同问题：

  1. 非合并：lane 请求无法落入少量连续 sector；
  2. 不对齐：访问范围跨越 sector/cache-line 边界；
  3. 不连续：stride 扩大地址跨度，搬回的数据中只有一部分真正被使用。

判断访问效率时，应结合 Nsight Compute 的 sector 数、请求数和 DRAM bytes，而不是只看 C++ 下标。

3.2 向量化访存的真实作用

使用 int2、int4 或 float4，主要收益是一个线程用更少的 load/store 指令搬运更多字节：

1
2

reinterpret_cast<int2*>(d_out)[i] =
    reinterpret_cast<const int2*>(d_in)[i];

它可以减少指令发射压力，但不会自动修复不对齐或非合并访问。指针必须满足向量类型对齐，Warp 间访问仍要连续。

向量化访存效果
K20X 上 scalar、vector2、vector4 copy bandwidth 随工作集大小的变化

图中小工作集阶段曲线接近，是因为固定开销和缓存占主导；进入大工作集后，vector2/vector4 减少了访存指令数量，因此可能获得更高带宽。该结果来自 K20X，不能直接外推到 Hopper；现代架构应重新实测。

一个完整带宽 benchmark 还应做到：

  • 数组大于 L2 Cache；
  • warm-up 后计时；
  • 使用不同输入避免编译器或缓存产生特殊路径；
  • 分别报告 scalar/vectorized 的指令数和 DRAM bytes；
  • 验证输出内容。

4. Memory Ordering：保证 benchmark 测的是正确程序

并发 benchmark 如果存在数据竞争，跑得再快也没有意义。需要区分四个概念：

  1. Compiler barrier：限制编译器重排；
  2. Memory fence：约束内存操作的观察顺序；
  3. Barrier：让一组线程在执行位置上会合；
  4. Visibility：其他线程能否通过正确的缓存/原子路径看到数据。

Fence 本身不等于线程同步，也不单独保证可见性。CUDA 官方 partial-sum 示例同时使用可见的结果存储、device-scope fence、原子计数器和线程同步。

1
2
3
4
5
6
7
8
9

// result 使用能让其他 block 正确观察的存储方式。
volatile float* result = ...;
result[blockIdx.x] = partial_sum;

// 只建立顺序：先发布 partial_sum，再更新 count。
__threadfence();

// 原子操作作为发布完成状态的信号。
unsigned int old = atomicInc(&count, gridDim.x);

如果没有 fence，其他线程可能先观察到 count 更新，再读取尚未按预期发布的 result。但仅加 fence 也不够：消费者仍需通过正确的原子/同步协议读取。

Release/Acquire 可以更精确地表达发布与消费关系。Scope 则限定哪些线程参与同步，常见范围包括 block、cluster、device 和 system。生产代码应优先考虑 cuda::atomic / cuda::atomic_ref 的 memory order 与 thread scope，而不是只靠传统 intrinsic 猜测语义。

Memory Scope 示例
CUDA memory scope：同步关系必须覆盖通信双方

Hopper 的 TMA/WGMMA 还引入了 async proxy。普通线程的 load/store 属于 generic proxy，TMA 和部分异步矩阵操作属于 async proxy；跨 proxy 交接数据时，需要对应的 proxy fence 或 barrier 协议。

5. Inline PTX：控制测量边界，而不是替代验证

5.1 读取计时器

读取 SM 本地 cycle counter：

1
2

uint32_t cycles;
asm volatile("mov.u32 %0, %%clock;" : "=r"(cycles));

注意 inline asm 字符串中的 PTX 特殊寄存器要写成 %%clock。同一 SM 内测短依赖链可以使用 %clock / clock64()；跨 SM 比较时间点时，可使用 %globaltimer，但仍要校准读取开销。

SM Clock
读取 SM 本地时钟寄存器
Global Timer
读取跨 SM 的 global timer
Clock 转换
计时值与周期/时间单位的转换

对于完整 kernel，CUDA Event 通常更稳妥；线程内计时器适合隔离短指令序列。二者回答的问题不同。

5.2 volatile 与 "memory" clobber

编译器默认认为 asm() 除输出操作数外没有副作用。若指令不能被删除或移动，应使用：

1

asm volatile("mov.u32 %0, %%clock;" : "=r"(cycles));

若 asm 通过指针隐式读写用户内存，或需要阻止编译器在 asm 周围做内存优化，应声明：

1

asm volatile("..." : : : "memory");

"memory" 是编译器层面的 clobber，不会自动生成 GPU memory fence，也不能替代线程同步。

PTX Memory Clobber
Inline PTX 的 memory clobber：告知编译器 asm 可能读写未显式列出的内存

5.3 操作数约束

Inline PTX 的 = 表示 write-only output，+ 表示 read-write operand：

1

asm("add.u32 %0, %0, 1;" : "+r"(x));

如果 PTX 需要读取操作数原值，就必须使用 + 或拆成独立 input/output；这与“阻止物理寄存器被其他指令复用”无关。

PTX 寄存器类型
Inline PTX constraint 与 PTX register type 的对应关系
PTX 操作数修饰符
输出操作数 = 与读写操作数 +

指针还涉及 Generic、Global、Shared 等 state space。内联 PTX 无法自动推断任意指针的目标空间，使用 ld.global、ld.shared 等指令前必须确认地址空间。

Memory State Space
PTX state spaces：Generic address 与具体地址空间

6. Hopper 案例一：WGMMA 异步矩阵流水线

传统 warp-level MMA 常先把 fragment 搬入寄存器，再由一个 Warp 协作计算。下面的图片适合作为 fragment layout 背景，但它不是 WGMMA 本身。

Tensor Core MMA 矩阵布局
Warp-level MMA fragment layout：用于理解寄存器中的矩阵分片

Hopper 的 WGMMA 由 4 个连续 Warp，也就是 128 个线程组成一个 Warp Group，共同执行异步矩阵乘加：

$$ D=A\times B+C $$

典型控制顺序是：

1
2
3
4

wgmma.fence
wgmma.mma_async.sync.aligned
wgmma.commit_group
wgmma.wait_group N
  • wgmma.fence：排序此前对 accumulator registers 的访问与后续 WGMMA；若操作数由 generic proxy 写入 Shared Memory，还要按 PTX 要求使用对应的 async-proxy fence；
  • mma_async：异步发射矩阵乘加；
  • commit_group：把已发射操作提交为一组；
  • wait_group N：等待，直到最多还有 N 组操作 in-flight。

.sync 要求参与 Warp 在该指令处会合；.aligned 要求 Warp Group 的所有线程一致执行同一条 WGMMA。若条件分支让 Warp Group 内线程执行路径不一致，行为未定义。

WGMMA Sync & Aligned
PTX 对 WGMMA .sync 与 .aligned 的一致执行要求

WGMMA 的 accumulator D 位于寄存器；B 通过 shared-memory descriptor 提供；A 根据指令变体来自寄存器或 shared memory。异步接口的价值在于允许多组 MMA 与下一阶段的数据准备重叠，而不只是“单条 Tensor Core 指令更大”。

WGMMA Sparse
Hopper sparse WGMMA 的操作数与 accumulator 组织方式

测 WGMMA throughput 时应：

  • 让矩阵尺寸足以填满多个 Warp Group；
  • 使用多 stage buffer，避免只测到 shared-memory 等待；
  • 用 commit_group / wait_group 控制 in-flight 深度；
  • 计入真实 MMA FLOPs，并验证结果；
  • 同时观察 Tensor Core utilization、stall reason 与 shared-memory throughput。

7. Hopper 案例二：TMA 数据搬运流水线

TMA（Tensor Memory Accelerator）把大块或多维张量在 Global Memory 与 Shared Memory 之间异步搬运。与 Ampere 的 LDGSTS 路径相比，TMA copy 可由一个 elected thread 发起，再由 TMA 单元完成地址生成和传输，让其余线程专注计算。

TMA 异步拷贝概览
A100 LDGSTS 与 H100 TMA：地址生成和大块传输从线程侧转移到 TMA 单元

一个典型 Hopper pipeline 是：

1
2
3
4
5
6
7

Global Memory
      │ TMA prefetch
      ▼
Shared Memory stage N
      │ WGMMA
      ▼
Registers

计算 stage N 时，TMA 可以预取 stage N+1。性能来自 copy/compute overlap，而不是单次 copy latency 必然更低。

7.1 Descriptor 与对齐

多维 TMA 使用 Host 端创建的 tensor map descriptor，描述基地址、维度、stride、tile shape 和布局。Device 端只需给 descriptor 和坐标即可发起传输。

TMA Descriptor
TMA tensor map descriptor 保存张量维度、stride 与 tile 元数据

对于 cp.async.bulk 一类强制使用 TMA 的操作，Global/Shared 地址通常要求 16-byte 对齐，传输大小是 16 Bytes 的倍数，mbarrier 地址要求 8-byte 对齐。高级 CUDA API 在不满足条件时可能回退到同步路径；直接 PTX 则可能产生未定义行为，必须查对应指令版本的约束。

TMA Non-Tensor Alignment
Bulk asynchronous copy 的地址与大小约束
TMA Tensor Alignment
Tensor TMA 对 tensor map、stride 和 tile 的约束

7.2 完成机制取决于传输方向

Global → Shared 的 TMA copy 通常由 transaction-aware mbarrier 跟踪：先初始化 barrier 和 expected bytes，消费者等待对应 phase 完成后再读取 Shared Memory。

Shared → Global 的 bulk copy 通常通过 bulk async-group 提交和等待。不能把 mbarrier 和 async-group 当成可随意互换的两种风格；它们对应不同数据方向和完成语义。

TMA completion mechanisms
不同 source/destination memory space 对应的异步完成机制
TMA Mbarrier 机制
Global → Shared TMA copy：用 mbarrier 跟踪事务字节完成
Async-group 机制
Shared → Global bulk copy：用 async-group 提交和等待

TMA micro-benchmark 至少要分别测：

  1. 单次 copy latency；
  2. 大工作集 steady-state bandwidth；
  3. TMA 与 WGMMA 重叠后的端到端 tile throughput；
  4. 不同 tile shape、stage 数与 cluster multicast 的影响。

只测 copy kernel 无法证明流水线更快；最终应比较“搬运 + 计算”的总时间。

8. 附录：与主线相关但不属于 benchmark 核心的知识

8.1 DMA 与 IOMMU

DMA 允许设备直接访问内存，IOMMU 则负责设备虚拟地址到物理地址的映射与隔离。它们对 GPU Direct、统一虚拟地址和设备间传输很重要，但不直接决定前面单 GPU kernel micro-benchmark 的 ALU/DRAM 峰值。

IOMMU 与 MMU 区别
MMU 与 IOMMU 分别服务 CPU 和设备侧地址转换

8.2 CUDA 模板显式实例化

如果模板定义放在 .cu，调用点位于另一个翻译单元，编译器看不到对应模板定义时，需要在 .cu 中显式实例化，并在头文件中声明接口。这是工程组织问题，不是 GPU 性能机制。

Template .cu 示例
在 .cu 中定义并显式实例化 CUDA 模板
Template 头文件示例
在头文件中声明跨翻译单元使用的模板接口

总结

可信的 CUDA micro-benchmark 不是“写一个很小的 kernel 然后计时”，而是一套受控实验：

  1. 明确测 latency 还是 throughput；
  2. 控制时钟、缓存、输入规模和 warm-up；
  3. 防止编译器消除，并检查最终 SASS；
  4. 验证输出和统计口径；
  5. 用 profiler 解释测量值为何接近或偏离理论上限。

在 Hopper 上，WGMMA 与 TMA 又把问题推进了一步：性能不再只取决于单条指令吞吐，而取决于能否把 Global Memory、Shared Memory、Tensor Core 和同步机制组织成稳定的多 stage 异步流水线。


更新于 2026-07-14
 CUDA, GPU
返回 | 主页