Concurrency Response of Plain Global Loads on the NVIDIA H100

Abstract

The bandwidth a memory-bound GPU kernel sustains is set by how many bytes it keeps in flight. We use Little's Law here as throughput accounting, not as a measured hardware pool. CUDA fills that budget on Hopper through plain loads (ld.global) and asynchronous copies (cp.async), among other paths; we characterize their concurrency response with clean-room microbenchmarks on three H100 SXM5 dies. Our main result concerns the plain-load path: attained LDG bandwidth peaks at a small offered per-thread load (K ~ 2) and then declines, by about 35% from K=2 to K=8 at our primary configuration. The decline survives a fixed-work control matching total issued logical loads across K, ascending and reversed sweep orders, and replication on two dies with the same instrument (-35.0% and -35.2%). Separately profiled counters show DRAM bytes nearly constant over K=2->8 while L2-sector traffic rises, and a 40x nominal allocation-size sweep (512 MB to 20 GB, all above the ~50 MB L2; no address trace) leaves the decline essentially unchanged, disfavoring a simple allocation-size dependence. Because the L2 hit-rate nonetheless rises with K at every allocation, the aggregate request stream does change with K; we report K as offered software ILP and leave the hardware mechanism open. A preliminary survey adds a matched cp.async-versus-plain-load comparison (2.1-2.9x at high offered depth, two dies), a die-B same-CTA two-stream observation whose companion die-C check differs and is not pooled, and a cross-die primitive baseline.