Concurrency Response of Plain Global Loads on the NVIDIA H100
Somashekar Manjunath, Rahul Ramachandra M
cs.AR, cs.PF
2026-08-16
三颗 H100 微基准实测:普通全局加载带宽在每线程 2 条达到峰值,加到 8 条反而掉 35%,固定工作量与反向扫描对照都成立;cp.async 同场景快 2.1~2.9 倍。
访存受限的 CUDA kernel 想吃满 HBM 带宽,靠的是让足够多的字节同时在路上。Little's Law(吞吐 = 在途字节 ÷ 往返延迟)在 H100 上的含义是:延迟几百纳秒、带宽 3.35 TB/s,每个 SM 得维持约 10 KB 在途数据才能饱和。涨在途字节的软件手段有两个:多开 warp(thread-level parallelism),或者让每个线程每轮多发射几条互相独立的加载(instruction-level parallelism,ILP)。第二条路在直觉里是白赚的:寄存器够、指令槽够,多发几条加载似乎没有理由更慢。
这篇用微基准证明这个直觉在 Hopper 上是错的。普通全局加载(ld.global)的带宽对每线程负载量 K 呈先升后降的响应曲线,峰值出现在很小的 K≈2,继续加到 K=8 反而掉约 35%。GT200 到 Kepler 年代有人测过 outstanding-request 上限,Fermi 上推出每 warp 6 条、每 SM 64 条的推断;Hopper 世代有整体架构解剖和 TMA 路径刻画,但没人做「固定逻辑工作量下扫每线程 ILP」这个受控实验。CoreOptX 的两位作者补的就是这块。
仪器是从公开 ISA 文档干净搭出来的 CUDA kernel,过了 SASS 审计:每轮恰好发射 K 条独立加载,零寄存器溢出;加载结果 XOR 累加进一个编译器无法证明死亡的 sink,防止被优化掉。延迟探针用指针追逐链,构造上就不可能被消除。
设计上真正下功夫的是拆混淆:
K 的定义也讲得诚实:它是源代码发出的 software ILP,不是硬件里同时在飞的条数。作者反复声明所有结论都是操作性的(给定负载下硬件交付什么),不推断任何物理队列结构。环境是 H100 80GB HBM3 SXM5(132 SM、约 50 MB L2),die-C 控制实验在 1830 MHz 稳频、ECC 开、无降频标志下采集。
主格点(HBM 后端、.cg、16B 宽度、132 CTA × 32 warp):
| 实验 | K=2 (GB/s) | K=8 (GB/s) | 降幅 |
| die-B 基线(总流量随 K 涨) | 3111.0 | 1990 | −36.1% |
| die-B 固定工作量,正序 | 3058.8 | 1988.7 | −35.0% |
| die-B 固定工作量,反序 | 3059.2 | 1987.0 | −35.0% |
| die-C 固定工作量(n=10) | 3048.6 | 1976.7 | −35.2% |
四组降幅 35.0%36.1%,聚得很紧。分配量扫描里降幅在 34.95%35.30% 之间平着走,同时 K=8 的 L2 命中率稳定在约 56%、K=2 时不到 1%。计数器审计:K=2→8 之间 DRAM 字节近乎不变(−1.57%+1.25%),发射的 load sector 匹配,但 L2 sector 流量涨约 50%。覆盖面上,288 条资格曲线里只有 8 条饱和,179 条呈峰后下降,101 条到网格边缘还在涨;下降在 L2 后端比 HBM 更普遍,窄宽度低占用格点会消失。
配套两条支线:cp.async(异步拷贝,直写共享内存、无目标寄存器)对匹配的普通加载 kernel,单个 group 时两者差 4% 以内,加到 16 个 group 时 cp.async 快 2.12.9 倍(die-B 2.152.87×,die-C 2.102.81×),占用越高优势越小。延迟阶梯:单 SM 指针追逐测得 L1 命中约 17 ns、L2 驻留约 145 ns、中间平台约 269 ns、HBM 驻留约 353 ns;流式读带宽达到 3350 GB/s 数据手册的 91%94%。
对写访存密集 kernel 的人这是一条可以直接查表用的结论:在 H100 上,每线程独立加载数堆过 2 以后,继续加 ILP 会负收益,普通加载路径掉三分之一带宽。深度学习推理里手写的 GEMM、attention 里(prefill 阶段)的数据预取循环,如果为了延迟隐藏把每线程加载摊得很开,可能正在付这笔税。对策也现成:把加载并到 cp.async / TMA 这类异步路径,论文给的同场景倍率是 2.12.9 倍。当然用 occupancy 换 ILP 是另一条路,但 H100 上 warp 数往往被共享内存和寄存器压着,ILP 是更常见的选择,所以这条曲线值得放进性能调优清单。
对 GPU 模拟器和编译器作者,这条曲线是 Hopper 世代缺失的校准点;Accel-Sim 这类模拟框架的验证数据里没有这个受控轴。
作者自己列得很全,这里挑实的说。只有 4.2 节(普通加载下降)是固定工作量+计数器+分配扫描的完整证据链;cp.async 对比是两颗 die 的纯计时,双流实验是 die-B 单会话观察,die-C 复测形状不同、没有合并。机制没有识别:工作量和扫描次序排除了,简单分配大小依赖排除了,但 L2 命中率随 K 从 1% 涨到 56% 说明聚合请求流确实在变,交织效应和 L2 路径都还开着,定位需要他们没采的 stall-reason 计数器。n=10 里有两个异常 launch(一个慢 08%,一个全程约半速),时钟记录正常、无解释。单 SM 延迟测不出 L2 近远分区结构。跨 die 复测是同仪器两颗 die 加一颗只对形状,不是跨硅片世代的普适性,完整时间复现性他们说要至少四个启动会话。读下来最该补的实验他们也点到了:固定合并度与占用、只变 K 相关地址排序,并做地址 trace 验证覆盖。