04 — Quantization: weight-only-4bit vs W4A4 are different regimes

Scope. This doc explains the regimes — why weight-only 4-bit and W4A4 are different things. For the per-hardware decision (which format on which part, through which kernel and engine) read 23-quantization-selection.md, which is newer and written against a larger corpus.

The single most consequential distinction in this repo is that "4-bit" describes two regimes that share no arithmetic, no hardware path, no accuracy profile and no failure modes. Conflating them is, in the words of W4A4 (4-bit weights and 4-bit activations, general term), "the single most common error in quantization blog posts."

The category record W4A4 (4-bit weights and 4-bit activations, general term) splits "W4A4" into three further things, and the distinctions matter more than the headline:

"(1) Integer W4A4 with INT4 tensor cores — does not exist on shipping NVIDIA or AMD parts; Hopper and older have no INT4 path at all and Ampere-era 'INT4' marketing was sparsity, not 4-bit multiply. (2) FP4 W4A4 — this IS native, on Blackwell, and that is exactly what NVFP4 (Blackwell block-scaled E2M1 + E4M3) is. (3) Simulated/fake W4A4 in research papers — SpinQuant, QuaRot, QuIP# — which measure quality but buy no speed without custom kernels."

And the honest state of the art, in TensorRT-LLM's own naming: "it ships W4A16 and W4A8 variants of AWQ and GPTQ and no W4A4 row at all."

1. Regime A — weight-only 4-bit (W4A16)

1.1 What it actually is

AWQ (activation-aware weight quantization) states the definition crisply: "AWQ is w4a16. Weights are 4-bit; activations and the KV cache stay 16-bit, so there is no hardware 4-bit-activation path and nothing on an INT4 tensor core is being used. Group size 128 is the scheme's defining choice and the file layout - it is not a tunable." And the runtime mechanism: "Runtime support is entirely software dequant-to-fp16 inside a fused GEMM."

GPTQ (one-shot approximate second-order PTQ) adds the crucial framing: "GPTQ's contribution is the rounding criterion (approximate inverse Hessian from calibration), not a new number format - a GPTQ file is INT4/INT3 with per-group scales, the same layout AWQ produces, which is why the two share every serving kernel."

That shared-kernel fact is why group size dominates every "GPTQ is slow" report: GPTQ (one-shot approximate second-order PTQ) notes "Group size 128 is the conventional default; 32/64/1024 are legal and change both quality and kernel eligibility, which is a frequent source of 'GPTQ is slow' reports." There is a gotcha record for exactly this — The ROCm Triton w4a16 path asserts on GPTQ models with a non-zero-group qzeros layout (major): serving GPTQ-Int4 fails on the ROCm Triton w4a16 path with a qzeros.shape == (K // group_size, N // 8) assertion.

1.2 Why it works at all: the two kernels, and the arch split

This is the most decision-relevant piece of the whole regime, because it determines which NVIDIA part you want:

• Marlin weight-only INT4 GEMM: repacking so dequant is free and tensor cores stay saturated — "removes the dequantization tax by moving it offline and restructuring the kernel so dequant and MMA are ordered so neither pipeline stalls." Design target as a ratio: "as long as we perform less than 25-50 tensor-core multiply-accumulates per 4-bit quantized weight, the FLOP:byte ratio of the GPU permits near-ideal 4x speedup, and the theoretical benefit of weight-only quantization should extend to batch sizes 4-8x larger than prior kernels achieved." Hardware floor, verbatim: "requires CUDA >= 11.8 and 'NVIDIA GPU with compute capability >= 8.0 (Ampere or Ada, Marlin is not yet optimized for Hopper)'." So: Ampere/Ada → Marlin is the fast path. Hopper → Marlin is the fallback.
• Machete weight-only 4-bit GEMM: CUTLASS 3.x tiles plus an offline prepacked int4 B operand, Hopper only — the Hopper answer. It "keeps CUTLASS's Hopper WGMMA tensor-core mainloop and changes the B OPERAND LAYOUT instead." And it is emphatic about what this is not: "this is a tensor-core bf16/fp16 GEMM over int4 weights dequantized into the MMA fragments on the fly, NOT a hardware int4 tensor-core path - there is no int4 MMA instruction being used." Floor, hard-coded in vLLM: get_min_capability() returns 90 and can_implement() returns False with "Machete requires compute capability of 90 (Hopper)" plus "Machete only supported on CUDA". So "a Hopper box gets Machete, an Ampere or Ada box falls through to Marlin, a consumer Blackwell (SM120) is not SM90 so also falls through, and ROCm is excluded outright."
• llama.cpp MMQ: INT8 activation quantization with per-arch tile configs, no tensor cores — the third answer, and it is a different machine entirely: MMQ "quantizes the ACTIVATIONS to q8_1 blocks of 32 elements and performs the dot product in INT8 accumulating to INT32, using DP4A where available. So the multiply-accumulates are integer SIMT instructions, not tensor-core MMAs - the FLOP:byte ratio the tensor cores exploit is not available at all." Its arch reach is the widest: "per-arch config headers exist for Ampere, Blackwell, CDNA, GCN, RDNA2/3/4/5 and Pascal (DP4A)."

The record states the honest comparison directly: "on a consumer or datacenter GPU, llama.cpp reaches 4-bit throughput through integer SIMT dot products with a broad arch reach (including RDNA and CDNA), while Marlin and Machete reach it through tensor cores on a narrower arch list. Neither is strictly better - MMQ works on more parts, Machete reaches much higher tensor-core utilization where it applies - and any comparison must state which path was taken."

Three W4A16 arch paths, one per part generation. Any advice that does not name the kernel is incomplete advice.

1.3 Quality, where the records have controlled numbers

The Q4_K_M correction that matters more than any of the numbers above. The record's notes: "IMPORTANT CORRECTION to the common 'Q4_K_M mixes 6-bit and 4-bit sub-blocks' claim: it does not. All 256 weights in block_q4_K are uniform 4-bit. What is mixed is the METADATA: the 8 sub-blocks of 32 elements each have a 6-bit scale and a 6-bit min." The quality advantage over legacy q4_0 therefore "comes from 6-bit per-32 sub-block scales plus per-32 mins on an fp16 super-block scale, not from wider weights."

And "Q4_K_M IS A MIXTURE, not a type" — per src/llama-quant.cpp, base q4_K with attn_v and ffn_down promoted to Q6_K. The plain 'Q4_K' CLI name is an alias for Q4_K_M.

Which also means quant names are not comparable across models. See Two models quantized at 'the same Q4_K_M level' can land at 4.75 and 4.30 bpw, so named-quant comparisons across models are not controlled experiments (minor): "Standard llama.cpp quantization applies hardcoded rules ('use Q4_K_M, except bump some tensors up/down, except fall back if incompatible, except keep some tensors unquantized')" — so "normalize comparisons to a matched bpw budget rather than a matched quant name", which llama.cpp's --target-bpw now solves.

1.4 The lever nobody records: imatrix

GGUF imatrix (importance matrix) - the calibration file that determines WHICH weights get more bits is "not a format - a per-tensor activation-energy file", and its note calls it "THE MAJOR PRACTICAL LEVER NOBODY RECORDS". Mechanism: "a per-tensor table of sum-of-squared activations (Sigma(Act^2))... fed back into llama-quantize as --imatrix file_name to weight the rounding-error objective per weight, so weights with high activation energy get their bits spent more carefully." It applies across formats — "every ggml type, every bit width, legacy and k-quant and i-quant alike." Its own value: "its effect is to REMOVE part of the quantization delta of the format it is applied to," with official per-type ppl strings Q4_K_M +0.1754 and Q5_K_M +0.056.

1.5 Format family roundup with real bits-per-weight numbers

All K-quants share QK_K = 256 per block with 12 bytes of packed sub-block scales (GGUF Q4_K_M (4-bit K-quant, medium)). Measured average bpw on Llama-3.1-8B, from the same record family:

Five K-quant records are contested (gguf-q3-k-m, gguf-q4-k-m, gguf-q4-k-s, gguf-q5-k-m, gguf-q6-k) — because the measured bpw of the *_M mixture types does not match the pure-block struct comments. Both numbers are retained per record.

1.6 Two low-bit families that are not post-training quantization at all

• Ternary / BitNet b1.58 — Ternary / 1.58-bit quantization (general), BitNet b1.58 (ternary LLM). "The framing that matters: ternary LLM quantization is a TRAINING-TIME recipe, not a post-training one. There is no such thing as PTQ-ing an existing fp16 LLM to 1.58 bits and keeping its quality." The paper's "matches fp16" means "a ternary model trained from scratch on the same token budget matches an fp16 model trained from scratch on the same token budget" — a statement about the ternary set's expressiveness, not about quantization loss on a pretrained checkpoint. native_support is empty on both records.
• EXL2 — EXL2 (GPTQ with mixed per-tensor bitrate) (draft, 0.6). "EXL2 uses the same optimization method as GPTQ... but mixes quantization levels within a model - and, per the upstream README, can apply multiple quantization levels to each individual linear layer - to hit any target average bitrate between 2 and 8 bpw." Its quality_delta is empty and the record says why: the README publishes throughput (Llama2-7B EXL2 4.0 bpw at 185 tok/s on a 3090 Ti, 211 on a 4090) but no perplexity or benchmark deltas vs fp16, and none was found elsewhere.

2. Regime B — W4A4, and the only genuine hardware-native 4-bit path

2.1 NVFP4: the one native W4A4

NVFP4 (Blackwell block-scaled E2M1 + E4M3) is the record that "most needs the two distinctions kept separate":

(1) It is genuinely w4a4. "the element format is E2M1 (1 sign, 2 exp, 1 mantissa, values ~-6..6) and BOTH weights and activations are quantized to it, with a Blackwell hardware path - unlike MXFP4 and unlike every w4a16 scheme above."

(2) It is NOT the OCP microscaling format, despite the identical element layout:

• scale granularity: NVFP4 16 values, MX 32 values
• scale format: NVFP4 uses a "power-of-two-free E4M3 FP8 scale"; MX uses an "E8M0 power-of-two scale"
• NVFP4 adds a second per-tensor FP32 level: "s_global = global_amax/(448*6), s_block = (block_amax/6)/s_global"

The vendor's own quantitative reason for the E4M3 scale: "encoding the scale itself in E4M3 rather than E8M0 drops the scale-encoding MSE from 0.72 to 0.08."

Native on Blackwell only. "TensorRT-LLM's hardware matrix marks NVFP4 Y on Blackwell sm120, sm100/103 and Rubin, and no on Hopper/Ada/Ampere." native_support: NVIDIA B200, NVIDIA B300 (Blackwell Ultra), NVIDIA GB200 NVL72 (per-accelerator). emulated_support: NVIDIA H100 SXM, NVIDIA A100 80GB SXM4, AMD Radeon Pro VII.

Quality — with the baseline stated correctly. The record's quality_delta is careful: "FP8 vs NVFP4, seven benchmarks (vendor-reported)" on DeepSeek-R1-0528 — MMLU-Pro 85% vs 84%, GPQA Diamond 81% vs 80%, HLE 15% vs 14%, LiveCodeBench 77% vs 76%, SciCode 40% vs 40%, Math-500 98% vs 98%, AIME 2024 89% vs 91%. NVIDIA states "1% or less accuracy degradation versus FP8." And the record's caveat is the important one: "This is a vendor comparison against an FP8 baseline, not against fp16, and is vendor-claim-grade."

Its "vendor memory/energy claims (~3.5x vs FP16, 25x/50x energy efficiency vs H100) are in the blog and are marketing, not measurements."

2.2 MXFP4: native, but a different format with a different risk profile

MXFP4 (OCP microscaling FP4, E8M0 scale): "MXFP4 is the OCP microscaling 4-bit format: E2M1 elements under a shared power-of-two E8M0 scale for every 32-value block." Same native_support (Blackwell B200/B300/GB200-NVL72), same emulated_support set.

The vendor-preference statement is the one to carry: "NVIDIA's own comparison table rates MXFP4 as 'accelerated hardware scaling: yes' but 'risk of noticeable accuracy drop compared to FP8', where NVFP4 is 'lower risk of noticeable accuracy drop'... i.e. the vendor prefers its own format on accuracy grounds."

Its quality_delta is empty, and the record says why: "No inference fp16 delta recorded from a controlled benchmark. The closest primary evidence is a PTQ study of SmoothQuant+AWQ+ GPTQ applied to MX formats, which reports 4-bit weights with 8-bit activations in MXINT with 'negligible accuracy loss compared to the uncompressed baseline'." Do not quote an MXFP4 accuracy number from this record.

Note also that MXFP4 is not MX-INT4: on Blackwell, "blockwise FP8" effectively means MXFP8 (OCP microscaling FP8, E4M3 + E8M0) (E4M3 act/weight with UE8M0 scales on sm100/103/107), whereas the SM90 FP8 recipe is E4M3 with FP32 scales.

2.3 The rest of the block-scaled family

Block-scaled FP8/FP4 tensor cores: MXFP and NVFP formats inside the MMA instruction gives the mechanism and the exact PTX shapes: warp-level mma has block-scaled forms where "the multiplicand type is .e4m3/.e5m2/.e3m2/ .e2m3/.e2m1 times a .ue8m0 scale at .m16n8k32, and .e2m1 times (.ue8m0 or .ue4m3) at .m16n8k64", and wgmma mirrors it at warpgroup level with fp8 at .m64nNk32. The payoff: "the scale is part of the MMA OPERAND, not a separate kernel, so a 32-element scale block is applied inside the tensor-core pipeline and the activation quantization no longer needs a fused dequant pass."

Three Blackwell-only formats carry it in the data: NVFP4 (Blackwell block-scaled E2M1 + E4M3) (grp 16), MXFP4 (OCP microscaling FP4, E8M0 scale) (grp 32), MXFP8 (OCP microscaling FP8, E4M3 + E8M0) (grp 32), plus FP8 block scaling (blockwise, E4M3 with per-block scale) (grp 32). And sparsity composes with sub-byte: 2:4 structured sparsity (mma.sp): halving bytes by dropping exactly half of every four notes "For FP4 on Blackwell, CUTLASS lists 2:4 sparsity support for FP4 as a separate feature - sparsity and sub-byte precision compose, which is the largest single byte reduction available."

The portability tax, which bites exactly when you switch parts: Block-scaled FP8/FP4 tensor cores: MXFP and NVFP formats inside the MMA instruction and DeepGEMM grouped/masked GEMM and Mega MoE: shaping MoE expert GEMMs for decode and for NVLink overlap both record that "SM90 requires scaling factors in FP32 while SM100 requires them packed in UE8M0 with four packed into a single int, and the LHS scale must have a TMA-aligned transposed layout for both. So switching an FP8 model from H100 to B200 changes the memory layout the engine must produce, not just the kernel it calls."

2.4 W4A4 with no silicon: SpinQuant and friends

SpinQuant (learned rotations + W4A4KV4) is the cleanest illustration of the confusion and is worth reading as the cautionary record. It is genuinely W4A4 — "both weights and activations quantized to 4-bit" — and its accuracy claim rests on rotations: W4A4KV4 on LLaMA-2 7B gives an average zero-shot reasoning accuracy of 64.0, a 2.9-point gap from full precision, while LLM-QAT is 22.0 points behind and SmoothQuant 25.0 points behind under identical precision (benchmarks: SpinQuant W4A4KV4 8-task zero-shot accuracy delta vs fp16, LLaMA-2 7B, SpinQuant W4A8KV8 8-task zero-shot accuracy delta vs fp16, Mistral-7B, SpinQuant random-vs-learned rotation accuracy spread, Mistral-7B W4A4KV4).

And then the sentence that decides the hardware question: "there is NO hardware INT4 tensor-core path. Its W4A4 numbers come from research Triton kernels plus fake quantization, not from silicon. On current NVIDIA parts 4-bit x 4-bit is not a tensor-core op; compare nvfp4, which is W4A4 and native, on Blackwell."

The engine record for the harness says the same from the serving side (SpinQuant research harness (fake-quant + Triton INT4 kernels)): "NO hardware INT4 tensor-core path: W4A4KV4 accuracy is a fake-quant / research-kernel result, not a silicon result", plus "SpinQuant_had needs an online Hadamard kernel, ~8% of network latency overhead per the paper."

3. Native vs emulated: the per-hardware map

native_support / emulated_support are populated on the records. The counts below were computed by tallying those two fields across all format records at the snapshot — they are a summary of the records, not figures quoted from any single one. Where a cell is a bare count, cite the accelerator record and re-tally if the data has moved:

Read this table as the answer to "why did my format not speed anything up?"

3.1 What "native" means per format, and the field most likely to mislead you

GGUF has ZERO native support on every accelerator in the repo — and the record says so directly. GGUF IQ4_XS (non-linear 4-bit, 256-value super-block with 6-bit sub-block scales): "Weight-only by construction... so native_support is deliberately EMPTY and the whole accelerator ecosystem is emulated_support. No accelerator has an i-quant or k-quant path in silicon. This is the single most commonly botched field in this repo." GGUF TQ1_0 / TQ2_0 (ternarized quants, QK_K super-block, trit-packed) repeats it: "No accelerator has a GGUF quant path in silicon."

This is not a criticism of GGUF — GGUF's win is portability and density, not silicon support. But if you are choosing a format to buy hardware for, note the consequence: no GGUF quant will ever be faster because the GPU supports it. It is faster because there are fewer bytes, and llama.cpp's MMQ path reaches them without tensor cores (llama.cpp MMQ: INT8 activation quantization with per-arch tile configs, no tensor cores).

AWQ, GPTQ, EXL2 and NF4 all have empty native_support. Every one is AWQ (activation-aware weight quantization)'s "software dequant-to-fp16 inside a fused GEMM". The speed comes from kernel engineering (Marlin, Machete, MMQ), never from silicon.

FP8 is the genuinely native 8-bit regime. FP8 E4M3 (per-tensor / per-block scaled): "FP8 is genuinely w8a8 with fp32 accumulate, so the low-precision operands actually reach the tensor core. Hardware-native from Hopper on NVIDIA... and from CDNA3 on AMD — MI300X lists 2614.9 TFLOPs dense fp8, again 2x fp16." And the boundary: "NOT native on Ampere (the A100 accelerator record has no fp8 entry) so FP8 on A100 is emulation by upcast." That matches the emulated lists exactly: A100 and A10 carry fp8-e4m3 in emulated_support; L4, L40S, H100, H200-class and Blackwell carry it natively.

3.2 Per-hardware guidance

NVIDIA Hopper (NVIDIA H100 SXM, NVIDIA H100 PCIe, NVIDIA H100 NVL) — the richest native menu in the repo (13 / 3 / 5 formats) but none of the 4-bit-block-scaled formats. Concretely: FP8-E4M3 native (H100 lists 1979 TFLOPS dense fp8, 2x its fp16); INT8-W8A8 native; W4A16 via Machete (sm90, CUDA only). Blackwell-only formats NVFP4, MXFP4, MXFP8, FP8-blockwise and W4A4 are all emulated_support here. Note vLLM's caveat that FP8 "needs Ada or newer" and FP8 E4M3 (per-tensor / per-block scaled)'s "not native on Ampere".

NVIDIA Blackwell datacenter (NVIDIA B200, NVIDIA B300 (Blackwell Ultra)) — the only place W4A4 is real. All eight native formats are the block-scaled family (nvfp4, mxfp4, mxfp8, fp8-blockwise, fp8-e4m3, fp8-e5m2, modelopt, wq4a4). Critically, Machete does not apply (not sm90), so W4A16 on Blackwell is Marlin or a w4a16 misroute — which is a gotcha, not a preference (§4).

Consumer Blackwell (NVIDIA GeForce RTX 5090, NVIDIA RTX PRO 6000 Blackwell Workstation Edition) — the trap. NVFP4 (Blackwell block-scaled E2M1 + E4M3) says TensorRT-LLM's matrix marks NVFP4 "Y on Blackwell sm120, sm100/103 and Rubin", i.e. SM120 counts as native for the format. But Consumer Blackwell (SM120) vs datacenter Blackwell (SM100a): one architecture name, two incompatible kernel targets says sm100a-compiled kernels "are not compatible with RTX 50 series GPUs", and FlashInfer: a kernel-selection layer over FlashAttention, cuDNN, CUTLASS and TensorRT-LLM notes FlashInfer lists SM100/SM103 and SM120/SM121 separately and that FP4 GEMM (NVFP4/MXFP4) is "for Blackwell GPUs" without promising SM120. There are no native_support lists for RTX 5090 or RTX PRO 6000 at all — they appear only in emulated lists. So the format claim and the kernel reality do not line up, and the records do not resolve it.

NVIDIA Ada consumer (NVIDIA GeForce RTX 4090) — 6 native formats, all weight-only or int8 paths (autoround-w4a16, bitsandbytes-load-in-4bit, llmcompressor-quant, ort-matmulnbits-int4, ort-quant-int8, ortgenai-int4-fp8). No FP8 in native_support despite the 4090 record listing 330.3 dense FP8 TFLOPS — FP8-E4M3 is not on its native list. W4A16 goes through Marlin (Ampere or Ada, per Marlin weight-only INT4 GEMM: repacking so dequant is free and tensor cores stay saturated).

NVIDIA Ampere (NVIDIA A10, NVIDIA A100 80GB SXM4) — Marlin's home turf (compute capability ≥ 8.0). But FP8 is emulated, and both appear in emulated_support for fp8-e4m3/e5m2. Int8-W8A8 is native on A100, emulated on A10.

NVIDIA L4 (NVIDIA L4) — 8 native formats including fp8-e4m3/e5m2 and int8-w8a8. Its record warns the memory side anyway: "300 GB/s is roughly a tenth of H100 SXM, so it is a fit-many-small-models card, not a fit-one-big-model card."

NVIDIA L40S (NVIDIA L40S) — 12 native formats, the second-richest in the repo, and 48 GB. Its record also carries the clearest dense/sparse warning in the repo (§2.5 of 01-hardware-selection.md).

AMD Instinct (AMD Instinct MI300X, AMD Instinct MI250X) — 8 native formats on MI300X, all int8 or weight-only paths. The vLLM constraint from vLLM is the practical limit: "Marlin (GPTQ/AWQ/FP8/FP4) has no AMD, Intel or CPU path; AWQ and GPTQ have no AMD GPU path." So MI300X's native menu exists but the two most popular W4A16 formats are unavailable there. MI350X/MI355X are absent from every native_support and emulated_support list in the repo — a genuine gap, not an emulated finding (§5).

Intel Gaudi (Intel Gaudi 2, Intel Gaudi 3 (HL-325 OAM / HL-338 PCIe)) — Intel Neural Compressor (multi-technique PTQ/QAT toolkit; low-bit LLM path is AutoRound) is native on both and nothing else is, which inverts the pattern: "advanced low-bit LLM/VLM quantization (2-4 bit, MXFP4, NVFP4, FP8 KV cache) is not implemented in Neural Compressor — it is integrated FROM AutoRound". Its note warns: "Treating the two as one tool would misattribute the group_size and asymmetric-quantization restrictions." Gaudi3 has just 2 native and 2 emulated; its emulated list is MoE mixed-precision expert bit allocation (expert-level bit-width, not one fixed precision) and MoE routing-instability under quantization (VSRAQ / EAQuant routing-consistency alignment).

Intel Arc (Intel Arc B580, Intel Arc Pro B60) — 2 native (autoround-w4a16, ort-matmulnbits-int4) and 18/13 emulated. Combined with llama.cpp's Intel caveats and On Intel Arc dGPUs, llama.cpp agentic workloads stall and enabling FA can cut prompt processing ~15x, Arc is the thinnest native-quant story in the repo. Its record also notes the capacity wall: 12 GB at 456 GB/s, and PCIe x8 not x16.

Apple Silicon (Apple M3 Max (MacBook Pro 16-inch, Nov 2023), Apple M3 Ultra (Mac Studio, 2025)) — now native, and this entry previously said the opposite; the earlier text was false. MLX affine quantized linear (mlx-lm's default 4-bit weight format, group_size 64) lists both apple-m3-max and apple-m3-ultra in native_support (alongside M4 Max, M4 Pro, M5 Max, M5 Ultra and M5 Pro), with group 64 and emulated_support empty. MLC LLM compiler quantization (q4f16 group-quant, plus q3f16, AWQ, FP8 per-tensor and FP8 block-scale) remains in emulated_support. The selection step and its caveats are in 23-quantization-selection.md §0 and §9.9 — read that row rather than restating it here.

What still does not resolve on Apple, and it is the reason to read 23 before choosing. MLX's own formats (MLX-LM: mlx-quantized-4bit, mlx-quantized-8bit) are Metal-native by construction but are not in the quantization schema's enum, so those two engine labels still have no format record of their own; MLX affine quantized linear (mlx-lm's default 4-bit weight format, group_size 64) is the closest thing and is keyed on the affine mode at group 64. Three benchmark records still point format_id at mlx-quantized-4bit, which is not a record id — a dangling reference the format record's own notes name. And MLX ships four modes at four different group sizes (affine 64, mxfp4 32, nvfp4 16, mxfp8 32), so a benchmark recorded as "4-bit on Apple" may be any of a 64-group affine grid or a 32-group MX grid, and those differ in both memory arithmetic and quality. "Native on Apple" is now true and is still not sufficient to compare two Apple checkpoints. §5 records the schema gap.

The parts with no quantization story whatsoever — AMD Radeon Pro VII (1 native, 21 emulated), Cerebras WSE-3 (CS-3 system), Groq LPU (Language Processing Unit, Tensor Streaming Processor), Google Cloud TPU v4 / Google Cloud TPU v5e / Google Cloud TPU v5p / Google Cloud TPU v6e (Trillium), Tenstorrent Wormhole Tensix (n150 / n300), Tenstorrent Blackhole Tensix (p100a / p150a / p150b), AWS Inferentia2, AWS Trainium2, and every asia-market ASIC record — appear in no native or emulated list.

4. The gotchas, grouped by which regime they poison

Silent-corruption-first ordering; the full list is in 05-known-traps.md.

Silently wrong output, weight-only path:

• NVFP4 models with --dtype bfloat16 produce garbled output on any GPU without native FP4 — blocker. NVFP4 + --dtype bfloat16 on SM < 100 (4090 SM89, V100 SM70): loads, starts, serves, every response garbled.
• Setting VLLM_MARLIN_INPUT_DTYPE=fp8 enables Marlin W4A8-FP8 on sm_121a where it silently produces garbage — blocker. sm_121a + WNA16 INT4 MoE + VLLM_MARLIN_INPUT_DTYPE=fp8: coherent-looking repetition loop, and "the corrupted path is also ~2.5% FASTER end-to-end, so throughput benchmarking selects for the bug."
• FlashInfer's b12x MoE backend silently defaults quant_mode to W4A4, so a weight-only NVFP4 checkpoint is computed at the wrong activation precision — major. A checkpoint marked W4A16_NVFP4 served with --moe-backend flashinfer_b12x runs the FP4xFP4 micro-kernel. Silent wrong output, and the record says explicitly: "Do not accept a b12x trace showing an FP4 kernel for a W4A16 checkpoint."
• NVFP4 fused-MoE loading detects mismatched gate/up global scales, warns once, and then dequantizes the up half with the gate half's scale anyway — major. The fused-MoE weight processor "keeps only the gate per-tensor global scale. When a checkpoint ships different scale_2 values..." The check is cheap and needs no GPU: grep the startup log.
• Converting an NVFP4 checkpoint to GGUF silently drops most per-tensor scales, and the model then emits one repeated token for every prompt — major. NVFP4 → GGUF conversion "produces a GGUF that does not load".
• vLLM's 'Your GPU does not have native support for FP4' warning fires on weight-only NVFP4 checkpoints, and gets quoted as evidence the silicon lacks FP4 kernels — minor, but it is a measurement gotcha: "'Your GPU does not have native support for FP4 computation'..." is emitted for a weight-only checkpoint. "Do not read this warning as a hardware capability statement." Establish FP4 support with torch.ops._C.cutlass_scaled_mm_supports_fp4(<capability>).
• compressed-tensors MXFP4 weight-only checkpoints are dispatched to a W4A4 kernel on SM100+, costing ~7.6 average points across AIME/GSM8K/MATH-500 — major, and the one with the clearest measured cost. On B200, a compressed-tensors MXFP4 checkpoint with input_activations: null (i.e. W4A16) logs "Using..." and routes to a W4A4 kernel. Two benchmark records capture the difference on the identical model: Qwen3-8B MXFP4 W4A16 served on MarlinMxfp4Kernel (correct W4A16 path): AIME-24/25 + GSM8K + MATH-500 average accuracy 78.16 versus Qwen3-8B MXFP4 W4A16 served on the FlashInfer W4A4 kernel: AIME-24/25 + GSM8K + MATH-500 average accuracy 70.58, against the unquantized Qwen3-8B BF16 reference: AIME-24/25 + GSM8K + MATH-500 average accuracy (baseline for the MXFP4 W4A16 dispatch comparison) at 82.36. That is a 7.58-point accuracy drop caused purely by kernel misrouting — the best single illustration in the repo of why the W4A16/W4A4 distinction is not pedantry.
• A llama.cpp ROCm commit raised WikiText-2 perplexity from 7.72 to 3024 on a Q4_0 model with no error, and only llama-per perplexity caught it — major, silent. ROCm/gfx1151, llama-perplexity PPL jumps two orders of magnitude.
• On gfx1151 the llama.cpp HIP backend produces corrupt output where Vulkan is correct with byte-identical flags, so cross-backend comparison is the only clean detector — major. On gfx1151, "the HIP backend is wrong and Vulkan is right", same machine, same build.
• llama.cpp HIP on gfx1151 returns wrong-but-plausible logits for any prompt longer than n_ubatch, and the default n_ubatch makes it almost universal — major, silent. "Any prompt longer than n_ubatch produces badly wrong logits on HIP/ROCm on gfx1151 - no error, no warning, no crash, and the numbers look plausible. Default n_ubatch is 512."

Silent-corruption consequence of combining quantization with other features:

• Enabling MTP speculative decoding together with prefix caching silently costs ~20% accuracy — blocker. MTP speculative decoding plus prefix caching on a finetuned MoE: ~20% internal classification accuracy drop.
• Draft-model speculative decoding is not lossless at temperature 0 when the target model is quantized — major. Greedy sampling with spec decode produces different text on a quantized target. n-gram speculation is verified lossless.
• Qwen3.6-35B-A3B NVFP4 W4A16 MoE on the Marlin backend: GSM8K accuracy (reference W4A16 path) — the same checkpoint measured through both paths: Marlin 0.9136 vs FlashInfer b12x 0.9037 GSM8K. Both verified.

Build/toolchain:

• vLLM's compiled extension fails to import against a source-built PyTorch with a different CXX11 ABI setting — blocker, affects any custom-quant build.

5. What the records cannot answer

1. MI350X / MI355X have no quantization-support entry anywhere — no native list, no emulated list, despite having MLPerf records using WMXFP4 weights with FP8 KV (MI355X, MLPerf Inference v6.0 Offline: Llama 2 70B WMXFP4 on 8xMI355X, aggregate output throughput). This is the single most consequential gap in this document: the newest AMD datacenter part has no recorded quantization compatibility.
2. No native quantization record for any Apple format. MLX 4-bit/8-bit are Metal-native by construction but have no quantization record (the scheme enum has no MLX member). Any Apple quantization advice is currently inferred from the engine record, not from a format record.
3. SM120 vs SM100 for NVFP4 is unresolved. The format record says SM120 is native; the kernel records say sm100a kernels do not run on SM120. No record reconciles them, and there are no native_support entries for RTX 5090 or RTX PRO 6000 at all.
4. No quality delta for MXFP4 from a controlled benchmark — the record says so explicitly, and the only evidence is a PTQ paper with 8-bit activations, not W4A4.
5. Format quality deltas are overwhelmingly Llama-3.1-8B or Llama-2-7B. No format has a controlled delta on a modern MoE, which is exactly where quantization behaviour is hardest.
6. No cost-per-quality curve. Records give isolated perplexity deltas at isolated bit widths with no shared baseline methodology, so there is no defensible "best 4-bit format" ranking — the comparisons in §1.3 have to be quoted with their model names attached.
7. The quality-vs-FLOPs claim for W4A4 on Blackwell is vendor-only. The 3.5x memory and 25x/50x energy-efficiency figures in NVFP4 (Blackwell block-scaled E2M1 + E4M3) are labelled "marketing, not measurements" by the record itself.
8. EXL2 has no quality delta at all, and the format is draft at 0.6 with an archived upstream.

Related documents

• 23-quantization-selection.md — the per-hardware decision table built on these regimes.
• 01-hardware-selection.md — which part to buy for each precision regime.
• 02-flop-map.md §3.7 — why quantization speedups are batch-conditional.
• 03-engine-selection.md — which engine supports which format on which backend.
• 05-known-traps.md — every silent-corruption gotcha, severity-ordered.