Arithmetic intensity and the roofline

Here is the symptom that starts every serving investigation. You load Llama-3-8B in bf16 onto an 80 GB H100, send a single request, and watch it generate. You get a couple of hundred tokens per second. You then send thirty-two requests concurrently and get roughly an order of magnitude more in aggregate — many times the work for barely more wall-clock time per step. The GPU did not get faster. Something about the first measurement was structurally wasteful, and the waste is not in your code. §3 derives both of those figures as bandwidth floors, so you can check the claim before you believe it.

Profile first, then identify the limiting resource. Arithmetic intensity relates useful work to transferred bytes, but better kernels may reduce actual rereads, improve achieved bandwidth, fuse launches, or change the algorithm. The roofline rules out improvements from extra arithmetic peak alone on a bandwidth-bound kernel; it does not rule out software optimization.

Think of the GPU as a workshop with an enormously productive machine floor and a single loading dock. Work arrives as pallets of bytes. If a pallet keeps the floor busy for a long time, the dock idles and the floor is the limit; if it is consumed instantly, the floor idles and the dock is the limit. Arithmetic intensity is how much work comes on one pallet: FLOPs performed per byte fetched from HBM.

The roofline plot is the picture of that. Intensity runs along x and achievable throughput along y. On the left, throughput is a diagonal — you are bandwidth-limited, so doubling intensity doubles throughput. On the right it is a flat ceiling — you are compute-limited, and more intensity buys nothing. The corner between them is the ridge point, and where your kernel lands relative to it is the entire diagnosis.

Read Figure 1 once and the rest of the book gets easier. Prefill sits on the ceiling. Decode sits in the far bottom-left, two and a half orders of magnitude from the corner. They are not two settings of the same workload; they are two different machines running on the same silicon, which is why §1.1 gives them a chapter of their own.

Definitions

Let $\pi$ be achievable peak compute in FLOP/s for the dtype you are actually using, and $\beta$ achievable HBM bandwidth in bytes/s. For a kernel, let $F$ be the FLOPs it performs and $Q$ the bytes it must move across the HBM boundary. Then

Three things about $Q$ matter and are routinely got wrong.

• For the optimistic calculations below, $Q$ is compulsory traffic, not total traffic — each distinct byte that must cross the HBM boundary at least once, counted once. A value re-read from L2 or SRAM does not count again. So the $I$ computed below is an upper bound on what any kernel can achieve for that computation; a kernel that spills tiles back to HBM sits further left than the algebra says. FlashAttention (§3.2) is exactly the story of reducing materialized intermediates and HBM traffic; finite SRAM still requires tile rereads. An operational roofline can instead use measured total HBM traffic, counting every HBM reread.
• $Q$ counts reads and writes. Inputs, weights, and outputs.
• $\pi$ and $\beta$ are achievable, not advertised. Spec-sheet peak assumes perfect tiling, no throttling, and no tail effects. Every number in this chapter uses the spec sheet because that is all a machine without a GPU can honestly cite; lab 01 replaces both with numbers measured on your card, and every prediction in this book should be re-checked against them.

Composing kernels, not only totals

For sequential kernels, a tighter lower bound is $\sum_i\max(F_i/\pi_i,Q_i/\beta_i)$, not merely $\max(\sum_iF_i/\pi,\sum_iQ_i/\beta)$. A compute-heavy kernel taking at least 3 ms and a separate bandwidth-heavy kernel taking at least 4 ms need at least 7 ms if they cannot overlap, even when the aggregate roofline suggests 4 ms. Launch gaps and synchronization add further costs. Shared-prefix reuse, quantization metadata, tensor-parallel communication, and speculative multi-token verification each change the accounting assumptions.

The H100 ridge point

Hold that number: 295 FLOP per byte. To saturate an H100's bf16 tensor cores you must do 295 floating-point operations for every single byte you pull out of HBM. That is a lot of arithmetic per byte, and almost nothing in decode comes close.

(a) A prefill GEMM

Take one linear projection: activations $X$ of shape $[T, d]$ times weights $W$ of shape $[d, d]$, producing $[T, d]$, all in a dtype of $b$ bytes. $T$ is the number of tokens in the batch, $d$ the hidden size.

The factor 2 in $F$ is one multiply plus one add per MAC. In $Q$: read the activations, read the weights, write the result. For bf16, $b = 2$ and this collapses to $I = Td / (2T + d)$.

Two limits carry the meaning. As $T \to \infty$, $I \to d/2$ — 2048 FLOP/byte for Llama-3-8B's $d = 4096$, seven times past the ridge, because the weight matrix is read once and reused by every token. At $T = 1$, $I \to d/(d+2) \approx 1$, and the GEMM is not a GEMM at all.

Setting $I(T) = I^{*}$ and solving gives the batch at which a square projection becomes compute-bound:

Derived. A Llama-3-8B projection needs about 345 tokens in flight before an H100's bf16 tensor cores are the limit. A 1,024-token prompt is comfortably past it; so is every value vLLM resolves max_num_batched_tokens to on an H100 — the 2,048 declared on SchedulerConfig is a fallback, and EngineArgs.get_batch_defaults (vllm/engine/arg_utils.py:L2580-L2671) picks 8,192 for vllm serve and 16,384 for the offline LLM class on any device with at least 70 GiB that is not an A100. Note the denominator: if $d \le b\,I^{*} = 590$, no batch size whatsoever makes a square bf16 GEMM compute-bound on an H100. Small models are structurally memory-bound.

(b) A decode GEMV

Now set $T = 1$ — one token per sequence, batch 1, which is what a decode step is. Keep the general $[1, K] \times [K, N]$ shape:

For bf16, $I$ approaches 1.0 FLOP per byte — the fused QKV projection at $K = 4096$, $N = 6144$ gives $50{,}331{,}648$ FLOPs over $50{,}352{,}128$ bytes, an intensity of 0.9996. Notice what dropped out: $K$, $N$, the layer, the model, the head count. Only $b$ survives. Every dense weight matrix, read once and applied to one token, does exactly $2/b$ FLOPs per byte. You load two bytes of bf16 weight and do one multiply-add with it, forever.

On the roofline that is $\beta \cdot I = 3.35$ TFLOP/s against a ceiling of 989.4: a batch-1 bf16 decode step can use at most 0.34% of the H100's advertised tensor-core throughput. Quantizing weights to fp8 doubles $I$ to 2.0 and doubles achievable throughput — still 0.68% of peak, but the time halved, because in the memory-bound regime the ideal memory-time term is bytes over bandwidth. Conversion, scale metadata, launch gaps, and achieved bandwidth determine the realized speedup. That is the whole argument for weight-only quantization, made before opening §4.1.

(c) Decode attention against the KV cache

The other half of a decode step reads the KV cache. For one sequence at context length $s$, one layer, with $h$ query heads, $h_{kv}$ key/value heads and head dimension $d_h$: the kernel must read $K$ and $V$ for all $s$ positions, then compute $q K^{\top}$ and $\text{softmax} \cdot V$ for each of the $h$ query heads.

The context length cancels. Decode attention has a fixed arithmetic intensity that depends only on the GQA group size $g = h/h_{kv}$ and the dtype: $I = 2g/b$, which for bf16 is simply $g$. Llama-3-8B has $g = 32/8 = 4$, so $I = 4$ FLOP/byte — 74 times left of the ridge. Multi-head attention ($g = 1$) gives $I = 1$; multi-query ($h_{kv} = 1$, $g = 32$) gives $I = 32$. GQA and MQA are visible on the roofline as pure rightward moves, and §3.5 is where they get built.

So if intensity does not degrade with context, what gets worse? First, the bytes grow linearly in $s$ even though the ratio does not, so the attention part of the step grows linearly with context while the weight part stays flat. Second — and this is the structural asymmetry the whole book turns on — batching amortises weight reads but does not amortise KV reads. Every sequence has its own cache, so doubling the batch doubles both the FLOPs and the bytes of attention and $I_{\text{attn}}$ is invariant in $B$ as well as in $s$. Taking $B \to \infty$ at fixed $s$, the blended intensity of a whole decode step tends to

where $P_{\text{stream}}$ is the parameters actually streamed per token. For Llama-3-8B at $s = 8192$ this is $(15.01 \times 10^{9} + 4.29 \times 10^{9}) / (1.074 \times 10^{9}) \approx 18$ FLOP/byte. Derived. At 8k context, no batch size on earth makes a Llama-3-8B decode step compute-bound on an H100 — a ceiling of 18 against a ridge of 295. Long context is not a capacity problem batching solves; it is a bandwidth problem only fewer KV bytes solve. That is the roofline predicting prefix caching (§2.3), KV quantization (§2.5) and MLA (§7.2) before any of them are described.

Batching, and why it is the biggest lever

Put $B$ sequences in a decode step, one token each. The flattened activation is $[B, d]$ and hits the same $[d, d]$ weight — algebraically identical to the prefill GEMM with $T \to B$. So the weight-GEMM intensity of a decode step is $I = Bd/(2B + d)$ for bf16, and batching walks you up the diagonal of Figure 1.

Sixteen times the batch for sixteen times the throughput at almost the same step time. That is the single largest multiplier available in LLM serving, and it is exactly what continuous batching exists to harvest (§1.3) — which is also why that chapter is about admission and eviction rather than matmuls. The hard part is keeping $B$ high, not making the GEMM fast.

The decode step lower bound

Because decode is memory-bound, its step time has a hard floor set by bandwidth:

Work it for Llama-3-8B in bf16, $L = 32$, $d = 4096$, $h = 32$, $h_{kv} = 8$, $d_h = 128$, vocabulary 128,256. Per layer: fused QKV $4096 \times 6144$, output $4096 \times 4096$, gate+up $2 \times 4096 \times 14336$, down $14336 \times 4096$ — 218.1 M parameters, times 32 layers is 6.98 B. Add the 525 M-parameter output head and you stream 7.50 B parameters per token. The input embedding table is gathered, not streamed — one row of 8 KB, not 1.05 GB — so $\text{bytes}_{\text{weights}} = 15.01$ GB, not the 16.06 GB the full 8.03 B parameter count suggests. KV per token is $2 \cdot L \cdot h_{kv} \cdot d_h \cdot b = 2 \times 32 \times 8 \times 128 \times 2 = 131{,}072$ bytes, exactly 128 KiB.

All four derived from the formula above with the cited 3.35 TB/s. The first pair is the number to carry around: a single H100 cannot generate Llama-3-8B faster than about 223 tokens per second at batch 1 in bf16 under one full target-model forward per accepted token and the stated streamed-weight assumptions. Speculative verification can emit multiple accepted tokens per weight stream; it is outside this particular ceiling. Not because the kernels are bad — because streaming 15 GB of weights at 3.35 TB/s takes 4.48 ms. Measure 6 ms and you have 1.5 ms of slack to hunt for in launch overhead and CUDA graphs (§8.1); measure 4.6 ms and you are done, go quantize.

The last KPI is the cruel one. At 8k context, KV traffic per token is 1.07 GB, so aggregate throughput cannot exceed $3.35\times10^{12} / 1.07\times10^{9} \approx 3{,}120$ tokens/s regardless of batch size — and you cannot even reach that. Capacity stops you long before bandwidth does: §0.1 budgets the same card at $0.92 \times 79.65 - 14.96 - 6 \approx 52.32$ GiB of KV pool — note that the weight term there is the full 14.96 GiB of resident parameters, not the 15.01 GB streamed per step, because the embedding table occupies memory whether or not it is read — which is about 429k tokens at 128 KiB each, or roughly 52 sequences at 8k context. Well short of the 345 the ridge wants. Part 2 is that fight.

The negative results are as useful as the positive ones. At $I = 1$ a kernel with greater arithmetic throughput alone cannot raise the bandwidth ceiling — you would be making the idle half of the machine idler. A compute-bound kernel may see little benefit from reduced bytes, but a full step mixes kernels with different bottlenecks. Most disappointing optimisations in this field are optimisations applied on the wrong side of the ridge.

Neither engine contains a file called roofline.py. What both contain is machinery whose sole purpose is to raise $T$ — the number of tokens handed to one forward pass — because it improves weight reuse; reducing transferred bytes is another way to move rightward.

vLLM: one flat token budget, then one flat GEMM

The V1 scheduler runs a token budget over running and waiting requests, and hands the model runner a per-request token count rather than a batch of sequences.

        req_to_new_blocks: dict[str, KVCacheBlocks] = {}
        num_scheduled_tokens: dict[str, int] = {}
        token_budget = self.max_num_scheduled_tokens
        spec = self.vllm_config.speculative_config
        draft_slots = spec.max_num_new_slots_for_drafting if spec is not None else 0
        input_budget = self.scheduler_config.max_num_batched_tokens

            num_new_tokens = (
                request.num_tokens_with_spec
                + request.num_output_placeholders
                - request.num_computed_tokens
            )
            if 0 < self.scheduler_config.long_prefill_token_threshold < num_new_tokens:
                num_new_tokens = self.scheduler_config.long_prefill_token_threshold
            num_new_tokens = min(
                num_new_tokens, token_budget, input_budget - draft_slots
            )

Each admitted request draws down the budget at vllm/v1/core/sched/scheduler.py:L704-L710. The default budget is small and explicitly labelled as a testing convenience — real deployments set it much higher, which is the point:

    DEFAULT_MAX_NUM_BATCHED_TOKENS: ClassVar[int] = 2048
    DEFAULT_MAX_NUM_BATCHED_TOKENS_FOR_BATCHED_DP: ClassVar[int] = 256
    DEFAULT_MAX_NUM_SEQS: ClassVar[int] = 128
# ...
    max_num_batched_tokens: int = Field(default=DEFAULT_MAX_NUM_BATCHED_TOKENS, ge=1)
    """Maximum number of tokens that can be processed in a single iteration.

    The default value here is mainly for convenience when testing.
    In real usage, this should be set in `EngineArgs.create_engine_config`.
    """

The model runner then flattens every request's scheduled tokens into one contiguous token axis. This is where the roofline point is physically created — the $M$ dimension of every weight GEMM in the forward pass is total_num_scheduled_tokens. Two GPU runners ship at this SHA and a dense model such as Llama-3-8B actually runs on the newer one (vllm/v1/worker/gpu/model_runner.py, selected by VllmConfig.use_v2_model_runner); the older file quoted here does the same flattening in NumPy rather than Triton, which is why it is the readable one. §11.4 owns the selection policy and reads both.

        total_num_scheduled_tokens = scheduler_output.total_num_scheduled_tokens
        assert total_num_scheduled_tokens > 0
        num_reqs = self.input_batch.num_reqs
        assert num_reqs > 0

        # OPTIMIZATION: Start copying the block table first.
        # This way, we can overlap the copy with the following CPU operations.
        self.input_batch.block_table.commit_block_table(num_reqs)

        # Get request indices.
        # E.g., [2, 5, 3] -> [0, 0, 1, 1, 1, 1, 1, 2, 2, 2]
        req_indices = np.repeat(self.arange_np[:num_reqs], num_scheduled_tokens)

        # cu_num_tokens: [2, 5, 3] -> [2, 7, 10]
        # self.query_pos.np[:10]: [0, 1, 0, 1, 2, 3, 4, 0, 1, 2]
        cu_num_tokens = self._get_cumsum_and_arange(
            num_scheduled_tokens, self.query_pos.np
        )

The comment's example — [2, 5, 3] tokens for three requests — is the roofline in miniature: three sequences become one $M = 10$ GEMM. Chunked prefill (§1.5) exists so a long prompt can be split across steps to top up this $M$ while decode tokens ride along, keeping every step near the ridge instead of alternating between a compute-bound prefill and a memory-bound decode.

SGLang: the same budget, a different shape

SGLang carries the equivalent knob as max_prefill_tokens, resolved from the model worker at python/sglang/srt/managers/scheduler.py:L1032-L1042 and declared as:

    max_prefill_tokens: A[
        int,
        Arg(
            help=(
                "The maximum number of tokens in a prefill batch. The real bound "
                "will be the maximum of this value and the model's maximum "
                "context length." + f"\n\n{human_readable_int.__doc__}"
            ),
            type_parser=human_readable_int,
        ),
        NS("schedule"),

The admission loop lives in PrefillAdder (python/sglang/srt/managers/schedule_policy.py:L511) — SGLang's per-step admission controller, taken apart properly in §1.4. It tracks three separate budgets and reports why it stopped:

    def budget_state(self):
        no_token = self.rem_total_tokens <= 0 or self.cur_rem_tokens <= 0
        if not no_token and self.is_hybrid_swa:
            no_token = self.rem_swa_tokens <= 0
        # Gate new mamba slots separately: rem_total_tokens' full_evictable can't
        # cover a mamba slot, which needs mamba-recoverable bytes (see __init__).
        if not no_token and self.rem_mamba_slots is not None:
            no_token = self.rem_mamba_slots <= 0
        if no_token:
            return AddReqResult.NO_TOKEN

        if self.rem_input_tokens <= 0:
            return AddReqResult.OTHER

        if self.dllm_config is not None:
            if self.rem_dllm_tokens <= 0:
                return AddReqResult.OTHER
        else:
            if self.rem_chunk_tokens is not None and self.rem_chunk_tokens <= 0:
                return AddReqResult.OTHER

Where they differ. vLLM keeps one scalar token_budget plus an input_budget and mixes prefill and decode tokens in a single undifferentiated flat batch. SGLang separates rem_total_tokens (KV capacity — what stops you reaching the ridge), rem_input_tokens (the GEMM-shape budget — what gets you there) and rem_chunk_tokens, and returns a typed reason for refusal. Same two roofline pressures; SGLang makes the distinction between out of batch and out of KV memory explicit in the return value, which is why its scheduler logs tell you which one bit.

The microbenchmark that draws the curve

vLLM's kernel benchmarks contain a runnable version of §3(a). benchmark_fp8_gemm.py sweeps $M$ from 1 to 16,384 against fixed model weight shapes and reports TFLOP/s:

@triton.testing.perf_report(
    triton.testing.Benchmark(
        x_names=["batch_size"],
        x_vals=[1, 16, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384],
        x_log=False,
        line_arg="provider",
        line_vals=_enabled,
        line_names=_enabled,
        ylabel="TFLOP/s (larger is better)",
        plot_name="BF16 vs FP8 GEMMs",
        args={},
    )
)
def benchmark(batch_size, provider, N, K):
# ...
    to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)

2 * M * N * K is exactly the $F$ of §3(a) and x_vals straddles the $T^{*} \approx 345$ crossing, so the printed TFLOP/s column is the roofline's y-axis, sampled — the closest thing in either repo to a first-party roofline experiment.

SGLang's single-batch harness times prefill and decode separately — precisely the two-point measurement Figure 1 needs. Both engines' historical entry points are now shims: python/sglang/bench_one_batch.py:L1-L22 re-exports from sglang.benchmark.one_batch, and benchmarks/benchmark_latency.py:L1-L17 points at vllm bench latency. The call path is cli_main() → main() → latency_test() → latency_test_run_once(), and inside that last function the two phases are bracketed independently. Prefill:

    model_runner.synchronize()
    tic = time.perf_counter()
    next_token_ids, _, batch = model_runner.extend(reqs)
    model_runner.synchronize()
    prefill_latency = time.perf_counter() - tic
# ...
    throughput = input_len * batch_size / prefill_latency
    rank_print(
        f"Prefill. latency: {prefill_latency:6.5f} s, throughput: {throughput:9.2f} token/s"
    )

Then decode, one timed iteration at a time, with the median reported rather than the mean — the first decode step after a prefill is not representative:

    for i in range(output_len - 1):
        model_runner.synchronize()
# ...
        tic = time.perf_counter()
        next_token_ids, _ = model_runner.decode(next_token_ids, batch)
        model_runner.synchronize()
        latency = time.perf_counter() - tic
# ...
        tot_latency += latency
        throughput = batch_size / latency
        decode_latencies.append(latency)

Two details make this a valid roofline probe rather than a vibe. The model_runner.synchronize() calls on both sides of the timed region drain the asynchronous CUDA stream, so perf_counter measures completed wall-clock work including host dispatch and synchronization overhead, not just enqueue time. CUDA events delimit device-stream time when that is the desired quantity. And extend() (python/sglang/benchmark/one_batch.py:L487-L521) versus decode() (python/sglang/benchmark/one_batch.py:L526-L537) differ only in prepare_for_extend versus prepare_for_decode and the resulting token count — same weights, same kernels, different $T$.

Walk the bytes for one decode step at batch 1, context 1024: 15.01 GB of weights streamed once, plus $1024 \times 131072 = 0.134$ GB of KV, is 15.14 GB — floor 4.52 ms, about 221 token/s. At --batch-size 64 the weights are still read once and KV becomes 8.59 GB, so 23.6 GB — floor 7.05 ms, about 9,080 token/s. All derived; the harness prints the measured version and the residual is the lesson.

Three commands, in increasing order of what they tell you.

# vLLM's GEMM sweep: M from 1 to 16384 against real Llama-3.1-8B weight shapes.
# The TFLOP/s column is the roofline y-axis, sampled.
python benchmarks/kernels/benchmark_fp8_gemm.py \
    --models meta-llama/Llama-3.1-8B-Instruct --tp-sizes 1

# SGLang. Prints "Prefill. latency ... throughput" and "Decode. median latency ...".
# Defaults are batch_size=(1,), input_len=(1024,), output_len=(16,)
#   — python/sglang/benchmark/one_batch.py:L196-L200
python -m sglang.benchmark.one_batch --model-path meta-llama/Meta-Llama-3-8B-Instruct \
    --batch-size 1 64 --input-len 1024 --output-len 16

# vLLM. Defaults: --input-len 32 --output-len 128 --batch-size 8
#   --num-iters-warmup 10 --num-iters 30  (vllm/benchmarks/latency.py:L34-L54)
vllm bench latency --model meta-llama/Meta-Llama-3-8B-Instruct \
    --batch-size 1 --input-len 32 --output-len 128

The third one gives an end-to-end latency for 128 output tokens (vllm/benchmarks/latency.py:L125-L135 wraps llm.generate in perf_counter). Divide by 128 and compare against the 4.48 ms derived floor. The gap is your budget for CUDA graphs, sampling, detokenization, and scheduler overhead — every one of which gets a chapter later.

Lab 01, 01-measure-your-gpu-roofline, does this properly: it sweeps matmul shapes to find achieved TFLOP/s, runs a streaming kernel to find achieved GB/s, computes your card's real $I^{*}$, and plots Figure 1 with your numbers. Every prediction in this book should be re-derived against them.

• Williams, Waterman, Patterson, Roofline: An Insightful Visual Performance Model for Multicore Architectures, CACM 2009 — the original model. Everything here is that paper with $\pi$ and $\beta$ swapped for GPU numbers.
• NVIDIA H100 Tensor Core GPU datasheet and the NVIDIA H100 Tensor Core GPU Architecture whitepaper — the source of the 989.4 TFLOP/s and 3.35 TB/s used throughout. Read the sparsity footnotes.
• NVIDIA, Matrix Multiplication Background User's Guide — the arithmetic-intensity treatment of GEMM shapes, including why skinny $M$ is pathological.
• Pope et al., Efficiently Scaling Transformer Inference (2022) — the operational-intensity analysis of PaLM serving that made the memory-bound-decode argument standard.
• Dao et al., FlashAttention (2022) and FlashAttention-2 (2023) — the canonical example of moving actual traffic down to compulsory traffic. Built properly in §3.2.
• vLLM PR #3130 — chunked prefill and issue #3861, plus Agrawal et al., Sarathi-Serve (OSDI 2024): an unusually clear record of engineers arguing about where on the roofline a mixed batch should sit. Picked up in §1.5.