ML Interview Notes
30 min read12 sections
Part 4 · Quantization · 04-04

When quantization pays — and when it does not

Status
SOURCE PINNED
Primary sources
  • benchmarks/kernels/
  • python/sglang/benchmark/one_batch.py
Edition pins
vllm a556f3f · sglang 7d89325

You swapped a bf16 checkpoint for a 4-bit one, batch-1 latency fell by a factor of three, and aggregate throughput under load went down. Nothing is broken. You moved a workload across a boundary that the roofline predicted, and this chapter derives where that boundary is, what the other three boundaries are, and how to price the accuracy you spent getting there.

§1

The problem

Two deployments, one model, one card. Llama-3-8B on an H100 SXM.

Deployment A is an interactive coding assistant. Two or three concurrent streams, and the number the users feel is inter-token latency. Quantizing the weights to INT4 group-128 with Marlin takes the decode step from a bandwidth floor of 4.48 ms to about 1.15 ms (§0.4's floor, recomputed at 4.125 bits per stored weight). That is close to a pure win.

Deployment B is a batch summarisation job. Five hundred concurrent sequences, and the number that matters is tokens per second per GPU. Same checkpoint, same kernel, same card, and the step time is now worse than bf16 — by a fixed 0.45 ms per step, at every batch size above roughly 270. At this batch that is about 6%; the penalty in milliseconds never shrinks, only its share of a growing step does. The dequantisation you are paying for on every forward pass no longer buys any bandwidth back, because bandwidth stopped being the constraint.

Both statements come from the same two-term model, and neither is a measurement — I have no GPU. What follows derives the model, marks its boundaries, then separates it from the other reason people quantize: quantization frees VRAM, VRAM buys concurrency, and concurrency buys throughput on an axis where the step-time argument is silent.

Standing on

§4.1 owns the design space and $B^{*} = I^{*}b_w/2$. §4.2 owns GPTQ, AWQ and SmoothQuant. §4.3 owns the kernels and $M^{*} = (b_{\text{eff}}/8)\cdot P/(2\beta)$. This chapter owns the decision and the measurement methodology, and re-derives none of them.

§2

Mental model

A decode step does two things that can overlap: it streams weights out of HBM, and it issues math. Quantization changes both terms, but not in the same way. Weight-only schemes divide the memory term and add a constant to the compute term. Weight-and-activation schemes divide the memory term and the compute term, because they unlock a faster tensor core. So the two families have different shapes: weight-only wins big and then fades to a loss; W8A8 wins a constant factor everywhere and never fades.

Figure 1 — decode step time against batch size, Llama-3-8B on H100 SXM. Derived from $t = \max\!\left(W b_w/\beta,\; 2WB/P + W c/D\right)$ with $W = 7.50$ B streamed parameters, $\beta = 3.35$ TB/s, $P_{\text{bf16}} = 989.4$ TFLOP/s, $P_{\text{fp8}} = 1978.9$ TFLOP/s. Not measured. Log batch axis.

0 4 8 12 16 1 4 16 64 256 1024 batch size B (concurrent decoding sequences) step time (ms) bf16 — 4.48 ms floor W8A8 FP8 — 2.24 ms floor W4A16 g128 — 1.15 ms floor B = 118 W8A8 takes the lead B = 266 W4A16 becomes a loss vs bf16 bf16 + 0.45 ms 2x in ideal model REGIME I — weight-bandwidth-bound fewest weight bytes wins REGIME II REGIME III — compute-bound fastest tensor core wins

Read the right-hand edge. At batch 1024 the W4A16 curve sits above bf16 by a fixed 0.45 ms and stays there. That gap is the dequantisation work, which is proportional to the number of weights and therefore constant in $B$, while the useful math grows linearly. Weight-only quantization does not become slightly less good at large batch; it becomes a permanent tax.

§3

First principles: two terms and three regimes

Model one decode step as a memory pipe and a math pipe that overlap:

$$ t(B) \;=\; \max\!\left(\underbrace{\frac{W\,b_w}{\beta}}_{\text{weight stream}},\;\; \underbrace{\frac{2WB}{P}}_{\text{mma}} \;+\; \underbrace{\frac{W c}{D}}_{\text{dequant}}\right) $$

$W$ is the number of parameters streamed per token; $b_w$ bytes per stored weight; $\beta$ HBM bandwidth; $B$ the number of decoding sequences in the batch; $P$ the peak of the tensor-core instruction the kernel actually issues; $c$ the number of non-tensor instruction slots spent per weight converting it back to the mma's input type; $D$ the non-tensor issue throughput. Dequantisation sits inside the compute term because it runs on the CUDA cores, competes with the mma for issue slots, and — crucially — is hidden underneath the weight stream only while the weight stream is the longer of the two.

For Llama-3-8B, §0.4 established $W = 7.50$ B streamed parameters (the embedding table is gathered, not streamed), giving 15.01 GB in bf16 and the 4.48 ms floor. Everything below reuses that $W$.

One assumption, stated once

Every quantized figure in this chapter applies the scheme to all 7.50 B streamed parameters, output head included. Most published W4A16 checkpoints leave the head in bf16 (§4.1.3), which is 0.53 B of the 7.50: the W4A16 floor becomes 1.40 ms instead of 1.15, the 8B capacity rows below lose about 1.5 GiB of KV pool. Recompute every intersection: the fixed bf16 head cost does not shrink with the quantized-layer memory and dequantization terms. The 70B rows are insensitive to it — the head is 0.7% of that model.

Counting $c$ from the kernel, not from intuition

Marlin's INT4-to-BF16 conversion is two specialisations in dequant.h, and you can count instructions off the page:

csrc/libtorch_stable/quantization/marlin/dequant.h:L173-L199 vLLM
template <>
__device__ inline void dequant<nv_bfloat162, vllm::kU4B8.id(), true>(
    int q, nv_bfloat162* frag_b) {
  static constexpr uint32_t MASK = 0x000f000f;
  static constexpr uint32_t EX = 0x43004300;

  // Guarantee that the `(a & b) | c` operations are LOP3s.
  // clang-format off
  int lo = lop3<(0xf0 & 0xcc) | 0xaa>(q, MASK, EX);
  q >>= 4;
  int hi = lop3<(0xf0 & 0xcc) | 0xaa>(q, MASK, EX);
  // clang-format on

  frag_b[0] = *reinterpret_cast<nv_bfloat162*>(&lo);
  frag_b[1] = *reinterpret_cast<nv_bfloat162*>(&hi);
}

template <>
__device__ inline void dequant<nv_bfloat162, vllm::kU4B8.id(), false>(
    int q, nv_bfloat162* frag_b) {
  dequant<nv_bfloat162, vllm::kU4B8.id(), true>(q, frag_b);

  static constexpr uint32_t SUB = 0x43084308;

  frag_b[0] = __hsub2(frag_b[0], *reinterpret_cast<const nv_bfloat162*>(&SUB));
  frag_b[1] = __hsub2(frag_b[1], *reinterpret_cast<const nv_bfloat162*>(&SUB));
}

Four weights come out of one 32-bit q. Producing them costs two lop3s and one shift; removing the INT4 zero-offset costs two __hsub2; applying the group scale costs the three instructions of scale() at csrc/libtorch_stable/quantization/marlin/marlin_template.h:L104-L116 (a num2num2 broadcast plus two __hmul2). Eight non-tensor instructions per four weights: $c = 2$. A GPTQ-symmetric INT4 checkpoint with bf16 activations takes exactly this path — dequant_skip_flop at csrc/libtorch_stable/quantization/marlin/marlin_template.h:L348-L354 is false when there are no zero-points and the scalar type is nv_bfloat16.

For $D$, take the H100 SXM datasheet FP32 vector rate of 67 TFLOP/s. Counting an FMA as two FLOP, that is $D = 33.5\times10^{12}$ instruction-lanes per second, roughly 30 times slower than the bf16 tensor core's 989.4 TFLOP/s. Then

$$ t_{\text{deq}} \;=\; \frac{W c}{D} \;=\; \frac{7.50\times10^{9}\times 2}{33.5\times10^{12}} \;=\; 0.448~\text{ms}. $$
What this model assumes

That dequantisation issue slots and mma issue slots add rather than overlap, and that the FP32 vector rate is a fair proxy for the lop3/__hsub2/__hmul2 mix. Both are approximations that no arithmetic can settle. What is not an approximation is the sign: above the ridge, W4A16 issues the same bf16 mma instructions as bf16 and additionally issues dequantisation, so it cannot be faster. The magnitude of the penalty is what Lab 06 measures.

The three regimes

Instantiate the model for the three schemes on an H100. The memory terms are $15.01/3.35 = 4.48$ ms (bf16), $7.50/3.35 = 2.24$ ms (8-bit) and $7.50 \times 0.5156/3.35 = 1.15$ ms (INT4 group-128, $b_{\text{eff}} = 4.125$ bits). The mma slopes are $2W/P$: 15.16 µs per unit of batch on bf16 tensor cores and 7.58 µs on fp8 tensor cores. Setting curves equal gives the boundaries.

The three regimes for Llama-3-8B on H100 SXM. Derived from the two-term model above; speedups are ratios of derived step times, not measurements.
RegimeBatchBinding constraintbf16W4A16W8A8 FP8Winner
IB < 118weight bandwidth4.48 ms1.15–2.24 ms2.24 msW4A16 (up to 3.9×)
II118 – 266mixed4.48 ms2.24–4.48 ms2.24 msW8A8 (2.0×)
IIIB > 266tensor-core issue15.16B µs15.16B µs + 0.45 ms7.58B µsW8A8 (2.0×); W4A16 is a loss
3.9×
W4A16 over bf16 at B = 1
B = 118
W8A8 overtakes W4A16
B = 266
W4A16 falls behind bf16
2.00×
W8A8 over bf16, ideal model

Three consequences worth stating separately.

W8A8-FP8's speedup is batch-invariant in this idealized model. Halving both terms halves their maximum. Real kernels, mixed-precision layers, activation scaling, KV traffic and hardware support need not preserve that factor. Treat it as a candidate for throughput-oriented serving, gated by quality and measurement, not a universal default.

W4A16's advantage decays before its own ridge. Setting $2WB/P + Wc/D = Wb_w/\beta$ gives $B = 47$, not the $M^{*} \approx 76$ that §4.3 derives. The 29-sequence difference is the dequantisation constant eating into the memory-bound headroom; $M^{*}$ is the crossover with dequantisation assumed free, and it is a useful arithmetic reference. A real dispatcher still uses hardware/shape-specific support and measured kernel behavior; a deployment comparison needs end-to-end measurements.

The bf16 ridge at $B^{*} = 295$ is not where W4A16 dies. W4A16 stops paying at 266, before bf16 itself becomes compute-bound, because it is racing bf16's floor, not bf16's slope. Between 266 and 295 both curves are flat-versus-rising, and the quantized one is already above.

§4

The other axis: capacity, not latency

Everything above holds $B$ fixed. But the reason most people quantize a 70B model is that they cannot hold $B$ fixed — the weights leave no room for KV cache, so $B$ is capped at something useless. This is a different effect with a different unit, and conflating the two is the most common analysis error in this area.

Use §2.1's budget exactly: $C = 79.65$ GiB per H100, $u = 0.92$ (vllm/config/cache.py:L80-L87), and a 6.0 GiB allowance for activations, CUDA graphs and non-torch memory. That leaves 67.28 GiB for weights plus KV, per GPU.

Concurrency ceiling, not latency. Derived from the budget above; KV in bf16, 128 KiB/token for 8B and 320 KiB/token for 70B (divided by TP degree). Requests are 8k tokens. Not measured.
DeploymentWeight formatweights GiB/GPUKV pool GiBresident tokens8k requests
Llama-3-8B, TP=1bf1614.9652.32428,61352
Llama-3-8B, TP=1fp8 W87.4859.80489,87759
Llama-3-8B, TP=1INT4 g1283.8663.42519,55263
Llama-3-70B, TP=2bf1665.701.5710,3101
Llama-3-70B, TP=2fp8 W832.8534.43225,61227
Llama-3-70B, TP=2INT4 g12816.9450.34329,89840
Llama-3-70B, TP=8bf1616.4350.851,333,049162
Llama-3-70B, TP=8INT4 g1284.2363.041,652,638201

The 8B rows move by 21%. The 70B TP=2 rows move by a factor of 40. That is the whole capacity argument in two lines: quantization buys concurrency in proportion to how much of your VRAM the weights were eating. When weights are 22% of the budget (8B), freeing three quarters of them buys you a fifth more sequences. When weights are 98% of the budget (70B at TP=2), freeing three quarters of them is the difference between a server that cannot hold two requests and one that holds forty.

And the throughput consequence is not the step-time ratio. Going from $B=1$ to $B=40$ at Llama-3-70B multiplies decode throughput by close to 40, because §0.4's memory-bound step time barely moves with $B$ in Regime I. The per-step effect in §4.4.3 is worth 2–4×; this one is worth 40×, and it is invisible to a batch-1 latency benchmark.

Figure 2 — two independent axes. Left: step time at fixed batch, which does not care how much VRAM is free. Right: concurrency ceiling, which does not care how fast the GEMM is. Both derived above; a scheme can win on one and lose on the other.

AXIS A — per-step latency (Llama-3-8B, B = 1) AXIS B — concurrency ceiling (70B, TP = 2) bf16 4.48 ms W8A8 2.24 ms W4A16 1.15 ms Same 80 GB card either way. Freeing VRAM changes nothing on this axis: the step still streams the same bytes. bf16 1 request fp8 W8 27 INT4 40 A 4x faster kernel changes nothing on this axis. Only bytes of resident weight do. 40x more sequences in flight.
§5

The MoE special case: routing divides the batch

Everything so far assumed every weight sees every row of the batch. In a mixture-of-experts layer it does not, and that single fact makes expert weights the best quantization target in any large model.

Take the routed-expert share of the parameters, derived from the published configs. Qwen3-30B-A3B (num_hidden_layers 48, hidden_size 2048, num_experts 128, num_experts_per_tok 8, moe_intermediate_size 768) has $128 \times 3 \times 2048 \times 768 \times 48 = 28.99$ B expert parameters out of 30.53 B total — 95.0%. DeepSeek-V3 (61 layers, 3 dense, 256 routed experts, moe_intermediate_size 2048, $d = 7168$) has 653.9 B of 670.9 B — 97.5%. Both totals reconcile with the published 30.5B/3.0B and 671B/37B headline figures, which is the check that the arithmetic is right.

Now the part that matters. Under roughly uniform routing, a batch of $B$ decoding sequences sends $B k / E$ rows to each expert, where $k$ is num_experts_per_tok and $E$ the expert count. The expert GEMM's effective $M$ is therefore

$$ M_e \;=\; \frac{B\,k}{E}, \qquad\text{so the expert layer stays memory-bound until}\qquad B^{*}_{\text{MoE}} \;=\; M^{*}\cdot\frac{E}{k}. $$

With $M^{*} \approx 76$ for W4A16 from §4.3: Qwen3-30B-A3B has $E/k = 16$, so its expert layers stay in Regime I up to $B \approx 1{,}216$; DeepSeek-V3 has $E/k = 32$, so up to $B \approx 2{,}432$. The dense attention projections in the same models crossed at 76. Routing buys the expert weights a winning region 16 to 32 times wider than the dense layers get.

Figure 3 — why experts stay memory-bound. Qwen3-30B-A3B at a serving batch of 512. Derived from $M_e = Bk/E$ with $E = 128$, $k = 8$ read from the published config.

batch B = 512 hidden 2048 router k=8 expert 0 M_e = 32 rows expert 1 M_e = 32 rows expert 2 M_e = 32 rows ... 125 more ... expert 127 M_e = 32 rows 32 < M* = 76 memory-bound: W4A16 still wins dense attention projections M = 512 rows, all of them 512 >> 76

Both engines encode exactly this asymmetry as first-class configuration. In vLLM the online shorthand table splits the spec by layer kind, and two entries deliberately leave dense linear layers alone:

vllm/config/quantization.py:L138-L147 vLLM
    ),
    # INT8 weight-only on MoE; linear stays unquantized (no `linear` field).
    "int8_per_channel_weight_only": QuantizationConfigArgs(
        moe=QuantSpec(weight=kInt8StaticChannelSym),
    ),
    # Online NVFP4 on MoE with per-token dynamic activation scales (Blackwell +
    # FlashInfer TRTLLM only); linear stays unquantized (no `linear` field).
    "nvfp4_per_token": QuantizationConfigArgs(
        moe=QuantSpec(weight=kNvfp4Static),
    ),

The container is QuantSpec(weight=..., activation=...) at vllm/config/quantization.py:L65-L77, and QuantizationConfigArgs carries one for linear and one for moe. Leaving linear at None means "inherit the method's default", which for an online scheme is unquantized — so --quantization int8_per_channel_weight_only quantizes 95% of a Qwen3-30B-A3B's bytes and touches none of the layers that would have been pushed past their ridge.

SGLang reaches the same place by a different route: a config class per mixed scheme rather than a per-layer-kind spec. W4AFp8Config dispatches on the layer's Python type:

python/sglang/srt/layers/quantization/w4afp8.py:L93-L105 SGLang
    def get_quant_method(
        self, layer: torch.nn.Module, prefix: str
    ) -> Optional[QuantizeMethodBase]:
        from sglang.srt.layers.linear import LinearBase
        from sglang.srt.layers.moe.fused_moe_triton import FusedMoE

        if isinstance(layer, LinearBase):
            if is_layer_skipped(prefix, self.ignored_layers):
                return UnquantizedLinearMethod()
            return Fp8LinearMethod(self)
        elif isinstance(layer, FusedMoE):
            return W4AFp8MoEMethod(self)
        return None

FP8 (W8A8) on the dense linear layers — which are compute-bound at serving batch, and where faster tensor cores are the only thing that helps — and INT4 weights with FP8 activations on the experts, which are memory-bound at the same batch. That is Figure 1's regime table applied twice within one model. SGLang's registry at python/sglang/srt/layers/quantization/__init__.py:L71-L101 lists three such MoE-targeted mixed schemes (w4afp8, quark_int4fp8_moe, mxfp_w4a8). The difference is ergonomic, not architectural: vLLM lets you compose an arbitrary pair on the command line, SGLang ships pre-validated pairs.

§6

Measuring accuracy honestly

The performance side of this chapter is arithmetic. The accuracy side is not, and it is where most quantization decisions are actually made badly. This section is the book's methodology; other chapters point here.

Perplexity is the wrong statistic

Perplexity averages the log-likelihood over every token in a corpus. Quantization error is not distributed that way: it changes a small number of high-leverage decisions and leaves the rest alone. Accuracy is Not All You Need (arXiv:2407.09141) makes the sharper version of the point — even aggregate accuracy is too coarse. Their finding is that a compressed model can match the baseline's benchmark score while a substantial fraction of individual answers flip, correct-to-incorrect and incorrect-to-correct, and that these flips correlate strongly with KL divergence from the baseline's output distribution. Their conclusion, quoted from the abstract, is that "compression techniques should also be evaluated using distance metrics", and they propose KL divergence and flip rate.

The practical form: run the same prompts through both servers and report score delta, flip rate, and mean per-token KL. Zero score delta with a 12% flip rate has changed your product; a 1-point delta with a 2% flip rate mostly has not.

Degradation is task-dependent, and not in the obvious direction

Two published sweeps disagree with the folk model, which is worth knowing before you pick an eval. Exploring the Trade-Offs (arXiv:2409.11055), covering 1B to 405B across four quantization methods and 13 datasets, reports that quantized models "often struggle with instruction-following and hallucination detection", that an MT-Bench LLM judge "highlights significant performance declines in Coding and STEM tasks", and — the counterintuitive one — that "hard tasks do not always experience the largest accuracy losses, indicating that quantization magnifies a model's inherent weaknesses rather than simply correlating with task difficulty".

Quantization Hurts Reasoning? (arXiv:2504.04823) sweeps DeepSeek-R1-distilled Qwen and Llama from 1.5B to 70B on AIME, MATH-500, GPQA and LiveCodeBench, finds W8A8 and W4A16 near-lossless with severe degradation below that, and explicitly refutes a widely repeated claim: "Contrary to expectations, quantized models do not exhibit increased output lengths." Generation length is not a cheap proxy for reasoning damage.

The operational conclusion is not "reasoning degrades more". It is: your eval must be the task you serve. Quantization amplifies whatever your model was already weakest at, and no published table tells you what that is for your fine-tune.

How many samples make a difference real

Accuracy on $n$ questions is a binomial proportion. At $p = 0.80$ and $n = 1319$ (the full GSM8K test set), the standard error is $\sqrt{0.8 \times 0.2 / 1319} = 1.10$ percentage points. A 95% confidence interval on a single score is therefore $\pm 2.16$ pp. Comparing two independent runs inflates the error by $\sqrt{2}$: the smallest difference GSM8K can resolve at 95% confidence is 3.1 pp. To resolve 1 pp at 95% confidence with 80% power you would need about 25,000 questions per arm. GSM8K has 1,319.

Two escapes. First, evaluate paired — same questions, both models, McNemar's test on the discordant pairs. The standard error becomes $\sqrt{f/n}$ where $f$ is the flip rate; at $f = 0.10$ that is 0.87 pp, resolving a 1.7 pp difference. This is the same flip statistic arXiv:2407.09141 argues for, arrived at from the statistics side. Second, stop trying to resolve 1 pp. If your decision changes at 1 pp of GSM8K, the decision is not about accuracy.

What a project actually accepts

The only accuracy thresholds that ship inside either repository are vLLM's GSM8K CI gates. The TurboQuant KV-cache configs are the cleanest example, because the same model appears at four aggressiveness levels:

tests/evals/gsm8k/configs/Qwen3-4B-TQ-k8v4.yaml:L1-L5 + tests/evals/gsm8k/configs/Qwen3-4B-TQ-t3nc.yaml:L1-L5 vLLM
model_name: "Qwen/Qwen3-4B"
accuracy_threshold: 0.80
num_questions: 1319
num_fewshot: 5
server_args: "--kv-cache-dtype turboquant_k8v4 --max-model-len 4096"
# ---
model_name: "Qwen/Qwen3-4B"
accuracy_threshold: 0.75
num_questions: 1319
num_fewshot: 5
server_args: "--kv-cache-dtype turboquant_3bit_nc --max-model-len 4096"

Four configs, four floors: k8v4 and t4nc at 0.80, k3v4nc at 0.78, t3nc at 0.75. Going from 4-bit to 3-bit KV buys a lower bar by 5 points, and a maintainer wrote that number down. That is the concrete answer to "what does a project consider acceptable degradation" — but read what the harness does with it:

tests/evals/gsm8k/test_gsm8k_correctness.py:L176-L192 vLLM
        measured_metric = results["accuracy"]
        expected_metric = eval_config["accuracy_threshold"]
        tol = eval_config.get("tolerance", 0.08)
# ...
        assert measured_metric >= expected_metric - tol, (
            f"GSM8K metric too low: {measured_metric:.4f} < "
            f"{expected_metric:.4f} - {tol:.4f} = {expected_metric - tol:.4f}"
        )

The default tolerance is 0.08 — eight percentage points, which at $n = 1319$ is 7.3 standard errors. None of the TurboQuant configs override it. So the effective gate on t3nc is 0.67, not 0.75. This is a regression smoke test designed not to flake, not an accuracy measurement, and treating vLLM's CI as evidence that a scheme is lossless would be a misreading of its own source.

The eval's own shape limits it further. evaluate_gsm8k defaults to max_tokens: int = 256, temperature: float = 0.0, num_shots: int = 5 (tests/evals/gsm8k/gsm8k_eval.py:L210-L223), and scores by extracting the last integer in the completion (tests/evals/gsm8k/gsm8k_eval.py:L68-L77). A 256-token cap cannot observe long-generation drift, greedy decoding cannot observe distribution-tail damage, and last-integer extraction cannot observe a formatting regression. All three are real quantization failure modes, and this harness is blind to them by construction. It is a good gate for what it gates.

§7

Worked trace: two deployments through the procedure

Figure 4 — the operational procedure. Thresholds are the derived boundaries of §4.4.3 for an H100 SXM; $B^{*}_{\text{MoE}} = M^{*}E/k$ from §4.4.5. §4.1's tree selects from the design space; this one starts from a measured workload.

Loading…

Deployment A, the coding assistant. Steady-state concurrency measured at 3. Llama-3-8B bf16 fits with 52 requests of headroom, so the capacity branch is skipped. Not MoE. $B = 3$ is far below 118, so W4A16 group-128. In vLLM this is not an online shorthand — there is no INT4 entry in _ONLINE_SHORTHANDS — so it means loading a GPTQ, AWQ or compressed-tensors checkpoint. get_quantization_config maps the checkpoint's quant_method: "gptq" through vllm/model_executor/layers/quantization/__init__.py:L150-L152 to AutoGPTQConfig, whose AutoGPTQLinearMethod calls choose_mp_linear_kernel and lands on Marlin (§4.3 owns that walk). Predicted step time 1.15 ms against bf16's 4.48. The accuracy gate is a paired run on the team's own code-review prompts, not GSM8K.

Deployment B, the summarisation batch. Steady-state concurrency measured at 480. Llama-3-8B bf16 fits. Not MoE. 480 is above 266, so Regime III: W8A8-FP8. On an H100 the fp8 tensor cores exist, so the weights go through a native fp8 GEMM rather than through Marlin — MarlinFP8ScaledMMLinearKernel at vllm/model_executor/kernels/linear/scaled_mm/marlin.py:L29-L45 announces itself as the "FP8 Marlin kernel for GPUs that lack FP8 hardware support", so it is the fallback, not the path. Predicted step time 3.64 ms against bf16's 7.28, and the W4A16 checkpoint that made Deployment A fast would have cost 7.73 ms here.

That the engines think in these regimes is visible in the dispatch itself. The FlashInfer blockscale FP8 path branches on $M$ and registers both branches into the compiled graph so the choice can be made per call:

vllm/model_executor/kernels/linear/scaled_mm/flashinfer.py:L304-L316 vLLM
    condition = input.shape[0] < 32

    # PyTorch's torch.compile cannot handle input-dependent control flow in standard
    # Python conditionals. torch.cond() explicitly registers both code paths in the
    # computation graph, allowing torch.compile to capture both branches.
    # without torch.cond, the M < 32 condition won't be able to be captured by torch
    # compile
    return torch.cond(
        condition,
        run_flashinfer_deepgemm_swapAB,
        run_deepgemm,
        (input, weight, weight_scale),
    )
§8

Pitfalls and war stories

Benchmarking the wrong batch size. vllm bench latency defaults to --batch-size 8, --input-len 32, --output-len 128 (vllm/benchmarks/latency.py:L34-L52). Batch 8 is deep in Regime I, where every weight-only scheme looks excellent. If your service runs at concurrency 300, that benchmark is measuring a workload you do not have, and it will tell you to ship W4A16 into the one regime where it loses. Sweep the batch axis or do not run it.

Benchmarking the wrong phase. latency.py times llm.generate(...) end to end and reports percentiles over whole-batch completions (vllm/benchmarks/latency.py:L125-L167). Prefill and decode are on opposite sides of the roofline, so a single blended number is exactly the statistic that cannot answer this chapter's question. SGLang's harness separates them by construction:

python/sglang/benchmark/one_batch.py:L784-L807 SGLang
    model_runner.synchronize()
    tic = time.perf_counter()
    next_token_ids, _, batch = model_runner.extend(reqs)
    model_runner.synchronize()
    prefill_latency = time.perf_counter() - tic
# ...
    tot_latency += prefill_latency
    throughput = input_len * batch_size / prefill_latency
    rank_print(
        f"Prefill. latency: {prefill_latency:6.5f} s, throughput: {throughput:9.2f} token/s"
    )
    measurement_results["prefill_latency"] = prefill_latency
    measurement_results["prefill_throughput"] = throughput

and reports median_decode_latency separately at python/sglang/benchmark/one_batch.py:L855-L865, with --batch-size accepting a list (python/sglang/benchmark/one_batch.py:L215-L226). It drives ModelRunner directly — no HTTP, no scheduler — which makes it the closest thing either repository has to a clean measurement of Figure 1.

A stale error message that sends you in a circle. §4.3 found this and it is worth repeating here because it bites during exactly this kind of experiment. When a tensor-parallel split makes a shard awkward, Marlin refuses:

vllm/model_executor/layers/quantization/utils/marlin_utils.py:L178-L186 vLLM
    # Validate output_size_per_partition
    if output_size_per_partition % GPTQ_MARLIN_MIN_THREAD_N != 0:
        raise ValueError(
            f"Weight output_size_per_partition = "
            f"{output_size_per_partition} is not divisible by "
            f" min_thread_n = {GPTQ_MARLIN_MIN_THREAD_N}. "
            "Consider reducing tensor_parallel_size or running "
            "with --quantization gptq."
        )

The first suggestion works. The second does not, at a556f3f: vllm/model_executor/layers/quantization/__init__.py:L150-L152 maps "auto_gptq", "gptq" and "gptq_marlin" to the same AutoGPTQConfig, whose docstring is "Config class for AutoGPTQ quantization using Marlin kernels", and override_quantization_method (vllm/model_executor/layers/quantization/auto_gptq.py:L219-L238) accepts all four aliases including "marlin" and returns "auto_gptq" for every one. There is no distinct config class selected merely by switching those aliases. That does not eliminate non-Marlin kernels: the subsequent mixed-precision selector can choose Machete, ExLlama or Triton when supported. Inspect the selected kernel and refusal reasons rather than inferring dispatch from the config name. The message is a fossil from when gptq and gptq_marlin were separate classes — which they still are in SGLang (python/sglang/srt/layers/quantization/__init__.py:L86-L88 registers GPTQConfig and GPTQMarlinConfig separately), so the advice is correct there and stale here.

Note also that this error is now rarer than it looks: marlin_padded_nk (vllm/model_executor/layers/quantization/utils/marlin_utils.py:L221-L242) zero-pads awkward $(n, k)$ up to the cheaper thread-tile family and warns rather than refusing — "Marlin requires thread-tile padding for some weight shapes in this model. Activations and/or outputs of the padded layers are padded/sliced on every forward; performance may be degraded." A warning you scrolled past is a performance regression you will later blame on the quantization scheme.

Comparing against the wrong baseline. The only comparison that answers this chapter's question is quantized-versus-bf16 on the same engine, same SHA, same flags, same batch. Against a number from a blog post you are confounding the scheme with the kernel, the scheduler and the hardware. Even arXiv:2411.02355 is a prior, not a substitute.

Running the deprecated shim. benchmarks/benchmark_latency.py at the vLLM repo root prints "DEPRECATED: This script has been moved to the vLLM CLI" and exits 1 (benchmarks/benchmark_latency.py:L1-L17). The real code is vllm/benchmarks/latency.py, reached as vllm bench latency. SGLang has the mirror-image shim: python/sglang/bench_one_batch.py:L1-L19 re-exports sglang.benchmark.one_batch with a FutureWarning.

Download the CPU quantization and parallelism reference checks. Run python quant_parallel_checks.py: byte ledgers, GPTQ correction, scaling, nibble layouts, paired accuracy, TP projection algebra, PP schedules, routing and semantic-order checks. These test mathematical contracts, not CUDA kernels or engine performance.

§9

Hands-on

You need a GPU for any of this; I did not have one. Lab 06 is the structured version. The minimum honest experiment is three commands per scheme.

Sweep the batch axis and find your own crossover shell
# 1. Decode step time versus batch, prefill excluded. SGLang, direct ModelRunner.
python -m sglang.benchmark.one_batch \
  --model-path meta-llama/Meta-Llama-3-8B-Instruct \
  --batch-size 1 8 32 64 128 256 512 --input-len 512 --output-len 32
# read "Decode.  median latency:" for each batch size; that is Figure 1's y-axis.

# 2. Repeat with a W4A16 checkpoint and with online FP8.
python -m sglang.benchmark.one_batch --model-path <gptq-int4-model> --batch-size 1 8 32 64 128 256 512 --input-len 512 --output-len 32
python -m sglang.benchmark.one_batch --model-path meta-llama/Meta-Llama-3-8B-Instruct --quantization fp8 --batch-size 1 8 32 64 128 256 512 --input-len 512 --output-len 32

# 3. The capacity axis, which step time cannot show you. Read the KV-blocks line at startup.
vllm serve meta-llama/Meta-Llama-3-8B-Instruct 2>&1 | grep -i "GPU KV cache size"

# 4. Closed-loop serving at a fixed concurrency, both schemes.
vllm bench serve --model <model> --dataset-name random \
  --random-input-len 1024 --random-output-len 256 \
  --max-concurrency 256 --num-prompts 2000

# 5. The accuracy gate, run against both servers on the same questions.
python tests/evals/gsm8k/gsm8k_eval.py --port 8000 --num-questions 1319

Step 4's --max-concurrency is a client-side semaphore (vllm/benchmarks/serve.py:L1588-L1599): it caps in-flight requests independently of --request-rate. It does not set the exact scheduler decode batch: in-flight requests include queued, prefill and finishing work. Record actual iteration shapes and decode-batch distributions. Use a kernel benchmark for a controlled GEMM $M$, and a serving benchmark for latency and goodput.

Step 5 gives you a score, not a comparison. For the flip rate, keep the per-question preds array from both runs — gsm8k_eval.py already builds it before reducing to np.mean(np.array(preds) == np.array(labels)) at tests/evals/gsm8k/gsm8k_eval.py:L189-L200 — and count the positions where the two runs disagree.

§10

Exercises

  1. Read and answer. Open vllm/config/quantization.py and list every entry in _ONLINE_SHORTHANDS that sets moe but not linear. For each, say which regime of Figure 1 the dense linear layers must be in for that choice to be the right one.
  2. Derive. Recompute the three regime boundaries for Llama-3-70B at TP=8. Per GPU, $W = 8.75$ B streamed parameters. Does the ordering of the schemes change? Does $B^{*}$ change? Explain in one sentence why the answer to the second question is what it is.
  3. Derive. DeepSeek-V3 has $E = 256$ routed experts and $k = 8$. At what serving batch do its expert GEMMs leave Regime I for W4A16? Now suppose you deploy with expert parallelism across 16 GPUs, so each GPU holds 16 experts. Does $M_e$ per expert change? Does the answer to the first part change?
  4. Predict, then verify. Predict the sign and rough magnitude of the W4A16-versus-bf16 step-time difference at batch 512 for Llama-3-8B on an H100. Then run exercise 1's command list on real hardware and check. If your measured penalty is much larger than 0.45 ms, name two mechanisms in §4.3 that could explain it.
  5. Statistics. A colleague reports that the INT4 build scores 0.79 on GSM8K against the bf16 build's 0.81 and concludes the quantization cost two points. Both runs used the full 1,319 questions. What is the 95% confidence interval on that difference under an unpaired analysis? What extra number would you ask them for, and how would it change the conclusion?
Answers

1. int8_per_channel_weight_only and nvfp4_per_token (vllm/config/quantization.py:L138-L147); both carry only a moe field, and the comments say so explicitly. Leaving dense linear unquantized is right when those layers are in Regime III — compute-bound at the serving batch — where weight-only quantization is a net loss, while the experts remain in Regime I because routing divides their effective $M$ by $E/k$.

2. Memory terms scale down by 8 and so does the mma term, because both are proportional to $W$. The boundaries are ratios of the two, so every boundary is unchanged: 47, 118, 266. $B^{*} = I^{*}b_w/2$ contains no model term at all. Tensor parallelism moves the absolute latencies, not the crossovers — which is exactly why the crossovers are worth memorising and the latencies are not. (Communication cost, which Part 5 owns, is the term this model omits.)

3. $B^{*}_{\text{MoE}} = 76 \times 256/8 = 2{,}432$. Under expert parallelism each GPU holds 16 of the 256 experts, but the routing fraction per expert is unchanged — a token still picks 8 of 256 — so $M_e = Bk/E = B/32$ regardless of placement. The answer to the first part is unchanged. What EP changes is the all-to-all traffic and the load-balance variance, not the arithmetic intensity of an individual expert GEMM.

4. The illustrative additive model predicts a positive sign, roughly 0.45 ms; this is not a measured or universal penalty: W4A16 issues the same bf16 mma instructions plus dequantisation. A much larger measured penalty would point at (a) Marlin thread-tile zero-padding, which marlin_padded_nk warns about once at load and then pays on every forward, or (b) falling off Marlin onto a slower mixed-precision kernel entirely — check the "Using ... for AutoGPTQLinearMethod" log line to see which kernel was chosen.

5. Standard error per score $\approx 1.1$ pp; unpaired difference standard error $\approx 1.56$ pp; 95% CI on a 2 pp difference is roughly $\pm 3.1$ pp, i.e. $[-1.1, +5.1]$ — consistent with zero. Ask for the flip rate on the same 1,319 questions. With paired binary correctness, difference $\Delta=.02$ and discordance $f=.03$ give $SE\approx\sqrt{(f-\Delta^2)/1319}=.00474$, so the approximate 95% interval is $[1.07,2.93]$ percentage points, not noise. Small discordance can make a paired difference more precise. Report directional discordant counts and use a paired bootstrap or McNemar test; answer-string flips are not correctness flips. Larger discordance can hide offsetting changes, a distinction discussed in arXiv:2407.09141's finding.

§11

Key takeaways

  • Weight-only quantization divides one term of $\max(\text{bytes}/\beta,\ \text{FLOP}/P)$ and adds a constant to the other. W8A8 divides both. That structural difference, not any accuracy argument, is why W4A16 has a crossover and W8A8-FP8 does not.
  • Above roughly $B = 266$ on an H100, W4A16 issues the same bf16 mma instructions as bf16 plus about 0.45 ms of dequantisation per Llama-3-8B decode step. It is not "diminishing returns"; it is a permanent tax that grows no smaller as batch grows.
  • Capacity and latency are orthogonal. Freeing 49 GiB of weights on a 70B at TP=2 takes the concurrency ceiling from 1 request to 40 — a 40× throughput effect that a batch-1 latency benchmark cannot see, and that no kernel improvement can substitute for.
  • MoE routing divides the effective GEMM height by $E/k$, so expert weights stay memory-bound to batch sizes 16–32× beyond where dense layers cross over — while being 95–97.5% of the parameters. Both engines ship MoE-only quantization specs for exactly this reason.
  • GSM8K at 1,319 questions resolves about 3 percentage points between independent runs, and vLLM's CI gate adds another 8 points of tolerance on top. Any claim that a scheme is "lossless" based on that pipeline is a claim about CI stability, not about accuracy.
  • Report score delta, flip rate and KL together, paired on identical prompts. Equal scores with a high flip rate is the failure mode the literature was written to warn about, and it is invisible to every aggregate metric.
§12

Further reading

  • arXiv:2411.02355"Give Me BF16 or Give Me Death"? Accuracy-Performance Trade-Offs in LLM Quantization, Kurtic, Marques, Pandit, Kurtz, Alistarh. Over 500,000 evaluations across the Llama-3.1 family; the largest published sweep.
  • arXiv:2407.09141Accuracy is Not All You Need. The flips-and-KL argument, and the methodological source for §4.4.6.
  • arXiv:2409.11055Exploring the Trade-Offs: Quantization Methods, Task Difficulty, and Model Size. 1B to 405B, four methods, 13 datasets.
  • arXiv:2504.04823Quantization Hurts Reasoning? Also the source for the negative result on output length.
  • docs/features/quantization/online.md in the vLLM tree — the online-quantization schema referenced from QuantizationConfigArgs's docstring; the canonical description of the linear/moe split.
  • tests/evals/gsm8k/README.md — how to run vLLM's accuracy CI locally, including the standalone gsm8k_eval.py path against an already-running server.
  • SGLang's developer-guide page benchmark_and_profiling.mdx — the comparison table of bench_serving, bench_one_batch_server, bench_offline_throughput and bench_one_batch, with the note that the last is for "kernel-level latency profiling of a single static batch".
  • Lab 06 — quantization tradeoff. Every number in this chapter is derived; the lab is where they become measured.

The Inference Engineering Course · code read at vllm@a556f3fccb and sglang@7d893255c3, pinned 2026-08-21. See SOURCES for the full provenance record.

Explore the library

Reading preferences

Appearance
18 px