PaperlayerView exact revision on arXiv

Source-linked AI summary

Dissecting GPU Utilization for LLM Inference on Nvidia Hopper

Mohammad Siavashi, Gerald Q. Maguire Jr., Dejan Kostic, Marco Chiesa

arXiv:2609.12923v1cs.PFcs.ARcs.DCcs.LG

TL;DR

A single SM-utilization percentage can conflate distinct denominators and obscure how much useful LLM work is done. The paper profiles inference across regimes and models with eight counter-validated views, finding that decode underutilization reflects several mechanisms, including Hopper’s fragment floor. These views connect utilization readings to mechanisms such as fragment fill, occupancy limits, stalls, wave quantization, and kernel selection.

  • Problem

    SM utilization is an elapsed-cycle pipeline-throughput proxy, not occupancy or direct useful matrix-multiply throughput, so one percentage cannot explain the distinct denominators relevant to LLM inference.

    Source: §1 · p. 1

  • Method

    The authors profile vLLM with FA3 and cuBLASLt on an H100 NVL across four models, three serving regimes, and per-layer kernel roles, deriving eight views from raw NCU reports and pinning each to a counter or formula.

    Source: §1 · p. 2

  • Results

    Decode underutilization reflects resource limits, spatial wave quantization, and an instruction-set fragment floor; Hopper’s m64 floor limits small-batch decode dense-GEMM fragment fill to 1.6–12.5%.

    Source: §1 · p. 2, §8 · p. 11

  • Takeaways & Limitations

    The views distinguish decode bottlenecks with different fixes, including raising the row dimension M rather than retiling when the m64 fragment floor limits useful matmul work.

    Source: §1 · p. 2, §8 · p. 11

    The measurements use a single H100 SKU and TP=1, and the authors do not claim absolute values transfer unchanged to other H100 SKUs, tensor-parallel configurations, FP8 paths, or intermediate prefix-cache hit rates.

    Source: §6 · p. 10

Abstract

from arXiv · show

A single SM utilization percentage can make an LLM inference workload look compute-saturated while hiding how much useful work is being done. The problem is not that the counter is wrong, but that it collapses several different mechanisms into one number. This is most severe during decode, where each request contributes only one new token and dense projection GEMMs become small-row matrix multiplications. On Hopper, the bfloat16 GMMA path executes these operations in fixed 64-row matrix fragments, so small-batch decode can fill only a small fraction of each fragment with real token rows. In this paper, we profile vLLM with FlashAttention-3 and cuBLASLt on an H100 NVL across cold prefill, warm prefill, and decode, sweeping sequence length and batch size. We replace the usual single utilization number with eight counter-validated views derived from raw Nsight Compute reports, each pinned to an NCU counter or explicit formula. Together, these views map utilization gaps to concrete mechanisms - fragment fill, occupancy limits, stall signatures, wave quantization, and kernel selection - across four production models and six per-layer kernel roles.

1. Introduction

A single SM-utilization percentage is a pipeline-throughput proxy, not a measure of useful LLM work, and it can conflate distinct denominators. The paper maps utilization readings to mechanisms across inference regimes, models, and kernel roles using counter-validated views.

1. Introduction

  • In Qwen3-14B dense decode at batch size 32, GEMMs report ≈72% sm_busy_pct while only half of each Hopper m64 fragment holds real token rows.Long_scoreboard memory stalls still dominate the issue budget, and the m64 floor is architectural rather than repairable by tile retuning.Source: §1 · p. 1
  • The same projection role reports ∼97% sm_busy_pct in cold prefill but ∼9% in warm prefill, while occupancy reads ∼14% against the 64-warp ceiling and 75–95% against its resource cap.These readings use different denominators and are internally consistent.Source: §1 · p. 1
  • The study profiles four production models on an H100 NVL across cold prefill, warm prefill, and decode, sweeping context length, batch size, and per-layer kernel role.It computes eight utilization views from raw Nsight Compute reports, each pinned to a counter or explicit formula.Source: §1 · p. 1
  • Decode underutilization reflects multiple mechanisms—resource limits, spatial wave quantization, and an instruction-set fragment floor—rather than one bottleneck.The paper ties each mechanism to a different fix and argues that a single scalar cannot reveal them.Source: §1 · pp. 1–2
  • For small-batch decode, Hopper’s GMMA m64 fragment caps dense-GEMM fragment fill at 1.6–12.5%, making per-kernel compute-throughput readings overstate useful matmul throughput by 8–64×.The stated lever is to raise row dimension M, not retile.Source: §1 · p. 2

2. Background and Motivation

The H100’s architectural resources and kernel roles shape how inference workloads use the GPU. Cold prefill, warm prefill, and decode differ in GEMM row dimensions, while occupancy depends on which of two limits is used as its denominator.

2.1. The H100 microarchitectural budget

  • The H100 NVL has 132 SMs and 528 Tensor Cores; its BF16 GMMA instruction uses a fixed 64-row M-axis fragment, and its 60 MB L2 and 3.9 TB/s HBM3 define relevant hardware limits.The BF16 roofline ridge is about 214 FLOP/byte; kernels below it are HBM-bandwidth-bound and those above it are compute-bound.Source: §2.1 · p. 2

2.2. Kernel families in an LLM inference layer

  • Dense layers dispatch QKV, output, Gate/Up, and Down projections plus FA3 attention; FA3 uses a persistent one-block-per-SM kernel followed by a smaller combine kernel.For the MoE model, a fused grouped-GEMM processes the top-8 of 128 experts, preceded by a small router-gate projection.Source: §2.2 · p. 2, Table 1 · p. 2

2.3. Three operational regimes

  • Cold prefill processes full prompts with M ≈ S·B, warm prefill re-evaluates a short suffix with M near B, and decode adds one query token per request against a growing KV cache.The warm-prefill setup reuses cached prefixes and re-evaluates at most one trailing 16-token block; it represents a near-perfect-hit endpoint, not intermediate-hit workloads.Source: §2.3 · pp. 2–3

2.4. Occupancy has two denominators

  • Occupancy can be normalized against the 64-warp architectural ceiling or the kernel’s resource cap, typically 8–12 warps; FA3 main reaches 7.6 warps/SM, or 12% and 95% of those respective limits.Both percentages are valid but answer different questions.Source: §2.4 · p. 3

3. Methodology

The profiling pipeline aggregates per-kernel NCU measurements across model layers and re-reads raw reports for metrics that upstream exports rename or blend. The methodology keeps distinct utilization denominators explicit and estimates per-call memory footprint for L2-fit analysis.

3.1. Profiling pipeline

  • The study profiles kernels with Nsight Compute’s MemoryWorkloadAnalysis, LaunchStats, SpeedOfLight, WarpStateStats, and Occupancy sections, aggregating per-layer, per-role instances through NCU’s CSV export.It re-reads raw report files for stall counters, tensor-core counters, and the FA3 main-versus-combine split; utilization is averaged across matched calls, while time decomposition is summed across roles.Source: §3.1 · p. 3

3.2. Metrics and their denominators

  • Utilization metrics use different denominators: SM utilization is elapsed-cycle based, occupancy uses either the 64-warp architectural limit or the kernel resource cap, and device-wide compute SOL scales by SMs reached out of 132.Per-active-SM and device-wide compute SOL coincide only when a launch reaches all 132 SMs.Source: §3.2 · p. 3

3.3. Working-set model

  • The L2-fit analysis estimates each kernel call’s BF16 footprint as weights plus activation and KV-cache bytes read or written, then compares it with the 60 MB L2 capacity.The notation distinguishes context length S, batch size B, and T = S · B for prefill or T = B for decode.Source: §3.3 · p. 3

4. Workloads, Systems, and Measurement

The measurements use vLLM with BF16, TP=1, FA3, cuBLASLt’s Hopper GMMA path, and CUDA Graphs on an H100 NVL. The sweep covers model-specific regime and shape corners, with a separate B=32 validation arm.

Workloads, Systems, and Measurement

  • The profiling setup runs vLLM in BF16 with TP=1 on an H100 NVL, with FA3, cuBLASLt Hopper GMMA, and CUDA Graphs enabled.Three warmup iterations precede profiling; NCU-profiled iterations are uncaptured, and reported end-to-end timings do not come from profiled iterations.Source: §4 · p. 3
  • Cold prefill is swept densely across all four models, warm prefill samples one small and two large memory-limit corners, and decode sweeps B∈{1, 2, 4, 8}.Qwen3-32B is KV-memory-capped: it has no B=8 decode point, and long contexts allow only B≤4.Source: §4 · p. 3
  • The main profiled grid records each model, regime, S, and B corner, with per-role NCU lines aggregated across model layers; the B=32 validation arm adds four decode cells.The grid has seven dense-model roles or eight for Qwen3-30B-A3B, with warm-prefill rows not always a strict cells-by-roles product.Source: §4 · p. 3

5. The Eight-Metric Characterization

The characterization replaces a single utilization reading with complementary metrics and shows how regime, kernel role, and denominator change the interpretation of GPU activity. Decode reaches most SMs, but dense kernels use only a small fraction of the architectural warp ceiling while approaching their own resource caps.

5.1. SM time-utilization

  • 92% mean cold-prefill dense-GEMM sm_busy_pct (73–97% range) contrasts with 9.4% mean decode dense-GEMM utilization (6.5–12.7% range); warm prefill is 8–11% at (S=2048, B=1).Warm dense-GEMM utilization recovers unevenly to 14–73% at the larger corner, with a 33% mean; FA3 and MoE also vary by regime.Source: §5.1 · pp. 3–4
  • The eight-metric diagnosis distinguishes occupancy against the 64-warp architectural ceiling from occupancy against each kernel’s resource cap, and tracks per-reason stall budget and L2-fit and wave efficiency.The metrics use NCU counters or explicit formulas to characterize utilization from complementary views.Source: §5 · p. 3, §5.1 · p. 4
  • GMMA M-fragment fill is Mactual/(⌈Mactual/64⌉·64), measuring per-instruction M-axis tile fill against the BF16 WGMMA 64-row fragment floor.The fixed m64 minimum is an architectural property of the BF16 wgmma.mma_async instruction.Source: §5.1 · p. 4
  • The metric set also treats dispatch share as a lower bound and kernel-selection diversity across the sweep as a stability surrogate.Kernel-selection diversity counts distinct kernels chosen by the autotuner.Source: §5.1 · p. 4

5.2. SM coverage and warp-pool occupancy

  • Decode dense projections reach 80–132 of 132 SMs, yet achieved warps are only 12–15% of the architectural limit and 75–95% of each kernel’s resource cap.The resource caps are 12 warps for dense GEMMs and 8 for FA3 main; the denominator gap is about 5×.Source: §5.2 · pp. 4–5
  • QKV grid sizes differ by model, while o_proj and down_proj launch 128 CTAs and gate_up launches 92–132 CTAs in decode.Qwen3-14B QKV uses 132 CTAs; Llama-3-8B uses 96, while Qwen3-32B and Qwen3-30B-A3B use 80.Source: §5.2 · p. 4

5.3. Occupancy cap: registers and shared memory

Kernel resource caps are set by the selected tile’s register and shared-memory demands, not simply by the hardware’s architectural warp ceiling. BF16 cuBLASLt tiles are co-limited, while the tested FP8 configurations show both shared-memory-only and co-limited cases.

5.3. Occupancy cap: registers and shared memory

  • FA3 main’s decode heuristic selects a smaller tile that allocates 256 registers per MMA thread and caps occupancy at 8 warps/SM, half the dense-GEMM cap.Both register and shared-memory limits bind for this 256-thread block.Source: §5.3 · p. 5
  • The BF16 tile configuration is selected by a proprietary nvjet heuristic, so its resource caps are verified with NCU runtime occupancy-limiter counters rather than source inspection.This constrains how the BF16 tile-selection mechanism can be independently inspected.Source: §5.3 · p. 5
  • Every profiled dense BF16 kernel is capped at 12 warps/SM, with registers and shared memory co-limiting; the cap does not change with context length or batch size.The cuBLASLt nvjet heuristic selects a small set of balanced tile configurations across projection roles, models, and regimes.Source: §5.3 · p. 5
  • For the tested FP8 CUTLASS tiles, about 90% of instances are SMEM-bound at 8 warps/SM, while the remaining co-limited configurations reach 12.A one-sided SMEM relaxation raises the SMEM-bound configurations from 8 to 16 warps/SM, but not the co-limited BF16 GEMMs; other FP8 kernel paths were untested.Source: §5.3 · pp. 5–6

5.4. Stall-reason decomposition

Dense-GEMM decode and warm prefill are dominated by memory-dependency stalls, while cold prefill shifts toward GMMA and synchronization stalls. Decode also wastes useful matmul capacity inside fixed-size GMMA fragments, so high utilization readings can overstate useful work.

5.4. Stall-reason decomposition

  • Long_scoreboard stalls account for 70–84% of the dense-GEMM per-cycle stall budget in decode and warm prefill, but 19–36% in cold prefill.Cold-prefill stalls shift toward gmma, barrier, and mio_throttle; math_pipe_throttle is not a top-three stall reason in any of 360 dense-GEMM rows.Source: §5.4 · p. 6
  • Llama-3-8B QKV at S=2048, B=1 has nearly identical long_scoreboard shares in decode and warm prefill: 72.7% and 72.8%.At cold prefill S=8190, B=4, only 23.9% of 20.85 cycles per issued instruction are long_scoreboard; barrier and gmma contribute 27.6% and 23.3%.Source: §5.4 · p. 6, Figure 4 · p. 6

5.6. The GMMA fragment-fill ceiling on decode matmul

  • Decode fills only 1.56–12.5% of the BF16 GMMA M-axis fragment across batch sizes 1–8, versus 100% in cold prefill.Across 136 decode rows, mean fragment fill is 5.86%; warm prefill averages 70.2%.Source: §5.6 · p. 7
  • For Llama-3-8B decode QKV at S=2048, B=1, 6.9% device-wide compute SOL corresponds to only about 0.11% useful FLOP throughput after accounting for 1.56% fragment fill.The passage attributes the gap to padded m64 MMAs counted as work; it says this bandwidth-bound shape cannot recover the gap through tile selection.Source: §5.6 · p. 7

5.5. Matmul throughput against the whole array

  • Device-wide compute SOL averages 92.0% in cold prefill and 7.9% in decode, an 11.7× gap across the four dense GEMMs.The measure scales compute SOL by SMs reached; its dense-decode mean is about 7.9% while DRAM SOL is about 64%, consistent with active memory-bound execution.Source: §5.5 · pp. 6–7, Figure 5 · p. 7

5.6. The GMMA fragment-fill ceiling on decode matmul

  • Per-active-SM compute SOL is the same NCU counter as sm_busy_pct, while the device-wide form additionally accounts for the fraction of SMs reached.The dense-GEMM means are 92 / 25 / 9% for cold prefill, warm prefill, and decode, respectively.Source: §5.6 · p. 7

5.7. Wave-quantization waste

Wave efficiency measures how fully CTA launches use the H100’s 132-SM waves, while the L2 fit ratio compares each call’s working set with cache capacity. Decode incurs wave waste and projection-specific L2 spills, with the amount depending on model and kernel role.

5.7. Wave-quantization waste

  • Decode wastes 19.5% of dense-GEMM SM-wave slots, compared with 14.5% in warm prefill and 0.2% in cold prefill.Wave efficiency is the share of per-wave SM slots doing useful work; the remainder is idle trailing-wave capacity.Source: §5.7 · pp. 7–8
  • Decode-QKV launches use 96 CTAs for Llama-3-8B and 80 for Qwen3-32B and Qwen3-30B-A3B, yielding 27% and 39% waste, respectively.Qwen3-14B uses 132 CTAs with 0% waste; these grids remain invariant across the swept sequence lengths and batch sizes.Source: §5.7 · p. 8, Figure 7 · p. 8
  • None of cuBLASLt’s 59 producible Llama-3-8B decode-QKV algorithm/tile pairs yields a 132-multiple CTA grid.The passage identifies a structural tile-space constraint; alternatives with better wave fill can substantially reduce per-CTA work, while the heuristic optimizes throughput rather than wave fill alone.Source: §5.7 · p. 8, Figure 6 · p. 8

5.8. Memory hierarchy and L2 working-set fit

  • Decode FFN projections spill L2 by 2–4× across dense models, while Llama-3-8B decode QKV fits L2 at ρ≈0.80 and larger Qwen QKV spills by 1.2–1.7×.ρ is per-call working-set footprint divided by 60 MB; values above 1 indicate spill across calls.Source: §5.8 · p. 8
  • Across kernel, sequence-length, and batch-size combinations, L2 spills occur in 53.8% of decode cases, versus 90.0% in cold prefill and 92.5% in warm prefill.At batch size 1, the MoE call touches roughly 1/16 the dense-FFN working set; larger batches can raise its effective expert-weight footprint.Source: §5.8 · p. 8

5.9. Compute-bound crossover batch

Decode’s compute–memory crossover lies beyond the measured batch range for dense GEMMs, but B=32 measurements show that behavior depends on model, context length, and kernel role. Kernel choice is stable, while per-token costs and bottlenecks shift across roles and regimes.

5.9. Compute-bound crossover batch

  • Within B≤8, no decode dense GEMM crosses the SOL-balance proxy; extrapolated crossover batches range from ∼25 to ∼49 for dense projections.The flash_attn proxy is closer to balance at 0.84, with extrapolated B∗∼8; dense-projection ratios range from 0.12 to 0.26, and production SLOs rarely permit B≥25.Source: §5.9 · pp. 8–9

5.10. Empirical validation at B=32

  • At B=32, Qwen3-14B dense-GEMM SM-busy rises from ∼12% at B=1 to 71.9%, while per-token time falls from 10.10 to 2.38 ms (∼4.2×).FA3-fwd takes 49.5% of the decode iteration versus 44.0% for dense GEMMs; the measured cell is near the predicted QKV crossover B∗≈31.Source: §5.10 · p. 9
  • At B=32, Qwen3-30B-A3B shares decode time between routed FFN and FA3, with their time shares within 5 percentage points.The fused-MoE share is 36.2% versus 35.7% at B=1, while FA3-fwd rises from 15.2% to 31.6%.Source: §5.10 · p. 9
  • Qwen3-32B reaches 72.6% dense SM-busy at B=32, S=2048, but only 23.8% at B=4, S=16K, where KV capacity caps batch size.At B=32, its per-token time is 0.98 ms; the 94 GB memory budget requires S≲3K for B=32, while the long-context sweep is capped at B≤4.Source: §5.10 · p. 9
  • At B=32, dense GEMMs retain long_scoreboard stalls, whereas FA3-fwd is led by stall_wait, distinguishing memory-dependency stalls from barrier or fixed-latency waits.Dense-GEMM long_scoreboard is 11.5–26.4 cycles/issue; FA3-fwd stall_wait is 1.8 and long_scoreboard is 1.1–1.2.Source: §5.10 · p. 9
  • At B=32, KV-cache writes remain ≤2% of decode iteration time across all four validation cells.The reported reshape_and_cache_flash shares range from 0.69% to 1.55%, so the cost decomposition needs no separate KV-write term.Source: §5.10 · p. 9

5.11. Kernel-selection stability

  • Decode kernel selection is stable: cuBLASLt chooses one kernel name for dense GEMMs and two for FA3 across the sweep.Cold-prefill QKV reaches seven distinct selections, so decode inefficiency is not attributed to unstable kernel choice.Source: §5.11 · p. 9

5.12. Per-token time decomposition

  • Across the main B≤8 sweep, FFN averages 58–66% of per-token cost, making it the dominant cost share across models and regimes.Table 7 separates largely flat code-generation metrics from workload outcomes that shift by orders of magnitude from cold prefill to decode.Source: §5.12 · p. 9, Table 7 · p. 10

6. Discussion

The discussion maps decode bottlenecks to different serving, software, and architecture levers, while qualifying the scope of the measurements. In particular, proposed batching targets and warm-prefill results depend on the measured setup and assumptions.

6. Discussion

  • Batching toward the dense decode crossover B∗∈[25, 49] is a target, but continuous batching alone does not raise the GMMA fragment fill ηM.The discussion distinguishes launch-amortizing persistent kernels from increasing effective token rows, which requires speculative, tree, or multi-token decoding.Source: §6 · p. 10
  • Warm-prefill measurements use synthetic warmup with ≥99% prefix-cache hits, leaving realistic multi-turn intermediate hit rates uncharacterized.This boundary qualifies how directly the warm-prefill characterization represents production cache states.Source: §6 · p. 10
  • The measurements use one H100 SKU and TP=1; tensor parallelism changes per-GEMM shape and adds all-reduce, while HBM-bound values may vary across H100 SKUs.These constraints limit direct transfer of measured kernel behavior to other serving configurations.Source: §6 · p. 10
  • The linear B∗ extrapolation is a diagnostic rather than a predictor, and its B≥25 behavior is refined by the B=32 validation arm.The validation found Qwen3-14B crossing near B∗≈31, Llama-3-8B approaching, Qwen3-32B crossing only at short context, and a MoE/FA3 co-bottleneck.Source: §6 · p. 10
  • The analytical launch-overhead constant τ=5 µs is host- and driver-dependent and excludes runtime-path, replay, and scheduling costs.The working-set model is first-order, omits scratch and tile staging, and is within ∼10–20% of dram__bytes on cross-checked kernels.Source: §6 · p. 10

7. Related Work

Prior work covers end-to-end inference performance, isolated Hopper primitives, roofline analysis, and LLM benchmarks. This paper instead characterizes per-role kernels in an end-to-end serving stack and connects those measurements to microarchitectural counters.

7. Related Work

  • Prior A100/H100 inference work measures end-to-end performance, while Hopper microbenchmarking studies primitives in isolation; this paper maps serving-stack kernels to NCU counters.The profiled stack composes vLLM, FA3, and cuBLASLt behavior in an end-to-end setting.Source: §7 · p. 10
  • FlashDecoding++ [5] proposes a flat-GEMM kernel that pads M to 8; this work measures the stock library path and finds its row deficit persists across available tiles.The authors conclude that their enumeration explains why the problem must be addressed at the custom-kernel level.Source: §7 · p. 10
  • Unlike training-focused MoE studies and MLPerf’s end-to-end latency or throughput, this work reports per-role microarchitectural metrics.Its analysis also extends roofline [25] with wave quantization, kernel-selection stability, and cuBLASLt algorithm-space enumeration.Source: §7 · pp. 10–11

8. Conclusion

The conclusion argues that SM utilization alone does not diagnose useful LLM work because different metrics expose different bottlenecks. Decode makes the distinction concrete: Hopper’s BF16 GMMA fragment floor limits useful dense-GEMM work, while FA3 loses time to a different stall mechanism.

8. Conclusion

  • SM utilization can conceal whether performance is limited by tensor-pipe work, launch-grid coverage, occupancy headroom, or useful GMMA fragment fill.The conclusion recommends counter-pinned, per-role metrics rather than a single device-wide utilization percentage.Source: §8 · p. 11
  • Small-batch decode leaves most rows of Hopper’s BF16 GMMA m64 fragments empty, an instruction-set ceiling that cuBLASLt retuning cannot lift.The stated lever is raising effective row dimension M through cross-request token packing or persistent decode kernels.Source: §8 · p. 11
  • FA3 decode loses time to barrier-style waits rather than memory stalls, so its bottleneck differs from the dense decode GEMM row deficit.The conclusion contrasts these mechanisms to motivate different fixes.Source: §8 · p. 11
Loading 2609.12923v1…