Concurrency Response of Plain Global Loads on the NVIDIA H100
Somashekar Manjunath, Rahul Ramachandra M
cs.AR, cs.PF
2026-08-16
Clean-room microbenchmarks on three H100 dies show plain-load bandwidth peaks at 2 independent ld.global per thread and falls 35% by 8, surviving fixed-work and reversed-order controls; cp.async reaches 2.1-2.9x the matched plain-load kernel.
A memory-bound CUDA kernel saturates HBM bandwidth by keeping enough bytes in flight. Little's Law on an H100, with a few hundred nanoseconds of latency against 3.35 TB/s of peak bandwidth, works out to roughly 10 KB in flight per SM. Software raises that number two ways: more warps (thread-level parallelism) or more independent loads issued per thread (instruction-level parallelism, ILP). The second route looks free: registers and issue slots are available, so more loads per thread should never hurt.
This paper shows that intuition is wrong on Hopper. Plain global loads (ld.global) respond to per-thread offered load with a rise-then-fall curve: bandwidth peaks near K≈2 and drops about 35% by K=8. Outstanding-request limits were last inferred for Fermi and Kepler; Hopper-era studies dissected the architecture and the TMA paths but never varied per-thread ILP at fixed logical work. Two authors from CoreOptX built that controlled experiment.
The instrument is a clean-room CUDA kernel built from public ISA documents and SASS-audited to issue exactly K independent loads per iteration with zero register spills; loaded values are XOR-accumulated into a sink the compiler cannot prove dead. The real craft is in dismantling confounds:
K is defined honestly as offered software ILP, not the hardware in-flight count. Every claim is operational: what the hardware delivers for a given offered load, with no inference about physical queues. The hardware is H100 80GB HBM3 SXM5 (132 SMs, 50 MB L2); die-C runs held a steady 1830 MHz SM clock with ECC on and no throttling flags.
Primary cell (HBM backend, .cg, 16B width, 132 CTAs × 32 warps):
| Arm | K=2 (GB/s) | K=8 (GB/s) | Decline |
| die-B baseline (traffic scales with K) | 3111.0 | 1990 | −36.1% |
| die-B fixed-work, ascending | 3058.8 | 1988.7 | −35.0% |
| die-B fixed-work, reversed | 3059.2 | 1987.0 | −35.0% |
| die-C fixed-work (n=10) | 3048.6 | 1976.7 | −35.2% |
The four declines cluster at 35.0–36.1%. Across the 40× allocation sweep the decline stays within 34.95–35.30% while the L2 hit rate at K=8 holds near 56% (under 1% at K=2). Counters: DRAM bytes nearly constant over K=2→8 (−1.57% to +1.25%), issued load sectors matched, L2-sector traffic up about 50%. Coverage-wise, only 8 of 288 qualification curves saturate; 179 show the post-peak decline, 101 are still rising at the grid edge. The decline is more prevalent on the L2 backend than HBM and vanishes at some narrow-width, low-occupancy cells.
Two side results: cp.async sustains 2.1–2.9× the bandwidth of a matched plain-load kernel at 16 offered groups (die-B 2.15–2.87×, die-C 2.10–2.81×), with the gap shrinking as occupancy rises; the two agree within 4% at a single group. The latency ladder on one SM reads L1 17 ns, L2-resident 145 ns, an intermediate plateau 269 ns, HBM-resident 353 ns, with streaming bandwidth at 91–94% of the 3,350 GB/s datasheet.
For anyone writing memory-bound kernels, this is a lookup-table conclusion: on H100, pushing independent loads per thread past 2 costs you a third of the plain-load bandwidth. Hand-written GEMM and attention prefill loops that spread loads wide for latency hiding may be paying this tax right now. The fix is ready-made: move loads onto cp.async/TMA-style asynchronous paths, where the paper measures a 2.1–2.9× advantage in the same setting. Trading ILP for occupancy is the alternative, but warps on H100 are usually capped by shared memory and registers, which makes ILP the common lever, so this curve belongs in the tuning checklist.
For GPU simulator and compiler authors, the curve is a missing Hopper-generation calibration point; Accel-Sim-class frameworks have no validated data on this controlled axis.
The authors are unusually candid. Only Section 4.2 (the plain-load decline) carries the full fixed-work, counter-checked, allocation-swept evidence chain; the cp.async comparison is timing-only across two dies, and the two-stream observation is a die-B single-session result whose die-C check differed and was not pooled. The mechanism is unidentified: work volume, sweep order, and simple allocation-size dependence are excluded, but the L2 hit rate rising from 1% to 56% with K proves the aggregate request stream changes, so interleaving and L2-path explanations remain open; pinning them needs stall-reason counters that were not collected. Two of ten primary launches were anomalous (one 0–8% slow, one at roughly half speed) with normal recorded clocks and no explanation. Single-SM latency cannot resolve near/far L2 structure. The cross-die replication is two dies with one instrument plus a third matching on shape only, not cross-silicon generality; a full temporal reproducibility bound would need at least four boot sessions. The discriminating experiment they name, varying the K-dependent address ordering at fixed coalescing, occupancy, and logical bytes with address-trace verification, is left to future work.