When quantization pays — and when it does not
benchmarks/kernels/python/sglang/benchmark/one_batch.py
a556f3f · sglang 7d89325You 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.
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.
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.
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.
First principles: two terms and three regimes
Model one decode step as a memory pipe and a math pipe that overlap:
$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$.
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:
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
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.
| Regime | Batch | Binding constraint | bf16 | W4A16 | W8A8 FP8 | Winner |
|---|---|---|---|---|---|---|
| I | B < 118 | weight bandwidth | 4.48 ms | 1.15–2.24 ms | 2.24 ms | W4A16 (up to 3.9×) |
| II | 118 – 266 | mixed | 4.48 ms | 2.24–4.48 ms | 2.24 ms | W8A8 (2.0×) |
| III | B > 266 | tensor-core issue | 15.16B µs | 15.16B µs + 0.45 ms | 7.58B µs | W8A8 (2.0×); W4A16 is a loss |
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.
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.
| Deployment | Weight format | weights GiB/GPU | KV pool GiB | resident tokens | 8k requests |
|---|---|---|---|---|---|
| Llama-3-8B, TP=1 | bf16 | 14.96 | 52.32 | 428,613 | 52 |
| Llama-3-8B, TP=1 | fp8 W8 | 7.48 | 59.80 | 489,877 | 59 |
| Llama-3-8B, TP=1 | INT4 g128 | 3.86 | 63.42 | 519,552 | 63 |
| Llama-3-70B, TP=2 | bf16 | 65.70 | 1.57 | 10,310 | 1 |
| Llama-3-70B, TP=2 | fp8 W8 | 32.85 | 34.43 | 225,612 | 27 |
| Llama-3-70B, TP=2 | INT4 g128 | 16.94 | 50.34 | 329,898 | 40 |
| Llama-3-70B, TP=8 | bf16 | 16.43 | 50.85 | 1,333,049 | 162 |
| Llama-3-70B, TP=8 | INT4 g128 | 4.23 | 63.04 | 1,652,638 | 201 |
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.
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
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.
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:
),
# 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:
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.
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:
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:
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.
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.
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:
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),
)
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:
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:
# 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.
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.
# 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.
Exercises
- Read and answer. Open
vllm/config/quantization.pyand list every entry in_ONLINE_SHORTHANDSthat setsmoebut notlinear. For each, say which regime of Figure 1 the dense linear layers must be in for that choice to be the right one. - 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.
- 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?
- 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.
- 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.
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
mmainstructions 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.
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.09141 — Accuracy is Not All You Need. The flips-and-KL argument, and the methodological source for §4.4.6.
- arXiv:2409.11055 — Exploring the Trade-Offs: Quantization Methods, Task Difficulty, and Model Size. 1B to 405B, four methods, 13 datasets.
- arXiv:2504.04823 — Quantization Hurts Reasoning? Also the source for the negative result on output length.
docs/features/quantization/online.mdin the vLLM tree — the online-quantization schema referenced fromQuantizationConfigArgs's docstring; the canonical description of thelinear/moesplit.tests/evals/gsm8k/README.md— how to run vLLM's accuracy CI locally, including the standalonegsm8k_eval.pypath against an already-running server.- SGLang's developer-guide page
benchmark_and_profiling.mdx— the comparison table ofbench_serving,bench_one_batch_server,bench_offline_throughputandbench_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.