Quantization Formats & Pipelines - imp

Sources of truth: docs/QUANTIZATION.md (formats, choosing a quant), docs/internals/QUANT_PIPELINE.md (files, GEMM-dispatch registry, boundary rules). This skill carries the gotchas only.

GGUF is legacy (docs/DESIGN_DECISIONS.md, 2026-10-30): fix load blockers and wrong output, no new GGUF kernels, routes or speedups. Perf and quality work goes to SafeTensors NVFP4 / FP8.

The two worlds

• GGUF -> NVFP4 decode cache at init; nvfp4_beneficial weights skip the FP16 cache; GDN/SSM projections excluded (quality lock) and served by the FP8 GDN projection sidecar (always on, native F16 only); GGUF Q8_0 projections decode in place via gemm_q8_0_dp4a_dense at every M (#2688: the FP8 copy cost 5280 MiB = 66 recurrent slots on Qwen3.8-27B Q4_K_M).
• gpt-oss native MXFP4 experts convert to NVFP4 at load.
• wcache_.nvfp4 is NOT a GGUF predicate (native models fill it too); GGUF = registry scan on dequant_gpu_supported(source) (#1792 trap).

StorageTier is the dispatch contract (src/core/storage_tier.h)

Undefined (FATAL if dispatched), FP32, FP16, FP8 (E4M3 + per-tensor scale), NVFP4 (two-level micro-scale, decode-GEMV), CUTLASS_NVFP4 (block-scaled, CUTLASS sm_120 fast path), MXFP4 (CUTLASS FMHA).

• ONE decision point since #621: plan_storage() in src/runtime/storage_planner.h; caches are RAII-owned by the executor. Extend the planner's rules, never add downstream overrides.
• NVFP4 != CUTLASS_NVFP4 at large M: plain NVFP4 falls to gemm_nvfp4 (dequant -> cuBLAS). SfAtom scales are built at load (src/exec/pre_dequant_*.cu); convert_scales_sfatom is a load-time artifact.
• M<=32 since #1766: the small-M kernel (#1788, always on) runs mma.sync.kind::mxf4nvf4 on the PLAIN packed layout (src/quant/nvfp4_gemm_smallm_v2.cu, zero extra VRAM; the CUTLASS SF atom static-asserts 128 rows, so no CTA_M<128 tile exists). +16.0% at 32 streams, +36.0% at 8. At M=1 the GEMV family stays (#1789).
• StoragePlanner prices zero-copy native-NVFP4 entries at incremental cost since #1795 (a ~15 GiB phantom demand printed "plan failed" before).

MoE

Per-expert NVFP4 tensors are packed into one contiguous [ne, N, K_packed] buffer per layer/projection at load (cache_moe_native_nvfp4); that is what makes graph capture possible (per-step dequant + cuBLAS is 5-17x slower, no graphs). Grouped-GEMM prefill class is closed at ~60% of the weight floor (sm120-cuda-expert).

KV cache dtypes

kv_cache.dtype = auto | fp16 | fp8 | int8 | int4 | nvfp4 | mxfp4 (CLI --kv-fp8 etc.).

• auto (default) has three upgrade arms in src/runtime/engine_init_resolver.cpp gated by src/model/model_arch.h lists (kv_nvfp4_default_safe, kv_fp8_hint_default_safe, kv_fp8_no_hint_default_safe); an arch missing from them silently gets FP16 KV.
• QWEN35 family (Qwen3.5/3.6/3.8 dense hybrids) resolves to NVFP4 since #1750: Qwen3.8-27B max_model_len 48512 -> 131072 for +0.29-0.35% PPL; FP8 86 848 -> NVFP4 126 432 tokens (+45.6%). MoE siblings keep the old default.
• NVFP4 KV decode was 13.5% slower than FP8 until #1817 (byte loads: 20 LDG.E.U8 per iteration); now 2.3% faster at 77k (74.1 vs 72.4 tok/s). PagedOracle.HD256_Sweep covers the shipped 24q/4kv shape across all 7 dtypes.
• --kv-fp8 / IMP_KV_FP8=1 on that family DOUBLES KV bytes (correct when auto meant FP16); the pin logs its context cost since #1823 (kv_pin_context_cost_factor, kv_dtype_is_explicit_pin in model.cpp).
• Sparse decode attention (attention.sparse_topk_tokens) covers F16/FP8/NVFP4 KV since #1818; the NVFP4 branches had to consume the compacted block table (a dtype gate alone = silent dense no-op). Budget arithmetic in src/exec/sparse_attn_geometry.h uses the resolved block size (32 on n_kv_heads <= 4, #1819).
• kv_cache.swa_snapshot_mb below the snapshot size switches prefix caching OFF (#1092).
• GDN recurrent state: gdn.state_bf16 (default on since #1778): BF16 storage, FP32 register arithmetic, +12.5% at 32 streams, +0.21% PPL, state pool 4848 -> 2544 MiB; resolver keeps FP32 under gdn.ref_kernel or unsupported HD/SS (WARN on stderr). FP16 state REFUTED (subnormal truncation ~6e-5).

Judging quantization quality

• Corpus tools/analysis/ppl_corpus_45k.txt (13 537 tokens, rebuild with tools/analysis/make_ppl_corpus.sh); the 199-token tools/analysis/ppl_corpus.txt inverts verdicts (+42%/+57% vs +25%/+19%; +1.0% vs -0.03% on FP8 SSM).
• --set runtime.deterministic=true both arms (implies runtime.deterministic_gemm; 0.35% run-to-run otherwise); --set speculative.mtp_k=0 (--perplexity otherwise loads the MTP head, +0.79 GiB, floors the 35B KV pool).
• Qwen3.6-35B PPL moves +-0.2..0.5% between fp32-equivalent kernels (routing flips): >1% = broken, below = no verdict. Numerics judge: Qwen3.8-27B-NVFP4-vllm (deterministic, fused GDN 4.6283).
• 35B recipe when model + caches fill 30 GB: corpus in ~1k-token slices (3200 chars), --max-seq-len 1280, both arms per slice, token-weighted NLL; warm-cache mount (/root/.cache/imp/warm) makes runs ~7 s; do NOT mount the library-reserve measurement (4399 MiB whole-init worsens the plan).
• Gemma-4-it scores text only inside its own turn framing (#2044): raw prose reads 19 312 in imp on the GGUF, the same prose framed as <|turn>user\n...<turn|>\n<|turn>model\n<|channel>thought\n<channel|> + text reads 18-27. Judge Gemma-4 with tools/analysis/ppl_corpus_gemma4_turn.txt, never with ppl_corpus_45k.txt or bare prose; the self-consistency oracle (teacher-force the model's own greedy answer, expect top-1 > 90 %) separates a scorer bug from a format prior in one 15 s run.
• PPL runs prefill; --prefill-chunk-size 1 runs every corpus token at M=1 through the decode GEMV path (teacher-forced decode PPL): judges decode-only sidecars (fp8_attn_proj, NVFP4 decode cache on GGUF). Qwen3-30B-A3B Q4_K_M 10.9443 (chunk 0) vs 10.9451 (chunk 1); nvfp4_decode_all 11.3011 (+3.25 %). PPL alone misses degenerate low-PPL models: run tools/analysis/degen_suite.py on a server too (check-degeneration).
• Per-block attribution: diagnostics.dump_hidden_dir + tools/analysis/layer_ab_diff.py (added divergence per block).

Gotchas