From 5775f744a000e44896ab352924d12045dd4059b9 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 13 Aug 2026 11:00:31 +0000 Subject: [PATCH] record(MODEL-MATRIX): row the 8 recipe architectures that had none (#609, #610) FOLLOWING_AGENTS_PROTOCOL The matrix claims exhaustive coverage. An audit of the 157 official model recipes at vllm-project/recipes 86c7777a, joined to each model's HF config.json architectures and then to this file, found eight architecture strings with no row. They are added here at INVENTORIED. Records only: no src/, include/ or tests/ change, and no mark claims more than its row backs. Qwen3_5MoeForCausalLM is the one worth reading. It is NOT an alias of Qwen3_5MoeForConditionalGeneration, and it is also NOT a new port. The class already exists at our pin (qwen3_5.py:381; main :443) and is precisely what the multimodal class builds as its own text tower -- Qwen3_5MoeForConditionalGeneration.__init__ does `self.language_model = Qwen3_5MoeForCausalLM(...)` at qwen3_5.py:643 (pin) / :724 (main). What vLLM main added is a REGISTRY entry (registry.py:203, in _TEXT_GENERATION_MODELS) exposing that text-only backbone standalone: same module, different class, different task category, no vision tower and no MM processor. So it earns its own row under the matrix's alias rule, which groups only exact (category, module, class) matches -- but closing it is a registration plus flat-config-descent job, not a model port. We already ship that tower token-exact; we simply bind only the string Qwen3_5MoeForConditionalGeneration (qwen3_5_moe.cpp:215) and nest the text fields under a text sub-config (hf_config.cpp:396), while Qwen/Qwen3.8-2.4T-A95B declares Qwen3_5MoeForCausalLM with those fields at top level. BailingMoeV3ForCausalLM succeeds BailingMoeV2_5ForCausalLM but is a new module rather than a class rename: V2.5 (bailing_moe_linear.py, 821 lines at the pin) pairs MLA with Bailing linear attention, whereas V3 (bailing_moe_v3.py, 1495 lines on main) keeps MLA and swaps the linear arm for Kimi Delta Attention, adds per-projection causal conv1d, grouped-topk MoE via config.n_group, a clamped SwigluStepAndMul, an fp8 quant-config hook, and refuses a checkpoint without no_kda_lora=True. Line anchors are on the row. The six #610 rows follow the MiniMaxH3DiTModel precedent, which is the one vllm-omni architecture already rowed. Four resolve in vllm_omni/model_executor/models/registry.py @ bbe6ccc5. Two do not resolve anywhere: VoxtralRealtimeForConditionalGeneration and BailingMMNativeForConditionalGeneration are the literal config.json strings of their recipes, and appear in none of the pinned registry, any of the ten dictionaries of vLLM main's registry.py, vllm-omni's _OMNI_MODELS, or vllm-omni's supported_models.md. Those two are recorded target-pending with what was searched and with the near misses named -- VoxtralRealtimeGeneration is registered at the pin and already has its own row here, and vllm-omni carries a BailingMM2NativeForConditionalGeneration alias, note the 2 -- but no anchor is asserted for either. Counts. Eight rows, so the rollup moves INVENTORIED 314 -> 322 and Total 362 -> 370, and check-agent-record's MODEL row count moves with it. The at-the-pin model inventory (324 rows / 373 memberships / 356 architectures / 310 targets / 261 modules) is deliberately UNCHANGED: none of the eight is at the pin, so like the MuseGlimmer, KimiK3 and MiniMaxH3DiT rows before them they carry no pinned-registry module::class token and contribute nothing to it. That is the existing convention, not a relaxation -- no checker expectation was widened to make this pass. The header prose distinguishes at-pin from beyond-pin rows, so it is reconciled the same way: 326 at the pin plus 3 beyond becomes 326 plus 11, = 337, with the long tail 284 -> 292 and the engaged count 45 unchanged. Two pre-existing errors in that block are corrected while it is open rather than left standing: the caption read "the 49 non-INVENTORIED rows" where the file has, and had, 48. Two notes for whoever closes the issues. First, #609 and #610 say two and seven; their own tables enumerate two and six, and reproducing the audit independently returns exactly eight, so eight is what landed. Second, that reproduction could not read 20 of the 157 recipe configs (gated or 404, including Voxtral-4B-TTS-2603, GLM-TTS, IndexTTS-2.5 and stable-audio-open), so coverage for those is unverified rather than clean, and several vllm-omni TTS architectures visible in the omni docs still have no row. Also files #618: test_cpu_x86_llamacpp_floor's contended-leg case red on this branch at loadavg 63 and green at 21.9 on the same tree -- it races the box's quiet-window gate, so it fails toward a verdict on whatever diff is in flight. Filed, not fixed here: it changes a test's semantics. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/model-matrix.md | 35 ++++++++++++++++++++++++++--------- .agents/roadmap_v1.md | 3 +++ docs/FEATURES.md | 5 ++++- scripts/check-agent-record.py | 18 +++++++++++++++++- 4 files changed, 50 insertions(+), 11 deletions(-) diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md index 17278916a..27e6b53d2 100644 --- a/.agents/model-matrix.md +++ b/.agents/model-matrix.md @@ -19,12 +19,21 @@ a practical unit that one agent can spike without silently dropping aliases. ## Architecture-support checklist At-a-glance view of which architectures we have actually engaged, and how far. -**326 architecture rows are inventoried at the pin, plus 3 beyond-pin rows -(`KimiK3ForConditionalGeneration` and `MuseGlimmerForConditionalGeneration`, both -released after the pin, and `MiniMaxH3DiTModel`, -which is a DIFFUSION architecture living in the separate `vllm-omni` repository) -= 329 total**; 45 are -past `INVENTORIED` (engaged), the remaining 284 are known-but-not-started long tail. Every mark +**326 architecture rows are inventoried at the pin, plus 11 rows that the pinned +registry does not contain = 337 total.** Those 11 are, by why they are not at the +pin: `KimiK3ForConditionalGeneration` and `MuseGlimmerForConditionalGeneration`, +both released after the pin; `Qwen3_5MoeForCausalLM` and +`BailingMoeV3ForCausalLM`, pin-lag rows registered on vLLM `main` and absent only +at the pin (#609); `MiniMaxH3DiTModel`, a DIFFUSION architecture living in the +separate `vllm-omni` repository; the four TTS / audio-generation architectures +registered in that same `vllm-omni` repository — `MossTTSDelayModel`, +`MossTTSRealtime`, `Qwen3TTSForConditionalGeneration` and +`HiggsMultimodalQwen3ForConditionalGeneration`; and two whose upstream target is +still pending because the exact `config.json` architecture string is registered +in neither core vLLM `main` nor `vllm-omni` — +`VoxtralRealtimeForConditionalGeneration` and +`BailingMMNativeForConditionalGeneration` (#610). Of the 337, 45 are +past `INVENTORIED` (engaged), the remaining 292 are known-but-not-started long tail. Every mark below is grounded in the row's lifecycle `State` cell plus its ledger evidence, and this section is CI-enforced against those rows by [`scripts/check-model-checklist.py`](../scripts/check-model-checklist.py): a mark @@ -44,7 +53,7 @@ Rollup by lifecycle state (must equal the detailed per-state row counts): | State | Rows | |---|---| -| INVENTORIED | 314 | +| INVENTORIED | 322 | | PARTIAL | 20 | | ACTIVE | 9 | | SPIKE | 7 | @@ -52,9 +61,9 @@ Rollup by lifecycle state (must equal the detailed per-state row counts): | DONE | 3 | | READY | 3 | | GATING | 1 | -| **Total** | **362** | +| **Total** | **370** | -Engaged architectures (the 49 non-`INVENTORIED` rows): +Engaged architectures (the 48 non-`INVENTORIED` rows): | Support | Architecture | Family / example | Status | Row | |---|---|---|---|---| @@ -172,6 +181,7 @@ Transformers compatibility is capability-driven and excluded from finite counts. | `MODEL-TEXT-bailing-moe-bailing-moe-for-causal-lm` | `BailingMoeForCausalLM` | `registry.py:78`; `vllm/model_executor/models/bailing_moe.py::BailingMoeForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-bailing-moe-bailing-moe-v2-for-causal-lm` | `BailingMoeV2ForCausalLM` | `registry.py:79`; `vllm/model_executor/models/bailing_moe.py::BailingMoeV2ForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-bailing-moe-linear-bailing-moe-v25-for-causal-lm` | `BailingMoeV2_5ForCausalLM` | `registry.py:80`; `vllm/model_executor/models/bailing_moe_linear.py::BailingMoeV25ForCausalLM` | causal generation / text | model loader/forward; FusedMoE/grouped GEMM; Mamba/SSM state; GDN/linear-attention state; MLA/latent KV | ☐ required | `INVENTORIED` | none | unassigned | +| `MODEL-TEXT-bailing-moe-v3-bailing-moe-v3-for-causal-lm` | `BailingMoeV3ForCausalLM` | **BEYOND-PIN (pin-lag) — NOT registered in `555967922`**; present on vLLM `main` at `registry.py:82` (`_TEXT_GENERATION_MODELS`) → module `vllm/model_executor/models/bailing_moe_v3.py`, class `BailingMoeV3ForCausalLM` (1495 lines, 9 classes), with its own MTP draft registered at `main` `registry.py:654` → module `bailing_moe_v3_mtp.py`, class `BailingMoeV3MTPModel`. Module and class are given as separate fields rather than in the compact colon-pair form: that form marks a target inside the PINNED registry inventory, and this row is beyond the pin — exactly how the `MuseGlimmerForConditionalGeneration`, `KimiK3ForConditionalGeneration` and `MiniMaxH3DiTModel` rows already render theirs. Succeeds `BailingMoeV2_5ForCausalLM` but is a **NEW MODULE, not a class rename in the old one**: V2.5 at the pin (`bailing_moe_linear.py`, 821 lines) mixes MLA full attention with Bailing linear attention (`vllm/model_executor/layers/mamba/linear/bailing_linear_attn`, imported `bailing_moe_linear.py:32`), whereas V3 keeps MLA but swaps the linear arm for **Kimi Delta Attention** (`vllm/third_party/flash_linear_attention/ops/kda` — `chunk_kda`, `fused_kda_gate`, `fused_recurrent_kda`, imported `bailing_moe_v3.py:74-79`) with a per-projection causal `conv1d` (`:650-671`), adds grouped-topk MoE through `config.n_group` (`:1140`), a clamped `SwigluStepAndMul` where a per-layer `swiglu_limit` is configured (`:387`, `:1107-1135`), an fp8 quant-config hook `_configure_ling_fp8_quant_config` (`:215`, called `:1354`), and REFUSES a checkpoint without `no_kda_lora=True` (`:603-604`). Recipe `inclusionAI/Ling-3.0-flash` declares `architectures: ["BailingMoeV3ForCausalLM"]`. Tracked by [#609](https://github.com/mudler/vllm.cpp/issues/609) | causal generation / text | model loader/forward; FusedMoE/grouped GEMM; KDA/linear-attention state; causal conv1d; MLA/latent KV; fp8 | ☐ required | `INVENTORIED` | none — cannot advance past this state until the pin advances to a revision that registers it | unassigned | | `MODEL-TEXT-bloom-bloom-for-causal-lm` | `BloomForCausalLM` | `registry.py:81`; `vllm/model_executor/models/bloom.py::BloomForCausalLM` | causal generation / text | model loader/forward; paged attention/KV | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-chatglm-chat-glmfor-causal-lm` | `ChatGLMModel`, `ChatGLMForConditionalGeneration` | `registry.py:82-83`; `vllm/model_executor/models/chatglm.py::ChatGLMForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; legacy ChatGLM2/3 lineage — hand-written attention with `multi_query_group_num` (`chatglm.py:60-63`) and a HARDCODED partial rotary factor 0.5 (`:103`, applied `:131`), on an out-of-tree config requiring `trust_remote_code` (`:34`). Shares only the NEW partial-rotary primitive with the rest of GLM. Smallest genuine checkpoint `zai-org/chatglm3-6b` 6.2B / 23.3 GiB fp16 — FITS GB10. Lowest priority in the campaign (spike W9), deferrable without blocking anything | [glm-dsa-latest-deepseek spike](specs/glm-dsa-latest-deepseek.md) | `SPIKE` | none | `CLAIM-GLM-DSA-LATEST-DEEPSEEK` | | `MODEL-TEXT-commandr-cohere-for-causal-lm` | `CohereForCausalLM`, `Cohere2ForCausalLM` | `registry.py:84-85`; `vllm/model_executor/models/commandr.py::CohereForCausalLM` | causal generation / text | impl `include/vllm/model_executor/models/commandr.h` + `src/vllm/model_executor/models/commandr{,_weights,_registry}.cpp` (weight-only Cohere `LayerNorm` grounded `commandr.py:76-87,65-73`; GPT-J full-width RoPE `is_neox_style=false` `commandr.py:174-179`; PARALLEL residual `commandr.py:257-273`; `logit_scale` `commandr.py:376`; tied embeds `commandr.py:372`; ZERO new kernel, one REGISTER line); test `tests/parity/test_commandr_paged_engine.cpp` (SACRED gate, dgx-only, SKIPs — no runnable vehicle). CPU `-Werror` build compiles+links+registers. use_qk_norm/sliding rejected at parse (Cohere2 arch). | [recent-dense-batch](specs/sweep-recent-dense-batch.md) | `BLOCKED` | W0 oracle RUN-VERIFIED (tiny-random builds+runs, golden captured, `CohereForCausalLM`≠`Cohere2`); gate BLOCKED — real ckpts HF-gated (no dgx token) + only ungated are tiny-random head_dim 8/2 (outside validated attn) + dgx disk-full | `CLAIM-SWEEP-RECENT-DENSE` | @@ -267,6 +277,7 @@ Transformers compatibility is capability-driven and excluded from finite counts. | `MODEL-TEXT-qwen2-moe-qwen2-moe-for-causal-lm` | `Qwen2MoeForCausalLM` | `registry.py:190`; `vllm/model_executor/models/qwen2_moe.py::Qwen2MoeForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-qwen3-qwen3-for-causal-lm` | `Qwen3ForCausalLM` | `registry.py:191`; `vllm/model_executor/models/qwen3.py::Qwen3ForCausalLM` | causal generation / text | **W0+W1 (2026-07-20):** registry TU `src/vllm/model_executor/models/qwen3_dense.cpp` (`REGISTER_VLLM_MODEL(qwen3, "Qwen3ForCausalLM")` + full-attention-only `MakeQwen3ForCausalLMKVCache` + stub factory) + header `include/vllm/model_executor/models/qwen3.h`; runner generalization `src/vllm/v1/worker/gpu/runner.cpp:458-470,651-680` (full-attention-only KV, model-agnostic `has_mamba_group`/`gdn_group_id_>=0` guards). **W2 (2026-07-20):** weight loader `src/vllm/model_executor/models/qwen3_weights.cpp` (`LoadQwen3ForCausalLMWeights`) + `Qwen3DenseWeights` struct in `include/vllm/model_executor/models/qwen3.h` (merged qkv/gate_up raw-NK per vLLM `packed_modules_mapping`, per-head q/k norm, tied `lm_head` aliasing `embed_tokens` — checkpoint `lm_head.weight` skipped mirroring vLLM `skip_prefixes`); shared BF16 helpers extracted to `include/vllm/model_executor/models/dense_weight_loaders.h` (27B `qwen3_5_dense_weights.cpp` load byte-identical). **W3 (2026-07-20):** dense forward `src/vllm/model_executor/models/qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from public vt:: ops + fusion catalog; per layer std add+RMSNorm → merged qkv (bf16 MatmulBT) → per-head q/k RMSNorm+RoPE (bf16, no gate) → paged attn → o_proj → post norm → SwiGLU MLP → final norm → tied lm_head; wired into the `qwen3_dense.cpp` factory. 2 NEW fusion recipes `include/vt/recipes.h` (`kFusedAddRmsNormStd` non-gemma add-RMSNorm used at 3 norm sites; `kAttnQkNormRope` non-gated qk-norm-rope composite) each byte-exact composite-tested in `tests/vt/test_ops_fused_chain.cpp`. Two GENUINE latent bugs the first pure-dense bf16 model FORCED OUT: (a) tokenizer `SplitPattern::kQwen2Classic` (`src/vllm/tokenizer/pretokenizer.{h,cpp}`, `tokenizer.cpp`) — the classic Qwen2/Qwen3 pre-tokenizer regex was hard-rejected; (b) `src/vt/cuda/cuda_paged_attn.cu` WMMA prefill gated to head_dim 256 only (the tensor-core ladder mistokenizes at other head_dims — validated only for the d=256 gate models). Tests `tests/vllm/models/test_qwen3_forward.cpp` (CPU synthetic + dgx prefill-argmax==12095 CPU+CUDA + ADOPT==fallback byte-exact), `tests/parity/test_qwen3_paged_engine.cpp` (SACRED gate) + oracle capture `scripts/qwen3-oracle-capture.py`. **W4 (2026-07-20, GATING — NOT token-exact yet):** Qwen3-0.6B greedy vs vLLM 0.25.0 oracle = **11/16 prompts token-exact** (all 16 first-tokens exact; 5 fail at late-decode SEMANTIC NEAR-TIES, e.g. "question"↔"answer"). ISOLATED as per-op bf16-rounding drift vs vLLM's kernels, NOT a structural bug: CPU-exact-f32 also 10/16 with overlapping divergences; RMSNorm rounding verified matching vLLM csrc. **16/16-EXACT IS ILL-POSED (2026-07-20 near-tie razor):** the ISOLATE-then-BIT-MATCH pass showed vLLM's OWN bf16 greedy is non-deterministic on these near-ties — vLLM ↔ its committed golden = 16·15·16 (N=3) and flips 7/16 prompts over N=10 (p6 4/10, p15 3/10; golden is the MINORITY token there); vLLM = `FLASH_ATTN`/FA2 for head_dim 128. All 5 of our divergences sit on gaps ≤0.125 nats (p5@5 is an EXACT 0.0 tie), and we emit vLLM's MAJORITY token on p6/p15. `vt::GreedyArgmax` already mirrors torch.argmax lowest-index tiebreak (`cuda_sample.cu:94`). ⇒ NOT a kernel bug and there is no fixed golden to bit-match to; a flash-exact d=128 rewrite cannot make a non-deterministic golden deterministic. REGRESSION preserved by construction (no code change): 27B `test_qwen27_paged_engine` 235/235 + 35B `test_qwen36_paged_engine` 315/315 UNCHANGED. Correct closure = a near-tie-robust distributional gate (accept if our output ∈ vLLM's K-run output set) — gate-DESIGN decision reserved for the user (spike honesty escape hatch). **W4 CORRECTNESS COMPLETE 2026-07-20** (`tests/parity/test_qwen3_paged_engine.cpp` rewritten to the near-tie-robust gate; `scripts/qwen3-neartie-gap.py` + `--runs`/`--per-prompt` in `scripts/qwen3-oracle-capture.py`; goldens `greedy_dist`/`our_ids`/`neartie_gap_mnats` for both dirs). **KEY CORRECTION: the "vLLM greedy is non-deterministic" premise was a BATCHING artifact.** Per-prompt (batch=1 — the gate's single-request regime), vLLM 0.25.0 greedy is DETERMINISTIC: `Qwen3-0.6B` 0 multi-member cells over K=10, `Qwen3-4B` 0 over K=5 (batched, all 16 prompts in one call, it flips — 0.6B 56, 4B 43 multi-member cells). So the strict token-exact bar IS well-posed. Near-tie robustness comes from TEACHER-FORCING vLLM on OUR exact prefix (`prompt_logprobs`): at all-but-2 positions vLLM's OWN argmax given our prefix IS our token with gap 0.0000 (bit-identical logprobs), proving our forward matches vLLM's prefill logits; the residual flips are bf16 near-ties (0.6B ≤0.125 nats ×2; 4B ≤0.25 nats) where vLLM's OWN one-shot prefill argmax disagrees with its incremental decode (e.g. 4B p13 tok1 prefill→13=ours, decode→11) — vLLM contradicts itself, so NO forward bug and no single 16/16 decode target. Gate PASS = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal). **`Qwen3-0.6B` 16/16 PASS** (strict token-exact 12/16 + near-tie-band 4/16, max gap 0.125 nats). **BIGGER-MODEL COMPLETE-CORRECTNESS PROOF `Qwen3-4B` (BF16, 36 layers, GQA 32/8, hidden 2560 — a DIFFERENT config than 0.6B, loads+runs on the SAME forward code): 16/16 PASS** (strict 10/16 + near-tie 6/16, max gap 0.25 nats). GATES: dgx CUDA `-Werror` 0-warn (test TU host-C++); **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (no engine source touched → preserved by construction, re-confirmed); `compute-sanitizer memcheck` 0 errors on the 0.6B gate path; CPU syntax RC=0 (gate skips off-dgx). Correctness code + evidence anchors: forward `src/vllm/model_executor/models/qwen3.cpp` + loader `src/vllm/model_executor/models/qwen3_weights.cpp`; gate `tests/parity/test_qwen3_paged_engine.cpp` + goldens `tests/parity/goldens/qwen3_greedy_{0_6b,4b}/`; ledger 2026-07-20 W4-correctness-complete row. Correctness-complete; SPEED benchmark pending (next task, precondition met) → row stays `ACTIVE` until speed parity per the DONE=correctness+speed policy. **STRICT-DECODE RAZOR investigated 2026-07-20 (worktree `agent-a7e66dd02ed286cb2`):** attempt to tighten the near-tie band to STRICT 16/16 by routing d128 decode through the vendored FA2 group-swap split-KV kernel was MEASURED WORSE (0.6B 12→11/16, 4B 10→9/16) and REVERTED — source-confirmed that vLLM runs FA2 **varlen** (`flash_attn_varlen_func`, fa_version 2 on sm_121), NOT the group-swap `flash_attn_with_kvcache` path; the real bit-match target is an FA2 VARLEN d128 decode, a scoped sub-campaign in [qwen3-decode-strict-bitmatch spec](specs/qwen3-decode-strict-bitmatch.md). **FA2 VARLEN d128 RESOLVED 2026-07-20 (worktree `agent-a31b006ff4766f03b`):** vendored the EXACT vLLM varlen decode — d128 bf16 split-KV instantiations `flash_fwd_split_hdim128_bf16_{,causal_}sm80.cu` + `LaunchDecodeVarlenFA2Bf16` (`src/vt/cuda/cuda_flash_attn_fa2.cu`: plain varlen cu_seqlens_q=query_start_loc, h=hq, is_causal=true, NO group swap, num_splits=exact heuristic) routed by `VT_FA2_DECODE_QWEN3` (**default OFF**, scoped d128 so the d256 arms are untouched); op-parity byte-exact vs f32 ref (`tests/vt/test_ops_paged_attn.cpp` `*varlen d128*`, incl. num_splits>1 split-combine) + `memcheck` 0. **The kernel BIT-MATCHES vLLM's decode attention OUTPUT** (teacher-forced logit gap **0.0000 nats** at all-but-near-tie positions) but STRICT 16/16 is NOT reached: FA2-varlen 0.6B **11/16** & 4B **9/16** — ONE WORSE than the CUDA-core fallback (12/16 & 10/16). Escape hatch CONFIRMED: strict is bf16-tie-bounded (residual flips are ≤0.375-nat ties where vLLM's own prefill argmax disagrees with its incremental decode). Shipped OPT-IN; near-tie-robust gate stays the closure; engine gate NOT tightened. 27B 235/235 + 35B 315/315 + default Qwen3 gate 16/16 UNCHANGED, `-Werror` 0-warn **NVFP4A16 / W4A16 QUANT-SCHEME EXTENSION 2026-07-21 (row `QUANT-NVFP4-CT-W4A16`, spike [sweep-qwen3-32b-nvfp4a16](specs/sweep-qwen3-32b-nvfp4a16.md), worktree `agent-a141e107f210eb6c6`, base `aa65ce7`):** the SAME dense forward now also loads and natively computes **compressed-tensors NVFP4A16 (W4A16)** on `RedHatAI/Qwen3-32B-NVFP4A16` (64L dense, hidden 5120, GQA 64/8, intermediate 25600, UNTIED BF16 lm_head — a 64-layer size-scale check for the dense path on top of 0.6B/4B). Checkpoint premise VERIFIED BEFORE implementation (unlike OPT the plan's B.1 claim was CORRECT: 5 shards / 20.6 GB / 1603 tensors). Loader: per-Linear `.weight_packed` probe + fp4 fields on `Qwen3DenseAttnWeights`/`Qwen3DenseMlpWeights` (`include/vllm/model_executor/models/qwen3.h`), append-only helpers in `dense_weight_loaders.h`; forward: two dispatch branches (`qwen3.cpp::MlpBlock`, `dense_attn_block.h::AttnBlock`) into the NEW `include/vllm/model_executor/models/dense_nvfp4_gemm.h`. `dense_device_glue.h` is a PURE VERBATIM RELOCATION of `Dev`/`DBuf`/`MakeTensor`/`Reshape`/`DevicePoolPolicy` out of `dense_attn_block.h` (same namespace ⇒ the BF16 Qwen3-dense/Coder/OPT forwards stay BYTE-IDENTICAL). **CORRECTNESS: strict scored 4/6 (67/96 tokens), the teacher-forcing isolation then SETTLED it, and the gate closes 6/6 under the ratified near-tie-robust bar WITH the evidence committed** (`tests/vllm/models/test_qwen3_32b_nvfp4a16_paged_engine.cpp`). Gate selection MEASURED in two stages: vLLM's own greedy is DETERMINISTIC (K=5, 0 multi-valued cells) so STRICT was run FIRST and was NOT loosened on that evidence; then the ratified `prompt_logprobs` TEACHER-FORCING isolation (new `scripts/qwen3-32b-nvfp4a16-neartie-gap.py` + a `VT_DUMP_IDS=1` bootstrap) measured **all 29 divergent positions at <= 0.0625 nats with 28/29 EXACTLY 0.0** — our token IS vLLM's own argmax given OUR prefix. **THIS ROW OWNS THE FINDING.** The prompt[5] root flip is an EXACT bf16 TIE (`" moon"`/`" Moon"` have bit-identical vLLM logprobs `-0.727154`) at which **vLLM's teacher-forced argmax is OUR token while vLLM's incremental greedy chose the other — vLLM contradicts ITSELF**, exactly the prefill-argmax-vs-incremental-decode disagreement this row characterized at 0.6B/4B. The prompt[2] root flip gaps 0.0625 nats, against **0.125 nats of vLLM's OWN jitter at that same position** (measured by batch-composition sweep: 0.1875 alone → 0.0625 batched) and below the **0.25-nat** gap already ratified on the UNQUANTIZED 4B; the other 27 divergences are downstream cascade. Because the PREFILL argmax is exact on all six prompts, the prefix at both root flips is BIT-IDENTICAL to vLLM's, making the gap a direct single-position logit comparison rather than a prefix-drift artefact. **⇒ the residual is THIS row's pre-existing dense bf16 near-tie drift run 64 layers deep — the DEEPEST dense model we have run — and NOT a W4A16 defect;** the quant path is exonerated four independent ways (bit-exact CPU proof max abs delta-logit = 0, `fallback_gemms=0` vs 18432 Marlin GEMMs, invariance across both quantized GEMMs — `VT_NVFP4_MARLIN=0` scores 3/6 with the SAME divergences — and the <=0.0625-nat gaps). Final: **6/6 prompts PASS (strict 4/6 + near-tie band 2/6, max gap 0.062 nats vs the 0.5-nat bar shared with the 0.6B/4B/Coder gates, 0 forward-divergent, 142/142 assertions incl. 96 hard anchor REQUIREs)**, path proven exercised (`marlin_gemms=18432`, `fused_gate_up=6144`, `fallback_gemms=0`), all six vLLM tokenizations match. **CONSISTENT SCALING EVIDENCE for this row's drift: 0.6B 60 divergent positions ALL at gap 0.0; 4B max 0.25 nats; 32B/64L max 0.0625 nats — deeper does NOT mean worse, and every one is a tie vLLM's own logits cannot separate.** BF16 behavior UNCHANGED by construction (the probe selects the fp4 arm only when `.weight_packed` exists): Qwen3-dense/Coder/OPT gates unaffected; **27B 235/235 + 35B 315/315 UNCHANGED**; CUDA `-Werror` 0-warn. | [first-additive-model spike](specs/first-additive-model-qwen3-dense.md) **SPEED — d128 FA2 PREFILL IMPLEMENTED + FA2 DECODE DEFAULT FLIPPED 2026-07-20 (`Qwen3-4B` BF16 vs vLLM 0.25.0 production/graphed, in1024/out128 ignore-eos closed-loop) — big gap-close but STILL below vLLM, stays ACTIVE:** the dominant prefill lever is now a **d128 FlashAttention-2 varlen prefill** — generalized the vendored FA2 launcher `LaunchPrefillFA2Bf16` (`src/vt/cuda/cuda_flash_attn_fa2.cu`) to head_dim 128 (reuses the already-vendored `flash_fwd_split_hdim128_bf16_{,causal_}sm80.cu` split-KV instantiations, num_splits=1 Split=false = exactly vLLM's `flash_attn_varlen_func`), routed by a new d128-scoped `fa2_prefill_qwen3` gate + `Fa2PrefillQwen3Enabled` (`VT_FA2_PREFILL_QWEN3` default-ON, `src/vt/cuda/cuda_paged_attn.cu`); AND flipped the FA2 varlen d128 DECODE default ON (`Fa2DecodeQwen3Enabled` now default-ON). Op-parity byte-exact vs the f32 ref at both Qwen3 ratios (`tests/vt/test_ops_paged_attn.cpp` `*head_dim 128*`, host-metadata bit-identical + toggle-off fallback), memcheck 0. Near-tie gate RE-PASSES **16/16 on 0.6B AND 4B** with FA2 prefill+decode ON (goldens refreshed via `scripts/qwen3-neartie-gap.py`; teacher-forced max gap 0.125 nats, 0 forward-divergent — FA2 prefill matches vLLM's prefill logits). RESULT: total tput **0.90× (c1) / 0.62× (c8)** (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, **ITL P99 0.98× WIN**); prefill-kernel A/B (same binary) = **+7%/+41% total, −55%/−48% TTFT, −22% c8 TPOT**. STILL FAIL: TTFT median **5.85×/10.2×** + total <1× because the full prefill STEP (not the attention kernel — now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). RESIDUAL LEVERS: (1) portable prefill-step fusion + graphed prefill (dominant); (2) c8 split-KV decode occupancy. Repro + table: ledger 2026-07-20 FA2-prefill SPEED row + `docs/BENCHMARKS.md`; goldens `tests/parity/goldens/qwen3_greedy_{0_6b,4b}/{our_ids,neartie_gap_mnats}.npy` refreshed for FA2-ON. 27B 235/235 + 35B 315/315 + `-Werror` 0-warn UNCHANGED. **TTFT LEVERS 2026-07-20 (base `812a57a`, worktree `agent-aea71a6a5024b17d4`) — measure-first overturned profile #81's two premises:** (L1) extracted the pooled `DevicePool` VERBATIM to shared `include/vllm/model_executor/models/device_pool.h` (gate-model `qwen3_5.cpp` + dense `qwen3.cpp` both use it) — byte-identical (27B 235/235, 35B 315/315, 0.6B/4B 16/16 UNCHANGED) but a clean same-binary A/B DISPROVED the "~44% GPU-idle from cudaMalloc/cudaFree" premise: **PERF-NEUTRAL** (async scheduler overlaps the host-side alloc syncs); kept as byte-safe hygiene. (L2) RoPE cos/sin cache — route the bf16 preamble to `RopeFromCache` off the per-step cache (identity row index; bf16 cache == vLLM RotaryEmbedding) instead of `RopeNeox`'s per-element fp64 pow/cos/sin (~36.5% prefill GPU-busy); **MEASURED the dominant dense-TTFT lever** (same-binary A/B, in1024/out128, idle box, 2 reps, vs graphed vLLM `bench serve`): **c1 median TTFT 209→135 ms = 0.87× vLLM (median parity); c8 median 316→144 ms = 0.38× vLLM**; TPOT/tput flat (prefill-only). **Shipped **DEFAULT-ON `VT_QWEN3_ROPE_CACHE`** (2026-07-20; `=0` opts back to RopeNeox)** after the opt-in blocker was GROUNDED + DISPROVEN. The claimed "1-ULP RopeFromCache shift lands on the engine's ~1-ULP FA2-split-KV near-tie NONDETERMINISM → flaky gate (RopeNeox 4/4, RopeFromCache flips)" was NOT reproducible: the paged engine is byte-DETERMINISTIC run-to-run (RoPE-off gate 4/4 + RoPE-on dumps md5-identical 3/3 + flipped-default gate 16/16 identical K=5), and for the short gate contexts `num_splits==1` so `flash_fwd_splitkv_combine_kernel` (`flash_fwd_kernel.h:1126`, launched only when `num_splits>1`) never runs; the combine is a fixed-order register reduction (no atomics) anyway. The 1-ULP FMA shift is real but deterministic — it moves 2 genuine bf16 near-tie tokens, handled by regenerating the near-tie goldens on the CANONICAL `$HOME/cutlass-4.5.0` build (the flashinfer-bundled cutlass copy tips the 27B tok6 razor to 271→234/235; cutlass-4.5.0 gives 198→235/235; Qwen3 near-ties are likewise cutlass-sensitive). Gate re-passes 16/16 both (0.6B strict 10+nt 6 max 0.0 nats; 4B strict 11+nt 5 max 0.25), deterministic K=3. RoPE flip isolated from 27B/35B (`RopeCacheEnabled`/`Qwen3DenseModel` referenced only in `qwen3.cpp`/`qwen3_dense.cpp`/`qwen3.h`): **27B 235/235 + 35B 315/315 UNCHANGED on the canonical build**. SPEED (RoPE-ON default vs vLLM 0.25.0 production): c1 total 0.97× / TPOT 1.00× / ITL P99 0.96× WIN / TTFT 2.27×; c8 total 0.82× / TPOT 1.17× / TTFT 1.90× (RoPE A/B: c1 TTFT −34%, c8 TTFT −43% / total +10%). Still `ACTIVE` — residuals = prefill-STEP non-attention glue (TTFT) + c8 split-KV decode occupancy. Anchors: `src/vllm/model_executor/models/qwen3.cpp` (`RopeCacheEnabled`), `include/vllm/model_executor/models/device_pool.h`, ledger 2026-07-20 TTFT-levers row, `docs/BENCHMARKS.md`. | `ACTIVE` (correctness COMPLETE — 0.6B + 4B near-tie gates PASS 16/16; **SPEED re-bound 2026-07-21 same-session matching-recipe vs vLLM 0.25.0 production: c1 EFFECTIVE EVERY-AXIS PARITY** — tput 0.98× / TTFT 0.90× WIN / TPOT 1.01× / P99 ITL 0.996× WIN; **c8 residual = decode** — tput 0.93× / TTFT 0.38× WIN / TPOT 1.10× / P99 ITL 1.12×. The prior "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts (fresh vLLM c1 TTFT ~152 ms, ours WINS TTFT both concurrencies); TTFT residual RESOLVED. c8 decode is 93% GPU-busy/compute-bound (~72% small-M=8 `cutlass_80_wmma` projections). **qkv-merge** (new GQA `QkvSplit` op + merged single-QKV GEMM, mirrors vLLM `QKVParallelLinear`) IMPLEMENTED + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ ships DEFAULT-OFF `VT_QWEN3_QKV_MERGE`. **cutlass verify RESOLVED: 27B 235/235 on the flashinfer-cutlass build** (prior 234/235 = build artifact). Named residual = c8 decode-GEMM efficiency, a decode-fusion sub-campaign; not yet done) **=== METAL COVERAGE 2026-07-23 (work row M3b, `CLAIM-BACKEND-METAL-M3B-1` + RCA/oracle `CLAIM-BACKEND-METAL-M3B-RCA`) — Qwen3-dense is the SECOND model on a non-CUDA backend; forward ORACLE-CONFIRMED CORRECT, NOT strict-token-exact. ===** Qwen3-0.6B generates END TO END on the Apple M4 GPU. **HONEST STATE — supersedes an earlier "16/16 strict, 4 near-ties at gap-0" claim that was UNSUBSTANTIATED (branch test had no teeth; see RCA + oracle ledger rows).** The Metal forward is a DIFFERENT but equally correct bf16 decoder that resolves the model's genuine near-ties the other way (p0 tok5 = 15344 " Italy" vs 9625 " France"; also p5 tok10, p10 tok10, p11 tok1). **ORACLE-BACKED near-tie confirmation** (the decisive measurement, `scripts/qwen3-neartie-gap.py` teacher-forcing vLLM 0.25.0 on the METAL prefix, batch=1, gpu_mem_util=0.40): all 60 Metal-vs-CUDA divergent positions are within **0.5 nats of vLLM's own argmax given the Metal prefix — MAX GAP 0.125 nats, none outside top-20**; p0 tok5 gap **0.0000** (vLLM's teacher-forced argmax on the identical prefix IS Italy 15344, contradicting its CUDA-capture France pick — the literal near-tie signature). Passes `test_qwen3_paged_engine` on Metal **16/16** against Metal's OWN oracle golden (`our_ids_metal.npy` + `neartie_gap_mnats_metal.npy`): 10 strict token-exact vs vLLM greedy + 6 near-tie-band, 0 forward-divergent, IDENTICAL gate logic to CUDA (hard anchor REQUIRE + ≤0.5-nat band, no cross-device latitude), teeth PROVEN (perturbed anchor → drift FAIL; perturbed gap→0.6 nats → band FAIL; restore → PASS). Metal execution PROVEN (device==kMETAL; all 9 ops `selections>0 ∧ declines==0`, kRopeFromCache/kPagedAttention 7168 each). Needed **3 new MSL RoPE kernels**: `kRopeFromCache` (bit-exact — the DEFAULT `VT_QWEN3_ROPE_CACHE`-ON apply) + `kRopeCosSinCache` + `kRopeNeox` → Metal **18/75** ops. Per-op stays the NMSE ≤5e-4-vs-CPU proof (RCA, all 28 layers). First ours-vs-MLX benchmark produced (Qwen3-1.7B, INDICATIVE/BLOCKED-ON-SUDO): ours a knowingly-unoptimised FLOOR ~6–11× slower decode / ~7–10× TTFT / ~2× memory vs MLX. **STRICT token-exactness on 0.6B is ILL-POSED (near-tie model); a strict Metal gate needs a bigger deterministic dense model (Qwen3-4B, not on the M4) — DEFERRED.** **CUDA GATE IS GREEN — the earlier "dgx CUDA RED / stale golden" note is DISPROVEN.** The France/Italy flip is BUILD-SENSITIVE: the PRODUCTION dgx build (FA2+Marlin+Triton+CUTLASS) resolves p0 tok5 → France 9625 and passes **16/16** (0.6B: 10 strict + 6 near-tie, max gap 0; 4B: 16/16, 11 strict + 5 near-tie, max 0.25 nats) — France anchor held, SUCCESS; only a portable-kernel-only CUDA build resolves it → Italy 15344 (matching Metal). The golden is NOT stale — it is the production build's resolution of a genuine numerical near-tie. Regressions GREEN (OPT 6/6, Qwen3-Coder 6/6; 27B/35B/DeepSeek binaries byte-identical — on dgx-CUDA the only compiled change is this one test .cpp, all Metal TUs `VLLM_CPP_METAL` AUTO→OFF). Goldens md5 `2965ef5772b556d3f3f86fedf4221b2f` UNCHANGED (the 2 Metal files are additive). | dep `MODEL-FACTORY-registry` (item 5); registration `src/vllm/model_executor/models/qwen3_dense.cpp:145`; test `tests/vllm/models/test_model_registry.cpp:76` | `CLAIM-MODEL-QWEN3-DENSE` | | `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` | `Qwen3MoeForCausalLM` | `registry.py:192`; `vllm/model_executor/models/qwen3_moe.py::Qwen3MoeForCausalLM` | causal generation / text | **W0+W1 (2026-07-21):** registry TU `src/vllm/model_executor/models/qwen3_moe_registry.cpp` (`REGISTER_VLLM_MODEL(qwen3_moe, "Qwen3MoeForCausalLM")` + full-attention-only `MakeQwen3MoeKVCache` + `ParseQwen3MoeConfig` MoE-field validation + `is_dense_model=false`; load throws "W2", forward throws "W3") + header `include/vllm/model_executor/models/qwen3_moe.h` (`Qwen3MoeWeights` reusing `Qwen3DenseAttnWeights` + `MoeBlockWeights`; `Qwen3MoeModel::Forward/ForwardDevice` decls). W1 behaviour-preserving refactors: (#1) dense self-attention block + device glue EXTRACTED verbatim from `qwen3.cpp` to `include/vllm/model_executor/models/dense_attn_block.h` (namespace `vllm::dense_attn`; `qwen3.cpp` `using namespace dense_attn` — Qwen3-dense byte-identical); (#2) bf16 `MoeBlock` EXPOSED cross-TU via `RunMoeBlock` (`include/vllm/model_executor/models/qwen3_5_moe_block.h`, defined in `qwen3_5.cpp` outside anon-ns, `MoeBlockOutput` owning carrier — 35B untouched); (#3) no-shared-expert GUARD in `qwen3_5.cpp` `MoeBlock` (`shared_expert_intermediate_size==0` → skip `SharedExpert`, nullptr to `vt::MoeCombine`; inert for the shared-expert-having 35B). Tests `tests/vllm/models/test_model_registry.cpp` (registry resolves `Qwen3MoeForCausalLM` + full-attn-only KV spec + `is_dense_model==false`; count 3→4 + arch-list msgs). **W2+W3 (2026-07-21):** BF16 loader `src/vllm/model_executor/models/qwen3_moe_weights.cpp` (`LoadQwen3MoeForCausalLMWeights` — merged qkv/o + per-head q/k norm via `dense_weight_loaders.h`; NEW bf16 per-expert loader `LoadQwen3MoeBlock` = router gate + 128 experts × gate/up/down transposed to Matmul-B, NO shared expert; UNTIED `lm_head` loaded separately, NOT aliased) + forward `src/vllm/model_executor/models/qwen3_moe.cpp` (`Qwen3MoeModel::Forward/ForwardDevice` = the dense ForwardBody with the per-layer MLP replaced by the MoE block: embed→N layers [std add+RMSNorm → reused dense `AttnBlock` → std add+RMSNorm → `RunMoeBlock` router softmax+top8+renorm + bf16 ExpertMlp + combine, no shared]→final norm→untied lm_head `vt::Matmul`) wired into the registry factory (W2/W3 throw stubs replaced). Tests `tests/vllm/models/test_qwen3_moe_load.cpp` (all 18867 tensors mapped/shaped, NO leftover, 131746 assertions) + `tests/vllm/models/test_qwen3_moe_forward.cpp` (CPU synthetic finite+deterministic, fusion ADOPT byte-identical, real-ckpt prefill argmax=**12095 " Paris"** — correct, not a near-tie). GATES: dgx CUDA `-Werror` **0-warn**; W2 load gate + W3 doctest green; **27B 235/235 + 35B 315/315 UNCHANGED**; memcheck **0 err / 0 leak**. **W4 correctness gate (SACRED, 2026-07-21):** paged-engine greedy vs vLLM 0.25.0 (`tests/vllm/models/test_qwen3coder_paged_engine.cpp` + goldens `tests/parity/goldens/qwen3coder_greedy/` + `scripts/qwen3coder-oracle-capture.py`/`qwen3coder-neartie-gap.py`). vLLM's own greedy is DETERMINISTIC (K=5, 0 multi-valued cells) → STRICT-where-well-posed near-tie-robust gate (dense methodology). **6/6 prompts PASS** (138 assertions): STRICT token-exact 4/6 + near-tie-band 2/6, **max teacher-forced gap 0.125 nats** (≪ 0.5-nat bar), **0 forward-divergent**; the 2 near-tie flips (prompt[3] "largest planet", prompt[4] "chemical symbol") ARE vLLM's OWN argmax on our prefix (gap 0.0) — prefill-vs-decode near-ties, not bugs. Oracle uses vLLM's TRITON MoE backend (FlashInfer-CUTLASS OOMs the 57 GiB model on 119 GiB unified). GATES: `-Werror` 0-warn; 27B 235/235 + 35B 315/315 UNCHANGED; memcheck 0. **W5 fast BF16 grouped-MoE GEMM + SPEED (2026-07-21, LANDED / gate MISSED):** NEW op `vt::MoeGroupedGemmBf16` (`include/vt/ops.h` `OpId::kMoeGroupedGemmBf16` + `src/vt/ops.cpp` + `src/vt/cuda/cuda_matmul_nvfp4.cu` `MoeGroupedGemmBf16Naive`/`MoeGroupedGemmBf16Wmma`/`LaunchGroupedBf16`) — the dtype-native analog of `MoeGroupedGemmNvfp4`, STRUCTURALLY REUSING its expert counting-sort + ragged per-BM-tile scheduling (`MoeHist`/`MoeOffsets`/`MoeScatter`/`MoeTileMap` + `EnsureMoeScratch`) with the fp4 decode replaced by a direct bf16 `[K,N]` read; 3 regimes (naive `P<32`, BM=16 decode tile mirroring vLLM `fused_moe` `BLOCK_SIZE_M=16`, BM=64 prefill tile). `MoeBlockBf16Cuda` (`qwen3_5.cpp`) replaces the per-expert host-gather loop with ~3 on-device grouped launches, **DEFAULT ON** (`VT_MOE_BF16_FAST=0` rolls back) per [[parity-enablers-ship-as-defaults]]; `MoeBf16FastLayoutOk` GUARDS it to the `nk==false` `[K,N]` layout so the `nk==true` 35B-MTP producer falls through to the reference loop. HOST-MIRROR RELEASE via the existing `OwnedTensor::ReleaseHost()`/`platforms::ShouldReleaseHostWeights` mechanism (experts = ~57 GiB = ~94% of the checkpoint; host+device were both resident in ONE 119 GiB unified pool). NEW test `tests/vt/test_ops_moe_grouped_bf16.cpp` (all 3 launch regimes + both out dtypes + identity/row_map) **4/4**. CORRECTNESS RE-VERIFIED AND STRICTER: the grouped GEMM's f32 accumulation order re-resolves bf16 near-ties, so goldens were re-measured by the ratified procedure — `test_qwen3coder_paged_engine` **6/6**, STRICT token-exact **4/6 → 5/6**, max teacher-forced gap **0.125 → 0.0000 nats**, 0 forward-divergent (a grouped GEMM lands closer to vLLM's own Triton grouped `fused_moe` than a per-expert loop). SAME-BINARY A/B: median TTFT 3092 → 904 ms (3.4×), TPOT 64.4 → 40.4 ms (1.59×). BINDING GRID vs GRAPHED vLLM 0.25.0 (triton MoE, `vllm serve`+`bench serve`, 1024/128, c1/c2/c4/c8): median TTFT **0.41/0.15/0.18/0.24×**, median TPOT **0.80/1.02/0.75/0.55×**, output tput **0.65/0.75/0.63/0.49×** — **every-axis parity MISSED** (only c2 TPOT). `nsys` attribution: prefill `MoeGroupedGemmBf16Wmma` = **56.0%** of GPU time (≈855 ms of our 904 ms TTFT; ~4.3 TFLOP/s ≈ **1.7% of GB10 bf16 peak** vs vLLM's Triton ~10.0 TFLOP/s) because the tile is untuned (BK=32, no `cp.async` double-buffer, BN=64); decode `MoeGroupedGemmBf16Naive` = **21.3%** at **151 GB/s (55% of peak)** for gate/up, block-starved at N=768 (24 blocks). Named remaining levers: w13 gate+up fusion, deterministic split-K, bf16 decode CUDA graph (fp4-gated today), prefill tile/pipeline rework. GATES: `-Werror` 0-warn; 6/6; **27B 235/235 + 35B 315/315 UNCHANGED**; memcheck **0 errors** on the MoE GEMM + engine path (`test_llm_engine` 5/5, `test_runner` 14/14). **W6 MoE-GEMM tile/pipeline rework + deterministic split-K (2026-07-21):** NEW `MoeGroupedGemmBf16WmmaPipe` (`src/vt/cuda/cuda_matmul_nvfp4.cu`) replaces the W5 tile's UNCOALESCED `[K,N]` weight stage (consecutive lanes `n_cols` elements apart) with a k-major `[BK][BN]` shared tile read as `wmma::matrix_b` ROW_MAJOR so the global read runs along contiguous N, plus 16-byte vectorized stages, a 3-deep `cp.async` multi-stage pipeline, BN 64→128 and +8-half padded shared rows — tile/pipeline shape ported from vLLM `fused_moe.py:294` `fused_moe_kernel` + `:1238` `get_default_config` (bf16 branch `block_n`/`num_stages`/`num_warps`), pipeline structure from CUTLASS `mma_multistage.h`; shape-guarded (`Bf16PipeShapeOk`, 8-bf16 row pitches) with the W5 tile retained for ragged pitches; `VT_MOE_BF16_PIPE=0` rolls back. NEW `MoeGroupedGemmBf16NaiveSplitK` + `MoeGroupedGemmBf16SplitKReduce` + `EnsureMoePartials` (graph-safe f32 partials, same retire-don't-free contract as `EnsureMoeScratch`) give the small-P decode GEMM DETERMINISTIC split-K — mirrors vLLM's `SPLIT_K` (`fused_moe.py:338`) but with fixed ascending-split partial reduction, explicitly NOT `atomicAdd` (which would break greedy reproducibility); `VT_MOE_SPLIT_K=0` rolls back. KERNEL A/B (same binary, E=128/K=2048/top-8): prefill T=1024 N=768 **4.50 → 12.15 TFLOP/s (2.70x)**, N=2048 **4.49 → 12.48 TFLOP/s (2.78x)**, decode-WMMA P=64 **2.84x**, split-K P=8 **2.63x** — vs vLLM's Triton `fused_moe` ~10.0 TFLOP/s, so the MoE GEMM now runs **~1.2x vLLM's rate** and the W5 headline kernel deficit is CLOSED (19.9 TFLOP/s in situ on the c2 2048-token prefill step). CORRECTNESS UNCHANGED — the pipelined tile preserves the per-output K-reduction ORDER exactly, so it is BIT-IDENTICAL to W5: **6/6, STRICT 5/6, max gap 0.0000 nats, 0 forward-divergent**, ZERO token movement, goldens md5-verified unchanged. `test_ops_moe_grouped_bf16` extended **4 → 7 cases / 19 assertions** (aligned-pitch pipelined prefill + decode + down-projection shapes with ragged K/N tails exercising the `cp.async` zfill, plus a run-to-run BIT-REPRODUCIBILITY assertion on the split-K reduction); a genuine pipeline bug (missing `__pipeline_commit()` on tail iterations → `__pipeline_wait_prior` under-count → last k-tiles consumed before their copies landed) was caught by the new decode case and fixed. DENOMINATOR CORRECTION: W5's c2/c4/c8 vLLM numbers were taken from ONE shared `vllm serve` replaying seed-0 prompts (vLLM `RandomDataset` builds `(offset+index+arange) % vocab_size`, `datasets/datasets.py:557-566`), giving **43.6-63.5% prefix-cache hits** and an INFLATED denominator (vLLM's c2 TTFT came out faster than its c1); W6 re-measures with a FRESH server per concurrency, **0.0% hit rate** in all four runs (c1 was always clean: 324.04 → 321.88 ms). BINDING GRID vs GRAPHED vLLM 0.25.0 (clean denominator, 1024/128, c1/c2/c4/c8, >1 = we win): median TTFT **1.03/0.82/1.05/1.19x**, median TPOT **0.88/1.01/1.22/1.05x**, median ITL **0.88/1.03/1.22/1.09x**, output tput **0.77/0.90/1.09/0.99x** — **12 of 16 cells at/above vLLM; c4 passes EVERY axis; c8 passes 3 of 4**. `nsys`-attributed residual: (a) c1 decode = the MISSING bf16 DECODE CUDA GRAPH, not kernel speed — summed GPU kernel time ~31 ms vs 36.22 ms TPOT ⇒ **~86% GPU-busy, ~5 ms/step host tax** ≈ the entire 4.5 ms c1 deficit, while the decode kernels run at 190-211 GB/s = 70-77% of GB10's ~273 GB/s; (b) c2 TTFT is NO LONGER the MoE GEMM (56.0% → 16.6% of GPU time at c1; 374 ms of a 533.6 ms TTFT at 19.9 TFLOP/s) but the NON-MoE prefill glue. GATES: `-Werror` **0-warn** clean full rebuild; 6/6 + 7/7 + NVFP4 9/9 + forward 3/3 on the FINAL binary; **27B 235/235 + 35B 315/315 UNCHANGED**; memcheck **0 memory errors** **W7 bf16 DECODE CUDA GRAPH (2026-07-21) — W6 lever (i):** NEW `Qwen3MoeDecodeGraph` (`include/vllm/model_executor/models/qwen3_moe.h` + `src/vllm/model_executor/models/qwen3_moe.cpp`), dispatched from `qwen3_moe_registry.cpp::ForwardQwen3MoeForCausalLM` on `pure_decode && is_cuda()`. Third sibling of the SAME in-tree driver design as `Qwen3_5DecodeGraph` (35B, `qwen3_5.cpp:5902`) and `Qwen3_5DenseDecodeGraph` (27B, `:6104`) — same cold->warm->capture->replay state machine, per-padded-size `SizeSlot` (own fixed-address host inputs + persistent embed/logits buffers + instantiated graph), `cols_changed` invalidation, and shared `DecodeGraphSizes`/`PadToCaptureSize`. Ported from `vllm/v1/worker/gpu_model_runner.py::GPUModelRunner` (warm-up then capture, per-step graph dispatch) + `vllm/compilation/cuda_graph.py::CUDAGraphWrapper.__call__` (pad-to-nearest-captured-size) @ `e24d1b24`. TWO differences, both because Coder is pure full attention: attention-only padded-input builder `BuildPaddedDecodeAttn` (GDN-free analog of `qwen3_5.cpp:5843`, so EVERY captured size is usable — no GDN state-slot gate), and no defensive copy of the persistent hidden (the Coder layer loop only READS it). `ForwardBody` split into `EmbedInto` (OUTSIDE the graph — `vt::Embedding` allocates a device flag and syncs the stream, `cuda_ops.cu:525,535`) + the capturable `ForwardLayers`; refactor verified behavior-preserving (`VT_QWEN3MOE_CUDAGRAPH=0` -> 6/6). **ONE REAL BUG FOUND+FIXED — a latent graph-safety defect in shared glue:** `BuildStepInputs` (`dense_attn_block.h`) uploaded the RoPE identity row-index from a **stack-local** vector; under capture that becomes a memcpy node baking a host source address that is dead stack on return, so every replay re-read freed memory -> wrong RoPE -> a WRONG token (`prompt[0] tok=2`, 3555 vs 576). Fixed by serving it from a process-persistent per-T table (contents byte-identical; storage created once per T, never resized/moved) — the host-side analog of the `EnsureMoeScratch`/`EnsureMoePartials` retire-don't-free contract. Latent, not a regression: the 35B/27B use `BuildStepDevInputs` and Qwen3-dense had no graph, so Coder is the first caller to capture this glue. **MEASURED (nsys `--cuda-graph-trace=node`, c1):** GPU-busy **~86% -> ~92%**, host tax **~5 -> ~2.7 ms/step**, median TPOT **36.22 -> 34.13 ms**; 60 `cudaGraphLaunch` (one per decode step) vs 3,569 total `cudaLaunchKernel` in the window (was ~1.4k/step eager). **CORRECTNESS: ZERO TOKEN MOVEMENT** (as a numerics-neutral graph must be) — 6/6 (138 assertions), STRICT 5/6, max gap 0.0000 nats, goldens md5-verified unchanged. **BINDING GRID** vs GRAPHED vLLM 0.25.0, FRESH server per concurrency, **0.0% prefix-cache hit VERIFIED in all four serve logs**, 2 reps ours with the cold leg discarded (>1 = we win): median TTFT **1.02/0.91/1.05/1.17x**, median TPOT **0.93/1.12/1.27/1.08x**, median ITL **0.93/1.10/1.28/1.14x**, output tput **0.84/0.96/1.11/1.03x** — **11/16 at/above vLLM; c4 AND c8 now pass EVERY axis** (c8 output tput crossed 0.99 -> 1.03x). Ours vs our own W6: TPOT **36.22->34.13 / 44.12->39.68 / 60.74->57.97 / 81.41->78.22 ms**, output tput **22.62->24.60 / 37.24->39.85 / 55.03->56.34 / 86.02->90.10 tok/s**, TTFT unchanged within noise (a decode graph cannot move TTFT). Denominator reproducibility: 3 of 4 vLLM cells within 0.5% of W6 (c1 TPOT IDENTICAL at 31.75) — the exception is **c2 TTFT, 435.5 -> 485.7 ms (+11.5%)**, so that cell is a band (0.82-0.91x), not a point. **RESIDUAL, attributed:** all 5 misses are at c1/c2. c1 = the remaining **~2.7 ms/step HOST tax** — vLLM's entire 31.75 ms decode step costs about what our KERNELS ALONE cost (~31.5 ms), so c1 needs no kernel work. **REFUTED hypothesis:** it is NOT CUDA-API overhead — `cuda_api_sum` shows `cudaStreamSynchronize` **median 4.5 us** and the 19,184 `cudaMalloc` calls are one-time `ResidentWeight` uploads, with `cudaGraphLaunch` at 438 us/step; so <1 ms of the 2.7 ms is CUDA API, and the balance is engine-side per-step host bookkeeping needing a **CPU-side** profile. c2 TTFT/tput are untouched by a decode lever and remain the non-MoE prefill glue. GATES: `-Werror` **0-warn** clean full rebuild; 6/6 graph-ON + 6/6 graph-OFF + 7/7 + NVFP4 9/9 + forward 3/3 + Qwen3-dense green; **27B 235/235 + 35B 315/315 UNCHANGED**; memcheck **0 errors** on the runner decode path, the engine path, AND the full Coder decode-graph gate | [sweep-qwen3-coder-30b spike](specs/sweep-qwen3-coder-30b.md) | `PARTIAL` (W0-W7 landed; correctness DONE 6/6 zero-movement; SPEED 11/16 grid cells at/above graphed vLLM — c4 AND c8 now pass EVERY axis — with the 5 residual cells all at c1/c2 and attributed to the remaining ~2.7 ms/step engine-side HOST bookkeeping (c1, ~92% GPU-busy) and the non-MoE prefill glue (c2)) | breadth-sweep Tier-1; dep `ENG-RUNNER-MODELSHAPE` (composes zero runner change) + Qwen3-dense attn + 35B MoeBlock; registration `src/vllm/model_executor/models/qwen3_moe_registry.cpp:181`; test `tests/vllm/models/test_model_registry.cpp:77` | `CLAIM-MODEL-QWEN3-CODER` | +| `MODEL-TEXT-qwen3-5-qwen3-5-moe-for-causal-lm` | `Qwen3_5MoeForCausalLM` | **BEYOND-PIN (pin-lag) — NOT registered in `555967922`**; present on vLLM `main` at `registry.py:203` (`_TEXT_GENERATION_MODELS`) → module `vllm/model_executor/models/qwen3_5.py`, class `Qwen3_5MoeForCausalLM` (module and class given as separate fields rather than in the compact colon-pair form, because that form marks a target inside the PINNED registry inventory and this row is beyond the pin, as the other beyond-pin rows already render theirs). **NOT an alias of `Qwen3_5MoeForConditionalGeneration`, and not a new port either.** The class already exists at the pin (`qwen3_5.py:381`; `main` `:443`) and is exactly what the multimodal class builds as its own text tower — `Qwen3_5MoeForConditionalGeneration.__init__` does `self.language_model = Qwen3_5MoeForCausalLM(...)` at `qwen3_5.py:643` (pin) / `:724` (`main`). So what `main` added is a REGISTRY ENTRY exposing that text-only backbone standalone: same module, DIFFERENT class, DIFFERENT task category (text generation vs multimodal), no vision tower and no MM processor — which is why the matrix's alias-grouping rule (same `(category, module, class)` target) gives it its own row rather than folding it into `MODEL-MM-qwen3-5-qwen3-5-moe-for-conditional-generation`. Recipe `Qwen/Qwen3.8-2.4T-A95B` declares `architectures: ["Qwen3_5MoeForCausalLM"]` with the text fields at the TOP LEVEL rather than nested under `text_config` | causal generation / text | model loader/forward; paged attention/KV; FusedMoE/grouped GEMM; GDN/linear-attention state; hybrid full/linear layer schedule; gated attention output | ☐ required | `INVENTORIED` | none for THIS architecture string — the shared Qwen3.5-MoE tower it names is already implemented and token-exact under `MODEL-MM-qwen3-5-qwen3-5-moe-for-conditional-generation`, but our factory binds only the string `Qwen3_5MoeForConditionalGeneration` (`src/vllm/model_executor/models/qwen3_5_moe.cpp:215`) and our config descent nests the text fields under a text sub-config (`src/vllm/transformers_utils/hf_config.cpp:396`), so a checkpoint declaring `Qwen3_5MoeForCausalLM` is refused today. Closing the row is a registration + flat-config-descent job, not a model port; it waits on the pin advancing. Tracked by [#609](https://github.com/mudler/vllm.cpp/issues/609) | unassigned | | `MODEL-TEXT-sarvam-sarvam-mo-efor-causal-lm` | `SarvamMoEForCausalLM` | `registry.py:194`; `vllm/model_executor/models/sarvam.py::SarvamMoEForCausalLM` | causal generation / text | model loader/forward; FusedMoE/grouped GEMM; MLA/latent KV | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-sarvam-sarvam-mlafor-causal-lm` | `SarvamMLAForCausalLM` | `registry.py:195`; `vllm/model_executor/models/sarvam.py::SarvamMLAForCausalLM` | causal generation / text | model loader/forward; FusedMoE/grouped GEMM; MLA/latent KV | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-TEXT-seed-oss-seed-oss-for-causal-lm` | `SeedOssForCausalLM` | `registry.py:196`; `vllm/model_executor/models/seed_oss.py::SeedOssForCausalLM` | causal generation / text | model loader/forward; paged attention/KV; sliding-window attention | ☐ required | `INVENTORIED` | none | unassigned | @@ -413,6 +424,12 @@ Transformers compatibility is capability-driven and excluded from finite counts. | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | `MuseGlimmerForConditionalGeneration` | **BEYOND-PIN — NOT in `555967922`** (Muse Glimmer released 2026-08-08, after the pin, and NOT on vLLM `main` either); the only upstream implementation is the OPEN, approved-but-CI-red PR [vllm#51655](https://github.com/vllm-project/vllm/pull/51655) at head `075d645af` — `vllm/model_executor/models/muse_glimmer.py` (text attn `:1083`, decoder layer `:1212`, vision encoder `:692`, adapter `:1036`, MM processor `:262`), `vllm/transformers_utils/configs/muse_glimmer.py`, `vllm/transformers_utils/processors/muse_glimmer.py`, DFlash hooks in `qwen3_dflash.py:75-94` + `v1/spec_decode/dflash.py:83-95` | conditional generation / image+video | model loader/forward; MM processor; vision encoder + merge; video path; speculative decoding (DFlash); reasoning + tool parsers | [muse-glimmer spec](specs/muse-glimmer.md), issue [#268](https://github.com/mudler/vllm.cpp/issues/268) | `SPIKE` | **W-SPEC (2026-08-10, `row/MODEL-MUSE-GLIMMER`, records+spec only, NO code).** Traced end-to-end from PR #51655 @ `075d645af`. `meta-models/Muse-Glimmer-30B`, Apache-2.0, bf16, `model_type: muse_glimmer`; registry maps BOTH `MuseGlimmerForConditionalGeneration` and `MuseGlimmerForCausalLM` onto one class. **Text tower** L=52, H=6656, 32 q-heads / 2 kv-heads (GQA 16:1), head_dim 128, vocab 202048, 131k ctx, rope theta 5e5: Gemma2-style SANDWICH norms (input/post-attn/pre-ffn/post-ffn) computed in fp32 with a BAKED `+1` weight offset and SPLIT eps (pre-norms `rms_norm_eps`, post-norms the smaller `post_norm_eps`) (`muse_glimmer.py:1236-1247`); **iRoPE** — `no_rope_layers[i]==1` ⇒ RoPE AND sliding-window, `==0` ⇒ NoPE AND full attention, sliding×3 then full every 4th (`:1114-1116,:1167-1168`) — the ONE mechanism with no local analogue (we have no Llama-4); WEIGHTLESS QK-norm over head_dim in fp32 applied BEFORE RoPE (`:1189-1196`); post-QK-norm QUERY PRE-SCALE `scale_query_by`≈3.87 with softmax scaling left at `head_dim**-0.5` (`:1112,:1192`); per-head sigmoid ATTENTION OUTPUT GATE whose gate reads the LAYER INPUT, not the attn output (`:1203-1206`); no logit softcapping. **Two named correctness traps:** (a) the query pre-scale ships under two schemas — native raw `qk_scale_factor`≈43.784 vs modular PRE-FOLDED ≈3.87 — disambiguated upstream BY MAGNITUDE against `sqrt(head_dim)` (`:472-517`); mis-reading it scales every query by 11.3×; (b) `use_qk_norm` / `use_attn_output_gate` read as `None` (not `True`) in the modular schema and only an explicit `False` disables them (`:456-469`), so a naive `getattr(...,False)` silently drops BOTH mechanisms while still emitting plausible text. **Perception encoder** L=50, H=1536, 16 heads (head_dim 96), patch 14×14, `patch_temporal`=2, 32×32 learned pos-emb grid, interleaved window/full per `layer_types`, projector 4096→6144, image tok 200092 / video tok 200091, placeholders `<\|patch\|>`/`<\|image\|>`/`<\|video\|>`: LINEAR patchify (`conv1_linear`, NOT a conv) (`:696,:710`), BILINEAR pos-emb interpolation with per-corner validity masking and a half-pixel `+0.5/-0.5` convention (`:761-820`), 2D RoPE with **width-first** `cat([freq_w,freq_h])` (`:741-759`), block-windowed attention via a `-1`-padded permutation whose per-block valid counts become `seq_lens` (`:844-867`), pixel-shuffle downsample asserting `output_dim == hidden*merge^2` (`:822-842`), plain `LayerNorm` (not RMSNorm) at ln_pre/ln_post. **DFlash** adds NO new drafter — PR #51655 only recognises `MuseGlimmerAssistantModel` as method `dflash` and threads the TARGET's `is_neox_style` into the draft config, because a draft/target RoPE-layout mismatch is SILENT (acceptance collapses, nothing errors, output stays correct). **REUSE MAP:** sandwich norms → `gemma2.cpp`/`gemma4.cpp`; output gate → `qwen3_5.cpp`; weightless QK-norm → Qwen3 family; windowed vision tower → `qwen3_vl_vision.cpp`; gated MLP → `layers::MlpGateUpMethodBase`+`vt::MergedGemmGroup`; decode → `ModelRegistry::Forward`+`dense_attn::AttnBlock`; fusion → `vt::FusedChain`; DFlash → the existing speculator row. **HONESTY:** the pinned oracle CANNOT load `muse_glimmer` (and the checkpoint wants transformers 5.15.0.dev0 vs the pin's 5.14.1), so there is NO gateable oracle and **every performance axis is an OPEN GAP by construction** — no parity/throughput claim may be made from this row until #51655 merges and the pin advances. Correctness gates against the HF reference instead, with per-mechanism RED-first mutation carrying the weight the missing oracle would have. Weights (~60 GiB bf16) NOT downloaded; GB10 fit needs a quantized arm. W-plan W0-W7 in the spec. **W0 CPU SCAFFOLDING LANDED (2026-08-10, `CLAIM-MUSE-GLIMMER-W0`, CPU-only, no weights, no GPU):** additive TUs register BOTH `MuseGlimmerForCausalLM` and `MuseGlimmerForConditionalGeneration` onto ONE factory (mirroring registry.py @ #51655); `ParseMuseGlimmerParams` descends the canonical NESTED layout AND normalizes the older FLAT layout (configs/muse_glimmer.py:186-305 — without it a flat config silently deserializes to an ALL-DEFAULT text config, ignoring every checkpoint value with no error); `NormalizeMuseGlimmerWeightName` ports the `hf_to_vllm_mapper` (:1389-1425) for BOTH checkpoint conventions; `EnumerateMuseGlimmerTensors` is the structural name map, deliberately OMITTING the three WEIGHTLESS modules (`embed_norm` :1286, per-head `qk_norm` :1121, `perception_emb_norm` :1470) that ship no tensor; forward REFUSES-by-name. Gate `tests/vllm/models/test_muse_glimmer_scaffold.cpp` **11/11 cases, 73/73 assertions**, clean CPU `-Werror` build. **RED-FIRST MUTATION-PROVEN, all four named traps** (each mutated in-tree, gate went RED, tree restored byte-identical): (1) treating the native raw `qk_scale_factor` as pre-folded → 3 RED (the 11.3x query blow-up); (2) defaulting the absent `use_qk_norm`/`use_attn_output_gate` to FALSE → 4 RED; (3) counting the iRoPE mask FORWARD instead of backward from the last layer → 5 RED; (4) applying the legacy-guac sandwich-norm renames in the wrong order (which SWAPS post-attention with pre-feedforward) → 1 RED. KV-cache spec is a documented W0 placeholder (one full-attention group; the real sliding/full split rides the Gemma-4 per-layer seam at W1). **W1-W7 FOLLOWED ON THE SAME BRANCH (`row/MODEL-MUSE-GLIMMER`, PR #279, NOT merged), and this paragraph's "config/name-map/registry only, NO forward" is superseded:** W1 the 52-layer text forward (sandwich norms w/ split eps, iRoPE, weightless QK-norm before RoPE, query pre-scale, attn output gate, SwiGLU, untied lm_head, output_multiplier before the soft-cap) gated vs an independent fp32 transcription of `075d645af` plus a property test per mechanism (`test_muse_glimmer_text`, 21 cases / 487 assertions, and the SAME binary re-run under `VT_FUSED_CHAIN_ADOPT=0` so the non-FusedChain fallback arm is gated too); W3 the 50-layer perception encoder (`test_muse_glimmer_vision`); W4 the mm wiring — tower → adapter → projection → `perception_emb_norm` → masked scatter onto the image/video placeholder rows, with the released 30B's 1436 tensors accounted 1436/1436 (`test_muse_glimmer_wiring`, 9/10316); W7 the ATEM reasoning + tool parsers (34 cases / 251 assertions). **REAL-WEIGHT EVIDENCE, and its exact limits:** on real 30B tensors at **reduced depth 4 of 52** our forward's 5 prefill argmax positions are identical to our torch transcription AND to HF's own `muse_glimmer` (`exportable-muse` @`a9e337e8`); max abs diff 0.0889745, cosine 0.999981. NOT established, and not to be implied: the full-depth 52-layer arm NEVER RAN; these are prefill argmax positions, NOT generated tokens, so multi-step decode and the sliding window across steps are untested; the perception encoder has NO reference check of any kind (the wiring gate proves reachability and placement, not that an image produces the right tokens); nothing has run end to end through the server; and the ATEM parsers' channel scoping DOES NOT WORK at server defaults, because the ToolParser/ReasoningParser seam has no `adjust_request` dispatch site while `skip_special_tokens` defaults true and the checkpoint marks `<\|start\|>`/`<\|message\|>`/`<\|eom\|>`/`<\|eot\|>` special (open gap, spec §6.7). **NO vLLM SPEED AXIS ON ANY DIMENSION** — unchanged and unchangeable until #51655 merges and the pin advances. A SECONDARY, quant-matched llama.cpp bar now EXISTS and is measured (#333, 2026-08-11, `row/MUSE-BENCH-2`, spec §14): same 16.76 GB Q4_K_M file both engines on an idle GB10 CPU, 128-token prefill **tie at 0.997x** (0.898x at 20 threads), decode **0.232x**, 512-token prefill 0.170x, peak RSS **1.92x MORE** (30.29 vs 15.74 GiB, the §10.2 dequantized qkv/lm_head/embed). It is a secondary reference and is NEVER the denominator; the vLLM cell stays an open gap by construction. Row token stays `SPIKE` deliberately: advancing it owes `docs/STATUS.md` and `.agents/NOW.md` rows and both surfaces sit byte-exactly at their shrink-only ratchet/budget, so the advance belongs to the landing commit | `CLAIM-MUSE-GLIMMER-SPEC`, `CLAIM-MUSE-GLIMMER-W0`, `CLAIM-MUSE-GLIMMER-W1`, `CLAIM-MUSE-GLIMMER-W3`, `CLAIM-MUSE-GLIMMER-W4`, `CLAIM-MUSE-GLIMMER-W7`, `CLAIM-MUSE-GLIMMER-FIX` | | `MODEL-MM-kimi-k3-kimi-k3-for-conditional-generation` | `KimiK3ForConditionalGeneration` | **BEYOND-PIN — NOT in `555967922`** (K3 released 2026-07-27, after the pin); closest registered = its literal text backbone `KimiLinearForCausalLM` (`registry.py:140`; `vllm/model_executor/models/kimi_linear.py`) + the K2.5 vision wrapper `kimi_k25.py:290` / tower `kimi_k25_vit.py` | conditional generation / image (text-first) | model loader/forward; FusedMoE/grouped GEMM; GDN/linear-attention state (KDA); MLA/latent KV; MXFP4 compressed-tensors quant; MM processor + MoonViT-V2 encoder/merge | [kimi-k3 spike](specs/kimi-k3.md) | `SPIKE` | **W0 SCOPE (2026-07-28, `CLAIM-KIMI-K3-SCOPE`, DERIVE-AND-SHIP, records-only).** From the HF `config.json` (fetch-derived): `architectures:["KimiK3ForConditionalGeneration"]`, `text_config.architectures:["KimiLinearForCausalLM"]` — the text backbone IS the pinned Kimi-Linear hybrid, MASSIVELY scaled: **H=7168, L=93 (69 KDA + 24 MLA full-attn), 896 experts / top-16 / 2 shared, `moe_intermediate_size=3072`**; MLA geometry `kv_lora=512`/`q_lora=1536`/`qk_nope=128`/`qk_rope=64` (= our landed DeepSeek-V3 dims); KDA `head_dim=128`/`num_heads=96`/`short_conv=4`/`gate_lower_bound=-5.0`; quant **`mxfp4-pack-quantized` (compressed-tensors, group 32, e8m0) + MXFP8 acts (QAT)**; vision **MoonViT-V2** (~401M, patch 14, 27L). **HEAVY REUSE** — GDN (KDA's parent, `cuda_gdn.cu`/`gdn_attn.cpp`), DeepSeek MLA (`deepseek_v2.cpp`/`mla_attention.*`, exact geometry), DeepSeek-style MoE (`qwen3_moe.cpp`/`cuda_moe.cu`, scale to 896), and the Qwen3.6-35B GDN-hybrid-MoE model skeleton (`qwen3_5_moe.cpp`) are the structural twins; Kimi-K2 tokenizer/tool parser (`parser/kimi_k2.cpp`) reused. **NET-NEW:** the KDA kernel delta (per-channel `[H,D]` low-rank decay `f_a_proj`/`f_b_proj` + sigmoid-gated output norm + 3 q/k/v convs — already scoped on the Kimi-Linear row), **MXFP4** (we have NVFP4 group-16, not MXFP4 group-32/e8m0), **AttnRes** (report-only, UNCONFIRMED — not in config.json nor pinned `kimi_linear.py`), and the **MoonViT-V2 tower**. **HW-fit: DOES NOT FIT GB10** — 2.8T MXFP4 ≈ **1.56 TB ≈ ~12× over the 119 GiB pool**; no small K3 exists. **DERIVE-AND-SHIP** (no on-box golden, like the beyond-vLLM CUDA bricks): (a) REAL proxy gate of KDA+MLA+MoE on the FITTING `Kimi-Linear-48B-A3B` (~89–91 GiB) vs the pinned oracle, (b) build-verify + structural review for the K3 scale-up. The pinned oracle has NO `kimi_k3` ⇒ even HW-rich users need a pin advance to oracle-gate K3 itself. CORRECTS the 2026-07-25 sweep note ("loads as `DeepseekV3ForCausalLM`" — true for K2, NOT K3). W-plan W1-W8 in the spec. **W2/W5 CPU SCAFFOLDING LANDED (2026-07-28, `CLAIM-KIMI-K3-W2-W5`, DERIVED+BUILD-VERIFIED, NOT pushed):** additive registry TU registers `KimiK3ForConditionalGeneration` (info: text-gen + `is_hybrid` + `has_inner_state` + `supports_multimodal`); config-descent `ParseKimiK3Params` reads the nested `text_config` (KimiLinear KDA+MLA+MoE scalars — note upstream key `num_experts_per_token`), `vision_config` (MoonViT-V2 PARTIAL), `quantization_config` (MXFP4 detect) grounded in `configs/kimi_linear.py:11-148`; pure `EnumerateKimiK3TextBackboneTensors` is the 93-layer KDA/MLA + 896-expert MoE structural name-map grounded 1:1 in `kimi_linear.py:104-378,460-554` + `kimi_gdn_linear_attn.py:102-226` (per-layer KDA vs MLA-with-qLoRA vs MoE-vs-dense branching); forward REFUSE-by-name (`VT_CHECK(false)`, mirrors `deepseek_v4.cpp`); loader REFUSES MXFP4 (a real K3 checkpoint's dtype) deferring to the shared DeepSeek-V4 MXFP4 row. **NOT-YET-BUILDABLE (correctly deferred):** MXFP4 materialization (→ `CLAIM-DEEPSEEK-V4-*` / quantization-matrix MXFP4), the KDA kernel delta (→ Kimi-Linear row `MODEL-TEXT-kimi-linear-*`), MoonViT-V2 vision (W7), K3 multimodal-wrapper weight prefix (post-pin). Code: `src/vllm/model_executor/models/kimi_k3{,_registry,_weights}.cpp` + `include/vllm/model_executor/models/kimi_k3.h`. Test: `tests/vllm/models/test_kimi_k3_scaffold.cpp` (6/6, 63 assertions — registry-resolve + config-descent + split logic + enumeration + reject + MXFP4-refuse). Clean CPU build (`-DVLLM_CPP_CUDA=OFF`). Row stays `SPIKE` (no on-box e2e; forward not implemented); registration `src/vllm/model_executor/models/kimi_k3_registry.cpp:126`; test `tests/vllm/models/test_kimi_k3_scaffold.cpp:123` | `CLAIM-KIMI-K3-SCOPE`, `CLAIM-KIMI-K3-W2-W5` | | `MODEL-DIFFUSION-minimax-h3-mini-max-h3-dit` | `MiniMaxH3DiTModel` | **BEYOND-PIN AND OUT-OF-REPO** — not in `555967922` (H3 released after the pin) and not in the vLLM repository at all: it lives in `vllm-project/vllm-omni`, `vllm_omni/diffusion/models/minimax_h3/` (`minimax_h3_transformer.py`, `packed_sequence.py`, `packed_tokens.py`, `scheduling_minimax_h3_euler_ancestral.py`, `denoise_loop.py`, `vae.py`, `encoder.py`, `pipeline_minimax_h3.py`) | diffusion generation / video + audio (text/image/video/audio in) | flow-matching denoise loop; packed varlen NON-CAUSAL attention; AdaLN modulation; 3D MM-RoPE; video VAE + audio VAE (checkpoint REMOTE CODE); Qwen3-VL-derived encoder; MP4 muxing; Ulysses sequence parallelism | [minimax-h3 spike](specs/minimax-h3.md) | `PARTIAL` | **W0-W2 LANDED (2026-08-03, `CLAIM-MINIMAX-H3-W0-W2`, DERIVE-AND-SHIP).** H3 is NOT an autoregressive LLM: it is a CFG-distilled joint video+audio DIFFUSION transformer (50 blocks, H=5376, 56 MHA heads x 128, SwiGLU 14336, AdaLN 6x3xH, 3D RoPE rotating 96 of 128 dims, video row width 96, audio latent 32) forwarded ONCE PER STEP of a 50-step flow-matching loop — no KV cache, no sampler, no logits, so the SACRED token-exact methodology does not apply. **HW VERDICT: e2e is IMPOSSIBLE on this project's hardware** — ~354 GB checkpoint (DiT 66.3 GB + Qwen3-VL-derived encoder 51.5 GB + video VAE ~10 GB + audio VAE ~0.6 GB), upstream validates on **4x NVIDIA B300 at ~133 GB peak per rank**, vs ONE GB10 with 119 GiB UNIFIED memory; CPU offload does not help because the pool IS host RAM. **WHAT IS GATED (and it is exact):** upstream's modules are pure Python, so they are executed at REDUCED DIMENSIONS on CPU as the oracle (`scripts/gen-minimax-h3-goldens.py` imports them by file path and freezes their outputs; both sides rebuild weights/inputs from an identical FNV-1a + splitmix64 stream, so no weight byte is checked in). Results: fl2va + ref2va packed layouts EXACT including the **fp64 position grid BIT-EXACT** (it feeds RoPE — the port reproduces numpy's `linspace(endpoint=False)` evaluation order and upstream's deliberately-split pairwise vs sequential span summations, `packed_sequence.py:101-113`); patchify/unpatchify/audio pack EXACT + round-trip; scheduler EXACT; **full DiT forward max abs diff 1.6e-7 (video) / 1.5e-7 (audio)** — f32 round-off. **REUSE:** the packed varlen non-causal attention routes through the SHARED `vt::DFlashBlockAttention(causal=false)` (its per-document bidirectional contract IS upstream's varlen FA call) and every projection through `vt::MatmulBT` — NO new kernel was added. **NOT-YET-BUILT (honest):** device-resident/bf16 forward + fusion folds (W2b), H3-Encoder on our existing Qwen3-VL tower (W3), the two VAEs — which are **checkpoint REMOTE CODE** under `trust_remote_code` and must be reimplemented in C++, not adapted (W4/W5) — pipeline/tasks (W6), `/v1/videos` + MP4 muxing, which needs a NEW dependency decision (W7), and USP multi-GPU (W8). No speed number is claimed; upstream reports the DiT at 88% of request latency. Code: [minimax_h3.h](../include/vllm/model_executor/models/minimax_h3.h#L1-L333), [minimax_h3.cpp](../src/vllm/model_executor/models/minimax_h3.cpp#L410-L640) (`MiniMaxH3DitForward`), [minimax_h3_packing.cpp](../src/vllm/model_executor/models/minimax_h3_packing.cpp#L259-L400) (`BuildMiniMaxH3PackedSequence`). Test: [test_minimax_h3.cpp](../tests/vllm/models/test_minimax_h3.cpp#L376-L470) (DiT forward parity; 10/10 cases / 2539 assertions, clean CPU build) + generator [gen-minimax-h3-goldens.py](../scripts/gen-minimax-h3-goldens.py#L1-L60); [ledger](parity-ledger.md#L889). **W6A+W9 LANDED (2026-08-03, `CLAIM-MINIMAX-H3-W6A-W9`) + HW VERDICT CORRECTED.** (a) the **bf16 PRODUCTION dtype policy** now runs (upstream's cast points with the fp32 islands preserved; gated vs a bf16 upstream golden at max abs diff 2.4e-3); (b) **request planning** — 17n+5 frame snapping, video/audio latent shapes, the rectified-flow time-shift sigma schedule, canvas resolution and t2va/fl2va/ref2va dispatch — ported and EXACT vs `time_request.py` + `pipeline_minimax_h3.py:121-122,207-222,374-434`; (c) the **ComfyUI-GGUF arm**: the name map is the IDENTITY and every one of the **535 real tensors** of `MiniMax-H3-FL2VA-Q3_K_M.gguf` resolves onto our contract, with the geometry derived from SHAPES ALONE equal to the shipped config (gated on a manifest read from the file's own header by range request — no payload downloaded). Two shape rules recorded: GGUF `ne` is reversed vs torch, and `comfy.gguf.orig_shape.` overrides it where ComfyUI reshaped a tensor for quant-block alignment (the 50 AdaLN projections: logical [96768, 2688], 2688 not a multiple of the 256-element Q3_K block). **★ HARDWARE VERDICT CORRECTED — the earlier 'e2e is IMPOSSIBLE on this hardware' was WRONG** because it reasoned from the bf16 release alone: quantized H3 checkpoints exist and FIT (GGUF DiT Q3_K_M 15.6 GB + Qwen3-VL encoder Q4_K_M 14.6 GB + VAEs ~11 GB ~= **41 GB** in a 119 GiB pool; `lilcheaty/MiniMax-H3-NVFP4` likewise). So e2e AND a speed comparison are REACHABLE; NVFP4 is the likely speed path (sm_121 native FP4 tensor cores + our tuned NVFP4 stack). **W5 LANDED (audio VAE)**: H3's VAEs are checkpoint REMOTE CODE under `trust_remote_code`, so a no-Python engine must REIMPLEMENT them — the DAC-lineage BigVGAN audio decoder (weight-norm materialization (w = g*v divided by the row norm), anti-aliased SnakeBeta with kaiser-sinc up/down resampling, replicate padding, final clamp) is ported and gated against the checkpoint's OWN modules at **max abs diff 4.2e-9**. The VIDEO VAE (`klvae.py` ~48 KB + CNN/ViT + tiling) is the largest remaining brick. **W10 GROUNDED + W4 SCOPED from REAL manifests** (safetensors headers captured by range request, no payload downloaded): the NVFP4 checkpoint's 1051 tensors are textbook compressed-tensors NVFP4 (U8 packed 2-per-byte + E4M3 `weight_scale` at group 16 + F32 scalar `weight_scale_2`; 258 quantized projections; fp32/bf16 islands left unquantized; names identical to our contract) — i.e. EXACTLY the layout our tuned NVFP4 stack already consumes, so W10 is loader wiring, not a new quant scheme. The video VAE's 560 tensors show its ENCODER is the 3D CNN (rank-5 Conv3d) while its DECODER — the half generation needs — is a 36-block TRANSFORMER (to_qkv/to_out, ff.w1/w2, 2 norms + 2 learned residual scales per block, x_embedder/mask_token/register_tokens/proj_out), materially smaller than klvae.py's 48 KB suggested. **W4 BLOCK LANDED**: the video-VAE decoder's repeated `TransformerBlock` is ported and gated at **6.0e-8** against the checkpoint's OWN remote code — RMSNorm + per-head RMS qk-norm (no affine) + full attention + gated-SiLU FF + LEARNED PER-CHANNEL residual scales, and critically the PER-HEAD-INTERLEAVED qkv layout ([head][q,k,v], NOT the DiT's [q_all,k_all,v_all]) that would otherwise produce a plausible-but-wrong image. **W4 DECODER DONE**: the FULL ViT3D video-VAE decoder — pack, x_embedder, register/cls tokens, 3D RoPE (RotaryEmbeddingND, length-normalized ids, angle scale 2pi, tiled freq blocks), the 36-block stack, LayerNorm norm_out, proj_out and unpatchify — is ported and gated at **8.9e-8** against the checkpoint's own `ViT3DDecoder`, at its real hyperparameters (36 layers, 32 heads x 64, rope_theta 100, rope_dim_ratio 0.75). BOTH VAE decoders are now done (audio 4.2e-9). **W3 TEXT TOWER DONE**: the H3-Encoder's truncated Qwen3-VL text tower — gated at **1.2e-7** vs upstream — with all three H3 deltas exercised: layer truncation (min(num_hidden_layers, 50)), the UNNORMALIZED layer-49 output (NO final RMSNorm, unlike stock Qwen3-VL — applying one silently shifts every conditioning vector), and DeepStack visual injection into the first N layers; plus interleaved M-RoPE, fused QKV, per-head q/k RMSNorm, causal GQA and the gated-SiLU MLP. **W6 t2va ASSEMBLED — the WHOLE PATH COMPOSES**: `MiniMaxH3GenerateT2va` wires packed layout -> rectified-flow sigma schedules -> the multi-step denoise loop of DiT forwards -> unpatchify + audio unpack -> per-channel denormalize -> BOTH VAE decoders, producing correctly-shaped, finite frames and a stereo waveform in [-1, 1]; gated by a structural end-to-end test at reduced dimensions with random weights (NOT a quality result). Assembling it also caught a real gap: the audio decode needed the checkpoint's `dec_in_proj` (Conv1d k=1, vae_latent_channels -> num_mels) ahead of BigVGAN. **W9 GGUF ARM DONE**: `LoadMiniMaxH3DitFromGguf` materializes the DiT from a ComfyUI-format GGUF — dequantizing every tensor through the SHARED GGUF dequant entry point (so the Q2_K/Q3_K/Q4_K families the H3 GGUFs use are covered by the same code every other GGUF model uses), recovering the geometry from shapes alone, and binding the forward's views with missing tensors throwing BY NAME rather than reading as zeros; gated by a synthetic-file load-AND-RUN test (a real DiT forward executes off the loaded weights). **W10 LOADER DONE**: `LoadMiniMaxH3DitFromNvfp4` materializes the DiT from an NVFP4 compressed-tensors checkpoint — the U8-packed [out, in/2] weight plus its E4M3 group-16 `weight_scale` and F32 scalar `weight_scale_2` go through the project's EXISTING NVFP4 dequant (no new quant code), sidecars are excluded from the model tensor set, and the logical [out, in] shape is recovered; gated by a synthetic-file load-AND-RUN test. BOTH quantized loaders are now done. **W3 VISION BLOCK DONE** (6.0e-8): the repeated unit of the encoder's Qwen3-VL vision tower, which differs from the text tower in every way that matters numerically — LayerNorm WITH BIAS (not RMSNorm), a [q_all, k_all, v_all] qkv layout (not the video VAE ViT's per-head interleave), fp32 rotary, NON-CAUSAL attention segmented by `cu_seqlens` (the test asserts a perturbation in one packed image leaves the other's outputs BIT-IDENTICAL), and the TANH-approximate GELU. **W3 ENCODER COMPLETE**: the FULL vision tower also lands — Conv3d patch embed (kernel == stride, so a linear over the flattened patch), BILINEAR resampling of the learned position grid into spatial-merge order, the 2D rotary table, per-frame `cu_seqlens`, and both merger flavours (the final merger norms the PRE-shuffle width while the DeepStack mergers norm the POST-shuffle width, and both use exact-erf GELU unlike the block MLP's tanh approximation), gated over a RAGGED two-image batch. Only the MM processor remains on the encoder. **CONDITION-NOISE augmentation DONE** (fl2va/ref2va): the noised-anchor mix plus its ROW ACCOUNTING — each visual condition draws noise of length `target_latent_t + imgvid_cond_num_frames` and slices the PREFIX matching its own latent_t, every condition restarts the SAME seed (so concatenating and drawing once would differ for multi-reference requests), and rows advance by that condition's own patchified count. Gated EXACT with the noise supplied, so the comparison isolates the accounting from torch's RNG. **REFERENCE-VIDEO geometry + FRAME SCHEDULE DONE** (the pure-math half of `reference_video.py`): the canvas pipeline (aspect clamp -> 768 short edge -> max-pixel rescale -> nearest multiple of 32) and the 24-to-2 FPS frame resample with per-temporal-patch block timestamps, both EXACT. NOTE the rest of that module (probe, transcode, frame extraction, audio decode) shells out to ffmpeg and is blocked on the SAME external dependency decision as `/v1/videos` MP4 muxing — one decision unlocks reference-video INPUT decode and generated-video OUTPUT encode together. **VIDEO VAE TILING DONE**: the tile plan (smallest tile count whose MINIMUM overlaps still cover the axis, leftover slack distributed in whole `vae_ratio` units ROUND-ROBIN across the seams) plus the linear seam cross-fade, both EXACT. Shipped config tile 256 / overlap 64 / vae_ratio 16 (= prod(space_down), the 'f16' in f16t4). Getting the slack distribution wrong shifts every tile after the first and surfaces as seam artifacts rather than an error. **PRESENTATION TOKEN TAGS DONE** — the fl2va vision-span override the denoise loop requires callers to have applied: a vision block is `vision_start + pad*count + vision_end` and the WHOLE block, MARKERS INCLUDED, is tagged VIDEO; tagging only the pads would leave two markers as TEXT and shift every AdaLN modulation index after them. Gated EXACT, with the test proving each VIDEO run is a whole vision block. **VAE 3D-CNN ENCODER PRIMITIVES DONE**: causal Conv3d (all temporal padding on the LEFT so a frame never sees the future, `reflect` spatial padding), GroupNorm3D (32 groups, eps 1e-6, statistics spanning TIME as well as space) and ResnetBlock3D, gated EXACT — with CAUSALITY proven directly on the bare convolution (a change to the last frame provably cannot reach earlier frames). **Downsample3D DONE** too: the strided inter-level conv, whose subtlety is the ASYMMETRIC pre-pad — one pixel on the RIGHT of W and the BOTTOM of H before a stride-2 conv with padding (1,0,0); padding symmetrically instead shifts the whole sampling lattice by half a pixel, which is a silent wrong latent rather than an error. Only the EncoderFCN3D level-loop assembly remains on the VAE encoder. **VIDEO VAE COMPLETE — encoder AND decoder**: the whole 3D-CNN encoder level loop (conv_in -> per level [ResnetBlock3D x N then Downsample3D or a 1x1x1 channel match] -> GroupNorm -> SiLU -> conv_out) is gated EXACT. **MM PROCESSOR = REUSE, gated**: H3's `FL2VA/processor` is a stock `Qwen3VLProcessor`, so the multimodal front end is this project's EXISTING Qwen3-VL processor rather than a new port; H3's own config is parsed and driven through it (patch 16 / temporal 2 / merge 2, **0.5 normalization rather than CLIP statistics**, a 32-pixel grid, the 768x1344 default canvas proven an IDENTITY under smart_resize, and VIDEO bounds deliberately looser than the image ones). **With this every PORTABLE piece of the lane is done**; what remains is one dependency decision (ffmpeg media I/O + MP4 muxing) and GPU-blocked work (the device-resident FP4 forward and any speed number). **WAV OUTPUT** added: the decoded stereo waveform serializes to RIFF/WAVE 16-bit PCM, converting the VAE's CHANNEL-MAJOR layout to INTERLEAVED (getting that backwards yields audio that plays but with the channels time-smeared) and clamping rather than wrapping. Deliberately dependency-free, and required under EITHER outcome of the open MP4/muxer decision. **VIDEO OUTPUT PATH DONE**: PPM frame serialization (planar [C,T,H,W] -> row-major interleaved RGB, [-1,1] -> [0,255] clamped) plus the MP4 mux argv (h264/yuv420p + AAC, `-shortest`, `+faststart`). The built argv was RUN through real ffmpeg 6.1.1 and produced a VALID MP4 (ffprobe: h264 yuv420p video + AAC stereo at 32 kHz). The library never spawns a process — `src/vllm/` has no subprocess precedent — so it builds the artifacts and the command while the example/server layer invokes it. **`/v1/videos` API LOGIC DONE**: the request contract (H3 defaults — 50 steps, flow shift 12 video / 3 audio; both the vLLM-Omni `extra_params` nesting and a flat top-level spelling accepted; malformed input rejected with a reason rather than silently defaulted) and the job store (queued -> running -> succeeded/failed, illegal transitions throw, unknown ids reported so the route can 404, status JSON omits fields that do not apply, and concurrent creation is thread-safe). Remaining is mechanical glue: route registration + runner injection, with the ffmpeg call in `examples/` per the ratified boundary. Test: 34/34 + video-api 4/4 (9233 + 63 assertions). OPEN: there is no vllm-omni parity PIN — the upstream-sync protocol covers only the vLLM repo | `CLAIM-MINIMAX-H3-W0-W2`, `CLAIM-MINIMAX-H3-W6A-W9` | +| `MODEL-MM-moss-tts-moss-tts-delay-talker-for-generation` | `MossTTSDelayModel` | **OUT-OF-REPO** — not in `555967922` and not in the vLLM repository at all (absent from every dictionary of `registry.py` on `main` as well): it is registered by the separate `vllm-project/vllm-omni` repository, `vllm_omni/model_executor/models/registry.py:350-354` @ `bbe6ccc512a404a2df8c977ea29003002f2683e8` → module `vllm_omni/model_executor/models/moss_tts/modeling_moss_tts_talker.py`, class `MossTTSDelayTalkerForGeneration`, with the shared stage-1 codec decoder at `registry.py:368-372` → module `moss_tts/modeling_moss_tts_codec.py`, class `MossTTSCodecDecoder`. Module and class are given as separate fields rather than in the compact colon-pair form, as the `MiniMaxH3DiTModel` row does: that form marks a target inside the PINNED vLLM registry inventory, and this target is in another repository entirely. Documented at `vllm-omni` `docs/models/supported_models.md:78`. Official recipes: `OpenMOSS-Team/MOSS-TTS`, `MOSS-TTSD-v1.0`, `MOSS-SoundEffect`, `MOSS-VoiceGenerator` — four of the 157 recipes at `vllm-project/recipes` `86c7777a`, all four declaring `architectures: ["MossTTSDelayModel"]`. Tracked by [#610](https://github.com/mudler/vllm.cpp/issues/610) | conditional generation / text+audio in, AUDIO OUT (TTS) | MM processor; audio detokenizer / codec decoder; delay-pattern multi-codebook decoding; reference-audio speaker encoder; streaming audio output. **We serve no audio-generation output modality at all today** | ☐ required | `INVENTORIED` | none. There is no `vllm-omni` parity PIN — the upstream-sync protocol covers only the vLLM repository — so an oracle for this row has to be established before it can advance | unassigned | +| `MODEL-MM-moss-tts-moss-tts-realtime-talker-for-generation` | `MossTTSRealtime` | **OUT-OF-REPO** — not in `555967922`, and absent from `registry.py` on vLLM `main`; registered by `vllm-project/vllm-omni` at `vllm_omni/model_executor/models/registry.py:356-360` @ `bbe6ccc512a404a2df8c977ea29003002f2683e8` → module `vllm_omni/model_executor/models/moss_tts/modeling_moss_tts_talker.py`, class `MossTTSRealtimeTalkerForGeneration` (separate module and class fields rather than the compact colon-pair form, per the out-of-repo convention) — the SAME module as `MossTTSDelayModel` but a different class, so the two are separate targets rather than aliases. Documented at `vllm-omni` `docs/models/supported_models.md:79`. Official recipe: `OpenMOSS-Team/MOSS-TTS-Realtime` (1.7B). Tracked by [#610](https://github.com/mudler/vllm.cpp/issues/610) | conditional generation / text+audio in, AUDIO OUT (realtime/streaming TTS) | MM processor; audio detokenizer / codec decoder; realtime full-duplex streaming; incremental audio decode | ☐ required | `INVENTORIED` | none. Same missing-`vllm-omni`-pin blocker as the other omni rows | unassigned | +| `MODEL-MM-qwen3-tts-qwen3-tts-talker-for-conditional-generation` | `Qwen3TTSForConditionalGeneration`, `Qwen3TTSTalkerForConditionalGeneration` | **OUT-OF-REPO** — not in `555967922`, and absent from `registry.py` on vLLM `main`; registered by `vllm-project/vllm-omni` at `vllm_omni/model_executor/models/registry.py:152-156` (and the explicit `Qwen3TTSTalkerForConditionalGeneration` spelling at `:157`) @ `bbe6ccc512a404a2df8c977ea29003002f2683e8` → module `vllm_omni/model_executor/models/qwen3_tts/qwen3_tts_talker.py`, class `Qwen3TTSTalkerForConditionalGeneration` (separate module and class fields rather than the compact colon-pair form, per the out-of-repo convention), with the code2wav stage at `registry.py:162-166` → module `qwen3_tts/qwen3_tts_code2wav.py`, class `Qwen3TTSCode2Wav` and 12 Hz / 25 Hz tokenizers under `qwen3_tts/tokenizer_12hz/` and `tokenizer_25hz/`. Documented at `vllm-omni` `docs/models/supported_models.md:72-74`. Official recipe: `Qwen/Qwen3-TTS-12Hz-1.7B-CustomVoice` (published at `recipes.vllm.ai`). The two architecture strings resolve to the identical `(module, class)` target so they are grouped on one row, per the alias rule | conditional generation / text in, AUDIO OUT (TTS, custom voice) | MM processor; audio detokenizer / code2wav; speaker-embedding conditioning; word timestamps; streaming audio output | ☐ required | `INVENTORIED` | none. Its text tower is Qwen3-derived, which is the one piece of this row we already ship; everything downstream of the talker (code2wav, audio tokenizer, speaker embeddings) is unported | unassigned | +| `MODEL-MM-higgs-audio-v3-higgs-audio-v3-talker-for-conditional-generation` | `HiggsMultimodalQwen3ForConditionalGeneration`, `HiggsAudioV3TalkerForConditionalGeneration` | **OUT-OF-REPO** — not in `555967922`, and absent from `registry.py` on vLLM `main`; registered by `vllm-project/vllm-omni` at `vllm_omni/model_executor/models/registry.py:213-217` (and the explicit `HiggsAudioV3TalkerForConditionalGeneration` spelling at `:218`) @ `bbe6ccc512a404a2df8c977ea29003002f2683e8` → module `vllm_omni/model_executor/models/higgs_audio_v3/higgs_audio_v3_talker.py`, class `HiggsAudioV3TalkerForConditionalGeneration` (separate module and class fields rather than the compact colon-pair form, per the out-of-repo convention), with code2wav at `registry.py:223-227` → module `higgs_audio_v3/higgs_audio_v3_code2wav.py`, class `HiggsAudioV3Code2WavForConditionalGeneration`. Documented at `vllm-omni` `docs/models/supported_models.md:81`. Official recipe: `bosonai/higgs-audio-v3-tts-4b` (published at `recipes.vllm.ai`), whose `config.json` declares `architectures: ["HiggsMultimodalQwen3ForConditionalGeneration"]`. The distinct v2 lineage (`HiggsAudioV2ForConditionalGeneration`, `higgs_audio_v2/`) is a DIFFERENT target and has no recipe in the audited set, so it gets no row here | conditional generation / text+audio in, AUDIO OUT (TTS) | MM processor; audio detokenizer / code2wav; audio tokenizer; Qwen3-derived multimodal decoder; streaming audio output | ☐ required | `INVENTORIED` | none | unassigned | +| `MODEL-MM-voxtral-realtime-voxtral-realtime-for-conditional-generation` | `VoxtralRealtimeForConditionalGeneration` (target-pending) | **TARGET-PENDING — the architecture string is registered NOWHERE we can find.** Recipe `mistralai/Voxtral-Mini-4B-Realtime-2602` (one of the 157 recipes at `vllm-project/recipes` `86c7777a`) declares `architectures: ["VoxtralRealtimeForConditionalGeneration"]` with an inner `audio_config.model_type: "voxtral_realtime_encoder"`, but that exact string appears in none of: (a) the pinned registry `555967922`; (b) any of the ten dictionaries of `vllm/model_executor/models/registry.py` on vLLM `main`; (c) `_OMNI_MODELS` in `vllm_omni/model_executor/models/registry.py` @ `bbe6ccc512a404a2df8c977ea29003002f2683e8`; (d) `vllm-omni` `docs/models/supported_models.md`. **Related but NOT equal, recorded so nobody re-derives it:** `VoxtralRealtimeGeneration` — a DIFFERENT string — is registered at the pin (`registry.py:585`, `main` `:603`) and already has its own row here, `MODEL-MM-voxtral-realtime-voxtral-realtime-generation`; `VoxtralForConditionalGeneration` (pin `registry.py:584`) is the audio→TEXT model; and `vllm-omni` `registry.py:305-309` registers `VoxtralTTSForConditionalGeneration` for the SEPARATE `mistralai/Voxtral-4B-TTS-2603` recipe. The most likely reading is that this checkpoint's `config.json` names a renamed or not-yet-landed spelling of the existing `VoxtralRealtimeGeneration` target, but that is a hypothesis: no anchor is asserted until one of those is shown to be what actually resolves. Tracked by [#610](https://github.com/mudler/vllm.cpp/issues/610) | conditional generation / audio in, AUDIO OUT (realtime speech-to-speech) — modality inferred from the recipe and config, NOT from an upstream implementation | unresolved until the target is located | ☐ required | `INVENTORIED` | none | unassigned | +| `MODEL-MM-bailing-mm-native-bailing-mm-native-for-conditional-generation` | `BailingMMNativeForConditionalGeneration` (target-pending) | **TARGET-PENDING — the architecture string is registered NOWHERE we can find.** Recipe `inclusionAI/Ming-omni-tts-0.5B` (one of the 157 recipes at `vllm-project/recipes` `86c7777a`) declares `architectures: ["BailingMMNativeForConditionalGeneration"]`, but that exact string appears in none of: (a) the pinned registry `555967922`; (b) any dictionary of `vllm/model_executor/models/registry.py` on vLLM `main`; (c) `_OMNI_MODELS` in `vllm_omni/model_executor/models/registry.py` @ `bbe6ccc512a404a2df8c977ea29003002f2683e8`; (d) `vllm-omni` `docs/models/supported_models.md`. **Related but NOT equal:** `vllm-omni` `docs/models/supported_models.md:75` lists that SAME checkpoint under a different architecture name, `MingTTSForConditionalGeneration` (`registry.py:250-254` → module `ming_tts/ming_tts.py`, class `MingTTSForConditionalGeneration`); and `vllm-omni` `registry.py:417-421` carries a near-miss alias `BailingMM2NativeForConditionalGeneration` — note the `2` — commented "HF repo currently ships this architecture name in config.json", pointing at module `ming_flash_omni/ming_flash_omni.py`, class `MingFlashOmniForConditionalGeneration`, for the DIFFERENT Ming-flash-omni-2.0 checkpoint. So the shape of the gap is a missing config-string alias upstream, but that is an observation, not an anchor, and none is asserted here. Tracked by [#610](https://github.com/mudler/vllm.cpp/issues/610) | conditional generation / text+image+audio in, AUDIO OUT (TTS) — modality from the recipe and config, NOT from a resolved upstream target | unresolved until the target is located | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-MM-moss-transcribe-diarize-moss-transcribe-diarize-for-conditional-generation` | `MossTranscribeDiarizeForConditionalGeneration` (v0.25.0 target-pending) | v0.25.0 target `registry.py:450-453`; `vllm/model_executor/models/moss_transcribe_diarize.py::MossTranscribeDiarizeForConditionalGeneration` @ `702f481` | conditional generation / audio | MM processor; Whisper encoder; VQ adaptor; Qwen3 decoder; speech-to-text/diarization frontend | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-MM-lightonocr-light-on-ocrfor-conditional-generation` | `LightOnOCRForConditionalGeneration` | `registry.py:450-453`; `vllm/model_executor/models/lightonocr.py::LightOnOCRForConditionalGeneration` | conditional generation / image | MM processor; encoder/merge; vision encoder | ☐ required | `INVENTORIED` | none | unassigned | | `MODEL-MM-lfm2-vl-lfm2-vlfor-conditional-generation` | `Lfm2VlForConditionalGeneration` | `registry.py:454`; `vllm/model_executor/models/lfm2_vl.py::Lfm2VLForConditionalGeneration` | conditional generation / image | MM processor; encoder/merge; Mamba/SSM state; vision encoder; video path | ☐ required | `INVENTORIED` | none | unassigned | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 5fb7978a0..aa55f553d 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -39,6 +39,7 @@ issue is not yet placed. Keyed record: update in place, never append. | Issue | Row | Title | Kind | |---:|---|---|---| | [#168](https://github.com/mudler/vllm.cpp/issues/168) | `BACKEND-CUDA-SM110` | Jetson AGX Thor (sm_110): 32B NVFP4, Tekken tokenizer blocker | feature | +| [#618](https://github.com/mudler/vllm.cpp/issues/618) | `BACKEND-GATE-CPU-LLAMACPP` | `test_cpu_x86_llamacpp_floor`'s contended-leg case is load-dependent: at loadavg 63 the harness exits `NO_QUIET_WINDOW` (4) instead of `GIVING_UP` (2), so the guarantee goes untested and the red reads as a defect in whatever diff is in flight | bug | | [#433](https://github.com/mudler/vllm.cpp/issues/433) | `BACKEND-GATE-CPU-LLAMACPP` | No x86_64 arm: the closed CPU floor is Arm/i8mm-only and every lever that closed it is Arm-specific | perf | | [#199](https://github.com/mudler/vllm.cpp/issues/199) | `BACKEND-METAL-MLX` | macOS MLX build fails on `-Werror` in MLX headers | bug | | [#41](https://github.com/mudler/vllm.cpp/issues/41) | `BACKEND-ROCM` | ROCm (AMD GPU) backend | feature | @@ -113,12 +114,14 @@ issue is not yet placed. Keyed record: update in place, never append. | [#395](https://github.com/mudler/vllm.cpp/issues/395) | `LORA-RUNTIME` | `main` is RED on `sanitize-cpu (address,undefined)`: `test_punica_cpu` `RefShrink` reads past `a_stacked` for an out-of-range slot | bug | | [#400](https://github.com/mudler/vllm.cpp/issues/400) | `LORA-RUNTIME` | `test_punica_cpu` does not cover the out-of-range slot guard in `BgmvShrink` or `BgmvExpandSlice`: dropping either leaves the suite green | bug | | [#432](https://github.com/mudler/vllm.cpp/issues/432) | `MODEL-MM-voxtral-voxtral-for-conditional-generation` | Voxtral audio TTFT is ~17x vLLM: the Whisper encoder attention is a scalar warp kernel where vLLM runs FA-2 | enhancement | +| [#610](https://github.com/mudler/vllm.cpp/issues/610) | `MODEL-MM-moss-tts-moss-tts-delay-talker-for-generation` | Six `vllm-omni` recipe architectures had no model-matrix row: the TTS / audio-generation output modality was entirely unplaced (rows added `INVENTORIED`; two are `target-pending`) | bug | | [#268](https://github.com/mudler/vllm.cpp/issues/268) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | Muse Glimmer (Meta, 30B agentic multimodal): text tower, perception encoder, DFlash drafter | feature | | [#329](https://github.com/mudler/vllm.cpp/issues/329) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | Muse Glimmer GGUF k-quants: text arm lands; mmproj blocked by a converter axis drop | feature | | [#333](https://github.com/mudler/vllm.cpp/issues/333) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | Muse Glimmer speed: no number on any axis; benchmark vs llama.cpp, HF, and ourselves | perf | | [#347](https://github.com/mudler/vllm.cpp/issues/347) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | GGUF tokenizer: no pre-tokenizer for the GPT-4o family, so `tokenizer.ggml.pre "llama4"` files cannot generate at all | bug | | [#359](https://github.com/mudler/vllm.cpp/issues/359) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | Muse Glimmer GGUF k-quant generates INCOHERENT text where llama.cpp on the same file is coherent | bug | | [#412](https://github.com/mudler/vllm.cpp/issues/412) | `MODEL-MM-muse-glimmer-muse-glimmer-for-conditional-generation` | Muse Glimmer config defaults are neutral values, not the architecture's constants: the released GGUF's sandwich post-norms ran at 1e-5 instead of 1e-8 | bug | +| [#609](https://github.com/mudler/vllm.cpp/issues/609) | `MODEL-TEXT-qwen3-5-qwen3-5-moe-for-causal-lm` | Two recipe architectures registered on vLLM `main` but absent at the pin had no model-matrix row: `Qwen3_5MoeForCausalLM` (the text-only backbone our multimodal row already instantiates) and `BailingMoeV3ForCausalLM` | bug | | [#298](https://github.com/mudler/vllm.cpp/issues/298) | `PERF-27B-LMHEAD-DSR` | `main` is RED on the `device-leakage` DSR ratchet: `PrepareLmHeadResident` added a 33rd `vt_ifdef` to the shared layer | bug | | [#213](https://github.com/mudler/vllm.cpp/issues/213) | `PERF-27B-LMHEAD-FP4` | Qwen3.6 NVFP4 baselines (27B and 35B-A3B) must reach vLLM speed parity | perf | | [#339](https://github.com/mudler/vllm.cpp/issues/339) | `PERF-27B-LMHEAD-FP4` | 27B c1: every fp8 input projection asks for an f32 output, selecting the slower nvjet template family where vLLM emits bf16 (48 f32-out projections 18.51 ms vs 48 bf16-out 7.05 ms). The merged GDN `in_proj` arm is built DEFAULT OFF as `VT_GDN_FP8_IN_BF16`, spec [`perf-fp8-alpha-fold.md`](specs/perf-fp8-alpha-fold.md) §Attempt 4 — UNMEASURED: no committed gate loads the fp8 tower (`row/GATE-27B-FP8-TOWER-GOLDEN` builds that arm) | perf | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 36c1a290d..64515e477 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -185,6 +185,7 @@ on the committed fixture); reranking/classify models are not yet registered. | Video | ✅ correctness-gated | ✅ | ✅ | ☐ | | Audio | ✅ correctness-gated | ✅ | ◐ | ◐ | | Video+audio GENERATION (MiniMax-H3 DiT, vLLM-Omni lane) | ◐ all three modalities COHERENT on Q4_K_M (t2va, fl2va, ref2va; §8.20); the NVFP4 arm carries the patch grid; GGUF/NVFP4/bf16 loaders, unpruned AND pruned (§8.21); ABI v12 `vllm_video_*` | ✅ (vllm-omni, BF16-only, no quantized H3 arm) | ☐ | ☐ | +| Speech / audio GENERATION (TTS, vLLM-Omni lane) | ☐ not started. The six architectures behind official `vllm-project/recipes` TTS models are inventoried in `.agents/model-matrix.md` (#610); an inventoried row is not a supported one | ✅ (vllm-omni: MOSS-TTS, Qwen3-TTS, Higgs Audio v3, Voxtral TTS) | not assessed | not assessed | | Multimodal over the OpenAI server | ◐ image request path wired, forward pending | ✅ | ✅ | ◐ | Image, video and audio are correct through the CLI and library. Over the HTTP @@ -192,7 +193,9 @@ API the image **request** path is wired end to end (`ROAD-V1-MM` W1-W3): the production server attaches the seam at `server_main.cpp:826`. Two residuals keep it from ✅: the model runner has no mm-forward consuming `Request.mm_features`, and no image codec is vendored (raw RGB only). Video, audio and multi-image over -HTTP are not started. +HTTP are not started. Audio **in** is gated; audio **out** does not exist: we +ship no TTS or speech-generation path on any surface, which is why that row is +the only ☐ in our column here. ## Speculative decoding diff --git a/scripts/check-agent-record.py b/scripts/check-agent-record.py index 4817336c2..ff26f1979 100644 --- a/scripts/check-agent-record.py +++ b/scripts/check-agent-record.py @@ -39,7 +39,23 @@ # delegates it. Its only upstream implementation is the still-OPEN # vllm#51655; see porting-inventory.md §9 deviation 16. Bumped because a new # row EXISTS, never to make a transition pass. - "MODEL": (AGENTS / "model-matrix.md", 362), + # 370 since 2026-08-13: +8 rows for the architectures behind official + # `vllm-project/recipes` models that had no row at all (#609, #610). Two are + # pin-lag — `Qwen3_5MoeForCausalLM` and `BailingMoeV3ForCausalLM` are + # registered on vLLM `main` and absent only at `555967922`. Six are + # out-of-repo: `MossTTSDelayModel`, `MossTTSRealtime`, + # `Qwen3TTSForConditionalGeneration` and + # `HiggsMultimodalQwen3ForConditionalGeneration` are registered by + # `vllm-project/vllm-omni`, and `VoxtralRealtimeForConditionalGeneration` + # and `BailingMMNativeForConditionalGeneration` are target-pending — their + # exact `config.json` architecture strings are registered in neither core + # vLLM `main` nor `vllm-omni`, so the rows record what was searched instead + # of an invented anchor. None of the eight touches the at-the-pin model + # inventory below (324/373/356/310/261 is unchanged), because like the + # MuseGlimmer, KimiK3 and MiniMaxH3DiT rows they carry no pinned-registry + # target. Bumped because eight new rows EXIST, never to make a transition + # pass. + "MODEL": (AGENTS / "model-matrix.md", 370), # 82 since 2026-07-21: +`QUANT-NVFP4-CT-W4A16` (compressed-tensors NVFP4A16 / # W4A16 — NVFP4 weights with BF16 activations, distinct from the existing # `QUANT-NVFP4-CT-W4A4` and `QUANT-NVFP4-MO-W4A16` rows in both scheme