FEATURES.md

Features

What vllm.cpp supports, next to the engines it is measured against. This page is a keyed table: one row per feature, kept current. It is not a changelog.

For measured speed see BENCHMARKS.md; for per-capability lifecycle conventions see Project status; for the agent-facing parity inventory with upstream file references see .agents/feature-matrix.md.

Legend. ✅ supported and gated. ◐ partial, usable with named gaps. ☐ not yet. n/a means the feature does not apply to that engine's design.

Reference versions: vLLM 0.28.1rc1.dev132 (e126687a9a, the parity pin since 2026-09-03), SGLang v0.5.15, llama.cpp b10451, MLX-LM as of 2026-07. Rows describing what vLLM has were read at the PRIOR pin 555967922 unless they say otherwise; the 290-commit-range PORT-NOW queue for the advance is classified and unworked (#2611). Competitor columns describe what those projects ship, and are our reading of their documented behavior, not measurements.

Three cautions about the citations and the results below, and they are different.

Citations. The issue tracker lost everything below roughly #1870. When this page was measured on 2026-09-03 it carried 157 distinct in-repo issue numbers, and 107 of them no longer resolved (queried against repos/mudler/vllm.cpp/issues/N; a dead number returns "Not Found", which is not a closed state). Those 107 have had their hyperlinks removed and keep the bare number as a historical marker, because a link a reader cannot follow costs a click and returns a 404. Where a dead citation carried a claim a reader might act on, it has been replaced by an anchor into this tree — a source path or a spec — which is what the reader wanted from the link in the first place. The recovered descriptions are archived in .agents/completed/issue-index.md. Tracked by #2825.

A result is only true of the tree it ran on. Twice on 2026-09-03 a landed commit silently invalidated a measurement recorded here, and nothing flagged either. A row that states a run therefore names the SHA it ran on, so a reader can ask git merge-base --is-ancestor rather than assume the result survived. The GlmMoeDsaForCausalLM ROCm arm is the worked example: the run was real at 9f3e6e223, and #2511 later withdrew the premise it depended on.

Results. Every "vs vLLM" figure on this page was captured at the prior parity pin 555967922 and has not been re-validated at the current pin e126687a9a, which advanced on 2026-09-03. .agents/oracles/vllm.md states in its own words that the pin advance "does NOT say any gate in this tree has been run against it", and .agents/NOW.md records "NO gate has run at it". The rows below name vLLM 0.25.0 where that is the version they were measured against. Tracked by #2794 (goldens predate the pin) and #2817 (the advance).

At a glance

	vllm.cpp	vLLM	SGLang	llama.cpp
Language	C++20	Python + CUDA	Python + CUDA	C/C++
Runtime deps	none	PyTorch	PyTorch	none
Install size	66 MiB	9.1 GiB	comparable to vLLM	comparable to us
Embeddable behind a C ABI	✅	☐	☐	✅
Weight formats	Safetensors + GGUF	Safetensors	Safetensors	GGUF
Correctness gate	token-exact vs vLLM	reference	own	own
Architectures	43 registered, 27 gated	130+	100+	100+
Downloadable server binaries	✅ v0.0.2: eight indexed archives with checksums, provenance, manifests, and SBOMs. Windows ZIP downloads do not exist; native CPU/Vulkan lanes await hosted runtime, dry-run, prerelease, and authenticated audit gates	✅ wheels/containers	✅ wheels/containers	✅ host-specific binaries
Native Windows builds	◐ CPU/Vulkan: /MT /W4 /WX, central NOMINMAX, UTF-8, aligned allocation, C++20 std::numbers pi, runtime ISA dispatch. Local closure includes the float-domain DeepSeek probe; hosted compile/runtime/release pending	✅	✅	✅

Serving and scheduling

Feature	vllm.cpp	vLLM	SGLang	llama.cpp
Continuous batching	✅	✅	✅	◐
Chunked prefill	✅	✅	✅	☐
Automatic prefix caching	✅, and mutually exclusive with DFlash speculative decode in effect: a request served from the cache DROPS its draft and decodes on the target alone, because the draft keeps a private context store the target's skipped prefill never fills. It buys that request TTFT and costs it draft acceptance, so on a shared-system-prompt workload speculation is off for most requests and output throughput can fall. Tokens are unchanged (the verify is lossless). Moving the draft context into the paged allocator removes the trade, and is owed (#2042, #1919)	✅	✅ (radix)	◐
Preemption and recompute	✅	✅	✅	☐
Priority scheduling	◐ gating (--scheduling-policy priority reaches the server; scheduler-unit tests only, no engine-level priority-vs-FCFS gate exists yet, #534)	✅	✅	☐
LPM cache-aware admission	✅	☐	✅	☐
In-batch prefix de-prioritization	✅	☐	✅	☐
Async / overlap scheduling	✅ default on (UAF-safe drain; device token-ids mirror on gate + classic-dense; the decode graph declines while the mirror is live (#323 fix, eager fallback); opt-in VT_ASYNC_EXECUTOR out-of-capture H2D staging)	✅	✅	☐
CUDA graph decode capture	◐ per-family	✅	✅	✅
Partial-prefill concurrency	☐	✅	✅	☐
Cascade attention	☐	✅	◐	☐

KV cache and memory

Feature	vllm.cpp	vLLM	SGLang	llama.cpp
Block-paged KV with refcount and LRU evict	✅	✅	✅	◐
Hybrid KV groups (full attention + GDN/Mamba)	◐ GDN gate activation resolved from the checkpoint's output_gate_type (silu/swish/sigmoid; anything else refused at load, #489)	✅	◐	◐
Sliding-window and chunked-local attention	◐	✅	✅	✅
fp8 KV cache	◐ --kv-cache-dtype fp8 halves the block, so a fixed --kv-cache-memory buys 2x the blocks and the DEFAULT 256-block path halves the pool bytes instead. Costs the bf16-native FA-2/WMMA/vector kernels (net UNMEASURED on most archs; the ROCm GQA4 f32-Q decode arm recovers parity, fork issue #7). 16 archs and MLA refuse before any write; only 1 arch names fp8 back. C ABI: vllm_model_params.kv_cache_dtype (ABI v24). CUDA UNRUN (spec)	✅	✅	✅
KV offload to host memory	✅	✅	✅	☐
External KV provider ABI (LMCache)	☐	✅	◐	☐
KV events (block create / evict publish)	◐ no transport	✅	☐	☐
Prefix-cache matching unit	◐ resolver only	✅	☐	☐
Compute directly on quantized blocks	✅	☐	☐	✅
Scratch allocator keyed by device (two backends, one process)	✅ since #516; a pool is bound to one backend and refuses any other, and a backend with no registered platform is refused rather than given another's residency cap	✅ device is field 0 of the allocation handle	✅	✅
Automatic memory sizing (no hand-tuned budget)	☐ hand-typed block count	☐ percent, hand-tuned	☐	◐
Memory cap with a pre-flight error instead of an OOM	☐	◐ KV pool only	◐	☐
Routed-expert weight streaming from disk	◐ default OFF (VT_MOE_EXPERT_STREAM=1), keep-quant/keep-f16 towers only — GgufExpertTowersReachSlotLane refuses every other residency (src/vllm/model_executor/model_loader/gguf_device_fit.cpp:127, was #1378); bounded slot cache; refuses unfittable slices by name. c1-c4 capacity, not throughput. CPU; staging device DECODES and the declared token gate G0-CORRECT FAILS (expert-stream-device-slots.md §Gates, #2838, was #1299)	☐ blanket cpu_offload_gb, not expert-granular	☐	◐ mmap only
Hybrid CPU/GPU expert placement (routed-expert compute on the CPU, attention and dense layers on the GPU)	◐ Six architecture families. RunMoePlaced routes Qwen3-MoE, Qwen3.5/3.6, Nemotron-H, DeepSeek-V2, Kimi-Linear and dots3-note through one shared placement seam. The vllm_cpp.placement object in --offload-config maps llama.cpp's -ot, -cmoe, -ncmoe, and --fit controls to per-layer decisions. The CPU is the only placement target. Accelerator targets are refused. So is the fp4-resident expert arm, because its device residents are built eagerly at load, so placing it would upload every expert and then compute across the bus, which is a defect a token gate cannot see. Laguna and Gemma4 need different forward interfaces before they can use the seam, and the reasons differ: Laguna runs its expert GEMMs on the DEVICE but presents a per-token host-float FFN boundary, so it has no [T,H] block to hand the seam (#2050), while Gemma4's expert path is a void Expert...Accum accumulate shape rather than a -> DBuf one. DeepSeek-V4 runs its experts on the host from host weights, so a placement has nothing to move. dots3-note joined the seam at W5 (#699): Dots3NoteMoeBlock routes through RunMoePlaced, so the seam now carries six architecture families, though no dots3-note checkpoint fits any host this project reaches. GLM-5-Next and Kimi-K3 have no reachable MoE forward yet, and refuse by name. qwen4_exp acquired one at W5f (#2031) — its layer loop calls RunQwen4ExpMoeBlock, which enters the shared sparse-MoE seam — and it is NOW on the PLACEMENT seam too (#2424): the loop calls RunMoePlacedPair, with no placeable refusal owed because this architecture has no fp4 expert arm. and REACHED: the unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S arm (67.564 GiB) executes this call site on dgx:gpu0, and a placed-vs-unplaced comparison of the MoE block at each arm's FIRST placed layer agrees within the 5e-04 cross-device bar: cpu_moe 48/48 layers at NMSE 5.239e-06, --fit 23/48 at 3.569e-05, both reproduced across two consecutive runs. The earlier 369 GiB checkpoint was the wrong artifact; this one fits. The layer matters and is not a detail: comparing at layer 0 is VACUOUS for --fit, which places TRAILING layers, and comparing deep is CONFOUNDED by accumulated divergence across the layers between — only each arm's first placed layer has an identical input on both sides. Until #2314 the resolved plan was never INSTALLED — SetActiveMoePlacementPlan had no caller in src/, so the seam read a global nothing wrote and no expert was ever placed, while the loader still PRINTED the resolved placement on stderr. LoadedEngine::FromModelDir now installs it on both the GGUF and safetensors paths ahead of all weight I/O, and test_placement_reach holds that call site by mutation. A model has now run end to end with placement in force: Qwen3.6-35B-A3B bf16 on GB10, --offload-config '{"vllm_cpp":{"placement":{"cpu_moe":true}}}', announced as 40 layers run their routed experts on cpu, the rest on cuda, both arms exit 0 and both completions are coherent. They are NOT token-identical, and cannot be: the round trip that is byte-identical is the DATA MOVEMENT, whereas a placed layer computes its experts with the CPU MoE kernels instead of the CUDA ones, and this project's own cross-device bar for arithmetic is NMSE <= 5e-4 rather than bitwise equality (tests/vt/test_backend_cross_device.cpp). Greedy decode amplifies any perturbation inside that tolerance into a different token, so a placed-vs-unplaced TOKEN gate asserts something unachievable and llama.cpp's -ncmoe diverges the same way; the correctness gate owed is an NMSE bound on the MoE block outputs, not token equality. The speed gate needs a discrete CPU/GPU system (#149, #2026)	☐ CPU MoE kernels exist, but selection requires the whole model to use the CPU platform	☐ CPU selection also uses module-level platform checks, with no per-layer device override	✅ -cmoe and -ncmoe select buffer types from tensor-name patterns, so compute follows weight placement

Quantization and weight formats

Format	vllm.cpp	vLLM	SGLang	llama.cpp
NVFP4 (W4A4 and W4A16 Marlin)	✅	✅	✅	✅ in GGUF, not safetensors (#979). Was wrongly ☐: GGML_TYPE_NVFP4 = 40 (ggml.h:430), CUDA MMQ and the ModelOpt repacking converter are UPSTREAM at pin b10451, the sm_121a GEMMs fork-local
NVFP4 dense sinks take vLLM's dense Marlin, not the single-expert MoE route	✅ VT_MARLIN_DENSE (single projection, efa6e40d) + VT_MARLIN_DENSE_PAIR (fused shared-expert gate_up), both default-ON; the pair sink measured +1.31% at c8 / +1.38% at c4 on 35B-A3B, SACRED 315/315 + 235/235	☐	☐	☐
Dense W4A16 MLP runs ONE merged gate_up Marlin GEMM (vLLM's MergedColumnParallelLinear topology)	✅ VT_DENSE_MARLIN_GATEUP, default ON (opt out =0): the A/B measured +2.12% c1 / +1.70% c8 on the 27B, arms separated, tokens identical (#365). Replaces the split pair's 193 Marlin calls/step vs the oracle's 129	✅	☐	☐
NVFP4 shared-expert down_proj kept bf16 (no f32 round-trip)	✅ VT_SHARED_DOWN_BF16 default-ON; bit-identical (both consumers widen bf16 in-kernel and re-round on store), SACRED 315/315 + 235/235 on BOTH arms with unchanged assertion counts; +2.05% c8 / +0.79% c4 on 35B-A3B	☐	☐	☐
NVFP4 lm_head kept packed (no dequant at load)	✅ VT_LMHEAD_FP4 default-ON, #213; CUDA-gated on nvidia@0893e160 (continuations byte-identical packed vs dequant, 235/235; RSS -1.70 GiB on CUDA, owed a re-measure; a no-fp4-GEMM backend keeps one bf16 operand too)	✅	☐	☐
GGUF k-quants and i-quants	✅ (CPU grouped keep-quant MoE bf16 regression in b4f5610a fixed 2026-08-06). CPU quant compute is ISA-tiered: Arm has i8mm + repack; x86_64 portable-only, MEASURED open on every axis (CIQ G5, #433). IQ4_NL (20) and Q5_0 (6) added for qwen4exp (#1989), decode bit-exact vs llama.cpp b10451; they are the ragged-K landing spots of upstream's own tensor_type_fallback (IQ4_XS -> IQ4_NL, Q4_K -> Q5_0, llama-quant.cpp:374). IQ2_XS (17) and IQ4_XS (23) added for the staged unsloth/GLM-5.3-Flash-GGUF UD-Q2_K_XL arm (#2240), decode bit-exact vs llama.cpp b10451 over REAL bytes from that checkpoint; "UD-Q2_K_XL" names a target average and not a format, and 82 of that file's 1412 tensors are IQ2_XS against two that are Q2_K. Both keep their blocks on the GEMM arm too as of #2247, which added their keep-quant vec_dot rows against the Q8_K activation — IQ4_XS pairs with Q8_K and not with the Q8_0 of its codebook sibling IQ4_NL, because its block is a 256-element super-block — gated BIT FOR BIT against the oracle's own kernels on those same real checkpoint bytes. Both also have a CUDA keep-quant kernel as of #2260, which is what stops these two encodings draining a CUDA stream to the host cores on every expert GEMM and stops the fused MoeGateUpSwiGLUGrouped seam refusing them outright; the CUDA arm is gated against the same oracle numbers on the same real bytes, and against a stream capture, because a dtype that falls back to the host is correct, slow and invisible to every value comparison. That is what takes the artifact from 426.72 GiB resident to 101.14 GiB, measured by driving the production residency decision over all 1412 of its tensors. IQ4_NL, Q5_0 and Q4_0 have a CUDA keep-quant kernel as of #2419, on a SECOND templated GEMM beside the Q8_K one: these are 32-element encodings that dot a BlockQ8_0 activation, so the 256-element super-block GEMM the twelve other dtypes share cannot take them at all. Before it the dense and grouped CUDA seams drained the stream and ran the CPU kernel -- MEASURED on GB10 to SEGFAULT rather than merely to be slow, because Backend::Alloc is a plain cudaMalloc and the CPU kernel dereferences a device pointer -- and the fused MoeGateUpSwiGLUGrouped seam threw by name. This is the released unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S expert path specifically, and the encodings are forced rather than chosen -- its expert_feed_forward_length is 640, so every routed expert row is indivisible by 256 and no K-quant can encode it, which also means the ffn_down_exps tower's K = 640 is refused by the Q8_K seams whatever dtype it carries. The device dots are bit-exact against this tree's CPU arm per block; the discriminating gate is a STREAM CAPTURE, because a host fallback returns the same numbers and no value comparison can see it. MXFP4 (39) is the one 32-element encoding still falling back, and is owed. IQ3_S (21) added for the unsloth "UD" dynamic-quant family (#2510), decode bit-exact vs llama.cpp b10451 over REAL bytes from Qwen3.8-27B-UD-Q4_K_M.gguf; it was the ONE hole in the reader's i-quant run 16..23, and four tensors out of that file's 866 refused the whole 16.4 GB artifact while every other encoding in it was already handled. It is the one FILE encoding here that DECODES and does not DOT: a gather table keeps its blocks, a GEMM weight EXPANDS to bf16 on every device, and the keep-quant vec_dot is owed jointly with the CUDA WType arm, because landing the CPU half alone routes every CUDA IQ3_S GEMM through the host fallback that was measured to segfault on a discrete card. On the measured artifact that expansion is 146.13 MiB of blocks against 680.00 MiB of bf16, 3.4 % of a 15.33 GiB file. Q4_1 (3) and Q5_1 (7) remain absent, so a -Q5_K_M build of a ragged-K model still refuses	☐	☐	✅
GGUF gather tables kept QUANTIZED (one row dequantized per gathered token)	✅ CPU, qwen35/qwen35moe/qwen3next/qwen4exp. vt::Embedding takes a block-quantized table, decoding one row per id — a port of llama.cpp's ggml_compute_forward_get_rows_q. Without it a 51.2 G-parameter n-gram table expands from 28.8 GB of IQ4_NL to 102.4 GB of bf16, which no device here has (#1989). This is a residency DEFAULT CHANGE on already-shipped GGUF models, not only a new arm: an existing qwen35 file with a quantized token_embd now keeps it compressed on CPU where it used to expand. Tokens do not move (every GGUF-path gather writes a bf16 output and the bf16 round is idempotent over the old expand-then-widen), so the change is memory-only today. deepseek4 and laguna are NOT reached: both consume token_embd as a flat host f32 array, so their loaders narrow the policy for that tensor and keep expanding it. A gather's admission is the ROW DECODER and not the vec_dot, so IQ2_XS and IQ4_XS joined this set the moment their decoders landed (#2240), a wave before they could keep on the GEMM arm; Q8_K is now the only encoding that still tells the two admission rules apart. CUDA joined it (KGATHER, spec). EmbeddingKernelCuda decodes a block row across all 19 encodings vt::cpu::BlockToFloat decodes (IQ3_S joined with #2510), so a gather table stays block-resident ON THE CARD. The gate is a REGISTRY QUERY and not a device list: vt::Embedding routes a block table to OpId::kEmbeddingQuant and DeviceQuantGatherSupported is OpRegistered(kEmbeddingQuant, dev), so a backend advertises the capability by registering the kernel and a device with no decoder refuses BY NAME. MEASURED ON A GPU, nvcc 13.0.88, 2026-08-31, and the two boxes measured DIFFERENT things: thor:gpu0 (Jetson Thor, sm_110) ran the branch HEAD through the production registration — 7 of 7 cases at 250 assertions, every encoding bit-exact against the CPU arm in f32 and bf16 out with i32 and i64 ids, against a pre-arm RED leg that failed by name at 32 assertions, and deleting the single production registration reds 5 of 7 cases, so the arm is reached and not merely present. dgx:gpu0 (NVIDIA GB10, sm_121a) then ran the SAME head leg green — 7 of 7 at 250, M5 red, every verdict matching thor — so the arm is proven on BOTH architectures and #2393 is closed. This is what makes a GPU arm of qwen4exp possible at all — its n-gram table is 26.822 GiB of IQ4_NL against 95.368 GiB expanded, on a box with ~119.6 GiB for everything — and it is a shape llama.cpp does NOT have: its CUDA get_rows dispatches the legacy quants only and aborts on every K-quant and every IQ type, so the llama.cpp column overstates the denominator for the encodings that matter here. Like the CPU arm this is a residency DEFAULT CHANGE on already-shipped GGUF models, and it reaches more than qwen4exp: a token_embd in a block encoding now keeps its blocks on CUDA where it used to expand, qwen35 GGUFs included. For a bf16 output the values do not move (both routes decode with the same scalar expressions and round once); for an f32 output the new path is strictly MORE precise than expand-then-widen, which is a behaviour change on already-shipped models and not merely a memory one. **Owed: the gather arms of METAL, VULKAN, ROCM and TENSTORRENT (#2394), which still refuse because each of their gather kernels asserts a float table by name. No throughput number is claimed	☐	☐	✅ get_rows for ~20 types, CPU and CUDA
GGUF F16 weights kept resident as F16 (no BF16 promotion)	✅ VT_GGUF_KEEP_F16 default-ON (CPU), the f16 GEMM computes on it directly. Default settled 2026-08-17, a memory-for-speed trade: 1.05 GiB less peak RSS for ~9% prefill and ~1.4% decode, tokens identical. 0 opts out	☐	☐	✅ ggml_vec_dot_f16
GGUF is a TWO-engine comparison at these pins (#979)	✅ text-only qwen35, no clip projector (#821)	☐ REMOVED from the tree in 6635279d8, now an unpinned out-of-tree vllm-gguf-plugin	☐ full stack present, qwen3_5 unreachable behind FOUR blockers, and the load path has NO completeness guard so a clean-looking load proves nothing	✅ native, LLM_ARCH_QWEN35
EXL3 trellis (exllamav3; codebooks 0 (3INST), 1 (MCG) and 2 (mul1), Hadamard-128 + sign vectors, NO scales)	◐ A stock EXL3 checkpoint GENERATES: turboderp/Llama-3.2-1B-Instruct-exl3 @ 3.0bpw loads through the shared dense container and emits coherent text from vllm-cli on a CUDA queue and on a CPU queue, which reaches LlamaForCausalLM and Qwen3-dense together. The scheme sits on vLLM's own LinearMethodBase seam, so it is no longer a DeepSeek-V4-private arm. bits and the codebook are both read PER TENSOR — the published 3.0bpw Llama has a 3-bit body and a 6-bit head, and ships no mcg marker, which means codebook 0 and not MCG. The DeepSeek-V4 arm is separate and unchanged: its rank-sliced SparkInfer artifact loads and executes end to end on a synthetic checkpoint. The rank-sliced 0xSero/deepseek-v4-flash-0731-spark routed-expert tower coalesces TP4->TP1 at load, the carried-* half (block-wise FP8 + BF16 + F32) is dequantized into the host-float tower the forward composes with, and DeepseekV4Model::Forward runs the whole model, dispatching one vt::Exl3MoeMlp per MoE layer over the routed experts (bits == 3, mcg codebook). The REAL artifact still does not run: its DSA compressor and indexer tensors are twice the width the host forward indexes and the loader refuses them by name, and its tokenizer is not read (#1924). The device half is now PARTLY VERIFIED (GB10 sm_121a, 2026-08-28): had_r_128 is BYTE-IDENTICAL CUDA-vs-CPU, exl3_gemm matches the f64 reference at rel_rms 5.538e-4 against a 1.0e-3 bound, and the m<=8 GEMV meets tier 3c at 5.160e-4 against 6.0e-3. What is STILL unverified on a device is the FUSED MoE arm, which cannot run on this code at all: it needs a device-resident tower and CudaBackend::DeviceMemoryIsHostAddressable() is false by design, so the routed-expert path executes on a CPU queue today. No speed number is claimed on any axis. Codebook 2 (mul1) and the 4- and 5-bit widths are now implemented so that Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw -- 409 trellis modules, 137 at 3 bpw, 270 at 4, one at 5 and one at 6, every one of them mul1-marked -- is no longer refused by name (#2495). The count read 272 until 2026-09-02 and was an UNDERCOUNT that omitted the 137 bits-3 modules, which is why (3, 2) was missing from the CUDA instantiation list and every MLP projection of that target refused on a CUDA queue. (3, 2) is now instantiated and gated on dgx:gpu0 (GB10, driver 580.173.02), agreeing with the CPU decoder at rel_rms 3.02544e-07 against a 1.0e-3 bound (#2574). cb 2 is a DIFFERENT decode and not a third multiplier: it sums the four bytes of the product into an fp16 bit pattern and maps it with a fused fp16 affine, and the host arm is gated against hand-computed upstream values rather than against itself. A ROCm arm now exists too (#2433, spec): Exl3Gemm on gfx1151 is a transcription of the PORTABLE CPU reference rather than a port of the CUDA kernel -- 90 KiB of requested shared memory does not fit a 64 KiB AMD LDS, and m16n8k16, ldmatrix and cp.async have no AMD form -- and because it uses no matrix cores it is gated at BYTE equality with the CPU arm on all three codebooks, not at a tolerance. It covers every width the host arm decodes, so the CUDA instantiation list below does not bound it. On the CUDA device arm (3,0), (3,1), (6,0), (3,2), (4,2), (5,2) and (6,2) are instantiated -- (3,2) is the width the undercount cost, added and device-gated by #2574; every other width and codebook and the device-resident tower are owed (spec). The m<=8 GEMV is a SEPARATE arm set from that list and a separate row (QUANT-EXL3-PERF, #2570): upstream instantiates (4,0) (4,1) (4,2) (2,1) (2,2) (3,1) (3,2) and this tree now carries (3,1), (3,2) and (4,2), so 407 of that checkpoint's 409 trellis modules reach the small-m fast path where before NO module of it could -- the only arm was (3,1) and the artifact contains zero (3,1) tensors. (3,2) was one template argument; (4,2), the 270-module arm, was a KERNEL port: a per-width LSTRIDE, two 24-lane guards that exist only because a 3-bit tile is 24 uint32 rather than 32, and a separate dq8_regs_4bits window read. (4,0), (4,1) and the whole 2 bpw arm stay OWED -- no artifact in this tree carries a 4-bit tensor at either of the first two codebooks, and nothing builds them, so that the kernel would compile for them is an inference and not a measurement. Instantiating an arm is necessary and not sufficient: Exl3GemvSelectConfig is upstream's envelope verbatim and DECLINES, and on Blackwell EVERY shape of this checkpoint, at both widths, is admitted only where size_n / 32 <= narrow_coresident, an occupancy query. The wide config is not an escape from it: its band needs size_k <= 4096 and the smallest 4-bit k here is 5120. The thresholds are >= 544 and >= 160 at the two bits-3 shapes and 32, 160, 192, 320, 384 and 544 across the six bits-4 ones, all pinned in tests/vt/test_exl3_gemv.cpp from both sides. The throughput effect is now MEASURED on GB10, and it is a NULL (#2570): three interleaved rounds on one binary read 17.20/17.18/17.09 tok/s with the arm off against 17.22/17.12/17.09 with it on, a 1.3% spread and no separation. nsys --cuda-graph-trace=node shows why, and it is a DECLINE and not an ineffective kernel: exl3_gemv_kernel appears 0 times at the default and 4 times only when forced. GB10 reports SM_COUNT=48 MAX_THREADS_PER_SM=1536, so the narrow config's 512-thread blocks ceiling narrow_coresident at 144, and only shapes needing <= 144 are admitted -- 34 of the checkpoint's 409 modules, 0.75% of its trellis bytes. The arms are correct, upstream-faithful and device-gated; on THIS part the envelope declines them. m == 8 is unmeasured, because the draft legs did not run. The FUSED MoE mgemm is a SEPARATE arm set from that list, and upstream's own bound on it is narrower: exl3_moe.cu:184 admits mcg and mul1 and refuses 3INST, so codebook 0 is not a MoE arm upstream either. It now instantiates bits 3, 4, 5 and 6 at codebook 1 where it carried bits 3 alone (#2756); before that a mcg expert tower at any other width could not run on a CUDA queue by ANY path, because the fused arm is default-ON and its refusal is not caught, and the VT_DSV4_EXL3_FUSED_MOE=0 rollback lands on a GEMM arm set with no (4,1), (5,1) or (6,1) either. Codebook 2 on the fused arm is owed as a LOADER slice and not a kernel one: five sites pin the codebook to 1 before the kernel is chosen, so a cb-2 instantiation added alone would be dead code. Its GDN linear-attention tower now loads and runs too (#2495 item 4, spec): 48 of that model's 64 layers are linear_attention, and the three projections of each were refused by name because nothing in the GDN forward consumed an Exl3Weight. The arm holds THREE trellises rather than one merged in_proj_qkvz owner, because the artifact ships in_proj_qkv and in_proj_z as two independently quantized tensors whose sign vectors are fitted per projection, so the bf16 merge has no trellis analogue. It also accepts the F16 in_proj_a/in_proj_b the artifact stores beside its BF16 conv1d, norm, A_log and dt_bias. The DFlash2 DRAFT loader now has an EXL3 arm too (#2495 item 7, spec): Mia-AiLab/Qwen3.8-27B-DFlash2-EXL3-5.0bpw ships all 36 of its modules at bits 5 with a mul1 marker, and LoadQwen3DFlash died on fc.weight, which an EXL3 module does not ship. The arm holds SEVEN trellises per layer rather than the merged qkv_proj and gate_up_proj owners the bf16 reader builds, for the same reason the GDN tower holds three. It also accepts the F16 candidate_selector.hidden_projection and *_conv.kernel_projection the repack left unquantized beside their BF16 base_kernel and codebooks — a second refusal that was invisible until the first was gone. Its mtp.* draft head loads and runs too, and all THREE readers of the target's trellis head now compute with it packed (#2495 items 5 and 6, spec): the MTP loader refused nine quantized mtp.* tensors as "expected BF16", and the MTP draft's ComputeLogits plus the DFlash/DSpark draft's SHARED head read knew only bf16 and NVFP4. The head is COMPUTED WITH rather than widened, which is what upstream does without a branch because its head is an nn.Module and _apply_head calls lm_head.quant_method.apply; a dequantized copy of the real 248320x5120 head would be 2.543 GB against 0.953 GB packed. The shared read also used to fall off the end of its bf16 loop and return SILENTLY with every owner empty when no arm matched, which is exactly what this target does -- it ships no lm_head.weight at all -- and it now refuses by name (#2569). The checkpoint's real inventory is 409 quantized modules, not the 272 this row previously recorded, measured from both downloaded shards' headers: bits 3 x137, 4 x270, 5 x1, 6 x1, every one mul1, and no tensor in the file has scale in its name. (bits 3, codebook 2) is now instantiated on the CUDA GEMM and gated on a device, so those 137 modules no longer refuse by name (#2574, owned by QUANT-EXL3-MUL1, not by this row). THE 27B CHECKPOINT NOW RUNS ON CUDA AND GENERATES. On dgx:gpu0 (GB10, driver 580.173.02, sm_121a) the published Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw loads and emits Paris. The capital of Germany is Berlin. The capital of Italy is at rc=0, with both shard sha256 values recomputed on the device. Warm decode is 16.7 tok/s target-only and 48.7 tok/s with its DFlash2 EXL3 draft at k=7, interleaved on one binary, and the two arms are token-identical (qwen38-27b-exl3-gb10). That is NOT a reproduction of the upstream README's 47.5 tok/s, and must not be quoted beside it: theirs is HumanEval-style at T=0.6 with acceptance 4.43, this is greedy T=0 on a prompt that continues into a list of capitals, the context is 8192 rather than 262144 with no NVFP4 KV cache, and the shipped paged draft route had to be disabled (#2274). A matched HumanEval-style run at T=0.6 with acceptance reported is owed	☐ no EXL3 at the parity pin	☐	☐
AWQ	◐ CPU dequant	✅	✅	☐
GPTQ	◐ CPU dequant	✅	✅	☐
MXFP4 compressed-tensors	◐ W4A16 Marlin, mem 2.63x less. gate_up FUSION + decode-graph default-ON; #44 3/3, 32B 6/6. VT_MARLIN_DENSE DEFAULT-ON (KERNEL-MARLIN-DENSE-EXEC): dense marlin 48-CTA, byte-faithful, beats MoE (c8 0.969)	✅	✅	☐
Compressed-tensors mixed-precision (config_groups, ordered regex targets)	◐ scheme read from the config, never from a dtype probe; an arm with no loader is REFUSED BY NAME. On unsloth/Qwen3.8-27B-NVFP4 the W4A4 group loads, the FP8 W8A8 group and the kv_cache_scheme are refused (#821)	✅	✅	☐
ModelOpt MIXED_PRECISION (quantized_layers, exact module names)	◐ the declared quant_algo is CROSS-CHECKED against the shipped operand names at load; a disagreement, an unimplemented algo or an unknown operand family is REFUSED BY NAME. r0b0tlab/...-MTP-sm121 loads (#821)	✅	✅	☐
ModelOpt MIXED_PRECISION: the cross-check moves no arm	✅ nvidia/Qwen3.6-27B-NVFP4 @0893e160 is NOT refused, and one case says so: same format, wildcard exclusions, an input_scale on every NVFP4 module, a kv_cache_scheme shipping zero scales (#821)	✅	✅	☐
ModelOpt MIXED_PRECISION: STATIC per-tensor FP8 + W4A16_NVFP4	◐ LOADS, never RUN. r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121 carries 208 input_scale, F32 scalar weight_scale and "dynamic": false, so the three unsloth blockers are absent (#1574, #821)	✅	✅	☐
fp8 weights, per-tensor scale	✅	✅	✅	☐
Block-wise (fine-grained 128x128) FP8, the weight_scale_inv layout	◐ RUNS on CPU and on CUDA sm120 (#1189 M4/M6): 10 projections as 7 GEMMs, gate_up+QKV merged; N,K %128==0 only. TOKEN GATE PASSED on Qwen/Qwen3.8-27B-FP8 vs the pinned oracle on GB10, 7 prompts x 16 tokens greedy, 6 strict + 1 in-band, weights resident at 1 byte/element; no speed number (gate, kernel)	✅	✅	☐
Per-tensor FP8 W8A8 linear is a shared seam any model can bind	✅ models/dense_fp8_gemm.h + layers::Fp8W8A8LinearMethod (#940), bound via layers::MakeLinearMethod. One definition, CUDA only (spec)	✅ Fp8LinearMethod	✅	☐
FP8 W8A8 works on a CUDA arch without cutlass-fp8	✅ vt::QuantFp8Static registers from an unconditional TU (#960); sm_110 measured (spec)	✅	✅	☐
fp8-tower GDN in_proj emits bf16, unlocking packed GDN decode	◐ VT_GDN_FP8_IN_BF16 + VT_GDN_PACKED_DECODE_FP8_TOWER (inert alone), both default OFF, ungated (#339) (spec)	✅ bf16 out_dtype	☐	☐
Merged fp8 projection folds per-column alpha in the GEMM epilogue	◐ VT_FP8_ALPHA_VEC_EPILOGUE, CUDA only, default off, ungated; refuses split-K under a bf16-D equivalence claim (claims_splitk1_premise, default off)	n/a	n/a	n/a
Per-tensor FP8 CUTLASS GEMM picks a decode-sized tile at decode-sized M	◐ VT_FP8_CUTLASS_SMALL_M (default on, =0 rolls back), CUDA sm_12{0,1}a only, UNGATED on a device: upstream's four-way sm120 M ladder (M<=16 16x64x128 with a 16x32 EpilogueTile, M<=32 32x64x128 with 32x32, M<=256 64x64x128, else 128x128x128) instead of the two-way ladder this tree shipped, where every M<=256 took the 64-row tile (#1866). Reached on VT_DENSE_CUBLASLT_FP8=0 and as the cuBLASLt lane's no-heuristic fallback; the DEFAULT fp8 arm is still cuBLASLt, so a stock decode step does not take it. Ladder gated by value on the host tier; the CUDA compile is measured green by the cuda-fat-build job on head d9bf525c0 (run 32802716762, whose own conclusion is cancelled for an unrelated cancelled leg), and no compiled CUDA has changed since, while the current head's own cuda-fat-build is PENDING; the token gates and the arm A/B are owed (spec)	✅ cutlass_gemm_sm120_fp8_dispatch	☐	☐
vt::MulColVecF32 carries a bf16 store width	✅ f32 arm byte-identical; bf16 arm rounds once; CPU + CUDA	n/a	☐	☐
bf16 / fp16	✅	✅	✅	✅
Safetensors direct load, no conversion	✅ at ANY tensor byte offset: the format aligns nothing, so no loader forms a typed pointer into the mapping. Last three fixed by #772; a checker is still owed on #627	✅	✅	☐
Weights uploaded straight from the file mapping (no host copy first)	◐ verbatim tensors only (37.8% of 27B BF16); arbitrary-offset reads are defined, including Laguna graph staging. Merged/transposed and merged FP4 weights still copy	✅	✅	✅ mmap

Model coverage

The supported set is exactly what the C++ registry registers: every architecture self-registers via REGISTER_VLLM_MODEL, and scripts/check-supported-models.py gates this list against the source so it cannot drift. Today that is 43 registered architectures. Each row names the checkpoint it was gated against and the verdict; caveats are in Project status, agent detail in .agents/model-matrix.md. A mergeable gate/up MLP routes through one shared merged-GEMM method, so a tuned arm added once reaches every such arch; Command-R, GLM-4, MiniCPM, MiniCPM3 and Phi-3 joined on 2026-08-10 (#299), and scripts/merged-gemm-consistency-allowlist.txt lists the rest with their blocker.

Gate words: strict is token-for-token identical to the vLLM oracle; near-tie is the ratified distributional gate used where vLLM's own greedy is bf16-non-deterministic; scaffold means registered and config/loader-gated but the forward is not yet a real-checkpoint run. Speed is a separate bar (match or beat the reference on every axis); most rows are correctness-complete and speed-pending, which BENCHMARKS.md tracks.

Registered architectures

Architecture	Tested checkpoint(s)	Correctness gate	Speed vs reference
Qwen3_5ForConditionalGeneration	Qwen3.6-27B NVFP4 (unsloth @890bdef7, nvidia @0893e160); Qwen3.5-4B BF16; Qwen3.8-27B BF16 @1d4bf0f2	27B strict 235/235 text + 32/32 image/video; 4B cached 3/3; Qwen3.8-27B 4/7 strict, 3 exact fp32 ties in band (#915)	unsloth 27B at/above vLLM, ModelOpt 0.85x; 4B 1.021x; 3.8-27B c4 0.963x, c1/c8 absolutes (#915). Loads BF16/per-tensor FP8/NVFP4 (CT+ModelOpt); modelopt_mixed FP8 tower NATIVE (#164), GDN qkvz merged. CUDA/CPU
Qwen3_5MoeForConditionalGeneration	Qwen3.6-35B-A3B (NVFP4 text; published BF16 text + vision tower)	NVFP4 strict 315/315 vs vLLM 0.25.0; published BF16 6/7 prompts strict 16/16 vs the pin, 7th an exact tie (#910). Image/video IMPLEMENTED, NOT GATED (#891): the tower loads and runs, mm gate OWED	gate model: 0.93x to 1.03x grid; NO BF16 or mm speed claim
Qwen4ExpForConditionalGeneration	GGUF (qwen4exp) — LOADS, DECODES, AND SERVES ONE SEQUENCE AT A TIME ON --device cpu (W5a, W5f, W5g, W5k, W5L, #2031)	THE GGUF LOAD IS GATED, AND SO IS THE LAYER LOOP; WHAT DOES NOT EXIST IS A SECOND STEP. A qwen4exp file reaches the architecture's own config builder through the GGUF dispatch, the registry resolves the class, and load_weights materializes the whole text tower — **on CPU and, since KGATHER, on CUDA (spec). The 51.2 G-parameter n-gram table would expand from 26.822 GiB of IQ4_NL to 95.368 GiB of bf16 on a device that cannot gather blocks, which the on-disk #1123 device-fit guard cannot see, so the load REFUSES BY NAME on such a device ahead of any tensor I/O (#2083). DeviceQuantGatherSupported is no longer a device list at all: it asks OpRegistered(kEmbeddingQuant, dev), CPU and CUDA register that op, and METAL, VULKAN, ROCM and TENSTORRENT do not and are still refused (#2394). So this row's earlier clause that the gate 'is true for CPU alone' is FALSE and is replaced — every convert-time transform inverted (the +1 fold on every norm gamma with ssm_norm the one exception, ssm_a back to log(-x), and the V-head reorder on every Gated DeltaNet tensor), gated in both directions against a committed 1224-tensor manifest of the shipped unsloth/Qwen3.8-Flash-Next-GGUF UD-IQ1_S and value-wise against a synthetic file. Qwen4ExpTextModel::Forward now exists (W5f, #2336) and ModelRegistry::Forward reaches it: the 48-layer loop composes all four block seams and is gated END TO END against the lane-pinned transformers 5.16.0 oracle at a tiny config, max	diff
Qwen3_5ForCausalLM, Qwen3_5MoeForCausalLM	none: no text-only Qwen3.5 checkpoint fits this hardware	NO RUN GATE, OWED. Gated on test_qwen3_8_text_only.cpp; NO token claim. Loader reads stacked BF16 experts (#740) plus BF16 towers, shared expert and lm_head (#864), so both published indices satisfy the load plan	not measured
Qwen3ForCausalLM	Qwen3 dense 0.6B/1.7B/4B/32B, NVFP4A16	near-tie strict 16/16 vs vLLM 0.25.0	c1 every-axis parity, c8 decode residual
Qwen3MoeForCausalLM	Qwen3-Coder-30B-A3B	strict 6/6 vs vLLM 0.25.0	11/16 grid cells at or above graphed vLLM
Qwen3VLForConditionalGeneration	Qwen3-VL-4B-Instruct (image + video)	image strict 32/32, video near-tie vs vLLM 0.25.0	vision tower 0.57x vs vLLM encode; umbrella pending
LlamaForCausalLM, InternLM3ForCausalLM	Llama-3.2-1B, 01-ai/Yi-Coder-1.5B-Chat, internlm3-8b-instruct	strict 16/16 each vs vLLM 0.25.0	pending
InternLM2ForCausalLM	internlm2-chat-1_8b	near-tie 16/16 vs vLLM 0.25.0	pending
MistralForCausalLM	Mistral-7B-v0.3	strict 16/16 vs vLLM 0.25.0	pending
OPTForCausalLM	facebook/opt-125m	strict 6/6 vs vLLM 0.25.0	pending
PhiForCausalLM	microsoft/phi-2	near-tie 16/16 vs vLLM 0.25.0	pending
Phi3ForCausalLM	microsoft/phi-4 (14B), Phi-3	strict 16/16 vs vLLM 0.25.0	pending
GemmaForCausalLM	google/gemma-1.1-2b-it, unsloth/gemma-2b	near-tie 48/48 vs vLLM 0.25.0	pending
Gemma2ForCausalLM	google/gemma-2-2b-it	near-tie 48/48 vs vLLM 0.25.0	pending
Gemma3ForCausalLM	google/gemma-3-1b-it	strict 48/48 vs vLLM 0.25.0	pending
Gemma4ForConditionalGeneration	Gemma-4 multimodal (unsloth/gemma-4-E4B-it)	text strict, image mm near-tie; audio pending	pending
Gemma4UnifiedForConditionalGeneration	Gemma-4 "unified" HF export (google/gemma-4-12B-it), no-PLE dense layout	shares the Gemma-4 text+mm forward; loads on the same factory (contributor #140); no separate oracle gate for this arch name yet	pending
GraniteForCausalLM	ibm-granite/granite-3.3-2b-instruct	near-tie 16/16 vs vLLM 0.25.0	pending
StableLmForCausalLM	stabilityai/stablelm-2-1_6b	near-tie 16/16 vs vLLM 0.25.0	pending
MiniCPMForCausalLM	openbmb/MiniCPM-2B-sft-bf16	strict 16/16 vs vLLM 0.25.0	pending
MiniCPM3ForCausalLM	openbmb/MiniCPM3-4B (MLA)	near-tie 16/16 vs vLLM 0.25.0	pending
Olmo2ForCausalLM, Olmo3ForCausalLM	allenai/OLMo-2-0425-1B; OLMo-3 (Olmo2 factory alias)	OLMo-2 strict 16/16; OLMo-3 oracle-blocked (vLLM 0.25.0 cannot build it)	pending
DeepseekV2ForCausalLM	DeepSeek-V2-Lite (MLA)	SACRED 8/8 cases vs vLLM 0.25.0 — 5/8 strict, 3/8 near-tie, 92/128 tokens strictly exact, max gap 0.25 nats (arch-fusion-fold-plan-2026-07-30.md §A2+A5). The 8/8 counts cases that PASS, including the near-ties; it is not 8/8 strict	speed short, attributed
DeepseekV4ForCausalLM	DeepSeek-V4-Flash GGUF (ds4 q2-imatrix, UD-IQ2); the SAFETENSORS arms now get past the tokenizer (#1924)	coherent near-tie vs ds4 oracle (vLLM cannot fit one GB10). Tokenizer ids are exact vs HF tokenizers on the checkpoint's own 6.4 MB tokenizer.json, and the GGUF arm's joyai-llm pre no longer resolves to an APPROXIMATION	decode beats ds4 1.144x, default on, via the deepseek-v4-gen CLI; the registered engine publishes DeepSeek-V4's real seven-group / 167-entry cache topology (#1973) and the runner now ALLOCATES all 167 of them (#2068), handing them to the forward keyed by the name each was published under; the FORWARD then refuses, because no registered forward consumes a cache set keyed that way yet (W5). At the default --block-size 32 a run reads the factory's own refusal first, since a compress-ratio-128 page needs 128 or 256. So the engine still cannot serve, one seam further along than it was
Glm4ForCausalLM	GLM-4-9B-0414	near-tie 16/16 vs vLLM 0.25.0	pending
Glm4MoeLiteForCausalLM	zai-org/GLM-4.7-Flash (31.2B, MLA MoE)	⚠️ the shipped 8/8 is not a measurement. Against the committed vLLM 0.25.0 golden the tokens match 69/128 positions (53.9%), 1 of 8 prompts exact. The only vs-vLLM assertion tests gap > 500 mnats against neartie_gap_mnats.npy, which is identically zero at all 128 positions, so it cannot fail (#2839)	pending
Glm5NextForConditionalGeneration	GGUF: unsloth/GLM-5.3-Flash-GGUF UD-Q2_K_XL @ d425e572f, 101.2535 GiB in four shards — LOADS on --device cpu, and the engine's multi-KV guard no longer refuses above the model's forward (W5b-2c, #2348; W5c, #2242; KV-cache spec + MoE W5, #2223; the forward W5b-2b, #2337). MEASURED on dgx:gpu0 2026-08-30 (#2343): all four shards load and the engine sizes its caches -- max_model_len auto-fits from 1048576 to 8192 against 256 blocks of 32 tokens, and max_num_seqs drops from 32 to 1 because one 4,390,912-byte GDN state fills a unified page -- in under 26 minutes wall, which is a DURATION and not a throughput number. At THAT change the first step then threw at the multi_kv guard at the TOP of ModelRegistry::Forward; W5b-2c (#2348) is the consuming forward that guard was waiting for, so it no longer fires for this model. THIS ARTIFACT GENERATES COHERENT TEXT ON dgx:gpu0 as of #2241. MEASURED 2026-08-30 on GB10 in the SHIPPED configuration (no diagnostic env set), vllm-cli --device cpu --max-tokens 2 at the prompt The capital of France is: it emits Paris., rc=0, prompt_tokens=5 completion_tokens=2, with peak RSS 104,792,300 kB = 99.94 GiB (VmHWM, polled) against the 99.47 GiB the broken binary read on the same box. The mechanism and the per-layer bisect come from two instrumented thor:gpu0 runs the same day, where four tokens read Paris. Paris is, the prefill top-5 is Paris (16.427) ahead of one, located, known and a at a margin of 1.279, and no NaN appears in any of 180 per-layer readings across four steps. The first generation attempt, at W5b-2c on dgx:gpu0, emitted token id 0 eight times because the loader repacked all 346 of this file's q8_0 tensors into the block_q8_0x4 i8mm interleave while the host bridge decoded them as plain blocks -- see the spec's ## Owed O30 for the bisect. NO SPEED NUMBER IS CLAIMED and the earlier ones are VOID: the 73 s/token of the W5b-2c run came from an all-NaN forward whose degenerate expert selection is not this model's work. the GB10 arm is the one measured above	THE WEIGHT TOWER IS PORTED AND THE FORWARD READS THE ENGINE'S PAGED CACHES. The config resolves and validates against transformers v5.16.1, the only revision of any admissible oracle that implements glm5_next — vLLM implements it at NO revision, and vllm#53906 is open and therefore inadmissible. All five upstream validate_architecture rejections are implemented, and both sources — a config.json and a GGUF — descend through ONE parser. The GGUF arm of load_weights now returns a real LoadedModel: the KDA layer with its three separate depthwise convs, the NoPE MLA with the two SPLIT absorbed halves, the DSA k-pool indexer, the flat mHC pair, the 288 stacked routed experts plus one shared, and the dense MLP on the leading three layers. The name map is gated against the REAL 1412-tensor artifact with no asset, in both directions, and blk.45 — the multi-token-prediction block — is read, counted and NOT built as a decoder layer. ModelRegistry::Forward DISPATCHES to the model as of W5b-2b (#2337), which is what discharges the six "gated but reached by nothing" debts this row carried, and the ENGINE path now REACHES that dispatch as of W5b-2c (#2348), which writes the forward the multi_kv guard at the top of the same function was waiting for (#2343, #2068): each DSA layer's MLA latent and indexer side cache are found BY NAME on MultiKvCacheIndex and read out of the engine's own pages, the 34 KDA states come off gdn_state positionally because that channel carries no names, and each step's new rows are written back into those pages rather than kept on the model. The tower stays block-resident exactly as loaded, ONE decoder layer at a time is bridged to host f32 and dropped, and only the 8 of 288 experts a token selects are decoded. That is arithmetic and not preference — a float tower is 426.72 GiB and the 42 sparse layers' expert banks alone are 1,134 GiB, against ~119.63 GiB usable on the largest box this project reaches, while the streamed forward's f32 peak is under 0.75 GiB. The vision tower and the safetensors arm still REFUSE BY NAME, as do a multi-request step (this forward is single-sequence and ragged batching is owed) and a non-CPU queue (every primitive here is a host f32 reference and the device arm is owed). The KV-cache spec no longer does (W5, #2223): make_kv_cache publishes three groups -- an MLAAttentionSpec at head 512 for the 11 DSA layers, ONE MambaSpec for the 34 KDA layers, and a second MLAAttentionSpec at head 257 for the indexer side cache -- and it is REACHED through the production factory hook. W5's 288+1 expert MoE block (glm5_next_moe) is now REACHED by the forward, along with W2's KDA arm, W3's DSA indexer, W4's mHC bricks and W5b-1's attention; deleting the production call site in the registry hook reds the focused gate. Use --device cpu, and as of #2260 the reason is no longer the quantization: the artifact's 82 IQ2_XS and 3 IQ4_XS tensors now HAVE a CUDA keep-quant kernel, so the expert GEMM no longer drains the stream to the host — which was measured on GB10 to SEGFAULT, not merely to be slow, whenever the tensors came from the ordinary CUDA device allocator — and the fused MoE seam no longer throws — and W9c-3a (#2464) then tried to spend that discharge and FAILED, so --device cuda is still not a path to use. The routed-expert device arm exists and is CUDA-gated at the unit level (NMSE 3.833e-15 against the host arm on dgx:gpu0), but driven end to end on the 101.24 GiB artifact BOTH --device cuda legs died with SIGSEGV (rc=139) emitting no token, reproducibly (spec O46). The cause is now diagnosed and it is NOT the mixed residency O46 inferred: the forward dies in StoreCaches, host-storing into the runner's cudaMalloc KV pages after the forward has already returned, which is a defect older than this arm and unreachable while the non-CPU refusal above it stood -- spec O49 and #2480, which owns the fix. The split is therefore OPT-IN and defaults OFF (VT_GLM5_NEXT_DEVICE_EXPERTS=1, for debugging that crash and not for serving), and the default behaviour of --device cuda is the refusal it always was. Use --device cpu, which emits Paris. on that artifact. State that precisely, because it is what is measured and no more -- no token has yet come out of this model on a GPU, and none is claimed; the end-to-end --device cuda leg on the 101.24 GiB artifact is queued on dgx:gpu0 and an untaken device gate is PENDING, never a pass. What IS gated, on x86_64: the forward admits a CUDA-typed queue instead of throwing, and the routed-expert arm's device path -- residency, operand construction, arm selection, the fit guard -- runs and agrees bit-for-bit with the host arm on a CPU-backed Dev. The forward SPLITS its queue -- it interposes a CPU queue for the host-reference arms and hands the caller's device to exactly ONE consumer, the routed-expert keep-quant GEMM, whose banks dense_attn::ResidentWeight uploads once per model and keeps in their blocks. READ THAT AS ONE ARM OF ELEVEN. The KDA recurrence, the DSA k-pool indexer, the eager MLA attention, both mHC sites, the router, the combine, the dense and shared MLPs, the embedding gather and the chunked lm_head ALL STILL RUN ON THE HOST, and the row's spec records that as O43 with #2410 owning the rest; the remaining port is priced at 2,500-3,500 lines off the two siblings that carry a device arm (kimi_linear_device.cpp 2,539, nemotron_h_device.cpp 2,144). A device that is neither CPU nor CUDA is still refused by name, and so is a CUDA queue in a build with no CUDA backend. A one-line stderr announcement names the device the expert GEMM ran on, because the two arms compute the same block and no logit can say which ran. NO end-to-end token gate exists or can exist on this fleet and that is a measured fact, not a schedule: no oracle registers this architecture at any revision it can also RUN here	none, and no speed claim is admissible from this row until a correctness gate exists
GlmMoeDsaForCausalLM	zai-org/GLM-5.3 (753.33B, DSA sparse MLA MoE), HF revision 935644c05e76fc198714f4cca449fd8b970ff6d7 — IT LOADS AND IT GENERATES TEXT, on two different backends, as of 2026-09-02. The architecture resolves, its config parses from a config.json and from a glm-dsa GGUF header through one validator, and the glm-dsa row of the GGUF dispatch table is reached from LoadedEngine::FromModelDir (W2, #2214). ROCm gfx1151 (strix:gpu0), unsloth/GLM-5.3-GGUF UD-IQ1_S, 201.83 GiB in six shards: the load completes with 78 layers' routed experts placed on the host CPU and all 228 expert towers kept COMPRESSED, 11.620 GiB paged in, and vllm-cli --prompt 'The capital of France is' --max-tokens 4 --temperature 0 emits Paris, which is at rc=0, reproduced n=2 on a separately rebuilt binary. That needed the placed-expert residency fix (#2516, #2517) and the ROCm kFusedNormRope kernel (#2564) — the ONE op blocking generation, because it is the only one on this path consulted through vt::OpRegistered BEFORE the call, so a miss was a refusal rather than a slow path. NO SPEED NUMBER IS ADMISSIBLE from that arm and none is claimed. That run observed five ops on the portable CPU reference tier, which docs/ROCM.md disqualifies: ConcatAndCacheMla, ConcatMlaNopeRope, MlaPrefillAttention, BatchedMatmul, MlaDecodeAttention. Three of the five — ConcatAndCacheMla, ConcatMlaNopeRope, BatchedMatmul — now have a native kROCM kernel, gated op by op against the CPU oracle on gfx1151 (#2715 W1, spec). The model arm has not been re-run since, so no hit count and no number is claimed for it. And that run's PREMISE has since been withdrawn on this board, which is a stronger statement than "disqualified". It ran at 9f3e6e223, and #2511's 6b97a6800 — which narrowed managed allocation to PageableMemoryAccess == 1 — is NOT an ancestor of it (git merge-base --is-ancestor 6b97a6800 9f3e6e223 returns false), while it IS an ancestor of current main. A probe on strix:gpu0 returns UnifiedMemory=0 DeviceMemoryIsHostAddressable=0 ReferenceTierEligible=0, and vt::ReferenceTierEligible (src/vt/op_provider.cpp:904-935) gates on exactly the second of those. With the tier ineligible, MaybeInstallReferenceTier installs no CPU fallback (:205), so GetOp REFUSES BY NAME for every op with no native kROCM kernel. The four MLA/DSA ops still unregistered on ROCm — MlaDecodeAttention, MlaPrefillAttention, DsaIndexerLogits, DsaTopkSelect — therefore BLOCK GENERATION on this part at current main, rather than running slowly on the reference tier (#2715 owns the four; W1 landed the other four of the seven). The distinction the earlier text drew — one op "consulted through vt::OpRegistered BEFORE the call" refusing while the rest merely fell back — holds only where the tier exists, and here it does not. What is NOT re-verified is the LOAD: the 228 compressed towers and 11.620 GiB were paged in under managed allocation, and whether plain hipMalloc still carries them on this part has not been measured. It is not asserted either way. GB10 (dgx:gpu0) with EXPERT STREAMING LIVE: the same checkpoint runs with [expert-stream] ON slots=8192 slot_bytes=6684672 resident=51.00 GiB, slot reuse climbing hits 0 -> 789 -> 1791 -> 3156 across four decode steps at exhausted=0, 29.73 GB moved, zero reference-tier hits. Read that as a HYBRID, not a pure streamed arm: VT_CPU_MOE was unset, but the engine's own --fit placed 33 of 78 layers' routed experts on the host to bring a 216433205760 B footprint under a 128452960256 B budget. There is no valid device-residency figure on a unified part (VmHWM peaked at 72.57 GiB, which is page-cache pressure; nvidia-smi reads 0 MiB). That run also CONVICTED #2544 for this model — the forward never read device_token_ids, so with the async mirror at its CUDA default every decode step after the first embedded token id 0 and emitted plausible garbage at rc=0; fixed in #2596, landed but NOT yet driven on the artifact. The safetensors arm still refuses by name (703.74 GiB across 141 shards, no streaming loader, no MoE-expert block-fp8 rung)	NONE, and none is reachable on this fleet. vLLM implements this architecture at the pin 5559679229bc961848b121ccdeaa8fa5d79bec98 (registry.py:117) and cannot fit it on any device this project reaches, so no wave may promise a token-exact number against it	It runs and emits text on ROCm gfx1151 and on GB10 with expert streaming; NO SPEED NUMBER IS PUBLISHED on either. The ROCm arm no longer merely DISQUALIFIES a number: on gfx1151 the reference tier is withdrawn (#2511), so the four unregistered MLA/DSA ops refuse and the arm does not generate at all (#2715); the GB10 arm is a hybrid, and no oracle that fits this fleet can gate its correctness
LagunaForCausalLM	poolside/Laguna-S-2.1-NVFP4, GGUF-Q4_K, Laguna-XS	⚠️ no golden, no token gate. tests/parity/goldens/ holds no Laguna entry and no shipped test compares tokens or distributions against vLLM. The strongest recorded vs-vLLM result is first 2 tokens match, then near-tie divergence against the MARLIN golden, and the one distributional item on the row (VT_LAGUNA_KV_BF16) is recorded UNRATIFIED. "byte-exact" elsewhere on this row means byte-exact against our own prior output under a lever A/B, never against vLLM (#2841)	vLLM parity+ 1.03x (.agents/benchmark-record.md CLAIM-LAGUNA-LONGCTX-LEVERS), default on (VT_LAGUNA_RESIDENT_BF16W, src/vllm/model_executor/models/laguna.cpp:141), via the laguna-gen CLI. ModelRegistry::Forward reached the f32 reference rather than the quantized forward and SIGSEGV'd on the NVFP4 arm (#2618; repair in #2835)
KimiLinearForCausalLM	Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE)	Folded onto the shared paged runner (ROW 7 §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128, WHICH IS A FAILING STRICT GATE, NOT A PROFILE — the oracle's greedy decode is bit-identical across K=3 (kimi-linear.md §589-591), so CLAUDE.md §Gates licenses no distributional bar; the shipped test asserts a FLOOR (matched*128 >= total*122) behind doctest::skip unless VT_KIMI_MODEL_DIR is set (#2840); FA2 paged MLA default-ON; SACRED post-fold green	Served via vllm_engine_load + vllm_complete_tokens (ABI v13); server 19.0 tok/s wall vs vLLM ~21 (~0.90×), speed residual open
KimiK3ForConditionalGeneration	Kimi-K3 (2.8T MoE)	scaffold: registry+config+enumeration gated, forward refuses	HW-infeasible (~1.56 TB); no run
Dots3NoteForCausalLM	dots-studio/dots3-note-prev @1e1e7b0c (280B-A16B multimodal MoE, 576,886,825,984 bytes bf16; the -fp8 sibling dots-studio/dots3-note-prev-fp8 @7c14222e is 298,673,280,504 bytes = 298.67 GB decimal / 278.16 GiB binary). Headers only — no tensor byte downloaded	W1+W2 scaffold: registry + config gated off the REAL released config.json, with one assertion per §4 config trap (ungrouped 1/1 router, GPT-J indexer RoPE, one nextn layer, the two LoRA rescales, the sliding theta); name map accounted 38006/38006 over the WHOLE released index — at W2 that split read 35381 language / 2195 vision / 430 audio, and W5c re-bucketed it (see below) — with the two tower files carried as named W6/W7 deferrals rather than dropped; W4a+W4b-2 put BOTH attention geometries on the DECODE PATH, reached through ModelRegistry::Forward: the 13 full-attention layers with the two LoRA rescales, k_rope_only_layernorm and the headwise gate, and the 33 sliding-window layers over a PADDED 1088-wide MLA cache row that each layer narrows to its own logical width on read; vt::MlaDecodeAttention and vt::MlaPrefillAttention grew an optional window whose absent state is bit-identical to no window; W4b-3c put the DSA lightning indexer's SELECTION on the same path, so a SINGLE-SHOT prefill longer than index_topk is now served sparsely instead of refused — vt::MlaDecodeAttention grew an optional selected-slot arm whose absent state is bit-identical to no selection and whose FULL selection reproduces the dense answer byte for byte, beside a new vt::DsaIndexerLogits / vt::DsaTopkSelect pair on CPU and CUDA. A STEP in which any request has CACHED CONTEXT and any request is past index_topk still REFUSES BY NAME — the sparse route is a property of the step, not of one request — because the indexer's own key cache is a second attention group owned by KV-DSV4-MULTICACHE (#1925). W5 put the 45 MoE layers on the same path through Dots3NoteMoeBlock over the shared RunMoePlaced seam — the UNGROUPED (n_group=1 / topk_group=1) noaux_tc router at 256/8 with the F32 e_score_correction_bias feeding the SELECTION and the unbiased scores feeding the routing weights, plus one shared expert at moe_intermediate_size * n_shared_experts = 1536 and NOT intermediate_size = 13824 — with no vt op changed. W5c removed the last refusal, and that one was a DEFECT rather than a gap: the nextn branch was STRICTER THAN UPSTREAM, which drops model.layers.46.* and model.mtp.* from the main model instead of refusing (utils.py:542 -> deepseek_v2.py:1618-1620; models/dots3_note/nvidia/model.py:624), so those 19 tensors are now a NAMED W10 deferral with their own accounting bucket (#2176) and the split reads 35362 language / 19 nextn / 2195 vision / 430 audio. So Dots3NoteDeviceRefusal is EMPTY for the released config.json, and that is representable rather than runnable: the MoE alone is 545.82 GB of a 576.89 GB checkpoint (94.62%), which no host this project reaches can hold. What still refuses BY NAME is the blockwise-FP8 sibling (W9 — weight_block_size [128, 128] with a weight_scale_inv per projection, which this port's bf16 loaders cannot read), GGUF k-quants (W9); the VISION tower no longer does, and since W7a the AUDIO tower no longer does either. W6a (#2512) puts the DENSE half of the VISION TOWER on a SERVED request: an image_url chat request travels ApiServer::handle_chat_completions -> the architecture-dispatched multimodal chat seam (a REGISTER_VLLM_MM_CHAT translation unit of its own, #2481) -> GPUModelRunner::execute_mm_encoder (#2398) -> ModelRegistry::EmbedMm -> ModelRegistry::Forward, and kDots3NoteFactory now carries encode_mm and embed_mm. The arm is patch_embed -> blocks 0-24 (fused attn.qkv with no bias, per-head q_norm/k_norm applied BEFORE a 2-D vision RoPE, bidirectional attention, and the three-tensor SwiGLU routed through layers::MlpGateUpMethodBase) -> post_trunk_norm -> the patch_merger adapter folding 4x1536 to 5120, plus a dots3 image processor with its own per-channel normalization, its own resized_size and its own marker. W6b (#2613) lifted the last refusal in front of the RELEASED checkpoint's tower: its 17 pyramid MoE blocks (608 routed experts, 1960 of the 2195 vision_encoder.* tensors) now COMPUTE, so all 2195 load and an image request against the real checkpoint is served instead of answered HTTP 400. The routed arm is mlp.gate_weight + the F32 mlp.router_bias — NOT the language tower's mlp.gate.weight + mlp.gate.e_score_correction_bias, which lives in the same checkpoint — through vt::MoeRouterTopK at one expert group, a per-expert SwiGLU on layers::UnquantizedMlpGateUpSplitMethod (split rather than merged, because merging 608 experts would cost 7.9 GiB of resident copy), and vt::MoeCombine carrying upstream's self-normalizing divide as a constant routed_scale. W6c (#2537) made a REAL image servable: ProcessImage used to refuse any image whose sides were not already multiples of patch_size * merge_size (28 on the released geometry) because the resample upstream always performs was not ported, so almost no photograph or screenshot cleared it. PilResizeBicubicRgb (src/vllm/multimodal/pil_resize.cpp) is Pillow's ImagingResampleInner 8bpc path and not a four-tap cubic: the filter support is scaled by max(1, in/out), so a downscale reads a much wider window and behaves as a weighted area average; samples are centred at (xx + 0.5) * scale; weights are normalized per output pixel; and the two passes run in 22-bit fixed point over a uint8 INTERMEDIATE that saturates between them. It is PIL's and NOT torchvision's, so it does not discharge Qwen3-VL's own deferred resize. Its gate is the row's two-arm reference — exact equality against a fixed-point reference (0 of 255 levels over 18 non-square cases) and 1.08 of 255 against a continuous one whose bound is 2 — plus a constant-field arm that only per-row normalization can pass and asserted-reached 0 and 255 saturation. What it does NOT wire is resized_size's detail / image_details / target_height / target_width arguments, which are request parsing and stay owed to W8 (#2645). W6b also lifted four config arms W6a had deferred — adapter_type = pixel_shuffle_mlp, post_norm = false, use_qk_norm = false, is_causal = true — and left three refusing BY NAME: use_bias = true (#2616) and the softmax / top-k-below-2 router arms (#2615), none of which any published checkpoint selects. Because top-k selection is a DISCRETE choice whose error is bimodal, the gate asserts selection-SET equality per token against the reference's own scan and prints the minimum decision margin (1.26e-2 on the fixture at token 8, over a spread that touches all 4 experts and all 6 possible pairs); a relative tolerance alone cannot see a selection flip, and a routed-arithmetic defect that preserves the selection set (spec §4.12.9's M6) is caught by the tolerance and NOT by the served suite. Ported from vLLM nvidia/vision.py (677 lines) / nvidia/vision_attention.py (477 lines) / common/processor.py read at 9035151d6, which is BEYOND our pin 5559679229; every anchor names that SHA because vision_attention.py is already 494 lines at vLLM main 7a100bb61. The gate is a CONSISTENCY gate (spec §6.4 option B): an independent in-test double-precision reference sharing no helper with the implementation, plus a served-request reachability suite whose load-bearing case is two DIFFERENT images compared on logprobs. W7a (#2703) puts the AUDIO tower on a SERVED request, through the same two hooks: an input_audio chat part reaches the 32-layer dots speech encoder and audio rows land in the prompt embeddings. The arm is a three-layer stride-2 Conv2d stem — composed as im2col + vt::MatmulBT as ONE EXACT TRACKED EXCEPTION, because src/vt/cpu/cpu_conv2d.cpp:111 is the only OpId::kConv2d provider in the tree and a stem routed through the seam would resolve no op on a CUDA queue (#2709) — masked at FOUR stages from valid_mel_lens = samples // 160, then 32 Whisper-convention blocks (k_proj alone carries no bias) with partial RoPE over half of each head, and a GELU adapter. Two shared seams grew to reach it: vllm::multimodal::MelFilterBankSlaney is Parakeet's filterbank EXTRACTED rather than a sixth mel copy, and it reproduces the committed third-party voxtral_mel_filters_f32.bin BIT FOR BIT (0 of 25728 values differ), which is the only assertion on this row whose other side this repository did not also write; and layers::MlpGateUpMethodBase gained a fourth member with a BIAS arm, reached by the audio fc1 [10240, 1280] + [10240] this checkpoint ships. That does NOT close #2616, whose subject is the VISION tower's use_bias and whose other two blockers this does not lift. W7b (#2797) lifted the chunk_seconds refusal, so a real recording is served rather than a clip: Dots3NoteAudioProcessor::ProcessWaveform now slices the waveform into chunk_seconds segments, takes a log-mel PER SEGMENT — the -8 floor is a GLOBAL max over whatever tensor it is handed, so one pass over the whole waveform would shift the quietest chunk — and Dots3NoteAudioForwardChunks runs each padded chunk at ITS OWN valid sample length and concatenates the first token_len * merge_factor rows of each IN ORDER (nvidia/audio.py:193-234 @ 9035151d6). A loop over the single-chunk tower IS upstream's one batched varlen call, because the pack builds each chunk its own cu_seqlens window and restarts its rope positions at 0 (nvidia/audio_encoder.py:664-685, :711-719), so no chunk can reach another; at one chunk the path is W7a's, byte for byte. NumAudioTokens is UNCHANGED — upstream's (n-1)//1280 + 1 and this port's ceil(n/1280) are the same function for every n >= 1 and differ only at n == 0, where a literal C++ transcription would invent a token. What W7b adds is the invariant a reader cannot infer: the tower's per-segment sum equals the prompt side's single ceil(total/stride) (common/processor.py:771) for every waveform exactly when chunk_samples is a whole number of token_strides. The released config satisfies it (960000 = 750 * 1280) and so does every EVEN chunk_seconds at 16 kHz; a config that does not is refused PER REQUEST past one chunk, at the chat seam and not inside the engine loop, and a clip inside one chunk is served there either way. W7c-1 (#2813) lifted the MONO restriction: a multi-channel PCM16 WAV already at 16 kHz is served, reduced by the per-sample mean over its channels — upstream's own load_audio(..., mono=True) -> np.mean (vllm/multimodal/media/audio.py:207-208, :220 @ 9035151d6), which dots3-note selects through AudioSpec(target_channels=1, channel_reduction=MEAN) at vllm/models/dots3_note/common/processor.py:523-525. The mean is accumulated in int32 — exact across the whole uint16 channel domain — and the served answer is the correctly-rounded float of that exact mean; it is bit-identical to upstream's float32 mean for every power-of-two channel count up to 512, C=1 and C=2 included, and that bound is TIGHT (at C=1024 the exact sum passes 2^24 and stops being a float32, and a sequential float32 accumulator then agrees on 126 of 2000 near-full-scale draws). W7c-2 (#2828) lifted the 16 kHz restriction: a PCM16 WAV at any sampling rate is served, resampled to audio_config.sampling_rate by ResampleAudioScipy. That is a RECORDED DIVERGENCE and the first on this row. Upstream's default resampler is pyav — libswresample — which is not bit-identical to itself across CPU dispatch on one binary and one input (ffmpeg 6.1.1: 24691 of 32000 samples differ, worst 9.686e-08), so no bit-exact gate against it can exist in principle. What is implemented is resample_audio_scipy (vllm/multimodal/audio.py:232-250 @ 9035151d6), ANOTHER ARM OF UPSTREAM'S OWN SWITCH, which vLLM ships in production for phi4mm (phi4mm.py:580); measured distance from the real default at 44100 -> 16000 on a 0 -> 7500 Hz sweep — content that reaches the output Nyquist, which is what a speech encoder sees — is 51.78 dB, against soxr 44.63 and torchaudio 29.02, and the probe is part of the number: on content well below the new Nyquist the ordering inverts and soxr wins by 30 dB. The gate is a CONSISTENCY gate against scipy.signal.resample_poly at its defaults, with committed goldens for four rate pairs plus a fifth carrying a tone above the output Nyquist, a length assertion, a lower bound and a difference assertion against a nearest-sample decimation — because a resampler that returns its input, or returns zeros, passes a tolerance alone. Refused BY NAME, and both bounds are deliberate divergences because the rate is named by the REQUEST: a non-positive rate; a reduced polyphase ratio whose max(up, down) exceeds 100000, which is the FILTER, at 20 * max(up, down) + 1 taps; and a reduced up/down above 8, which is the OUTPUT (found by the fresh review of PR #2842). The second bound is not implied by the first: up is target_sr / gcd and can never exceed 16000, so a fmt chunk declaring 1 Hz passes the filter bound and then asks for 16000 output samples per input sample — a measured 1220.7 MB from a 40 KB upload, twice per request, and a std::bad_alloc that would be answered HTTP 500 for a property of the request. 8 is four times the largest ratio this row serves (telephony's 8000 -> 16000) and admits any source rate down to 2000 Hz. Any container but RIFF/WAVE PCM16 is refused too, and that one is NOT this row's — mp3/flac/ogg need a demuxer this tree does not vendor, five surfaces refuse compressed media for the same missing brick, and #2814 owns it. W8a (#2860) makes ONE request carry MORE THAN ONE ITEM and MORE THAN ONE MODALITY. Each modality's placeholder used to be expanded in its own pass over the id stream, each pass reporting offsets into the vector IT built, so the second span was short by every earlier expansion and a mixed image+audio request was refused BY NAME rather than served wrong. vllm::multimodal::ApplyPromptReplacements replaces both passes with upstream's own planner, apply_token_matches (processing/processor.py:944-957 @ 9035151d6), narrowed to the non-empty REPLACE subset every dots3-note rule uses: it walks the prompt ONCE, takes the earliest match across every rule that still has items, breaks a tie by the rule that appears earlier, emits spans in ID-STREAM order — which is what GetMmFeaturesInWindow's two binary searches are a precondition of — and THROWS BY NAME when a rule's items are not all consumed, because dropping one silently is the outcome that answers fluently from media the model never saw. It also keys on the whole [start, pad, end] target rather than the pad id alone, which is upstream's key, so a bare pad typed into a user's text is now an ordinary token instead of a 400. The declared per-prompt limits move with it, from {"image": 1, "audio": 1} to {"image": 512} plus {"audio": 128} when the checkpoint carries an audio_config — upstream's own numbers (common/processor.py:527-534 @ 9035151d6) — and a user --limit-mm-per-prompt can still only LOWER them. video stays ABSENT although upstream declares {"video": 1} beside them, so a video part is still refused at limit 0 with upstream's own message: the decode it would promise needs a container demuxer, an H.264/VP9/AV1 bitstream decoder and a JPEG codec, none of which this tree vendors, and W8a widened #2814 to hold it. supports_multimodal went TRUE -> FALSE at W5 and FALSE -> TRUE at W6a, and the trail is the honest record of what this port could back at each point; what W8 still owes is the video front end, and the include/vllm.h multimodal request path, which LEFT this row to #2862 because it must serve Qwen3-VL as well	No oracle, on any host we own (298.67 GB fp8 against a 122 GiB ceiling), so NO number is claimable on any axis and the e2e gate is an open gap by construction (spec §6.4, #699)
NemotronHForCausalLM	Nemotron-3.5-Lightning-30B-A3B-NVFP4 (nvidia @29f2d174)	config+enumeration+KV-shape gated; hybrid forward COMPUTES; loader materializes 18487/18487 as SHIPPED; A3 e2e gate 96/96 STRICT PASS on GB10 at 0ea5d249f (#1221); NO run against current main	PAGED (#810 A2-P): K/V go to the runner's pages; conv+SSM rows carry at the metadata's state indices. G-SAFE: num_reqs <= 1. Device lm_head (A2-Q2b), UNMEASURED. Owed: FP8 mamba (A2-Q1), MTP, GGUF
MuseGlimmerForCausalLM	real tensors, bf16 depth 4/52 only: 5 prefill argmax positions match a torch transcription of vllm#51655 and HF. GGUF full depth generates coherently (#347, #359) but is NOT token-exact	text forward + loader vs an fp32 reference, per-mechanism property tests, scaffold 11/11, GGUF gate 17/17. An ABSENT config key now takes the architecture's constant (#412): GGUF post-norms ran at 1e-5, not 1e-8	no vLLM denominator (pin cannot load it); SECONDARY llama.cpp, same GGUF, GB10 CPU: prefill tie 0.997x, decode 0.232x, RSS 1.92x (#333)
MuseGlimmerForConditionalGeneration	vision: no reference run of any kind; enumeration gated vs the released 30B index (1436/1436). Image/video need bf16 safetensors: mmproj-kquant.gguf is refused by name	perception encoder loaded and wired, so an image or video prompt runs; perception_emb_norm now armed by default (#405). Reachability plus placeholder scatter only, no image or video correctness	not measurable; anchored to open vllm#51655
LlamaModel	landed tiny synthetic embedding fixture (engine path == direct pooler path, identical vectors; f64 LAST+normalize reference); real checkpoint (e5-mistral class) is a NAMED residual	pooling/embed only, text paths refuse by task; vllm_embed + /v1/embeddings	n/a (CPU correctness-grade embeddings)
ParakeetForCTC, ParakeetForRNNT, ParakeetForTDT	nvidia/parakeet-ctc-0.6b/-1.1b, -rnnt-0.6b, -tdt-0.6b-v3 (transcribed, ids exact vs HF generate(), P4/P6 2026-08-07; not retained) + committed synthetic fold fixture	ASR transcription-only (SupportsTranscription mirror; text paths refuse by task); fold gate byte-identical to the pre-refactor pipeline	n/a (CPU correctness-grade ASR via vllm_transcribe + /v1/audio/transcriptions)
CohereForCausalLM	Command-R / Cohere (and Cohere2)	scaffold: W0 tiny-random oracle run-verified; real-checkpoint gate blocked	no run

Standalone and non-registered lanes

These run through dedicated forwards, not the REGISTER_VLLM_MODEL registry, so they sit outside the gated list above. One caveat the LTX-2.5 row is too narrow to carry: its text tower's prompt tokenization mirrors upstream only while the checkpoint's tokenizer post_processor adds nothing. The shipped one is MEASURED empty, so this port's plain encode plus an explicit BOS prepend matches upstream's add_special_tokens=True today; a checkpoint with a non-empty post_processor would tokenize differently here, and Ltx2TokenizeGemmaPrompt in ltx2_text_encoder.cpp is the call that would have to change.

Lane	Tested checkpoint(s)	Correctness gate	Speed vs reference
Voxtral audio (VoxtralForConditionalGeneration)	Voxtral-Mini-3B-2507	near-tie-robust 16/16 vs vLLM 0.25.0	decode 0.97x (beats vLLM); encoder FORWARD 15.90x of vLLM's whole TTFT (pin 46.02 ms), or 2.89x with opt-in VT_WHISPER_ENC_FA2=1 (costs 3 near-tie divergences vs 0). Not a TTFT ratio. Pending
Whisper audio encoder	openai/whisper-small; whisper-large-v3 (Voxtral cfg)	encoder tower 77/77; large-v3 tower 203/203	pending
MiniMax-H3 DiT (MiniMaxH3DiTModel, vllm-omni lane)	MiniMax-H3 (33.1B video+audio)	portable 79/79; all three modalities COHERENT on Q4_K_M (§8.20); PRUNED ckpts run, Q8_0 seam 0.9941 (§8.21); ref2va grid was NVFP4 quant error, §8.9 REFUTED; GGUF/NVFP4/bf16 shards stream	FP4/Marlin landed; speed pending; no bf16 render yet. Render from the Q4_K_M GGUF, not the NVFP4 arm. Krea 2 text-to-image (roadmap C11) is scoped to reuse these DiT seams
LTX-2.5 DiT (LTX2VideoTransformer3DModel, Lightricks lane)	LTX-2.5 (21.00B video+audio)	SPIKE. DiT, VAEs+ENCs, cond, pipeline, quant loaders gated, reduced dims. Prompt AdaLN host+dev; Gemma-4->xattn FIXTURE-gated. Img chain PPM->resize->encode->place->noise. Temporal x2 ups DRIVEN. Render OWED	ltx-2.5/ltx2-gen. NVFP4 ~29 GB/GB10, FP8 ~44, bf16 42.0; +24 GB tower. FP8/torchao/NVFP4/bf16; kf abs-pos ported; ALL 3 load, NO allow_unported. IMG+LAST kf crf=0, A2V WAV+LoRA; --lora REPEATABLE and N adapters fuse in order, the 2nd onward through upstream's addmm_ form (#932); a 2nd adapter REFUSES on a2vid_two_stage/ti2vid_two_stage/keyframe_interpolation, whose stage 1 needs a proper subset this engine cannot hold; DiffVAE/ref refused. PENDING
MiniMax-Music3 (MiniMaxMusic3ForConditionalGeneration, diffusers lane)	MiniMax-Music3 (8.6B Qwen3 LLM + 0.646B RVQ decoder + 2.4B fp32 DiT + DAC Flow-VAE); diffusers arm, ~28.5 GB	ACTIVE. Loader 1413/1413; AR, acoustic and the 8.6B LM forward all gated vs real weights; SpeechRegistry + vllm_speech_* v21 + /v1/audio/speech; GGUF Q4_K depth decoder value-gated. HTTP request OBSERVED (#852)	No reference number. Host kernels multi-core, same song bytes (§12). PARTIAL device arm (#672): 8.6B LM + 2.4B fp32 DiT staged once (§14); rest host. Denominator SGLang-Omni production
LTX-2.5 DFR base + generated keyframe slots	LTX-2.5 (21.00B video+audio)	gated vs EXECUTED upstream dfr_layout + 3 dfr_pipeline helpers @ fd4ded7f (test_ltx2_dfr 11/11, 652 assertions); canvas, tiles, stitch, carry-forward as EXACT index vectors, since each defect is plausible	--pipeline-kind dfr plus --lora, which is now REQUIRED (#1445, both stages); NO keyframe_slot_sft base is published either, so the arm is REFUSED in practice. Canvas PADS 9->25, trims back; x8-grid slots MARKED, read before trim. ROUNDS below; detail LoRA (#975) refused
LTX-2.5 DFR temporal rounds (temporal_upsample_rounds)	LTX-2.5 (21.00B video+audio)	7 test_ltx2_video cases via LoadVideoEngine/Generate; deleting the upsampler call site reds the reachability counter while frame counts stay green; the four per-tile guarantees are red-first by mutation	--temporal-upsample-rounds + --temporal-upsampler. FIRST caller of the temporal x2 ups. (n-1)*2**r+1 frames; cond fps capped 60, playback not. Tiles clamp to segments; unclamped arm ungated (#1493). Fixture-gated
LTX-2.5 checkpoint class	LTX-2.5 DiT	8 load-path cases; deleting validation reds 4 engine cases	--checkpoint-class: full, distilled, or keyframe_slot_sft. Required except for dmd2; declarations are checked, but matching bf16 headers prevent automatic detection (#1137)
LTX-2.5 tiled + streaming Conv VAE decode	LTX-2.5 video VAE	gated vs executed upstream ltx_core @ fd4ded7f (test_ltx2_tiling 10/10, 915 assertions); one-tile and untiled-spatial controls BIT-EXACT vs untiled on both causality arms; an untiled frames axis is REFUSED	Streams temporal chunks through upstream's AUTO layout (768/64 px, 80/24 frames); above one tile the pixel volume is never materialized. NO-OP below 768px and 81 frames; 81-120 IS tiled, differing 6.70% of range
LTX-2.5 Conv VAE decode arithmetic width	LTX-2.5 video VAE	test_ltx2_vae "the decode's convolution accumulates in f32", entering through Ltx2VideoDecodeStreaming; widening the accumulator to double, or deleting the production call site, each turns it RED	f32, the width F.conv3d uses at f32 AND bf16 (MEASURED). Was f64 at 8 sites (#1008). Conv sums BLOCKED per input channel, as torch's. STORAGE follows the checkpoint since A24 wave 3 (#2786); f32 is the PARITY arm
LTX-2.5 text tower arithmetic width	LTX-2.5 Gemma-4 text tower + both caption projections	test_ltx2_text_encoder's five ltx2 text bf16: cases against upstream EXECUTED in torch.bfloat16 at fd4ded7f, compared as raw 16-bit patterns; the production case enters through Ltx2EncodePromptToConditioning and asserts every returned value survives a bf16 round trip while the f32 arm FAILS the same predicate, so the gate is measured to discriminate rather than assumed to	bf16, the single dtype upstream resolves for the whole pipeline (distilled.py:109, handed to PromptEncoder at :111-113). The render loads the checkpoint's own 16-bit values through Ltx2TextProjectionsAsBf16 instead of doubling them: ~2.3 GB rather than ~4.6 GB for the two projections, and half of every full-width activation buffer. Upstream's TWO norms are NOT the same dtype — V1's returns f32 because its int64 denominator promotes, V2's is bf16 and rounds each square before the mean — and the epsilon that reaches either is bf16(1e-6). f32 stays as the PARITY arm the goldens are measured in (Ltx2WidenTextProjectionsToF32). A24 wave 1 (#2676); the other seven components stay f32 and are owed
LTX-2.5 Conv VAE decode STORAGE width	LTX-2.5 video VAE decoder + the kLtx2Vae CPU kernels + the tiled-decode buffer	test_ltx2_vae's bf16 arms against upstream EXECUTED in torch.bfloat16 at fd4ded7f: a SHALLOW two-convolution arm BIT-EXACT with three rejected rules 0.0039-0.0078 away, five per-kernel rules bit-exact each against the hypothesis it rejects, and the deep thirteen-convolution arm bounded by the chain's own measured one-ulp response (0.0117) which the two reachable defect rules exceed at 0.0352 and 0.0156. The production gate enters through Ltx2VideoDecodeStreaming and reads 0 of 27648 decoded pixels wider than bf16 where the f32 arm reads 27645, with the latent that ENTERS the decode reading 8 of 8 as the live control	bf16, the single dtype upstream resolves for the whole pipeline (distilled.py:109, handed to VideoDecoder at :148). Every volume carries an element width and the weights are no longer expanded through Bf16ToF32. SIX rounding rules were EXECUTED rather than read and three scalars in two files give three different answers: PixelNorm's epsilon narrows for its add, the noise blend's Python floats do NOT, and the registered per-channel statistics do. bf16 is CPU-ONLY and the refusal names why -- vt::Conv3d has no bf16 arm on CUDA (#1007) and every convolution goes through it -- which costs nothing because the render decodes on the CPU queue. The noise DRAW stays f32 (#2780). A24 wave 3 (#2786); the upsampler and duration head stay f32 and are owed
LTX-2.5 Conv VAE ENCODE arithmetic + STORAGE width	LTX-2.5 video VAE encoder (Ltx2ConvVideoEncode, Patchify, SpaceToDepthFold, SpaceToDepthDownsample) sharing the decoder's VaeStore and kLtx2Vae CPU kernels	test_ltx2_vae's "the video ENCODER's BF16 arm matches upstream ltx_core" case against upstream EXECUTED in torch.bfloat16 at fd4ded7f: BIT-EXACT, 0 of 64 words differing, and the case ASSERTS that both rejected rules (0.00390625 each) sit BELOW the chain's own one-ulp sensitivity (0.0078125), which is why a wave-3-shaped tolerance band would have been a mute switch here and a bit-exact gate is used instead. The production gate enters through the render path and reads vae_encode_not_bf16 0 of 20 where the f32 arm reads 20 of 20, with vae_encode_in_not_bf16 at 15352 of 15360 pixels as the live control. Deleting the kBF16 argument at the loader call site leaves the whole 3570-assertion value suite GREEN and reds only the render-path counter, which is the reachability half; carrying the group mean in f32 across the skip add, or taking the per-channel statistics off the bf16 grid, each reds the value case and leaves the counter green	bf16, the single dtype upstream resolves for the whole pipeline (distilled.py:109, handed to ImageConditioner at :120-125, which builds this encoder with it). Upstream's VideoEncoder.forward casts NOTHING: the pixels arrive already bf16 from load_image_and_preprocess, so the one narrowing sits at the boundary before any arithmetic. THREE rules were EXECUTED rather than read, and a fourth probe is reported as separating nothing rather than as a confirmation: the group mean widens internally and rounds ONCE on the store (sampling.py:50-51, and a probe at group_size == 2 gates it at zero, so the fixture is laid at 4); the skip ADD has no width question of its own, but carrying the mean's result unrounded ACROSS it is wrong on 18-24% of the block's output (sampling.py:63); and PerChannelStatistics.normalize NARROWS both registered buffers (ops.py:81-84), which is 39-47% of the encoder's last arithmetic and which no token gate can see, because the latent is a std::vector<float> on either arm. The encoder is HOST-ONLY on BOTH arms and this row does not change that: SpecOf is called without a queue, so it never reaches a device path and the decode's device refusal has no counterpart here; its residency is owed (ltx25-vae-device-residency.md). A24 wave 4 (#2850); the latent upsampler and the duration head stay f32 and are owed, the first because it routes through no shared seam at all and the second because nothing reaches it
LTX-2.5 Conv VAE decode threading	LTX-2.5 video VAE	test_ltx2_vae "the decode DISPATCHES its convolutions to the CPU threadpool" and "...BIT-IDENTICAL across thread counts", through Ltx2VideoDecodeStreaming; 34 golden margins UNCHANGED; TSan clean	Parallel over CONV output lines via vt::cpu::ParallelForRows (#1009). ~9x at 16-20 workers, contended box, 21-23% spread. Bit-identical at any count
LTX-2.5 Conv VAE decode DEVICE arm	LTX-2.5 video VAE	test_ops_conv3d 4/4 vs an independent scalar reference; test_diffusion_device_seam runs a non-CPU dispatch and requires byte-identical pixels	Routes through the new vt::Conv3d op on the queue the engine resolved at load (#1007). NO GPU has run the CUDA arm (#1452); other stages stay host-side (#1451)
LTX-2.5 retake (RetakePipeline, regenerate a time window)	LTX-2.5 DiT + video VAE encoder	test_ltx2_retake 4/4 (69 assertions) and 4 test_ltx2_video cases entering through Generate; mask, conform and the four-way plan pinned to upstream fd4ded7f	--pipeline-kind retake on ltx2-gen. Source is a frame_%06d.ppm DIRECTORY; a container is REFUSED (no demuxer). Geometry comes from the clip. A folder has no audio, so the soundtrack is generated
LTX-2.5 text-to-audio (T2AOneStagePipeline)	LTX-2.5 DiT + audio VAE, no video VAE	test_ltx2_video's ltx2 t2a: cases, entering through Generate; 18 mutations, 17 DETECTED (four by review of a conditional-only #1039 gate) and the 18th proven an identity, not a blind spot	--pipeline-kind t2a_one_stage. NO picture: 0 frames, no mux argv. The only AUDIO-ONLY guided arm (CFG + STG, 3 forwards/step), so it needs a text tower. CPU only; the device forward is refused by name
LTX-2.5 HQ preset (TI2VidTwoStagesHQPipeline, res_2s sampler)	LTX-2.5 DiT	6 test_ltx2_pipeline cases + 2 test_ltx2_video cases through Generate, vs UPSTREAM'S OWN loop run at fd4ded7f: video latents BIT-EXACT on 3 of 5 fixtures, 1 ulp on 2. 20 mutations, 18 DETECTED	--pipeline-kind res2s_two_stage plus --lora, which is now REQUIRED (#1445, both stages, at ONE strength rather than upstream's 0.25/0.5; #1144). 2.5 only. TWO denoiser calls per step plus a terminal one, and stage 1 is GUIDED at cfg 3.0 / modality 3.0, so 15 + 3 steps is 38 calls and 100 forwards. The preset IS the sampler
LTX-2.5 T2A guidance space	LTX-2.5 DiT (T2A arm)	test_ltx2_video "the guider is handed x0 predictions" through Generate, on all 3 arms plus the guider output and the Euler input; a seam case puts the two spaces 1.5e-07 apart at rescale 0 and 0.352 at 0.7	Combines denoised (x0), mirroring X0Model (model.py:590-604). Was velocity space, which agrees only at rescale 0 (#1039)
LTX-2.5 VIDEO guidance	LTX-2.5 DiT, joint video+audio	test_ltx2_video's ltx2 one_stage: cases through Generate; all FOUR arms carry the x0 invariant and the guider output replays EXACTLY	--pipeline-kind one_stage runs _guided_denoise: 4 forwards/step, combined per modality in x0. Was ONE unguided forward, every video_guidance field dead (#1092)
LTX-2.5 cross-attention perturbations	LTX-2.5 DiT	test_ltx2_video and test_ltx2_device each gate one direction ALONE, on a forward where the other stream is PRESENT but DISABLED. Swapping the two flags is RED on both arms	SKIP_A2V_CROSS_ATTN / SKIP_V2A_CROSS_ATTN, which modality_scale = 3.0 selects on every video row. On the DEVICE forward too since 2026-08-19 (#1092)
LTX-2.5 latent upsampler dims=2 arm	LTX-2.5 spatial upsampler checkpoint whose config sets dims: 2	ltx2 the latent upsampler reproduces upstream on the dims=2 arm (test_ltx2_pipeline) against the EXECUTED upstream module, plus a dims=2 upsampler checkpoint RENDERS (test_ltx2_video) through LoadVideoEngine/Generate; removing the fold reds the golden at max	diff
LTX-2.5 audio-to-video (A2VidPipelineTwoStage)	LTX-2.5 DiT + audio VAE encoder + spatial upsampler	test_ltx2_pipeline and test_ltx2_video's ltx2 a2vid: cases through LoadVideoEngine+Generate; the take's latent is bit-identical across SEEDS and moves with the WINDOW	--pipeline-kind a2vid_two_stage. Guided half-res stage 1, DERIVED schedule, plain Euler; distilled 3-sigma stage 2. --audio-path and --lora REQUIRED; the distilled adapter rides stage 2 ALONE (#1118). ANY sample rate: the mel front-end resamples (ops.py:36-42, #2583) and the soundtrack returns at the FILE's rate
LTX-2.5 keyframe interpolation (KeyframeInterpolationPipeline)	LTX-2.5 DiT + spatial upsampler	ltx2 keyframe: cases in test_ltx2_pipeline / test_ltx2_video via LoadVideoEngine+Generate: frame 0 APPENDS against a ti2vid_two_stage control, the x0 invariant on four arms, the 4096 anchor	--pipeline-kind keyframe_interpolation. No frame-0 special case, so --first-frame is guidance to interpolate FROM; stage 2's audio leaves. --lora REQUIRED. --last-frame new (#1191). CPU fixtures
LTX-2.5 two-stage text/image-to-video (TI2VidTwoStagesPipeline)	LTX-2.5 DiT + spatial upsampler	test_ltx2_pipeline and test_ltx2_video ltx2 ti2vid: cases through LoadVideoEngine+Generate; the x0 invariant on all FOUR arms, and the 4096 anchor read at two geometries against a res_2s control that moves	--pipeline-kind ti2vid_two_stage. Guided half-res stage 1 on the UNADAPTED model, plain Euler; distilled 3-sigma stage 2. --lora REQUIRED, no --audio-path; stage 1's audio leaves. CPU fixtures, Full-model run owed
LTX-2.5 guidance knobs	LTX-2.5 request surface	test_ltx2_video renders with an override and refuses one on a fixed recipe	Seven video/audio guider extras mirroring default_1_stage_arg_parser, plus a negative embeds pair for a tower-less engine. Refused whole on distilled_two_stage and retake, whose guidance is distilled in
MTP speculator	Qwen3.6-27B, Qwen3.6-35B-A3B	token-identical to vLLM mtp at c1	~4% faster c1; +16% output tput (MoE)
MTP speculation DEPTH (num_speculative_tokens > 1)	Qwen3.5/3.6 mtp.* heads	k=1..4 through the loader, greedy tokens unmoved, two witnesses per arm: the draft decode forwards the propose RAN, and whether the DELIVERED draft row varied with depth. test_mtp_depth 5/5, 63 assertions	Default stays k=1. NO speed claim at k>1. Drafts are proposed and verified, never ACCEPTED, and neither witness proves per-column provenance. Both await the owed DGX gate (#81)
DFlash block-diffusion	Qwen3 (DFlash draft)	near-tie e2e 27/27 vs vLLM	2.9x over spec-off, 1.003x vs vLLM DFlash-on
DFlash2 block-diffusion (dynamic conv + candidate selector)	Qwen3 DFlash2 draft, safetensors or GGUF (bf16 / Q8_0 / Q4_K_M)	Gated against vLLM: 4/4 token-exact, 45/47 draft blocks identical, acceptance identical per prompt. All 7 DFlash2 suites green on sm_121a, zero CUDA skips	GREEDY only. Speed 0.8017x vLLM: RECORDED, no floor, NOT a pass (#1562). A GGUF draft is dequantized to bf16 at load (#1314)
Async scheduling × speculative decoding	any Eagle-type speculator (mtp, dflash/DFlash2, dspark)	test_mtp_depth W7 cases: a spec engine resolves async ON and emits the sync scheduler's exact tokens through both engine fronts (depth-1 and the depth-2 batch queue); test_engine_core_proc pins the -1-placeholder / worker-fill contract	Mirrors vLLM's polarity (async disabled only OUTSIDE the Eagle-type family): drafts ride as -1 placeholders the worker fills from its own propose; host ngram and draft_model stay synchronous. VT_ASYNC_SCHED=0 rolls back. Spec steps keep the host sampler (device-resident spec sampling owed); the GPU TPOT A/B owed (#1824)
DFlash/DFlash2 shared lm_head kept PACKED	a DFlash or DFlash2 draft off an NVFP4 safetensors target	test_qwen3_dflash2_draft 36/36 (353): block logits BITWISE equal to Qwen3_5MTPModel::ComputeLogits on the same packed head, and FromModelDir loads and drafts off one	Widening a head stays refused by name: GGUF output.weight, FP8, W4A4. VT_LMHEAD_FP4=0 rolls back to the refusal. DSpark and the CUDA arm owed (#1628)
DeepSeek-V4 MTP	DeepSeek-V4-Flash (nextn head)	lossless 5/5; real-model weight-blocked	pending

Inventoried but blocked

Enumerated in .agents/model-matrix.md, not registered, no runnable GB10 gate:

Architecture	Model	Why blocked
DeepseekV3ForCausalLM, DeepseekV32ForCausalLM	DeepSeek-V3 / V3.2	671B, ~642 GiB fp8 vs 119 GiB unified; V3.2 also DSA-indexer dep-blocked
GlmMoeDsaForCausalLM	GLM-5 (DSA)	~1404 GiB bf16; dep-blocked (GLM-5.x is DeepSeek-V3.2 verbatim)
MiniMaxM2ForCausalLM	MiniMax-M2	~230B, ~428 GiB bf16, ~4x over the unified pool
Dots3NoteMTPModel	dots3-note nextn head (the target arch Dots3NoteForCausalLM IS registered; see the supported table above)	W10 owns it and it is deliberately NOT registered: a speculator that cannot propose makes the engine accept a speculative config it then dies on mid-run. The checkpoint ships exactly one nextn layer, and since W5c (#2176) its 19 tensors are a NAMED W10 deferral in the language tower's accounting rather than a refusal, which is what vLLM's own loader does with them. Blocked behind the target row: no oracle runs here, 298.67 GB fp8 against a 122 GiB ceiling, so NO number is claimable on any axis (spec, #699)

27 of the 39 registered text-generation architectures carry a passing correctness gate today; the rest are honestly marked scaffold or blocked above. (The 43 registered total also covers 3 Parakeet ASR entry points and the LlamaModel embedding arch, which are not text generation.) vLLM registers 130+ text architectures, so this is a curated, gated subset, not a breadth claim. The first EMBEDDING architecture is registered and live (LlamaModel, task=embed, LAST pooling, the as_embedding_model mirror, gated on the committed fixture); reranking/classify models are not yet registered.

Multimodal

Input	vllm.cpp	vLLM	SGLang	llama.cpp
Image	✅ correctness-gated	✅	✅	◐
Video	✅ correctness-gated	✅	✅	☐
Audio	✅ correctness-gated	✅	◐	◐
Video+audio GENERATION (MiniMax-H3 DiT, LTX-2.5 DiT)	◐ H3: all three modalities COHERENT on Q4_K_M (t2va, fl2va, ref2va; §8.20); the NVFP4 arm carries the patch grid; GGUF/NVFP4/bf16 loaders, pruned too (§8.21). LTX-2.5: a second lane, SPIKE, gated at reduced dims	✅ H3 (vllm-omni, BF16-only, no quantized arm); LTX-2.5 only through the generic diffusers adapter, no native recipe (vllm-omni#6066)	☐	☐
Speech / audio GENERATION (TTS, vLLM-Omni lane)	◐ IndexTTS-2.5: vllm_synthesize renders TEXT to AUDIO on real weights, and the reference clip CONDITIONS it -- CAMPPlus speaker vector into the talker's row 0 and the S2Mel style; two clips give different audio (rms 0.0064 vs rms 0.0956), same clip twice is bit-identical. STRUCTURE only: emotion conditioning is excluded and vLLM-Omni is unpinned, so nothing here is a correctness claim (#634, #633)	✅ (vllm-omni: MOSS-TTS, Qwen3-TTS, Higgs Audio v3, Voxtral TTS, IndexTTS-2.5)	not assessed	not assessed
MUSIC generation (MiniMax-Music3)	✓ every stage gated; an HTTP request observed e2e over a REAL SOCKET against a MUSIC-ONLY server (#852, #672, spec §10); adjacent caption italics match upstream (#1083)	☐ absent from the pin, from vLLM main and from vllm-omni	◐ SGLang-Omni serves the NATIVE layout; its 32 kHz resample and batching are OWED	☐
Multimodal over the OpenAI server	◐ an image chat request now reaches the MODEL FORWARD: the scheduler carries mm_features and schedules encoder inputs, the GPU runner runs the vision tower, gathers its rows, merges them and sets ModelForwardInput::mm (#2379). Gated on a CPU queue over the real ApiServer -> AsyncLLM -> Scheduler -> GPUModelRunner -> ModelRegistry::Forward chain with a SYNTHETIC tiny Qwen3-VL, so it is a REACHABILITY result and not a token one. Still open: the token-exact gate on real Qwen3-VL-4B weights (GPU), the PNG/JPEG codec, video and audio through the runner, and Qwen3-VL serving ONE sequence per step (its registered forward returns only the last row and refuses a batched step by name)	✅	✅	◐
Per-modality input LIMITS (--limit-mm-per-prompt, --language-model-only)	✅ limits, refusals, and the TOWER SKIP: a tower whose every modality sits at 0 is constructed but never loaded. Byte saving measured on Qwen3-VL-4B-Instruct only: 0.770 GiB of host RSS at load, --device cpu, threshold MET on both pairs, 2026-08-28 (#607). Not a general or a VRAM claim. An earlier 2026-08-24 run read 1.542 GiB, about half of which was our own bf16→f32 widening (#1359); its Qwen3-VL half landed and the rerun measured the 0.499x fall, which is correct rather than a regression. muse-glimmer-30b is still unmeasured (benchmark)	✅	☐	☐

Image, video and audio are correct through the CLI and library. Over the HTTP API the image path now runs from the request to the model forward: the production server attaches the seam in server_main.cpp, the scheduler carries mm_features to the worker and schedules encoder inputs against an encoder cache, and the GPU runner runs the vision tower, gathers its output rows, has the model merge them into the token embeddings and sets ModelForwardInput::mm (#2379, which re-files the invisible #2300). What is gated is that the request ARRIVES there, on a CPU queue over the real ApiServer -> AsyncLLM -> Scheduler -> GPUModelRunner -> ModelRegistry::Forward chain against a synthetic tiny Qwen3-VL; the token-exact gate against real Qwen3-VL-4B weights needs a GPU and a checkpoint and is still owed. Four residuals keep the row from ✅: no image codec is vendored, so the server accepts raw RGB only and refuses a PNG or JPEG data URI first; Qwen3-VL serves ONE sequence per step, because its registered forward returns only the last token's logits and now refuses a batched step by name; only the image modality reaches the runner; and no token number exists. Video, audio and multi-image over HTTP are not started, and the seam refuses a video_url or an input_audio part with HTTP 400. Audio in is gated. Audio out has a surface now (/v1/audio/speech, vllm_speech_* v20), but no family renders from a prompt: both refuse, naming what is missing.

Speculative decoding

Speculator	vllm.cpp	vLLM	SGLang
MTP (multi-token prediction)	✅ token-identical, ~4% faster at c1, depth k configurable with default 1	✅	✅
Draft model	◐ CPU brick	✅	✅
Medusa	☐ spike only	✅	✅
EAGLE / EAGLE3	☐	✅	✅
DFlash block diffusion	✅ 2.9x over spec-off, at/above vLLM DFlash-on	✅	☐
DFlash2 block diffusion (a SECOND DFlash architecture, not a change to DFlash)	◐ safetensors AND GGUF drafts DRAFT (bf16, Q8_0, mixed Q4_K_M: 45 Q4_K + 4 Q6_K), greedy, 0.8017x vLLM RECORDED not a pass (#1562), no published artifact LOADED yet (spec)	✅ BEYOND-PIN, vllm#52816	☐
n-gram / prompt lookup	✅ 27B 5/5 strict vs vLLM	✅	✅
DSpark (semi-autoregressive block drafter)	◐ both gate models (spec): token-identical to spec-off; T=1+k verify CAPTURED. Cross-engine ratio UNSETTLED (0.834x matched-and-warm); Marlin MoE CLEARED as the residual	✅	◐
DSpark draft routing (which draft the loader takes)	◐ Qwen3DSparkModel, Gemma4DSparkModel and (BEYOND-PIN, vllm#52197) DSparkDraftModel + qwen3 take the Qwen3 lane; DeepSeek-V4 DSpark is REFUSED by name (spec)	◐ at the pinned 555967922 that pair routes to DeepSeek-V4; ✅ only since vllm#52197, merged 2026-08-17	not assessed
Other methods (ngram-gpu, suffix, custom-class, dynamic-k, mlp-speculator)	☐ inventoried	✅	◐

Structured output and tool calling

Feature	vllm.cpp	vLLM	SGLang	llama.cpp
JSON schema constrained decode	✅	✅	✅	✅
Regex constrained decode	✅	✅	✅	✅
GBNF grammars	✅	☐	☐	✅
xgrammar backend	✅	✅	✅	☐
Jump-forward decoding	◐ knob only, NOT REACHED: the driver (DrainForcedTokens), the grammar forced-token hook, --[enable|disable]-jump-forward and the ABI v10 field enable_jump_forward are built and unit-gated, and no production path calls the driver, so setting the knob changes nothing. The scheduler splice that must recompute KV for the jumped tokens is owed (#2387)	☐	✅	☐
Tool-call parsers	✅ 38 families; hist. OpenAI args decode #526 (CPU child; live pending)	✅	✅	◐
Reasoning-content parsers	✅ 12	✅	✅	☐
Muse Glimmer ATEM parsers (muse_glimmer)	◐ UNIT-GATED ON STRINGS; CHANNEL SCOPING FAILS AT SERVER DEFAULTS: no adjust_request seam, so skip_special_tokens: true strips the framing. OPEN GAP, spec §6.7	✅	☐	☐
Custom logits processors	◐ CPU-verified	✅	✅	☐

Backends and hardware

Backend	vllm.cpp	vLLM	SGLang	llama.cpp
CUDA	✅ sm_80 to sm_121a	✅	✅	✅
CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open, and the closed 20-core floor ran a SUPERSEDED fork denominator rather than the stock b10451 pin, re-take owed #1003)	✅ CPU_ATTN registered (#1371/#1392, spec)	◐	☐	✅
Metal (Apple Silicon)	✅ builds under Apple Clang with project warnings promoted to errors, the Qwen3.5 MoE loader included; its layout-refusal path uses the same messages and behavior on every platform (#1054)	☐	☐	✅
Vulkan	◐ EXL3 is now NATIVE (#2530, spec): kCastF16 and kExl3Gemm were the only two ops an EXL3 checkpoint ran on the CPU reference tier on a Vulkan queue, MEASURED at two fallback notices and now at ZERO. kExl3Gemm is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone exceeds Vulkan's 16 KiB guarantee before one reaches mma.sync, ldmatrix, cp.async or a grid-wide barrier Vulkan has at no version -- so the gate is BYTE equality with the CPU arm on all three codebooks and every width, not a tolerance. Verified on llvmpipe with NO GPU and NO lease, which is what makes it a CI gate rather than a lease gate; no real Vulkan GPU and no checkpoint run, and NO speed number is claimed on any axis. kExl3MoeMlp and kExl3HadR128 are owed	☐	☐	✅
ROCm	W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. M4: Qwen3-0.6B/3.5-0.8B 16/16 (#41). M3: ROCM_ATTN registered (#1056/#1065, spec). EXL3 is now NATIVE on gfx1151 (#2433, spec): kCastF16 and kExl3Gemm were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. kExl3Gemm is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease (#2511). kExl3MoeMlp is owed. CPU parity open (#269)	52 registered ops: full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). APU managed allocation is now MEASURED and NARROWED (#2511): hipMallocManaged on a part reporting PageableMemoryAccess = 0 faulted the GPU in 17 of 21 legs on gfx1151, against 0 of 21 for plain hipMalloc, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and UnifiedMemory() follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and VT_ROCM_MANAGED_ALLOC=1. ROCm guide	✅	✅
XPU / TPU	☐	✅	◐	☐
Tenstorrent Blackhole	◐ ACTIVE, OPT-125m 6/6; Qwen3-0.6B wired; Mistral-7B-v0.3 16/16 on P150 (spec). 16x16 rerun and residual-RMS owed (spec)	✅	☐	☐
Tenstorrent host-free decode	◐ DEFAULT since #1604 (0 opts out): no per-step host readback; 2.1x default-leg tok/s; both golden pairs re-adjudicated, both paged gates 16/16. Capture opt-in only (#1625 hang); async off (#1627)	☐	☐	☐

CUDA runtime-verified on GB10 (sm_121a), Jetson Thor (sm_110) and Jetson AGX Orin (sm_87). sm_110 has no CUTLASS FP4 tensor-core kernels and no fp4-mma, so it stays a correctness venue for those; the one fast path it does get is the vendored Marlin NVFP4 W4A16 GEMM, enabled since 2026-08-11 and validated on Thor silicon (8.0x-29.0x per GEMM at M=1, e2e 16.61 to 81.63 tok/s at c=1 on Qwen3-1.7B-NVFP4A16). That is a kernel-level result, not a token-exact model-level gate.

Vulkan runs a model end to end: opt-125m greedy is STRICT token-exact, 6/6 prompts vs the vLLM 0.25.0 oracle, every op dispatched natively with zero provider declines. Qwen3.6-27B runs too, both GDN recurrences and the fused attention preamble native: decode 4.36 tok/s vs llama.cpp's 4.35, parity met narrowly, denominator SUPERSEDED (#1003), and prefill 21.5x, a SELF-ratio (GB10). A load keeps one copy of the weights, not two, and is 1.54x faster warm: 27B peak RSS 100.8 -> 53.4 GiB. Still partial at 25 natively registered ops of 112 (8 GDN), the rest on the portable CPU tier; quant/MoE/MLA have none. Build with -DVLLM_CPP_VULKAN=ON; off by default.

Serving, API and operations

Feature	vllm.cpp	vLLM	SGLang	llama.cpp
OpenAI-compatible /v1/chat/completions	✅	✅	✅	✅
Streaming (SSE)	✅	✅	✅	✅
Offline batch API	✅	✅	◐	☐
Prometheus metrics	✅ live per-step values on the serving path, not just the catalog; async detach and server teardown wait for the final fold	✅	✅	◐
Container images	◐ cuda/vulkan/cpu lanes build and gate from one Dockerfile (amd64+arm64, ENTRYPOINT vllm-server, ffmpeg included); nothing published to GHCR yet	✅	✅	✅
Graceful shutdown on SIGTERM	✅ clean exit in 0.25 s, including as container PID 1 (#312)	✅	✅	✅
Plugin / out-of-tree model registration	✅ in-tree factory DONE + plugin seam	✅	◐	☐
Fetch a checkpoint from Hugging Face by name	◐ vllm-server --model org/repo[:QUANT] fetches into the HF cache. HTTPS via system OpenSSL (VLLM_CPP_HF_DOWNLOAD, ON). The musl-static archive has no TLS. vllm-cli and the C ABI take a local path (#1280)	✅	✅	✅
A registered forward opens its OWN model type, not whatever it was handed	✅ all 35 entry points establish the concrete type first and refuse a mismatch by name (#775, swept in #847)	n/a	n/a	n/a
Multiple engines in one process (build, destroy, rebuild)	✅ resident device state is owned by the weights, so a new engine never inherits a freed one's pointers	✅	✅	✅
LoRA adapters	☐ CPU brick only	✅	✅	✅
Embedding / pooling endpoints	◐ /v1/embeddings live (task=embed; score/rerank/classify pending)	✅	✅	✅
OpenAI video generation /v1/videos (Sora shape)	✅ model/size/seconds aliases + GET /{id}/content; input_reference and metadata references condition the render; --video-family pins the family (default DETECT), --video-extra K=V carries family knobs	◐ (vllm-omni, its own request shape)	☐	☐
OpenAI speech generation /v1/audio/speech (createSpeech shape)	◐ route + ABI live, opt-in behind --speech-model; lyrics + description are extra named fields for a music family; 20 unsupported keys refused by name; every key read at the top level and under extra_params	◐ (vllm-omni)	☐	☐
logprobs / top_logprobs / prompt_logprobs on both generate routes	✅ generated-token logprobs on completions and chat; prompt_logprobs scores the prompt and reaches the client (per choice on completions, top-level on chat), with vLLM's three refusals. -1 (whole vocabulary) is SERVED here — there is no separate max_logprobs cap. OpenAI echo still does not prepend the prompt to the payload (#223)	✅	✅	◐
Flat C ABI for embedding in other languages	✅ versioned	☐	☐	✅
Request-length bound before tokenization	◐ /tokenize + both generate routes REFUSE a prompt above max_model_len x the longest vocabulary token, naming the limit, never truncating; other routes and the C ABI unbounded (#1541)	◐ header block only; a prompt COUNT bound and an audio-upload byte bound, no prompt byte bound	☐	☐

C-ABI capability coverage

• Which capabilities an embedder drives through the flat C ABI (include/vllm.h, the only installed header), gated by scripts/check-surface-coverage.py: a reachable row names an entry point that exists; an embedder-unreachable row is tracked in scripts/abi-capability-allowlist.txt against its fold row (ARCH-ONE-SURFACE). The ABI is text-generation-complete; the one embedder-unreachable row (multimodal input) is the open capability gap.

Capability	C-ABI surface	Embedder-reachable
Text completion (blocking + streaming)	vllm_complete, vllm_complete_stream	reachable
Pre-tokenized completion (token-id prompts, ABI v13)	vllm_complete_tokens	reachable
OpenAI chat (tools, streaming)	vllm_chat, vllm_chat_stream	reachable
Async request submission	vllm_request_submit	reachable
Structured output / grammars	structured_json, structured_grammar	reachable
Tool + reasoning parser selection	tool_parser, reasoning_parser	reachable
Speculative decoding config	speculative_config	reachable
Custom logits processor	vllm_logits_processor	reachable
Embeddings / pooling (task=embed)	vllm_embed, vllm_embedding_result_free (ABI v15; pooling checkpoints load via vllm_engine_load)	reachable
Audio transcription (Parakeet ASR)	vllm_transcribe, vllm_transcription_params_default, vllm_transcription_free	reachable
Video+audio generation (MiniMax-H3, LTX-2.5)	vllm_video_engine_load, vllm_video_generate, vllm_video_result_free, vllm_video_mux_argv, vllm_video_engine_family (ABI v18 family registry), vllm_video_last_phase_log (ABI v23 render phase table)	reachable
Explicit device selection (auto/cpu/cuda)	device field on vllm_model_params (ABI v14; 0=auto keeps the probe, explicit absent device fails loud)	reachable
Run the OpenAI server (server as a thin ABI client)	vllm_server_main (ABI v18)	reachable
Speech + music generation (MiniMax-Music3; the IndexTTS-2.5 seam)	vllm_speech_engine_load, vllm_synthesize, vllm_speech_result_free, vllm_speech_engine_family, vllm_speech_engine_sample_rate, vllm_speech_engine_requires_reference_audio (ABI v20)	reachable
KV-cache storage dtype (fp8)	kv_cache_dtype field on vllm_model_params (ABI v24; NULL/"auto" is byte-identical default)	reachable
Speculative acceptance telemetry	vllm_engine_spec_acceptance filling vllm_spec_acceptance (ABI v25; counts are cumulative over the handle, so a caller subtracts two reads). drafts_accepted EXCLUDES the bonus token, matching vLLM's spec_decode_num_accepted_tokens, so accepted/proposed is the acceptance rate; drafted_request_steps counts (request, step) pairs that carried a draft and not forward passes, matching vLLM's spec_decode_num_drafts and SGLang's spec_verify_ct, so 1 + accepted/drafted_request_steps is vLLM's mean_acceptance_length exactly and is NOT SGLang's accept_length, whose numerator also carries the prefill token. vllm-cli prints the per-leg delta on stderr	reachable
Multimodal input (image/audio/video)	none	embedder-unreachable

Parallelism and scale-out

Single-GPU today. Every mode below is scoped against one vt::Communicator abstraction, and world_size == 1 stays byte-identical.

Mode	vllm.cpp	vLLM	SGLang
Tensor parallel (TP)	◐ CPU-gated, no 2-GPU run; TP-W1 LANDED 2026-08-08 (rank-layout group table + per-rank handle); TP-W2..W4+W7 CPU-completable	✅	✅
Collective / process-group abstraction	✅ CPU + NCCL transport	✅	✅
Pipeline parallel (PP)	☐ spike written	✅	✅
Expert parallel (EP) + EPLB	☐ spike written	✅	✅
Data parallel (DP)	☐ spike written	✅	✅
Context parallel (PCP / DCP)	☐ scoped	✅	◐
Multi-node	☐ spike written	✅	✅
PD disaggregation	☐	✅	✅

CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the CPU supports them (SSE2 before), selected by a runtime probe, and can take a transpose-free [K,N] weight path via an opt-in load-time repack (VT_CPU_ELEM_KN_REPACK, CPU only, default off). Byte-identical to the portable tier either way.

Not supported yet

Gap	State	Detail
Kimi-Linear-48B-A3B (KDA + NoPE-MLA + MoE hybrid)	Runner fold LANDS (ROW 7 §21, #122): the ENGINE/SERVER surface serves Kimi at the 122/128 golden profile (engine==CLI 128/128); STRICT stays closed (intrinsic p7 near-tie)	server 19.0 tok/s wall / CLI 18.9 vs vLLM ~21 (~0.90×), speed residual named (§21)
Muse Glimmer 30B (Meta)	Text gated at reduced depth 4/52 only; vision wired but never reference-checked	spec / #268. Full depth, multi-step decode, image/video, server path and parser scoping open. vLLM speed OPEN GAP; llama.cpp bar #333
LTX-2.5 AUTO duration (the duration head)	Brick ported, never constructed	duration_head_path is REFUSED by name rather than accepted-and-ignored (#611); supplying a head cannot load one. Give num_frames or duration
LTX-2.5 arms a request CAN reach	Refused by name at the call site	The spatiotemporal latent upsampler (both flags set). Supplying that checkpoint names that arm, not the temporal one. Its operator is small; what keeps it out is that it returns [c, 2f-1, 2h, 2w] while both spatial call sites require the frame count back unchanged, so porting it would trade a named refusal for a shape complaint (#2584). The temporal-only x2 arm is ported, not refused, and so is dims=2; what stays refused in 2-D are its two contradictions, temporal_upsample and rational_resampler, neither of which upstream can run either
LTX-2.5 resolution	Off-grid sizes refused, naming the offending axis and a size you can actually pass; frames still round	--width/--height must divide 64 (two-stage) or 32 (one-stage), from the VAE factor times the phase downscale (#919). --frames rounds to 8k + 1. No size cap
LTX-2.5 arms nothing can request	Declared, not requestable	int8-convrot (ComfyUI-only), single-node multi-GPU, BetaScheduler (upstream selects no scheduler either). No flag or extra asks for these. multishot was RETIRED (absent upstream) and kLoraFusion too (now served)
Qwen3.8-27B quantized arms (Q4_K_M GGUF, its clip mmproj, the unsloth "NVFP4" = mixed-precision)	All three files ACCOUNTED against committed manifests (866 / 334 / 1968 names). Q4_K_M RUNS on CPU and its token gate vs llama.cpp b10451 FAILED (2026-08-23): tokenizer exact 6/6, generation divergent 5/6, every divergence a rank-2 loss under 0.18 logits over 288 steps, so a precision difference and not a wiring defect (evidence). No speed or memory number is admissible. NVFP4 still un-run (#1632)	BF16 gated (#915). --mmproj (ABI v22) loads a clip GGUF; a load refuses a file with a tensor no loader reads. "NVFP4" @7d6f8d4d: W4A4 loads, FP8 refused (spec)
Qwen3.8-27B, the SECOND NVFP4 artifact (r0b0tlab/...-MTP-sm121, a ModelOpt checkpoint)	2001 names ACCOUNTED per scheme against four committed manifests; LOADS, never RUN: no token gate (#1632)	@36f717a2: 208 per-tensor STATIC FP8 + 193 W4A16_NVFP4 modules, both halves load. Not the unsloth format (spec)
Multi-GPU execution	Hardware-blocked	TP proven equal to tp=1 on CPU; no 2-GPU box to run it
LoRA end to end	CPU brick landed	Unwired standalone; not usable through the server
Multimodal over HTTP	Image request served end to end on CPU; token gate, codec and batching pending	ROAD-V1-MM W1-W3 plus ENG-MM-INPUT-PIPELINE P2 (#2379) landed: Request.mm_features now reaches the forward. Open: no token-exact gate on real weights, no image codec (raw RGB only), no video/audio through the runner, and Qwen3-VL serves one sequence per step. Video/audio/multi-image still refuse with HTTP 400 rather than drop (#686)
Reranking / classify models	Engine side only	Embeddings are LIVE (LlamaModel, vllm_embed, /v1/embeddings); the classify/score heads are landed ops with no registered arch
ROCm	W0 community-verified on 5 gfx archs; classic-dense, GDN-hybrid and EXL3 e2e run all-native; correctness gaps remain	52 registered ops including the GDN state/conv/postconv/recurrence set, MoE combine/gate, keep-quant expert GEMM, the narrowing CastF16, and the EXL3 trellis GEMM (#2433); the APU managed-allocation branch is measured and now narrowed to devices that can take a recoverable page fault, which withdraws the CPU reference tier on gfx1151/gfx1103 (#2511). ROCm guide
Vulkan EXL3	Native, correctness only	kCastF16 + kExl3Gemm registered and BYTE-identical to the CPU arm on llvmpipe; zero reference-tier hits where S1 measured two (#2530). No real GPU, no checkpoint run, no speed number. kExl3MoeMlp needs a grid barrier Vulkan does not have and is owed
XPU, TPU	Not started	CUDA, CPU, Metal and Vulkan are the built backends
Custom logits processors on CUDA	Open, not root-caused	Segfaults in a CUDA build, 232/232 green on CPU
Memory budgeting (ROAD-V1-MEM, #83)	M1+M2 landed (absolute bytes)	--kv-cache-memory sizes the KV pool from an absolute byte budget (ABI v16, per-layer divisor since #1963 — the group-aware one counted placeholder names and overshot by the layer count); --num-blocks overrides; --gpu-memory-utilization needs the M3 profile run (dgx-gated). See specs/kv-sizing.md
Gemma4 MoE ROCm FP8 + SharedK-WMMA	Partial	Dual-GPU FP8 resident experts, SharedK-WMMA prefill (RDNA4); decode-graph and forward extract deferred. Env VT_GEMMA4_*/VT_ATTN_*, seam test_gemma4_rocm_fp8_seams. spec

How to read this page

A ✅ means the implementation has a named gate. For a model, that gate compares the same workload with the pinned oracle. For an engine feature, the gate is a named test in the tree. A ◐ means the path works only within the limits in its table row.

The marks describe support, not speed or current ownership. See Status for lifecycle state and Benchmarks for performance. An inventoried row is not a supported feature.