Track papers across labs, topics, and trends.
Search papers, labs, and topics across Lattice.
This study investigates the concurrency response of plain global loads on the NVIDIA H100 GPU, revealing that the achieved bandwidth peaks at a low per-thread load (K ~ 2) and declines by approximately 35% as the load increases to K=8. The findings are consistent across multiple configurations and dies, indicating that while the L2 hit-rate increases with K, the overall request stream changes, suggesting complex interactions in memory access patterns. The research utilizes clean-room microbenchmarks to provide precise measurements of bandwidth performance under varying loads and allocation sizes.
Bandwidth on the NVIDIA H100 drops by 35% when increasing the per-thread load from K=2 to K=8, challenging assumptions about scaling performance in memory-bound applications.
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.