Per-engine per-format NATIVE-vs-EMULATED quantization KERNEL matrix (which kernel actually executes on which arch, and which fallback it reaches when none is native) — the layer underneath the engine support matrices

GROUP SIZE LEFT NULL and NOT MEANINGFUL FOR THIS RECORD: columns span NVFP4 group 16, MXFP4/MXFP6/MXFP8 group 32, Marlin/Machete group -1/32/64/128/256. PROXY ENUM VALUE - DECLARED 2026-10-05: a MATRIX record filed under 'nvfp4'; use the nvfp4 record for NVFP4 itself. SIBLING OF w6q-vllm-kernel-capability-matrix (wave 4), NOT A REPLACEMENT. Wave 4 answered 'which format has a NATIVE GEMM on which GPU generation, for vLLM'. This record answers the question UNDERNEATH that one, which wave 4 could not close because it was scoped to one engine: (a) which kernel actually EXECUTES per engine x format x arch, including the case where the format is supported and the kernel is a weight-only substitute; (b) every CUDA/ROCm/FlashInfer/toolkit VERSION gate, which no support matrix expresses; (c) engines wave 4 did not cover; (d) verified negatives with code evidence. Three of its findings CORRECT or NARROW wave-4 records and are marked CORRECTION below. ================================================================================ PART 1 — THE FALLBACK LADDER, WHICH IS THE WHOLE POINT ================================================================================ Wave 4 framed vLLM's MX and NVFP4 fallback as 'Marlin OR Emulation'. The current vLLM registries show the Marlin rung is not an alternative to Emulation but its PREDECESSOR, which changes what a reader should expect to happen. On NVIDIA the CUDA candidate lists are ordered: MXFP4 = [FlashInfer, Marlin, Humming, B12x, Emulation]; MXFP6 = [Humming, Emulation]; MXFP8 = [FlashInferCutedsl, FlashInferCutlass, Marlin, B12x, Emulation, Humming, FlashInferTrtllm]; NVFP4 = [FlashInferCuteDsl, FlashInferCutlass, FlashInferB12x, Cutlass, FlashInferCuteDslW4A16, Marlin, FlashInferTrtllm, FlashInferCudnn, Fbgemm, B12x, Torch, Emulation, Humming] with the in-code comment 'Weight-only: a fallback on a W4A4-capable checkpoint, not a preference.' Selection is first-match-wins over that order. Because Marlin's own gate is only is_cuda() and has_device_capability(75), the DEFAULT outcome for an MXFP4, MXFP8 or NVFP4 checkpoint on ANY SM75-or-later NVIDIA part is the WEIGHT-ONLY W4A16 Marlin kernel. Emulation is the second-order fallback, reached only sub-SM75, on ROCm, or for MXFP6 on ROCm. This is a correction of emphasis rather than of fact: the degradation wave 4 identified is right, but it is reached more often and on more hardware than 'emulation' implies, and the two are materially different in cost. Marlin preserves the 4-bit weight storage and therefore the MEMORY saving, and loses the W4A4 compute saving. EmulationMxFp8LinearKernel additionally gives up the memory saving: its code says 'Dequantize MXFP8 -> BF16 ONCE here, at load time, so apply_weights runs a plain BF16 linear with no per-step dequant -- i.e. run as if from a BF16 checkpoint. The 1-byte MXFP8 weight is replaced by BF16 (2x its size...)'. So the correct one-line rule, and it is the sentence to carry away from this record: for MX and NVFP4 formats on NVIDIA, a supported-not-native cell means Marlin W4A16 in the common case and BF16 emulation in the residual case, and the two are distinguishable ONLY by which name appears in the startup log, because Emulation's is_supported() returns (True, None) unconditionally and is guarded by nothing but a logger.warning_once. All four Emulation classes (MxFp4, MxFp6, MxFp8, NvFp4) are unconditional; the NVFP4 one carries the comment 'Always available as a last-resort fallback.' PER-FORMAT, PER-ARCH KERNEL THAT EXECUTES (vLLM main @ 528772a4, 2026-10-05). NVFP4 W4A4. Native: sm100/103 via FlashInferCuteDslNvFp4LinearKernel (is_device_capability_family(100)) or FlashInferCutlassNvFp4LinearKernel (cutlass_fp4_supported() AND has_device_capability(100) AND has_flashinfer()) or CutlassNvFp4LinearKernel (cutlass_fp4_supported()) or TorchNvFp4LinearKernel (family 100 + torch._scaled_mm + torch.float4_e2m1fn_x2). sm120/121 via FlashInferB12xNvFp4LinearKernel (has_device_capability(120) AND has_flashinfer_b12x_gemm()) or B12xNvFp4LinearKernel (family 120 + pip install vllm[b12x]) or the same Cutlass path (cutlass_fp4_supported covers 120<=cc<130). Non-native: MarlinNvFp4LinearKernel (SM75+) on everything else, then EmulationNvFp4LinearKernel, then HummingNvFp4LinearKernel (SM75+). On ROCm the registry is [EmulationNvFp4LinearKernel] and nothing else — no AITER, no CUTLASS — so NVFP4 on AMD is unconditional BF16 emulation, re-confirmed on current main. MXFP4. Native: sm100/103 AND sm120 via FlashInferMxFp4LinearKernel, whose entire gate is `if current_platform.has_device_capability(100) and has_flashinfer_cutedsl()`. has_device_capability is a >= comparison, so this PASSES on sm120 as well as sm100 — the only thing separating a 5090 from a native MXFP4 GEMM is whether the optional flashinfer.cute_dsl module is importable. Also sm120 via B12xMxFp4LinearKernel (family 120 + vllm[b12x]); SM75+ via HummingMxFp4LinearKernel (optional package); SM75+ via MarlinMxFp4LinearKernel which logs 'MarlinMxFp4LinearKernel is a weight-only (A16) kernel; the requested activation quantization (%s) is ignored.' VERIFIED NEGATIVE, re-confirmed on current main: sm90 AND BELOW CANNOT REACH A NATIVE MXFP4 GEMM AT ALL. The candidate list is FlashInfer (cc>=100), B12x (cc family 120), Humming (cc>=75) and Emulation, so on Hopper the only survivors are Humming-without-the-package and Emulation. Contrast TensorRT-LLM, whose matrix marks MXFP4 an unconditional Y on sm120: same silicon, and in vLLM the Y is conditional on a Python module being present. On ROCm, AiterMxfp4LinearKernel gates on current_platform.supports_mx() AND a required quark package; supports_mx() is `any(gfx in _GCN_ARCH for gfx in ['gfx95', 'gfx1250'])`, i.e. CDNA4 only, so gfx942 (MI300/MI325) is NOT in the set. MXFP6. Two candidates on CUDA, both narrow: HummingMxFp6LinearKernel (SM75+ and optional package) and EmulationMxfp6LinearKernel. There is NO CUTLASS and NO FlashInfer MXFP6 path in vLLM main, so MXFP6 is emulation wherever Humming is absent. On ROCm the list is [EmulationMxfp6LinearKernel] alone. MXFP8. Native on sm100/103: FlashInferCutedslMxfp8LinearKernel (family 100 + cutedsl module), FlashInferCutlassMxfp8LinearKernel (cc>=100 + has_flashinfer()), FlashInferTrtllmMxfp8LinearKernel (family 100). Native on sm120/121: B12xMxfp8LinearKernel (family 120 + vllm[b12x]) and FlashInfer mm_mxfp8, which additionally requires N % 32 == 0 on SM12x with the reason string 'mm_mxfp8 on SM12x requires N to be divisible by 32'. Weight-only on SM75+: MarlinMxfp8LinearKernel, gated is_fp8_marlin_supported() = has_device_capability(75), and requiring K % MXFP8_BLOCK_SIZE == 0. Native on AMD: RocmDotScaledMxfp8LinearKernel, ROCm-only, gated on supports_mx() with the comment 'supports_mx() == gfx95x (CDNA4 native microscaling hardware). On other archs dot_scaled would upcast to BF16, so the kernel selector falls through to the BF16 emulation (hipBLASLt) path instead.' MXFP8 is the ONLY MX format with a native AMD path in vLLM, and even that is CDNA4-only. FP8, THREE SCALING MODES, THREE DIFFERENT SUPPORT SETS AND THREE DIFFERENT FALLBACK SHAPES. (i) Per-tensor: CUTLASS-gated at has_device_capability(89) via cutlass_scaled_mm_supports_fp8 — CUDA >= 12.4 on sm89, CUDA >= 12.0 on cc>=90, false below 89. CudapPlatform.supports_fp8() is the same 89 floor. So Ada and Hopper (and everything above, which inherits the cc>=90 branch) are native, and Ampere is NOT merely unlisted but has a NAMED fallback: MarlinFP8ScaledMMLinearKernel, whose module comment reads 'For GPUs that lack FP8 hardware support, we can leverage the Marlin kernel for fast weight-only FP8 quantization', gated only at has_device_capability(75) and additionally refused under VLLM_BATCH_INVARIANT. An FP8 checkpoint on an A100 loads and runs as weight-only FP8 through Marlin with FP16 activations: memory saving without the W8A8 compute saving. There is also a third rung for pure FP8 on CUDA that no wave-4 record lists: PerTensorTorchFP8ScaledMMLinearKernel and ChannelWiseTorchFP8ScaledMMLinearKernel (torch._scaled_mm) and HummingFP8ScaledMMLinearKernel (SM75+ optional). (ii) Block-scaled (1x128/128x128 DeepSeek layout): the CUDA list is [FlashInferFp8DeepGEMMDynamicBlockScaled, DeepGemmFp8BlockScaledMMKernel, CutlassFp8BlockScaledMMKernel, B12xFp8BlockScaledMMKernel, MarlinFP8ScaledMMLinearKernel, HummingFP8ScaledMMLinearKernel, TritonFp8BlockScaledMMKernel, BlockWiseTorchFP8ScaledMMLinearKernel] — a long ladder whose lower rungs are weight-only Marlin and then plain Triton. DeepGemmFp8BlockScaledMMKernel requires is_deep_gemm_supported(), whose failure reason is 'Currently, only Hopper and Blackwell GPUs are supported.' CutlassFp8BlockScaledMMKernel requires CUDA 12.0 on 90<=cc<100 and CUDA 12.8 on cc>=100. B12xFp8BlockScaledMMKernel is family 120 and requires group_shape exactly (1,128) activations and (128,128) weights. has_flashinfer_fp8_blockscale_gemm() requires is_device_capability(90) EXACTLY plus the attribute flashinfer.gemm.fp8_blockscale_gemm_sm90, so the FlashInfer rung is Hopper-only by exact equality even though the CUTLASS rung above it accepts Blackwell. BlockWiseTorchFP8ScaledMMLinearKernel is the sharpest one: its own comment says 'torch._scaled_mm DeepSeek-style (1x128, 128x128) block scaling is only implemented for SM90 (Hopper) on NVIDIA CUDA' with a link to PyTorch's _check_deepseek_support(), and it returns False for anything that is not exactly SM90 — so the generic torch rung of the block-FP8 ladder is Hopper-only while the CUTLASS rung is Hopper-and-Blackwell. (iii) Rowwise / per-token dynamic: THE SILENT DEGRADATION, and it is the cleanest example in this record of a loss that no support table can express. In Fp8LinearMethod.__init__ the activation quant key is chosen as `if self.act_q_static: kFp8StaticTensorSym; elif cutlass_fp8_supported(): kFp8DynamicTokenSym; else: kFp8DynamicTensorSym` — commented 'Use per-token quantization for better perf if dynamic and cutlass'. On a part where cutlass_fp8_supported() is false, an FP8 checkpoint that asked for per-token activation scaling SILENTLY GETS PER-TENSOR ACTIVATION SCALING INSTEAD, with no error and at most a logger line. The output is still correct, the memory saving is intact, and the accuracy headroom per-token scaling buys is simply gone. Same shape, different mechanism, on ROCm: RowWiseTorchFP8ScaledMMLinearKernel exists in the ROCm list but refuses float16 output ('hipblaslt rowwise _scaled_mm only supports BFloat16') and refuses per-tensor scales. W4A8 IS TWO FORMATS AND ONLY ONE OF THEM IS A HARD FAILURE. W4A8-FP8 (int4 weights, fp8_e4m3 activations) is Hopper-EXACTLY-90: CompressedTensorsW4A8Fp8.get_min_capability() returns 90 and CutlassW4A8LinearKernel.can_implement requires is_device_capability(90) exactly, act_type float8_e4m3fn, zero_points False, weight_type int4, group_size EXACTLY 128, in_features and out_features both divisible by 128, out_dtype bfloat16. The equality check means sm100/103/107 and sm120 are EXCLUDED from W4A8-FP8 even though they are newer and have strictly better FP8 hardware; the class constructor also rejects anything but 'W4A8 kernels require group quantization with group size 128'. W4A8-INT8 (int4 weights, int8 activations) is the opposite: the CUDA mixed-precision list [CutlassW4A8, Machete, Marlin, Conch, Exllama, TritonW4A16, Humming] rejects the config in all seven cases and choose_mp_linear_kernel raises ValueError('Failed to find a kernel that can implement the WNA16 linear layer. Reasons: ...') at load time. It IS available on CPU (Dynamic4bitLinearKernel with no x86 exclusion, and CPUWNA16LinearKernel). W4A16 (AWQ/GPTQ): the one family genuinely native from Turing up. MarlinLinearKernel.get_min_capability() = 75, MARLIN_SUPPORTED_GROUP_SIZES = [-1,32,64,128], and query_marlin_supported_quant_types returns [] below cc 75. MacheteLinearKernel requires is_device_capability(90) EXACTLY with query_machete_supported_group_sizes = [-1,64,128] for fp16/bf16 and [-1,128] otherwise, and query_machete_supported_quant_types = [uint4, uint8] with zero points or [uint4b8, uint8b128] without. ExllamaLinearKernel returns min capability 60 (Pascal) but requires float16 activations exactly and is 'only supported on CUDA and ROCm'. TritonW4A16LinearKernel.get_min_capability() returns 0 with the comment 'Triton handles capability checks itself' and group sizes [-1,32,64,128,256], making it the universal W4A16 fallback on CUDA, ROCm and XPU. ConchLinearKernel returns 80 (note: wave 4 and this record's wave-4 sibling describe Conch as group sizes [-1,128]; current main returns min capability 80) and needs the conch-triton-kernels package. CPUWNA16LinearKernel.get_min_capability() = -1. AMD, BRIEFLY, BECAUSE 'EMPTY' IN AN NVIDIA TABLE READS LIKE 'UNKNOWN'. Native MXFP8 on gfx95/gfx1250 (CDNA4) via RocmDotScaledMxfp8. Native MXFP4 on gfx95/gfx1250 behind AITER plus a required quark package. Native FP8 on CDNA and RDNA4 via AITER: AiterInt8ScaledMMLinearKernel requires ROCm, compute_capability >= 90 and VLLM_ROCM_USE_AITER=1 plus VLLM_ROCM_USE_AITER_LINEAR=1; AiterPreshuffledPerTokenFp8ScaledMMLinearKernel and AiterHipbMMPerTokenFp8ScaledMMLinearKernel require the same env flags plus aiter, and AiterPerTokenFp8ScaledMMLinearKernel additionally refuses any shape for which is_per_token_w8a8_gemm_tuned(N, K, fp8_dtype) is False with the reason 'requires a tuned configuration for N ... and K ... and fp8 dtype ...' — i.e. on AMD an untuned layer shape silently drops a rung down the FP8 ladder. AiterFp8BlockScaledMMKernel requires group_shape (1,128) activations and, on RDNA4 (gfx12), a per-(N,K) tuned Triton blockscale shape ('RDNA4 (gfx12) only has the aiter Triton blockscale backend, which needs a per-(N,K) tune. Reject untuned shapes so the dispatcher falls through to the generic backend.'). ROCmFP8ScaledMMLinearKernel requires CDNA3+ (gfx942/gfx950) or RDNA4 (gfx12x), VLLM_ROCM_USE_SKINNY_GEMM, and is NOT built on gfx1250. Native INT8 via AiterInt8ScaledMMLinearKernel and TritonInt8ScaledMMLinearKernel. W4A16 via RDNA3W4A16LinearKernel (requires on_gfx1100(), and requires torch.ops._rocm_C.gptq_gemm_rdna3 to exist or it returns False with 'torch.ops._rocm_C.gptq_gemm_rdna3 missing — rebuild C++ extension'), RDNAHybridW4A16LinearKernel (targets gfx11/gfx12 only), TritonW4A16LinearKernel and ExllamaLinearKernel. NOT NATIVE ANYWHERE ON AMD: Marlin (not in the ROCm list at all), Machete, CUTLASS, and NVFP4 (emulation only). Numerical caveat a capacity plan needs: on gfx94-class parts is_fp8_fnuz() is True and fp8_dtype() returns torch.float8_e4m3fnuz, not OCP e4m3fn, so an AMD FP8 number is not the same format as a Blackwell one. ================================================================================ PART 2 — EVERY VERSION GATE FOUND, ONE PLACE ================================================================================ These are the gates a support matrix cannot express, because they are properties of the RUNTIME, not the silicon. A format marked supported on hardware that is right still fails or degrades on a runtime that is too old. CUDA RUNTIME GATES, vLLM C++ (all re-confirmed on main @ 528772a4, unchanged from wave 4). cutlass_scaled_mm_supports_fp8: CUDA >= 12.4 for cc in [89,90); CUDA >= 12.0 for cc >= 90; FALSE below 89. Because the cc>=90 branch has no upper bound, sm100/103/107/120 inherit the Hopper rule — FP8 per-tensor above Hopper is CUDA-version-gated, not arch-gated. cutlass_scaled_mm_supports_block_fp8: CUDA >= 12.0 for cc in [90,100); CUDA >= 12.8 for cc >= 100; FALSE below 90. cutlass_scaled_mm_supports_fp4: cudaRuntimeGetVersion() >= 12080 RUNTIME (not compile-time CUDA_VERSION, unlike the other two — a real asymmetry worth knowing), then 100<=cc<120 or 120<=cc<130, and only if the matching ENABLE_NVFP4_SM100/SM120 was compiled in; its own comment states the reason: 'Only report support when the SM-specific kernel was actually compiled in, so the Python-side backend selector does not choose CUTLASS and then hit TORCH_CHECK_NOT_IMPLEMENTED (or worse, fall through to Marlin).' NEW, NOT IN WAVE 4: cutlass_group_gemm_supported — 'CUTLASS grouped FP8 kernels need at least CUDA 12.3 and SM90 (Hopper) or CUDA 12.8 and SM100 (Blackwell)': CUDA >= 12.8 for 100<=cc<120 under ENABLE_CUTLASS_MOE_SM100, CUDA >= 12.3 for 90<=cc<100 under ENABLE_CUTLASS_MOE_SM90, else false. So a MoE deployment needs a HIGHER CUDA floor than a dense one on the same Hopper part (12.3 vs 12.0), and the same 12.8 as native NVFP4 on Blackwell. FLASHINFER VERSION GATES. SGLang's MXFP4 MoE path on sm90 requires FlashInfer >= 0.6.11 by name: RuntimeError('moe_runner_backend=flashinfer_mxfp4 on SM90 requires the interleave_moe_{weights,scales}_for_sm90_mixed_gemm helpers from FlashInfer PR #3084 (>= 0.6.11). Upgrade flashinfer-python or pick a different backend (e.g. marlin / triton_kernel).') In vLLM the equivalent coupling is a RENAME, which is the same class of breakage wearing a different hat: has_flashinfer_b12x_gemm() accepts EITHER name, with the comment 'FlashInfer 0.6.11 renamed Sm120BlockScaledDenseGemmKernel -> Sm120B12xBlockScaledDenseGemmKernel (commit 223f2a49). Accept either.' A vLLM build pinned against the wrong FlashInfer minor silently loses every b12x kernel — which are the sm120 native NVFP4/MXFP4/MXFP8/FP8-block paths — and falls through to Marlin. has_flashinfer_cutedsl() is a module-presence test (importlib.util.find_spec('flashinfer.cute_dsl')), and it gates FlashInferCuteDslNvFp4LinearKernel, FlashInferCuteDslNvFp4W4A16LinearKernel, FlashInferMxFp4LinearKernel, FlashInferCutedslMxfp8LinearKernel, has_flashinfer_bf16_fp4(), has_flashinfer_b12x_gemm() and has_flashinfer_cutedsl_grouped_gemm_nt_masked(). has_flashinfer_bf16_fp4() additionally requires flashinfer.gemm to expose BOTH mm_bf16_fp4 and prepare_bf16_fp4_weights. vLLM's block-FP8 FlashInfer rung has a third gate on top: is_flashinfer_fp8_blockscale_gemm_supported() additionally requires the ENV VAR VLLM_BLOCKSCALE_FP8_GEMM_FLASHINFER to be set, so block-FP8 GEMM on Hopper can be switched off by configuration. SGLANG'S OWN ARCH PREDICATES ARE CUDA-VERSION GATES. In sglang/srt/utils/common.py, built on _check_cuda_device_version(device_capability_majors=[...], cuda_version=(M,N)): is_blackwell = majors [10,11,12] AND CUDA >= 12.8; is_sm120_supported = majors [12] AND >= 12.8; is_sm100_supported = majors [10] AND >= 12.8; is_sm90_supported = majors [9] AND >= 12.3; is_sm80_supported = majors [8] AND >= 11.0; is_ampere_with_cuda_12_3 and is_hopper_with_cuda_12_3 are the 12.3 versions. Plus exact-match predicates the major-based helper cannot express, with the comment 'GB10 (DGX Spark and OEM equivalents). Not expressible with _check_cuda_device_version, which only matches on the major.': is_sm120() == (12,0) exactly, is_sm121() == (12,1) exactly. CONSEQUENCE: in SGLang 'is Blackwell' is a CUDA-12.8-AND-Blackwell test, so a B200 on CUDA 12.6 is not Blackwell to SGLang's dispatch and takes a different branch than the same B200 on 12.8. SGLang's cutlass_fp8_supported() is the same predicate as vLLM's C++ in Python form: major >= 9 AND CUDA >= (12,0); major == 8 and minor == 9 AND CUDA >= (12,4); else False. TORCH GATE. SGLang's NVFP4 KV cache requires hasattr(torch, 'float4_e2m1fn_x2') else ValueError('--kv-cache-dtype=nvfp4 requires torch.float4_e2m1fn_x2 support. Please use PyTorch 2.8.0+ with CUDA 12.8+.'), and it also emits a request-time WARNING('%s KV Cache might lead to an accuracy drop!'). vLLM's NVFP4 KV gate adds two more: force_use_trtllm_attention() must be true, and it is DISABLED under VLLM_BATCH_INVARIANT. BUILD-TIME GATES, WHICH ARE THE DEADLIEST BECAUSE NOTHING WARNS. ENABLE_NVFP4_SM100 / ENABLE_NVFP4_SM120 decide whether native NVFP4 CUTLASS exists AT ALL; ENABLE_SCALED_MM_SM90 / SM100 / SM120 / C2X decide the FP8 arch dispatch; ENABLE_CUTLASS_MOE_SM90 / SM100 decide grouped FP8 MoE. All are compile-time macros, so a wheel built without them reports False from the Python predicate and the selector quietly takes the next rung. pip install vllm[b12x] is the same class of gate wearing a hat: B12x kernels are OPTIONAL and every B12x is_supported() ends in 'Install the B12X backend with `pip install vllm[b12x]`'. ML-LLM's cutlass groupwise GEMM is in the same class and worse, because it is a TVM compiled-in function rather than a package: BlockScaleQuantizeLinear.forward() probes `extern.get_store().cutlass_gemm and tvm.get_global_func('cutlass.groupwise_scaled_gemm_e4m3fn_e4m3fn', allow_missing=True) is not None` and otherwise falls back to dequantize-then-BF16. NON-FP4 HARD FAILURES THAT ARE NOT EMULATION. SGLang raises rather than degrades: 'ModelOpt NVFP4 native dense GEMM backends require SM100+. Use --fp4-gemm-backend marlin on SM80-SM90.' (this check sits AFTER the Marlin early-return, so the sequence is: marlin backend -> group_size must be 16 else ValueError; any other backend -> not is_blackwell else ValueError), and for MXFP4 MoE on anything that is not sm90/sm100/sm120, NotImplementedError('moe_runner_backend=flashinfer_mxfp4 requires SM90, SM100, or SM120.'). Marlin in SGLang additionally locks group size: 'NVFP4 Marlin requires group_size=16, got ...' and 'NVFP4A16 Marlin requires group_size=16, got ...'. Contrast vLLM, where the same checkpoint loads and runs degraded. Same hardware, same format, two engines, one fails loudly and one fails quietly. ================================================================================ PART 3 — ENGINES ADDED, AND WHAT EACH ONE'S KERNEL DISPATCH ACTUALLY LOOKS LIKE ================================================================================ Wave 4 covered vLLM, SGLang, TensorRT-LLM, LMDeploy, TGI, Triton, llama.cpp and added Xinference, TEI, Dynamo, Envoy. From this wave's read of source, the engines with quantization KERNELS worth recording — and the dispatch shape each uses, which is itself the useful taxonomy — are: (1) TensorRT-LLM: tile-candidate heuristics, not kernel registries. cutlass_heuristic.cpp is a per-arch candidate-tile generator: sm==90 && HOPPER -> get_candidate_configs_sm90; sm in [100,120) && BLACKWELL -> get_candidate_configs_sm100; sm >= 120 && BLACKWELL && (FP4_ONLY || FP8FP4_MIXED) -> get_candidate_configs_sm120; everything else -> get_candidate_tiles(sm, ...) with Ampere-style FP8 tiles gated `if (sm == 89 || sm >= 120)` and an explicit 'no valid ampere style fp8 configs for sm90' empty return. The hard structural gate is 'if ((FP4_ONLY) && !(BLACKWELL)) { // FP4 is only supported on blackwell return {}; }'. min_stages is 3 on sm89 and max_stages is 2 below sm80. CORRECTION: TensorRT-LLM NOW SHIPS A MARLIN NVFP4 KERNEL, contradicting the wave-4/docs-23 reading that the engine has no FP4 route off Blackwell. See the corrected claim below and w6q-trtllm-marlin-nvfp4-2026-10. (2) llama.cpp / GGUF: six kernel families by arch class, and the arch gate is a COMPILE-TIME __CUDA_ARCH__ constant rather than a runtime query. This is the only engine in the set whose kernel family is chosen by #if, which means a build compiled for the wrong arch silently gets a different kernel with no runtime check. (3) MLC-LLM: compiled-operator dispatch. Scheme dataclasses (GroupQuantize, BlockScaleQuantize) + a runtime probe for a compiled-in cutlass extern. No hand-written CUDA kernels, so its 'kernel matrix' is really a scheme-enum-and-build-flags matrix. (4) TGI (Gaudi backend): a closed string dispatch, exactly enumerable at four values. (5) LMDeploy: prose arch lists plus SmoothQuant-produced checkpoints and QLinear module replacement. (6) XPU inside vLLM: XPUMxFp4LinearKernel and XPUMxFp8LinearKernel gate on is_xpu() ONLY — no compute-capability or arch test whatsoever, and there is NO EmulationMxFp4LinearKernel in the XPU registry (XPU MXFP4 = [XPUMxFp4LinearKernel] alone) while XPU MXFP8 = [XPUMxFp8LinearKernel, EmulationMxfp8LinearKernel] and XPU FP8-block = [XPUFp8BlockScaledMMKernel, TritonFp8BlockScaledMMKernel, BlockWiseTorchFP8ScaledMMLinearKernel]. That asymmetry — MXFP4 on XPU has no fallback rung at all, MXFP8 has two — is a genuine gap and no record states it. (7) CPU inside vLLM: the FP8 story is AMX-shaped, and the fallbacks are named in the failure reasons. CPUFP8W8A8ScaledMMLinearKernel requires torch.cpu._is_amx_tile_supported() AND ops.cpu_has_amx_fp8() ('requires native AMX-FP8 hardware (Diamond Rapids); falls back to the Torch FP8 kernels otherwise'); CPUFp8W8A8BlockScaledMMKernel requires AMX tile AND cpu_has_amx_fp8 ('requires native AMX-FP8 hardware; falls back to CPUFp8BlockScaledMMKernel (W8A16) otherwise'); CPUFp8PerTensorScaledMMLinearKernel and CPUFp8BlockScaledMMKernel require only AMX tile ('requires AMX tile support (Sapphire Rapids or newer)') plus ops._supports_cpu_fp8_w8a16. So 'x86 CPU supports FP8' resolves to a two-level cascade: AMX-FP8 (Sapphire Rapids+) for W8A8, AMX tile (Sapphire Rapids+) for W8A16, and the shape constraints bite too — the AMX tinygemm path requires N % 32 == 0 ('The AMX tinygemm kernel tiles N in chunks of 32 and cannot handle a remainder tile smaller than that'), and the block kernels require weight group_shape col exactly 128 and row a positive multiple of 32. CONSIDERED AND DELIBERATELY NOT RECORDED AS KERNELS, with reasons, because recording them would be inventing a claim. MindIE and Firefly: closed-source Huawei distributions; no kernel source is readable, so nothing beyond the existing engine record can be asserted. PaddlePaddle / vLLM-Ascend: the Ascend NPU quantization path routes through CANN/Ascend-CANN custom ops and torch_npu, whose gates are not in the vllm-ascend tree and were not read; the repo already carries niche-vllm-ascend-repo and the SGLang npu_mxfp4.py module exists but its gate is a CANN operator check this wave did not verify. LocalAI, GPT4All, Jan, KoboldCpp: these are serving wrappers or CPU-first runtimes with no quantization GEMM kernel of their own (they delegate to llama.cpp, exllamav2 or a backend library), so they have no kernel matrix to record — the honest entry is that they inherit one. ExLlamaV2: a real kernel library, but its CUDA QDQ kernels (exllamav2_ext/cuda/quant/qdq_{2,3,4,5,6,8}.cuh) are dequantize-to-half dot-product kernels — qdq_util.cuh is literally dq(q, qzero, scale) = __hmul(__int2half_rn(q - qzero), scale) — so it is a W4A16-style dequant path with no arch gate in the header; recorded here as a classification, not as a gate. GGUF k-quant variants are covered in (2) via the MMQ family ladder. ================================================================================ PART 4 — VERIFIED NEGATIVES, WITH THE CODE EVIDENCE ================================================================================ N1. Native NVFP4 exists only on sm100/103 and sm120/121 AND requires CUDA 12.8+ — STILL HOLDS. cutlass_scaled_mm_supports_fp4 returns false unless cudaRuntimeGetVersion() >= 12080, then only for 100<=cc<120 (ENABLE_NVFP4_SM100) or 120<=cc<130 (ENABLE_NVFP4_SM120). Same asymmetry in TensorRT-LLM, where FP4 candidates require BLACKWELL and only Sm100/Sm103/SM120 launcher templates are instantiated. WHERE IT DOES NOT HOLD: llama.cpp, where GGML_CUDA_CC_BLACKWELL is 1200 not 1000, so 'Blackwell' in that engine means sm120-sm129 only and an NVFP4 GGUF on a B200 runs the AMPERE integer MMQ config. See N6. N2. Native MXFP4 does not exist on sm90 and below — STILL HOLDS in vLLM (FlashInfer needs cc>=100, B12x needs family 120, Humming needs the optional package, so Hopper's only non-emulation rung is package-gated) and INDEPENDENTLY CONFIRMED IN TensorRT-LLM, where cpp/tensorrt_llm/kernels/cutlass_kernels/fp4_gemm/ instantiates Sm100, Sm103 and SM120 launchers and there is NO Sm90 FP4 template instantiation at all. On sm100/sm120 it still sits behind an optional module in vLLM (flashinfer.cute_dsl, or vllm[b12x], or humming, or a required quark package on ROCm) — that half also still holds. N3. There is NO 2-bit KV dtype at all because CacheDType is a closed enum — STILL HOLDS and is now CORROBORATED BY A SECOND ENGINE. vLLM's CacheDType still reads exactly: auto, float16, bfloat16, fp8, fp8_e4m3, fp8_e5m2, fp8_inc, fp8_ds_mla, nvfp4_ds_mla, turboquant_k8v4, turboquant_4bit_nc, turboquant_k3v4_nc, turboquant_3bit_nc, int4_per_token_head, int8_per_token_head, fp8_per_token_head, nvfp4, nvfp4_4over6 — 18 members, lowest sub-8-bit width is 3 bits (turboquant_k3v4_nc, turboquant_3bit_nc) or 4 (int4_per_token_head), and there is no 2-bit member. SGLang's configure_kv_cache_dtype is SHORTER and reaches the same conclusion by a different route: it accepts only auto, fp8_e5m2, fp8_e4m3, mxfp8, bf16/bfloat16 and nvfp4/fp4_mx_block16, else raise ValueError(f'Unsupported kv_cache_dtype: ...'). SGLang has no int4_per_token_head, no int8_per_token_head, no fp8_per_token_head and no TurboQuant family at all. TWO ENGINEES, TWO DIFFERENT CLOSED ENUMS, ONE CONCLUSION. N4. Volta has no Marlin/AWQ because query_marlin_supported_quant_types returns [] below cc 75 — STILL HOLDS in vLLM (`if device_capability < 75: return []`) but the two engines now DISAGREE about the boundary: SGLang's equivalent returns [] below cc 80 (`if device_capability < 80: return []`). So on Turing (sm75) vLLM has Marlin and SGLang does not. Same kernel family, two floors, five points apart, and no record in the repo states the disagreement. N5. NEW — the NVFP4 registry in vLLM contains two kernels with NO architecture gate. FlashInferTrtllmNvFp4LinearKernel and FlashInferCudnnNvFp4LinearKernel both have is_supported() = `if has_flashinfer(): return True, None` and can_implement() = `return True, None`. Every other NVFP4 W4A4 kernel in the registry checks an arch family, a CUTLASS predicate, or a torch attribute. Selection is first-match-wins over an ordered list, so with FlashInfer installed these two entries are reachable on any NVIDIA part including a V100. WHAT THIS INVALIDATES: the repo's recommended probe torch.ops._C.cutlass_scaled_mm_supports_fp4(capability) remains the correct way to ask 'does CUTLASS have a native FP4 GEMM for this part', but it is NOT a complete answer to 'will vLLM run an FP4 tensor-core GEMM for my NVFP4 checkpoint', and it is not a complete answer to 'is this cell safe'. See w6q-vllm-nvfp4-ungated-flashinfer-2026-10. N6. NEW — llama.cpp's native FP4 MMQ path exists ONLY on cc 1200-1299, not on sm100/sm103. GGML_CUDA_CC_BLACKWELL is defined as 1200 ('While BW spans CC 1000, 1100 & 1200, we are integrating Tensor Core instructions available to 1200 family') with GGML_CUDA_CC_RUBIN 1300, and BLACKWELL_MMA_AVAILABLE requires __CUDA_ARCH__ >= 1200 && < 1300. Consequence: an MXFP4 or NVFP4 GGUF on a B200/B300 runs ggml_cuda_mmq_get_config_ampere (integer SIMT), while the SAME quant type on an RTX 5090 reaches ggml_cuda_mmq_get_config_blackwell. That is the inverse of vLLM and TensorRT-LLM. It also narrows an existing repo claim: 'No accelerator has an i-quant or k-quant path in silicon' remains TRUE for i-quants and k-quants and is now FALSE for MXFP4 and NVFP4 GGUF on sm120-sm129. N7. NEW — even on sm120+, llama.cpp's FP4 MMQ config is conditional on the ACTIVATION precision, not just the weight. The ampersand-branch is `if (prec_src1 != GGML_PREC_Q4 && (type == GGML_TYPE_NVFP4 || type == GGML_TYPE_MXFP4)) return ggml_cuda_mmq_get_config_ampere(...)`, commented 'only src1 at Q4 uses the native FP4 config, higher precisions keep src1 at Q8_1'. So at Q8 activations the FP4 weight type drops back to the Ampere config even on a 5090. N8. NEW — MXFP6 in vLLM is emulation by default, with exactly one real kernel. The CUDA list is [HummingMxFp6LinearKernel, EmulationMxfp6LinearKernel] and the ROCm list is [EmulationMxfp6LinearKernel]. No CUTLASS, no FlashInfer, no AITER MXFP6 path exists. Contrast MXFP8 and MXFP4, both of which have vendor kernels. N9. NEW — MLC-LLM's group-quant family is a closed enum over {int3, int4, int8} and its block-scale family is a closed enum over {float8_e4m3fn, float8_e5m2}, so MLC has no FP4 group-quant scheme at all and no non-FP8 block-scale scheme. Two schemes, two closed enums, no FP4 anywhere. N10. NEW — TGI's Gaudi backend supports exactly four quantize values (compressed-tensors, awq, gptq, fp8) plus None, and None maps to the FP8-capable HybridFP8UnquantLoader rather than to unquantized. There is no exl2/exl3 loader, no int8 W8A8 quantize value, and no nvfp4/mxfp4 value. ================================================================================ CORRECTIONS TO WAVE 4, STATED PLAINLY ================================================================================ C1. 'vLLM falls back to emulation' is the wrong default for MXFP4/MXFP8/NVFP4 on NVIDIA. Marlin W4A16 precedes Emulation in every one of those CUDA registries and its own floor is SM75, so the usual outcome on Turing-through-Blackwell is Marlin, and Emulation is the residual case (sub-SM75, ROCm, or MXFP6 on ROCM). Wave 4's degradation finding was right; its default-case attribution was not. C2. TensorRT-LLM DOES have an NVFP4 kernel off Blackwell now. isMarlinNvfp4SmSupported(sm) = sm >= 89 && sm < 100; isMarlinNvfp4DenseSmSupported(sm) = that OR sm == 120 OR sm == 121; the launch is TLLM_CHECK_WITH_INFO-gated on it with the message 'Marlin NVFP4 GEMM is only supported on SM89, SM90-99, and SM120/121; current SM = %d'. It is a W4A16 kernel, so it reproduces the same degradation — but the arch set is the COMPLEMENT of native FP4 on sm100/103, which is the elegant part and the reason it is worth recording. C3. docs/23's sm120-inversion finding stands and got sharper, not weaker: on sm120 TensorRT-LLM now has BOTH native CUTLASS FP4 templates AND a Marlin NVFP4 W4A16 kernel, and still no AWQ/GPTQ kernel in the sm120 candidate list. And vLLM on sm120 has native NVFP4 (b12x/FlashInfer, optional), native MXFP4 (conditional on flashinfer.cute_dsl), native MXFP8 (b12x), block-FP8 that CUTLASS REFUSES on sm120 (cutlass_scaled_mm_supports_block_fp8 returns false for cc<100 that is not >=90 — i.e. sm120 takes the else branch and returns FALSE) and a B12x FP8-block rung that requires family 120 plus group shapes (1,128)/(128,128). So on consumer Blackwell, FP4 is native, block-scaled FP8 is available through an optional package but not through CUTLASS, and AWQ/GPTQ has no vendor CUTLASS path — Marlin or Triton. Three different answers to three formats on one architecture family, and no support table in the repo can express all three. ================================================================================ WHAT A CAPACITY PLAN SHOULD DO WITH THIS RECORD ================================================================================ Ask the running process which kernel it chose, not what the hardware supports. Concretely: read the engine's startup log for the kernel-class name (vLLM logs 'Using %s for MXFP8 GEMM' and 'Using %s for CompressedTensorsW4A8Fp8'; SGLang exposes get_fp4_gemm_runner_backend() and get_moe_runner_backend()); and treat the presence of 'Marlin', 'Emulation', 'Triton' or 'torch' in that name as the fact that determines your capacity plan, regardless of what the format record says. The three checks that cost nothing remain: the kernel/backend line in the log; the scheme name in the config (nvfp4_w4a16 is the tell); and a direct probe of the CUTLASS predicates for the specific format — with the PART 2 caveat that a FALSE from cutlass_scaled_mm_supports_fp4 does not by itself tell you which kernel you will get, because of N5, and a TRUE does not tell you the build compiled the SM you need.