GPU Datapath: Where Every Byte Moves
in One LLM Inference Step

Five Things the Byte Ledger Says

Performance folklore about LLM inference is full of plausible claims that are rarely checked against hardware counters: fusion saves bandwidth, caching helps decode, attention is the bottleneck. Measuring the actual byte traffic at every level of the memory hierarchy lets each claim be confirmed, quantified, or retired. The five findings below are the ones that change how a serving system should be tuned.

The Hierarchy a Byte Must Climb

Every operand of every kernel travels the same physical ladder: HBM holds the model, the device-wide L2 caches it, each SM's unified SRAM splits into a hardware-managed L1 and an explicitly-managed shared memory, the register file feeds the compute pipes, and Tensor Cores consume operands from registers (Ampere) or directly from shared memory (Hopper's wgmma). What distinguishes the levels is not just speed — it is who manages them.

One Step = 394 Kernel Launches

Tracing a vLLM engine step for Llama-3.1-8B (32 layers) shows a fixed structural inventory: 12 kernels per layer plus ~10 framework kernels — 394 launches per step. Every one of them falls into one of the ten classes below; together the classes cover 100% of the step's GPU time.

The sweep resolved the last, arch-dependent class: the separate splitKreduce_kernel appeared only on A100 and L40S, and only for the down projection (K=14336, the deepest reduction) at 64–256 tokens — exactly where split-K theory predicts it, a deep dot-product with too few output tiles to fill the GPU. Hopper never dispatched it at any measured shape: its persistent segment_k schedules fold the reduction into the main kernel.

Measuring the Real Kernels, One at a Time

Attaching Nsight Compute to a live multi-process vLLM server is slow and fragile, so we use sentinel microbenches that dispatch the exact same kernels: GEMMs go through torch.matmul (the cuBLASLt path vLLM uses, with CUBLASLT_LOG_LEVEL=5 recording the dispatched tile and algorithm as evidence), attention calls vllm.vllm_flash_attn.flash_attn_varlen_func (FA3 on Hopper, FA2 elsewhere; decode uses the paged block_table path with shuffled pages), and the light ops call the registered vLLM CUDA ops directly through torch.ops._C. Each cell wraps a single call in cudaProfilerStart/Stop so ncu profiles exactly one launch after 10 warmups, with clocks pinned to base.

# one cell of the 39-cell sweep: GEMM/light ops at {1, 64, 256, 2048} tokens; attention at {1024, 2048, 4096} + paged decode
ncu --clock-control base --replay-mode kernel --profile-from-start off \
    -k 'regex:gemm|cutlass|Kernel2|wgmma|ampere|sm80|sm90|xmma|gemv' --launch-count 2 \
    --metrics dram__bytes_read.sum,dram__bytes_write.sum,lts__t_sector_hit_rate.pct,\
l1tex__m_xbar2l1tex_read_bytes.sum,l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum,\
smsp__inst_executed_op_ldsm.sum,sm__sass_inst_executed_op_ldgsts.sum,... \
    --csv --page raw \
    python experiments/ncu/profile_datapath.py --op gemm --proj qkv --m 2048

The two flags doing the real work are -k, a regex over demangled kernel base names that selects which launch to profile, and --launch-count 2, which also captures the secondary kernel some cells dispatch (the split-K reduce after decode GEMMs, the combine pass after split-KV attention). Each requested counter then maps onto one column of the ledger:

Six Datapaths, Measured

The ten kernel classes of the inventory collapse into six distinct movement patterns, and each pattern uses the hierarchy differently: some multiply bandwidth through reuse, some stream through it once, and some never leave the L2 at all. This section walks each pattern with its measured byte counts, beginning with the most reuse-rich kernel and ending with the most latency-bound one.

A · Prefill GEMM — the reuse cathedral

The QKV projection at M=2048 tokens (4096×6144 weight) is the canonical compute-bound kernel. Watch the byte counts grow as data climbs toward the SMs on H200:

Read path — operand supply

Write path — epilogue

H200, qkv_m2048: amplification 1.46 (114+21 MB moved vs the 92.3 MB theoretical min). Read supply 1.61 GB = 768 tiles × 2.1 MB panel fetches — every DRAM-delivered byte served ~14× (24× per unique byte: weight rows read by 16 CTA rows, activation columns by 48), 92.9% answered by L2 (the 61.8% sector hit rate ncu reports counts accesses, not bytes). Writes land in write-back L2 first; 20.9 of 25.2 MB sank to DRAM before kernel end. Grid 112 persistent CTAs, 0.85 waves, accumulators pinned in registers for the whole K-loop.

The same logical GEMM walks a different micro-datapath on each architecture, and the instruction counters prove it. On H200 the kernel issued only 768 cp.async and zero ldmatrix instructions — operands ride the TMA engine into shared memory and wgmma consumes them there directly, skipping the register file on the way in. On A100 the same cell issued 2.36 M cp.async and 6.32 M ldmatrix instructions: bulk copies into SMEM, then explicit fragment loads into registers for mma.sync. Two generations, one algorithm, two physically different byte routes.

Anatomy of the 113.9 MB, sweep by sweep

The persistent grid (28×4 = 112 CTAs) walks the 16×48 tile grid four rows at a time — four sweeps for M=2048, each touching all 48 weight panels once. Sweep 1 starts with a cold L2: every panel fetch is compulsory (W 50.3 MB as each column's cluster first touches it, plus X rows 0–3, 4.2 MB). But the kernel's full footprint (W+X+C ≈ 92 MB) far exceeds the 60 MB L2, and all 48 weight panels are re-touched every sweep: as each sweep's later fetches and dirty output compete for the same cache sets, LRU evicts the earliest weight panels before sweep 2's CTAs (rows 4–7) reuse them — so those panels miss and re-read from HBM. Measured: about 31% of W (15.7 MB) per sweep boundary.

HBM → L2 read pressure over kernel time (model anchored to the measured total)

Integral = 54.5 + 3 × 19.86 = 114.1 MB — exactly the measured DRAM reads. Peak pressure is 24% of H200's 4.8 TB/s, which is why the re-reads hide under compute (TC duty 95.5%): the cost is energy and power budget, not kernel time.

The model has exactly one fitted constant — the per-sweep eviction fraction e — and a dedicated 9-point M-sweep on H200 pins it: every M from 256 to 8192 dispatches the same 128×128×64 kernel, and the measured re-reads grow linearly with the sweep count as predicted (e = 0.365, R² = 0.978). The scatter around the line tracks cuBLASLt's launch-shape choice — the heuristic picks anything from 28×4 to 60×2, 120×1 and 14×8 across this sweep, and each shape walks the tiles in a different order with slightly different locality.

B · Decode GEMM & LM head — the weight stream

At M=1 (one new token), every projection degenerates to reading the entire weight matrix once to produce a vector. Amplification collapses to 1.00–1.07 across all four projections and the LM head on H100, H200 and A100 (L40S sits apart at 1.13–1.20 — see the callout below): there is no reuse to exploit, because there is only one row of activations. The LM head is the heaviest single stream — 1.05 GB per step, 250 µs on H200 at 4.2 TB/s (88% of peak bandwidth).

C · FlashAttention prefill — the quadratic that stays on-chip

Attention is the one operator whose intermediate state grows quadratically with sequence length, so its datapath is decided by a single question: where does the score matrix live? FlashAttention's answer — in shared memory and registers, tile by tile, never materialized — can be verified directly from the byte counters of the A100 cell at 2048 tokens:

Read path — Q/K/V tile supply

Write path — output O

A100-80GB, attnp_s2048 (FA2, causal, GQA 32:8): DRAM total 42.7 MB vs 41.9 MB minimum — amplification 1.02. Read supply 307 MB = 11.8× re-read of Q/K/V tiles, 91.5% answered by L2; the s² score matrix lives entirely in SMEM and registers. A100's 40 MB L2 evicts the output promptly, so 99% of O lands in DRAM within the kernel. At s=4096 the service ratio rises to 21.9× (1.15 GB served from 52.6 MB of DRAM reads).

A naive attention would materialize the 2048×2048 score matrix per head — hundreds of megabytes — write it to HBM after QKᵀ, read it back for softmax, and read it again for PV. The ledger shows none of that: DRAM traffic is Q+K+V+O and nothing else. The quadratic intermediate is born in Tensor-Core accumulators, renormalized in registers by online softmax, and dies there. This single design decision is why the measured 42.7 MB replaces what would otherwise be >1 GB of HBM traffic.

FA3 on Hopper sharpens the same picture. The H100 cell at s=2048 finishes in 87.7 µs — 3.9× faster than A100's FA2 at 342.6 µs — with DRAM reads of 25.25 MB, matching the 25.17 MB logical Q+K+V to within 0.3%. Its DRAM amplification even drops to 0.66, because 85% of the output write is still sitting in L2 when the kernel ends. H200 runs the same cell in 91.5 µs, and L40S's FA2 needs 366 µs at 78% Tensor-Core duty. The TMA + warpgroup pipeline does not change what moves; it changes how little time the movement takes.

D · Paged decode attention — the zero-reuse stream

Decode attention reads a paged KV cache through a block table (we shuffle the pages to reproduce vLLM's scattered layout) with one query token per sequence. At batch 64, context 2048 on A100: 537.55 MB read against a 537.92 MB minimum. Amplification 1.00. L2 service ratio 1.0. The split-KV kernel parallelizes along the context dimension to fill SMs, and the page-granular (16-token) reads remain fully coalesced — scattering pages costs nothing measurable at this granularity. There is simply no locality here to mine: each step touches every cached KV byte exactly once.

E · RMSNorm, RoPE, embedding — the L2 residents

The per-layer elementwise kernels read and write activations whose entire working set (16–50 MB at T=2048) fits inside L2. The ledger makes their datapath unmistakable: Tensor Cores 0%, shared memory essentially unused, and — the striking part — DRAM writes far below logical writes, because the write-back L2 holds the dirty lines for the next consumer.

At decode (T=1) the same kernels become almost weightless — RMSNorm launches a single thread block (one CTA on a 132-SM GPU, 4.3 µs) and the whole op is launch-latency, not bytes. Thirty-two layers of these per step is why kernel-launch overhead, not arithmetic, dominates the light-op budget during decode.

F · Greedy sampling — one block, one scan

The final argmax over the 128256-entry logits row reads 0.51 MB of fp32 and runs as a single-block reduction — pure latency, 0.01% Tensor-Core duty, one CTA: 45.2 µs on H100, 46.1 on H200, 59.8 on L40S and 69.3 on A100, none of them above 13 GB/s. At batch 64 the same kernel reads 32.9 MB at real bandwidth (740 GB/s on H200, 742 on H100, 628 on L40S, 481 on A100). Like everything else at batch 1, it is bound by being small, not by being slow.

Adding Up a Whole Step

Multiplying each measured cell by its structural call count (32 layers; 64 RMSNorms; one LM head, embedding and argmax) reconstructs the full step's HBM budget. The four numbers below pair each phase's measured traffic with its theoretical minimum — the gap between them is the entire optimization opportunity:

The decode floor is easy to sanity-check: 15.6 GB at H100's 3.35 TB/s is 4.7 ms of pure streaming; the measured kernel-time sum is 7.7 ms (base clock, no inter-kernel gaps counted), so the byte ledger explains the majority of decode latency from first principles. The complete H200 ledger gives the same shape — 15.59 GB against a 15.29 GB minimum, a 3.2 ms floor against 6.3 ms of kernel time — as do A100 (7.7 ms floor vs 13.0 ms) and L40S, where the floor explains fully 86% of measured kernel time (20.4 vs 23.8 ms). Prefill is different: H100's excess 10.7 GB comes almost entirely from the four projection GEMMs' amplification (×32 layers), the one place where a better tiling/persistent-kernel strategy has real bytes to recover.

Where the Removable Bytes Are

An optimization is only as large as the traffic it can actually remove, so the right way to rank candidates is to subtract each kernel's theoretical minimum from its measured bytes and see what remains. Applying that subtraction across the whole step ledger produces the table below — and it disagrees with intuition in instructive ways.

Reproducing the Measurements

Every number on this page comes from a 39-cell sweep that runs unattended as a single SLURM job per GPU and finishes in about seven minutes. The harness is three small files — a microbench multiplexer, a parser, and the sweep template — and the same recipe ports to any CUDA GPU with profiling-counter permission. Each run also archives its evidence: per-cell wide-CSV counters, the cuBLASLt dispatch logs, and three full Nsight Compute reports whose Memory Workload Analysis charts draw the same flows this page tabulates.

The fastest way to explore a single kernel is to capture one cell with the full section set and open the report in the Nsight Compute UI, whose Memory Workload Analysis chart draws the same per-level byte flow this page tabulates. The sanity-run line first confirms the op dispatches correctly before paying the profiling overhead:

# profile one op locally (any CUDA GPU with ncu counter permission)
python experiments/ncu/profile_datapath.py --op attn_decode --m 64 --ctx 2048   # sanity-run without ncu
ncu --clock-control base --profile-from-start off \
    -k 'regex:flash|fmha|attn|device_kernel' --launch-count 2 --set full -o attnd_b64 \
    python experiments/ncu/profile_datapath.py --op attn_decode --m 64 --ctx 2048
# then open attnd_b64.ncu-rep in the Nsight Compute UI and read the Memory Workload Analysis chart