Skip to content
Review

Concurrency Response of Plain Global Loads on the NVIDIA H100

Aug 2026 · 0 citations · 15 references
Computer Science

TL;DR

The 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 the authors' primary configuration.

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.

View source

Similar papers

Conference Aug 2026

Characterizing Predictability–Latency Trade-offs of KV-Cache SSD Offloading in LMCache for LLM Serving Systems

KV-cache offload is widely used to stretch GPU memory for LLM serving, but its storage behavior has not been characterized at the block-device level. In this paper, we study LMCache through realworld multi-session workloads that span same/different context $\times$ same/different prompt, using over 100 stateless requests per workload. Our comparison is observational rather than factorial: it contrasts two realizable deployment snapshots—LMCache 0.3.0 with buffered I/O and the kernel page cache, and LMCache 0.3.16 with prefix-aware deduplication and O_DIRECT. Key findings define the paper. First, the legacy stack appears almost read-free at the SSD layer in all four workloads, but a flush_ram experiment shows that this is conditional on page-cache warmth: once the cache is evicted, the same path issues about 1 GB of sequential reads in 683 ms. Second, deduplication delivers a dramatic write reduction only for the exact-repeat workload: one set collapses from 1.62 GB to 11.75 MB of writes, while the other three sets remain in the 2–5 GB range. Third, O_DIRECT converts a hidden and bimodal read cost into an explicit and stable one: on one set, warm TTFT rises from 27 ms to 93 ms because each of the 99 warm requests re-reads the same ~12 MB partial chunk from SSD. The main conclusion is therefore not “newer is better,” but a system trade-off: predictability versus average-case latency.

Ying He, Dingsen Shi, Yanbo Dai et al. · 0 citations
Preprint Aug 2026

Over the Memory Wall, Into the Instruction Wall: The New Bottleneck in GPU Data Processing

Datacenter GPUs have seen an order-of-magnitude increase in memory bandwidth with the adoption of newer generations of HBM. Meanwhile, GPU database systems are gaining traction, many building on cuDF, an open-source library of GPU relational operators. Previously, query performance was bound by memory bandwidth, but the increase in memory bandwidth has not resulted in a proportional speedup of cuDF kernels. To investigate why performance has not kept up, we built Valk, a performance analysis tool that combines data from multiple profilers. We profile cuDF running TPC-H in-memory on two extremes of hardware capability, the L4 and GH200 GPUs. The GH200 has 13.4$\times$ the memory bandwidth and 2.5$\times$ the instruction throughput of the L4, yet is only 5.2$\times$ faster in running TPC-H. Our analysis shows that when memory bandwidth is increased, kernels become compute bound. From our analysis, we make three recommendations to fully utilize the GPUs'potential for relational workloads when the memory wall is removed: kernels need to 1) make more efficient use of caches, and 2) increase occupancy and/or instruction level parallelism, and 3) execute fewer instructions per memory access.

S. Hepkema, Bowen Wu, Christos Kozyrakis et al. · 0 citations
Preprint Aug 2026

What Irregularity Costs: CUDA C++, Rust, and Triton on a Hash-Blocked GPU Workload

GPU language comparisons are almost always run on tiled dense linear algebra, where every toolchain is good and the differences are small. We implement the same hash-blocked TSDF fusion kernel in CUDA C++, in Rust through NVIDIA's cuda-oxide, and in Triton, and measure it on a workload with the opposite character: an open-addressed hash table with compare-exchange insertion, data-dependent per-lane probe depth, and contended scatter. The result is a split. On the regular stage, which walks a truncation band and accumulates, all three languages land within a small factor of each other. On the irregular stage, which probes and inserts, Rust stays close to hand-written CUDA C++ while Triton is more than an order of magnitude slower. Language choice is nearly free on the work that is usually benchmarked and expensive on the work that is not. We attribute both gaps to specific things the languages cannot express, not to ratios. Triton's cost follows from a probe loop that must run to a compile-time bound and from tl.atomic_cas taking no mask, which forces a scratch structure with no counterpart in CUDA. Rust's cost was invisible in every instruction count: its kernel issues fewer instructions, fewer compare-exchanges and fewer registers at identical occupancy, yet was slower. Hardware counters located it in L1 residency. A GPU-scope atomic load must be coherent across SMs, no NVIDIA L1 is, so the type-correct way to read a shared location bypasses the cache on every access. Triton's bounded probe is also a correctness problem for fusion: at load factors an ordinary depth trajectory reaches, it silently discards blocks and the reconstruction loses patches of surface with nothing reported. We also report a defect found and fixed in cuda-oxide itself, now merged upstream: its scoped atomic load and store could not be called at all in the build mode that produces real kernels.

P. Korolev · 0 citations
Sep 2026

Resource Efficiency and Performance Predictability in A Groupwise, Hardware-Prioritized Cache on NVMe SSDs

Thanks to notable performance and capacity advantages, NVMe SSDs promise an effective cache tier for alleviating the load pressure of back-end storage servers. Compared with DRAM, an NVMe SSD exhibits up to hundreds of times larger capacity but delivers two orders of magnitude smaller bandwidth per gigabyte. This paper reveals that these hardware characteristics challenge fundamental design goals of caches: (1) storing tons of objects in a large-capacity SSD easily induces severe, persistent fragmentation, resulting in low cache space utilization; (2) intensive front-end requests contend for limited SSD bandwidth, causing unpredictable cache lookup latency. To tackle these challenges, this paper presents Gemini. The core of Gemini is heat-informed, group-based object management with a hardware-assisted I/O prioritization mechanism. Specifically, we introduce a tree-based prediction model for agile object grouping and fast reaction to hotspot shift. It provides SSD-friendly, bulk object eviction with a zero-write, remapping-based migration mechanism. In addition, to ensure request performance, we propose a novel I/O model that mitigates bandwidth congestion and prioritizes their processing at the hardware layer. Evaluation results show that Gemini delivers up to 3.8<inline-formula><tex-math notation="LaTeX">${\boldsymbol{\times}}$</tex-math><alternatives><mml:math><mml:mrow><mml:mo mathvariant="bold">×</mml:mo></mml:mrow></mml:math><inline-graphic xlink:href="cai-ieq1-3711088.gif"/></alternatives></inline-formula> throughput improvement, 2<inline-formula><tex-math notation="LaTeX">${\boldsymbol{\times}}$</tex-math><alternatives><mml:math><mml:mrow><mml:mo mathvariant="bold">×</mml:mo></mml:mrow></mml:math><inline-graphic xlink:href="cai-ieq2-3711088.gif"/></alternatives></inline-formula> tail latency reduction, and 2.7<inline-formula><tex-math notation="LaTeX">${\boldsymbol{\times}}$</tex-math><alternatives><mml:math><mml:mrow><mml:mo mathvariant="bold">×</mml:mo></mml:mrow></mml:math><inline-graphic xlink:href="cai-ieq3-3711088.gif"/></alternatives></inline-formula> higher cache space utilization than state-of-the-art caches for both YCSB workloads and production applications.

Miao Cai, Junru Shen, Baoliu Ye · 0 citations
Preprint Jul 2026

Every Microsecond Matters: Achieving Near Speed-of-Light Latency in GPU Collectives

GPU collective communication is typically optimized for bandwidth, yet many emerging workloads are increasingly limited by latency. Long-context decode-heavy large language model (LLM) inference is a prime example, where serving large models requires multiple GPUs, and many small collectives lie directly on the critical path of token generation. Therefore, even microsecond of overhead can impact performance and cost. In this work, we study how to approach the hardware Speed-of-Light (SoL) lower bound for GPU collectives within a scale-up network. We identify key principles for near-optimal designs, including barrier-free synchronization and efficient use of symmetric memory and multicast. Building on NCCL's device-side API, we develop low-latency interfaces for constructing custom collective kernels and use them to implement new symmetric collectives in NCCL. Microbenchmarks show substantial latency reductions for small and medium messages, reducing overhead to within 7% of the absolute SoL lower bound. When integrated into real applications, these kernels improve inter-token latency and throughput in LLM inference and accelerate cuSOLVERMp, demonstrating benefits for both AI inference and traditional HPC workloads.

Siyuan Shen, Anton Korzh, J. Bachan et al. · 0 citations