Roadmap-v1 reality audit (2026-07-31, CLAIM-ROADMAP-V1-AUDIT, .agents/specs/roadmap-v1-audit-2026-07-31.md): an 8-lane code-grounded audit + adversarial DONE-verify (26 agents) vs the 18-row plan-of-record. Correctness + CPU-gated features are REAL; the dominant open gate across nearly all RI rows is every-axis SPEED; most GPU "DONE" proofs are DGX-recorded (not re-verifiable on the current box). 4 DONE claims refuted: MM-video is near-tie not STRICT; Gemma-4 image is 16/18 near-tie not STRICT 18/18 (audio unbuilt); ROAD-V1-A (async code DONE but both SGLang-floor arms never ran, 27B 114/124 coin-flip not clean); C6 async-serving (SERVE-ASYNC-LLM still GATING). D5 LoRA is an unwired standalone CPU brick. Fixed backend-matrix BACKEND-GATE-CUDA-VLLM staleness 49/124 → 114/124. (Follow-up: reworded a spec code-snippet check-agent-record misread as a dangling link — rc=0.)
This is the per-capability status ledger. README.md is the user-facing landing page and carries only headline state; every capability's current lifecycle stage, its active gaps, and its next gate live here.
Capabilities are labelled: correctness-complete (token-exact vs the vLLM oracle), speed-pending (correct, throughput work in progress), build-only (compiles for a target with no runtime proof here), and hardware-blocked (cannot run on the hardware available).
The forensic chronology and the raw evidence live in the append-only
.agents/state.md, the
parity ledger, the area matrices, and
docs/BENCHMARKS.md. This file keeps ONE binding current-state
line per capability, not a run-by-run log.
The reference numbers here are measured against the pinned vLLM oracle on GB10,
greedy, same workload, same tokens. Where a claim cannot be measured on the
hardware here, it is stated as such rather than implied. The parity pin
advanced 2026-07-26 to vLLM 0.26.0.dev0 (55596792) + transformers 5.14.1
(from 0.25.0); correctness was re-validated bit-identical on the new oracle
(zero golden drift; see .agents/specs/pin-advance.md),
so the token-exact claims hold against the new pin. Historical speed figures
citing "vLLM 0.25.0" are the last binding measurement against the prior oracle
(our engine is unchanged by the advance); a re-benchmark against 0.26 is pending.
GCC 12 production-library maintenance (2026-07-31): the two known -Werror
blockers in GGUF prefaulting and KV-offload temporary-file naming are fixed
without behavior or lifecycle changes. The production library and focused
tests build clean; the all-target build remains PARTIAL on two unrelated
test-only -Wrestrict diagnostics, recorded in BENCHMARKS.
vllm.cpp implements an intentionally focused subset of vLLM, held to token-for-token correctness against the pinned oracle.
| Capability | State | Notes |
|---|---|---|
| Qwen3.6-27B (NVFP4) text generation | Correctness-complete, at/above vLLM speed | Token-exact greedy on GB10; beats vLLM 0.25.0 total throughput at every concurrency (1.007-1.045x), effective parity 115/124 axes |
| Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete, decode at-parity, prefill speed-pending | Token-exact greedy; decode at or beyond vLLM, remaining gap is prefill TTFT |
| Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. D1 (2026-07-31, CLAIM-D1-BF16-MERGED-QKV): the bf16 merged-QKV path (Qwen3QkvMergeEnabled/VT_QWEN3_QKV_MERGE) is now default-ON — one vt::MatmulBT over the merged [qdim+2kdim,H] owner + a contiguous vt::QkvSplit (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit test_ops_qkv_merge byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed |
| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete, speed-pending | Direct ON and OFF are token-identical, and token-identical to the previous series; direct loading cuts peak/stable host PSS by 73.4%/91.1% and mean TTFT by 12.7%. Against an oracle built at the actual parity pin: 0.9970x total throughput, TTFT passes (0.773x), TPOT 12.4% high; re-validated unchanged (0.9972x / 1.1247x) after rebasing onto 139 upstream commits. The failing axis is the discrete-GPU async-overlap gap, scoped as ENG-ASYNC-SCHED W4 |
| Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense AttnBlock — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6 |
| Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling |
| Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) |
| OPT (learned pos-emb, cross-family) | Correctness-complete, speed-pending | Strict token-exact 6/6 (OPT-125m); additivity canary across model families |
| DeepSeek-V2 MLA | Correctness-complete, speed-pending | Token-exact 8/8 (DeepSeek-V2-Lite); 0.86-0.95x vLLM output rate, TTFT faster at c4/c8. A2+A5 MLA norm-rope fold landed 2026-07-31 (bit-exact; DGX SACRED gate 8/8 UNCHANGED, goldens md5 identical, 27B 235/235 + qwen3_dense 16/16+16/16 inert — see docs/BENCHMARKS.md): the kv_c(nope)+k_pe(rope) A-projections merge to ONE vt::MatmulBT and {kv_a_layernorm(latent); decoupled-k_pe RoPE} fold into ONE new shared vt::FusedNormRope (kFusedNormRope) launch — per-layer A-proj GEMMs 3→2 (q-lora) / 2→1 (no-q-lora) plus the norm+rope launch fold, default-ON (VT_MLA_FUSED_NORM_ROPE=0 = byte-exact rollback). kimi_k3 inherits it for free |
| GLM-4 dense (sandwich norms, partial rope) | Correctness-complete, speed-pending | Token-exact 16/16 (GLM-4-9B-0414); first GLM-family model; partial interleaved RoPE + Gemma2 sandwich norms + biased qkv |
| GLM-4.7-Flash (MLA + GLM MoE) | Correctness-complete, speed-pending | Token-exact 8/8 (GLM-4.7-Flash, 31.2B); reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora query branch + noaux_tc sigmoid router with routed-scaling |
| Gemma-3 dense (GeGLU, dual rope, sandwich norms) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-3-1b-it); first Gemma-family model; GeGLU (gelu_pytorch_tanh) + dual per-layer RoPE theta + Gemma-RMSNorm sandwich norms + sqrt(hidden) embed-scale + query_pre_attn_scalar scaling |
| Gemma-2 dense (attn + final logit soft-cap) | Correctness-complete, speed-pending | Near-tie-band 48/48 (gemma-2-2b-it): 44/48 strict on vLLM's greedy + 4/48 at 0.0-nat ties in vLLM's own logits; proves the attention + final logit soft-cap primitives (attn_logit_softcapping 50 + final 30); the inverse of Gemma-3 (both soft-caps, no QK-norm) |
| Gemma-1 dense (the original Gemma) | Correctness-complete, speed-pending | STRICT token-exact 48/48 greedy (gemma-2b); two fused norms/layer, head_dim scale, GeGLU + sqrt(hidden) embed-scale, tied lm_head; no soft-cap/QK-norm/sliding. D1 (2026-07-31): the whole Gemma family (1/2/3/4) folded to the default-ON bf16 merged-QKV descriptor (MergedQkvEnabled); re-gated Gemma-2 SACRED 48/48 (global+sliding) + Gemma-4 STRICT 32/32 — its existing gate held |
| OLMo-2 dense (pure post-norm, full-width QK-norm) | Correctness-complete, speed-pending | Token-exact 16/16 (OLMo-2-0425-1B); first OLMo-family model; ZERO new compute kernel (pure post-norm norm_after + full-width QK-norm reuse existing ops); real ByteLevel-tokenizer gate (no BOS). D3 (2026-07-31, CLAIM-D3-OLMO2-FULLWIDTH-QKNORM): the full-width q/k RMSNorm + RoPE preamble folds onto the SHARED FusedChain catalog — new sibling recipe kAttnQkNormRopeFullWidth (the norm_full_width shape-param variant of Qwen3-dense's kAttnQkNormRope; FStep.norm_full_width additive, default false ⇒ every existing recipe byte-identical). Composite-only realization (fast_op=kNoFastOp) is byte-exact to the pre-fold standalone RmsNorm(qdim)+RmsNorm(kdim)+RopeFromCache sequence — CPU RED-first test_ops_fused_chain 10/10·379 GREEN (byte-exact, full-width proven distinct from per-head). DGX SACRED (olmo2_greedy_1b 16/16) + qwen3/27B canaries OWED — this lane ran on a CPU-only box (no GPU/cutlass); bit-exact by construction on CUDA (composite = the identical device-agnostic walker), empirical GPU run still owed. See docs/BENCHMARKS.md |
| Phi-3 / Phi-4 dense (Llama subclass, LongRoPE) | Correctness-complete, speed-pending | Token-exact (near-tie-robust) 16/16 (Phi-4-mini-instruct: 7 strict + 9 near-tie, 0 divergent) by the ratified root-divergence gate; bigger-dense STRICT anchor phi-4 (14B) 14/16 fully token-exact + 2 exact-tie (0.0 nats); pre-fused qkv/gate_up loader + LongRoPE cache bit-identical to vLLM's |
| Granite-3 dense (Llama + 4 scalar multipliers) | Correctness-complete, speed-pending | Token-exact 16/16 (granite-3.3-2b-instruct); first IBM Granite model; ZERO new compute kernel; embedding/residual/attention/logits multipliers (attention scale 1/64, not 1/sqrt(head_dim)) threaded over the shared dense path |
| StableLM dense (LayerNorm, partial rope, qkv bias) | Correctness-complete, speed-pending | Token-exact 16/16 (stablelm-2-1_6b): 14/16 strict on vLLM's greedy + 2/16 at bf16 near-ties (max gap 0.438 nats), 0 forward-divergent; first Stability AI model; ZERO new compute kernel (nn.LayerNorm with weight+bias + partial NeoX rope 16/64 + merged qkv bias, all reuse existing ops) |
| OLMo-3 dense (dual rope, interleaved sliding window) | Implemented, oracle-blocked | Loads + runs in our engine (dual rope: plain sliding + YaRN full-attn, per-layer sliding window); no SACRED gate: vLLM 0.25.0 oracle cannot run OLMo-3-1025-7B (KeyError: 'rope_theta'; transformers 5.13.1 nests rope_parameters per layer-type, no flat rope_theta; run-verified W0 2026-07-26) |
Laguna-S-2.1 MoE (LagunaForCausalLM, 118B/8B) |
REAL vLLM BAR ESTABLISHED (2026-07-31, CLAIM-LAGUNA-VLLM-NVFP4): FIRST-EVER vLLM Laguna run — prior numbers (incl. the correctness oracle) were all llama.cpp, never vLLM. vLLM on official poolside/Laguna-S-2.1-NVFP4 (single GB10, greedy, eager, MARLIN backend forced via VLLM_TEST_FORCE_FP8_MARLIN=1 because the auto-default FLASHINFER_CUTLASS needs an absent nvcc): ~18.8 tok/s (64-tok steady) — a LOWER bound. Our GGUF-Q4_K engine = 7.7 tok/s (vLLM ~2.4×); llama.cpp GGUF = 27.8 (still fastest at batch-1). llama.cpp is now a labeled SECONDARY "beat best-in-class GGUF" note; vLLM-NVFP4 is the headline bar. TRUE apples-to-apple still owes OUR NVFP4 Laguna forward arm (same tensor-core path as 27B/35B) — bring-up W-plan SPEC'D in .agents/specs/laguna-nvfp4-arm-2026-07-31.md (~85% reuse of the 35B NVFP4 W4A4 MoE infra + a name-map; bf16 attn/dense + fp4 experts; N1-N5 bricks, DGX-gated). N1-scaffold LANDED (2026-07-31): additive LagunaMoeWeights.experts_{gate,up,down}_fp4 + shared_{gate,up,down}_fp4 (Nvfp4Weight, mirror qwen3_5), dead until the N1 loader; CPU build clean + test_laguna_scaffold 8/8·167 unchanged. N1b loader IMPLEMENTED (2026-07-31, build-verified): LoadLagunaForCausalLMWeights (laguna_weights.cpp) replaces the VT_CHECK(false) stub — resolver + per-layer LoadBf16Direct (attn/dense/norms/embed/lm_head/router/shared-expert) + F32 e_score_correction_bias + LnLoadCtNvfp4Raw W4A4 experts. Name-map + dtypes VERIFIED against the real poolside/Laguna-S-2.1-NVFP4 index (router mlp.gate BF16, bias F32, experts W4A4, shared-expert BF16). N1b RUN-VERIFIED (2026-07-31): test_laguna_nvfp4_loader 2/2·29 — a synthetic 2-layer NVFP4 checkpoint (L0 dense, L1 MoE 2 W4A4 experts) round-trips byte-identically through the real loader (BF16 weights, expert packed/scale byte-identical + scale2/alpha math + per-expert indexing distinct, F32 bias, shared-expert BF16; RED-first missing-tensor throws); test_laguna_scaffold 8/8·167 unchanged. PENDING: N2 forward-branch + N4 DGX correctness / N5 speed-vs-vLLM-W4A4 (the real bar, faster than the 18.8 MARLIN lower bound). See docs/BENCHMARKS.md CLAIM-LAGUNA-VLLM-NVFP4. Prior DECODE-SPEED ATTRIBUTED (W7 profile-only, 2026-07-31, CLAIM-LAGUNA-W7-SPEED): nsys of the W6 decode on the real UD-Q4_K_XL GGUF (GB10) shows the 0.66 s/tok (~1.5 tok/s, vs llama.cpp 27.8 tok/s on the identical bytes, ~15-18x) is HOST-ORCHESTRATION, not kernel compute: GPU active only 32.7% of the step, 67.3% host/idle; 22,115 cudaStreamSynchronize (~2,764/step, zero GPU overlap) from the ~1,795 per-GEMM DrainQueue sync in LagunaForwardGgufCached + scalar host glue; 39.4% of the GPU time is QuantizeQ8K activation-quant (re-quantized per GEMM), the weight GEMVs run un-grouped at ~22% of the 240 GB/s peak (llama.cpp ~76%). Ranked levers (all in-tree from ds4): device-resident decode (kill the syncs) 1.5->~5-7 tok/s, grouped-expert GEMM (MatmulBTQuantGrouped) +1.5-2x + dedupes the activation-quant, decode CUDA-graph, tuned MMVQ; + free host cleanups (LagunaEmbed copies the whole 1.23 GB embed table/token; per-token RoPE-cache rebuild). Honest reachable ~13-20 tok/s, 27.8 a stretch. NO code changed. See .agents/specs/laguna-s21-w7-speed-2026-07-31.md. Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS full-recompute — TOKEN-IDENTICAL (byte-equal ids, md5 match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real 3-shard UD-Q4_K_XL GGUF (GB10, --gpu, keep-quant), same "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of "cultural capital"…". LagunaKvCache (mirrors DeepseekV4KvCache, MLA-latent → GQA multi-head K/V) caches post-QK-RMSNorm/post-RoPE K + raw V at f32 (bit-exact by construction: RoPE/QK-norm are position-only and attention is causal). MIXED attention handled per-layer: 12 GLOBAL layers grow the cache unbounded (full causal); 36 SLIDING-WINDOW-512 layers EVICT the oldest rows beyond the 512 window (gemma2/3 is_sliding), capping their K/V. LagunaForwardGgufCached + shared LagunaAttention/LagunaFfnBlock helpers used by BOTH forwards (identical float ops — the recompute path's ids are unchanged after the refactor); examples/laguna_gen --stateless forces the W5 recompute for the A/B gate. No cache bug: bit-exact on the first run. Next speed: grouped-expert GEMM + device-resident decode (both in-tree from ds4). See .agents/specs/laguna-s21-w6-2026-07-31.md. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10, keep-quant). laguna-gen "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of "cultural capital" as developed by French soci…" — the FIRST token is "Paris.", matching the llama.cpp-Poolside reference on the identical bytes. Multi-shard GGUF reader (LagunaGgufCtx routes each of 814 tensors to its shard; shard-1 = header only) + keep-quant tower (attn/dense/shared/experts/lm_head stay Q8_0/Q4_K/Q5_K COMPRESSED, consumed via vt::MatmulBT; norms/router/bias/embed → f32) + LagunaForwardGguf (the f32 composition with the ~9 GEMM sites swapped to keep-quant Gemm/GemmRowSlice, ds4 precedent) + examples/laguna_gen. Real GGUF metadata verified: dual-RoPE freq_base 500000/10000, dims 64/128, YaRN factor 32, sigmoid ungrouped-noaux router (scale 2.5), per-layer Q-head [48 global/72 sliding], per-head softplus out-gate, QK-RMSNorm. Load 20.6s, peak 71 GiB (fits 119 pool). Prior W4 IN PROGRESS (2026-07-31): 73.4 GiB UD-Q4_K_XL GGUF FETCHED + read authoritatively (814 tensors); 3 CPU-verified fidelity corrections grounded in the real GGUF + llama.cpp — per-head QK-RMSNorm (attn_q/k_norm, the scope MISSED it), GGUF-authoritative dual-RoPE mscale (llama.cpp yarn_attn_factor·(1+0.1·ln(factor)), factor 32 not HF 128), separate ffn_gate/up_exps. Keep-quant tower materialization + ForwardGguf + the real-model greedy run vs llama.cpp-laguna same-quant oracle = W5 close. Prior: W3 REAL host-reference forward + 3 new ops (laguna_ops.cpp, CPU -Werror clean, test_laguna_scaffold unit-gated) |
Poolside Laguna: 48 layers (12 global + 36 sliding-window-512), 256 routed top-10 + 1 shared expert, per-head softplus attention output gate, sigmoid noaux_tc router, dual per-layer RoPE (YaRN full-attn / plain sliding), GQA 8 KV / 128 head-dim, 1M ctx. W3 (2026-07-31): the 3 genuinely-NEW small host ops landed in laguna_ops.cpp — per-head softplus attn out-gate (LagunaSoftplusHeadGate), ungrouped sigmoid-noaux router (LagunaUngroupedRouterTopK, ds3 noaux_tc MINUS the group step + tie-break razor), dual per-layer RoPE cos/sin builders (BuildLaguna{FullYarn,Sliding}CosSin, reusing the pinned YaRN inv_freq over the partial-64 dims); LagunaModel::Forward is now a REAL runnable host-reference composition (variable-Q-head GQA + dual RoPE + sliding-window mask + softplus gate + dense L0 / ungrouped-MoE L1..47 + untied lm_head) replacing the VT_CHECK(false) stub; test_laguna_scaffold 8/8·166 (softplus math, router selection+tie-break RED-first, dual-RoPE bit-match, variable-Q-head shapes, forward composition on synthetic weights), test_model_registry 24/24. W2 (2026-07-30): registered, ParseLagunaParams, GGUF blk.N.* name-map + UD-Q4_K_XL quant-mix (Q4_K/Q5_K/Q6_K/Q8_0 ALL already decoded → ZERO new kernel). W1 oracle DECISION: vLLM NATIVE laguna.py (in pin → config constructs); dual-oracle = vLLM-NVFP4/-FP8 (fits GB10 119 GiB; BF16 235 GiB does NOT) + llama.cpp-Q4_K token-exact. ~85–90% reuse (ds4-MoE + Gemma-sliding + OLMo-3-dual-rope + Q4_K keep-quant, ALREADY landed). DEFERRED (W4): GGUF keep-quant tower materialization + device/paged production forward (loaders still LOUDLY throw) + strict dual-oracle greedy gate on a fetched checkpoint + poolside_v1 parser. See .agents/specs/laguna-s21-w3-2026-07-31.md (+ W1/W2 laguna-s21-w1w2-2026-07-30.md, W0 laguna-s21-scope-2026-07-30.md) |
InternLM2 dense (fused-wqkv interleaved split) |
Correctness-complete, speed-pending | Token-exact 16/16 (internlm2-chat-1_8b): 12/16 strict + 4/16 bf16 near-tie (max gap 0.0 nats), 0 divergent; first InternLM model; ZERO new compute kernel (reuses the Llama dense forward; the only delta is a loader-side de-interleave of the fused wqkv, which packs q/k/v interleaved by KV-group) |
Command-R / Cohere dense (CohereForCausalLM) |
Implemented, gate-blocked | ZERO-new-kernel port grounded in vLLM commandr.py: weight-only Cohere LayerNorm + GPT-J full-width RoPE + PARALLEL residual + logit_scale + tied embeddings, all reuse; compiles, links, self-registers. No SACRED gate yet (real checkpoints HF-gated, ungated ones tiny-random, GPU box disk-full); oracle run-verified at W0. See docs/BENCHMARKS.md |
Phi-1 / Phi-2 dense (PhiForCausalLM, parallel residual) |
Correctness-complete, speed-pending | Token-exact 16/16 (microsoft/phi-2): 9/16 strict + 7/16 bf16 near-ties (max gap 0.25 nats), 0 forward-divergent; the OLDER Microsoft Phi arch, DISTINCT from Phi-3/Phi-4; ZERO new compute kernel (GPT-J parallel residual, LayerNorm-with-bias, biased qkv/dense, partial NeoX rope 32/80, non-gated NewGELU MLP reusing vt::GeluTanh, untied biased lm_head); F16 dtype-aware loader |
MiniCPM dense (MiniCPMForCausalLM, three scalars) |
Correctness-complete, speed-pending | Token-exact 16/16 (openbmb/MiniCPM-2B-sft-bf16): 10/16 strict + 6/16 bf16 near-ties (max gap 0.0 nats), 0 forward-divergent; first OpenBMB MiniCPM model; ZERO new compute kernel (the Llama/Granite dense forward plus three scalars: scale_emb, scale_depth/sqrt(layers) residual, dim_model_base logit scaling), tied lm_head; .bin-only weights converted to safetensors via trusted torch |
MiniCPM3 (MiniCPM3ForCausalLM, MLA + three scalars) |
Correctness-complete, speed-pending | Token-exact 16/16 (openbmb/MiniCPM3-4B): 13/16 strict + 3/16 bf16 near-ties (max gap 0.0 nats), 0 divergent; the first MLA-attention MiniCPM; ZERO new compute kernel (the MiniCPM three scalars with attention swapped GQA to DeepSeek-style MLA, reusing the landed DeepSeek-V2 MLA block; DeepSeek-V2 re-gated 8/8 unchanged); tied embeddings, .bin-only weights converted via trusted torch |
| Yi (Llama architecture) | Correctness-complete, speed-pending | Token-exact 16/16 (01-ai/Yi-Coder-1.5B-Chat): 13/16 strict + 3/16 bf16 near-ties (max gap 0.125 nats), 0 divergent; modern Yi adopted the Llama architecture (architectures: ["LlamaForCausalLM"]), so it runs on the existing Llama path with zero code changes; confirms a distinct non-Llama-branded checkpoint (64000 vocab, RoPE theta 1e7) loads and decodes correctly |
InternLM3 (InternLM3ForCausalLM, Llama alias) |
Correctness-complete, speed-pending | Token-exact 16/16 (internlm3-8b-instruct): 14/16 strict + 2/16 bf16 near-ties (max gap 0.0 nats), 0 divergent; a plain Llama architecture in vLLM 0.25.0 (registered as a one-line Llama alias; dynamic-NTK RoPE factor 6.0, GQA kv=2, untied lm_head), not InternLM2 + sliding window; zero forward or loader delta |
| Long-context RoPE + sliding-window attention | Correctness feature-positive on GB10, speed-pending | Shared scaled-RoPE (YaRN, Llama-3, Phi-3/4 LongRoPE, dynamic-NTK) + sliding-window attention, GPU-gated vs the vLLM 0.26 oracle: LongRoPE (Phi-4-mini) + llama3 (Llama-3.2-1B) + dynamic-NTK (InternLM2) 16/16, sliding-window (Gemma-2/Gemma-3) 48/48, operator local-mask kernel positive. YaRN and chunked-local model e2e are reachable-blocked (no cached vehicle); speed pending |
| LoRA / multi-LoRA adapters | In progress (W1 CPU brick), not yet usable end-to-end | Highest-demand missing feature. W0 spike + W1 CPU runtime brick landed (LORA-RUNTIME ACTIVE): the LoRALayerWeights container, the punica shrink/expand ops (-1-slot skip), AddLoraLinear, and a single-linear LoRALinear (ReplicatedLinear create/set/reset/apply, scaling folded at SetLora), RUNTIME-VERIFIED on CPU (test_punica_cpu 6/6, 101 assertions, vs double-precision references, RED-first). Packed/TP/merged layers, mapping metadata, adapter load, LRU multi-adapter manager, the load/unload endpoints, and the GPU kernels + model gate are named W2-W7 in .agents/specs/lora-adapter.md. No model can be served with an adapter yet |
| Safetensors loading | Supported | Both gate models plus every registered dense/MoE family |
| GGUF loading (F32/F16/BF16/Q4_0/Q8_0/Q2_K/Q3_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/IQ3_XXS/NVFP4) | Supported; compute-in-quant (keep-quant) on CPU AND now CUDA for the six K-block encodings PLUS Q2_K/IQ2_XXS/IQ3_XXS (DeepSeek-V4 W8, 2026-07-29 - the FIRST CUDA keep-quant GGUF k-quant GEMM KERNEL-QUANT-CIQ-GEMM-CUDA, MMVQ-style dequant-in-kernel, GB10-gated 92401/92401 vs the CPU oracle, so a CUDA runner keeps blocks compressed and dots them on the GPU instead of the ARM cores); NVFP4 now COMPUTES IN FP4 on CUDA for the dense-MLP and full-attention projections (2026-07-29, CLAIM-GGUF-NVFP4-COMPUTE), no longer materialize-only |
Weights in six block encodings stay compressed from file to matmul on CPU (no BF16 expansion). NVFP4 (ggml type 40) DEQUANTIZES, including the per-tensor (per-expert) <stem>.scale sidecar the container keeps outside the blocks; gated BIT-EXACT against the compressed-tensors NVFP4 path on real Qwen3.6-27B bytes from both containers. It no longer expands to bf16 on CUDA: an NVFP4 matmul/expert weight is REPACKED at load into the same (weight_packed [N,K/2], weight_scale [N,K/16]) operand pair the compressed-tensors path produces - a pure byte permutation, gated BYTE-IDENTICAL against that container - and the existing vt::MatmulNvfp4* kernels run on it, so no new kernel exists and no numerics are re-derived. Covers the dense MLP + full-attention q/k/v/o and the MoE shared/routed experts; the GDN in_proj_* family and ssm_out still expand (the V-head reorder rewrites their layout) and a CPU build still expands everything - the documented part subset. MEASURED on GB10 (2026-07-29), same-binary A/B, one flock, idle box, 2 reps per arm: peak RSS 50.8 -> 25.7 GiB, load-and-generate 1:58 -> 0:41, and the 256 projections that move cost 35 840 MiB expanded against 10 080 MiB fp4-resident (3.56x). The divergence against the safetensors sibling CLOSES: the fp4 arm is token-IDENTICAL over 24 greedy tokens where the bf16 arm of the same binary diverges at index 4, which retires the reading that that divergence was permanent. It is REPORTED, not gated: the two containers are not the same model - the GGUF NVFP4-quantizes 192 GDN in_proj tensors the safetensors keeps BF16 (mean relative weight error ~0.18) and their activation global scales differ - so identity is not guaranteed and a cross-container throughput arm is not valid. SACRED gates unmoved: test_qwen27_paged_engine 235/235, test_qwen36_paged_engine 315/315. The MoE (35B) stacked-expert arm is now HARDWARE-GATED too (2026-07-29), superseding the gap recorded here: the real 35B A3B NVFP4 GGUF loads and generates through the fp4 path, its 120 routed-expert stacks x 256 experts repack to the modelopt safetensors' own operands with ZERO differing bytes over 840 sampled (tensor, expert) slabs, and all 840 per-expert <stem>.scale[e] are bit-identical to that expert's weight_scale_2 - the per-expert scale INDEXING, mutation-proved against both a scales[0]-for-every-expert and an expert-0-slab-for-every-expert mutant. Same-binary A/B: peak RSS 68.5 -> 22.7 GiB (3.01x), load-and-generate 1:51.9 -> 0:28.8, tokens IDENTICAL (correct here, since the 35B routed experts run the W4A16 grouped GEMM in both arms). Recorded as OPEN, not smoothed over: this case's 24-token greedy stream is NOT run-to-run stable (one of three use_a16 runs and one of four safetensors-reference runs differed), so the binding results are the weight-level byte identity and the residency audit, not a token-exactness claim; test_qwen36_paged_engine is token-exact at ITS engine params, so the instability belongs to this case's configuration and attributing it is owed work. That run also found and FIXED a latent defect the MoE arm made reachable: the two fp4 fused MoE blocks issued the router GEMM assuming the safetensors [K,N] gate layout and threw matmul: inner dims mismatch on the GGUF's [N,K] one; MoeRouterLogits now branches on nk (inert for the safetensors path, SACRED gates unmoved). Q2_K (id 10) + IQ2_XXS (id 16) DEQUANTIZE (2026-07-29, CLAIM-DSV4-GGUF-LOADER): the ~2-bit types the single-Spark DeepSeek-V4-Flash-GGUF UD-IQ2_XXS/UD-Q2_K_XL vehicles use, ported 1:1 from llama.cpp ggml-quants.c (iq2xxs_grid codebook + signs; Q2_K nibble sub-scale/min), unit-gated on hand-derived bytes (test_gguf_dequant 15/15). Dequant-only (no vec_dot -> expand-bf16). A V4-GGUF model still cannot RUN: the V4-GGUF name map (tensor-manifest-blocked) + the V4 forward (W3-W8) remain |
| AWQ / GPTQ quantization | W0 spike + W1 CPU INT4 dequant primitive; not yet loadable end to end | INT4 unpack+dequant-to-bf16 for BOTH community formats, mirroring vLLM 1:1 (AWQ reverse-order awq_triton.py; GPTQ qdq_4.cuh with zero_offset v1/v2 + act-order g_idx). Unit-gated RED-first (hand-computed known bytes + double-precision roundtrip). NOT wired to a loader, no GPU Marlin compute, no model run yet: config recognizer (W2), Marlin GPU GEMM riding the vendored NVFP4 Marlin (W4), CPU e2e (W3), GPTQ 8/2/3-bit (W5) and MoE (W6) are named next bricks. See .agents/specs/awq-gptq-quant.md |
MXFP4 (compressed-tensors mxfp4-pack-quantized) |
Weight dequant landed (CPU), build-only; compute + e2e pending | The shared quant path DeepSeek-V4-Flash (MXFP4 experts) and Kimi-K3 (its real checkpoint) both need. CPU weight unpack + E8M0 dequant to bf16/f32 (2^(byte-127), group 32, no global scale — distinct from our NVFP4 group-16 fp8-scale path); unit-gated 5/5·1142 vs a double-precision dq_mxfp4_torch port (RED-first, bias-128 mutation fails 446 assertions). GPU W4A4 fp4 GEMM + Marlin W4A16 fallback + MoE expert path + e2e are named later bricks; the two model loaders keep their MXFP4 refusal until wired |
| CPU backend vs llama.cpp | At or ahead on every axis (GGUF) | Prefill 1.18x ahead, decode at parity, peak memory 1.01x, byte-identical greedy tokens. Single-stream only; no concurrent-serving comparison has been measured |
| Paged KV cache + prefix caching | Supported | Block-paged full attention, hybrid full-attention + GDN state groups, automatic prefix caching (APC) on by default for dense models (cache-ON gated end to end: token-identical output, cache hits, faster TTFT) |
fp8 KV cache (cache_dtype=fp8) |
In progress (W1 CPU brick), not yet usable end-to-end | HIGH-priority memory/throughput lever (halves the KV footprint). W0 spike + W1 CPU brick landed (KV-FP8 ACTIVE): fp8-e4m3 K/V STORE (Quantize(hp/scale)) + the paged-attention READ dequant (Dequant(fp8)*scale) + the cache_dtype config parse, all CPU-gated RED-first (test_ops_fp8_kv_cache 8/8·511; a wrong store direction fails 3/480). Storage is 1-byte fp8 (DType::kI8) + a Fp8KVCacheDataType interpretation enum, per-tensor k/v scales (mirroring vLLM BaseKVCacheMethod). The CUDA store + fp8 paged-attention read (the GPU memory-halving path, DGX-blocked), the runner/spec integration (half-sized KV blocks + checkpoint-scale threading + --kv-cache-dtype/--calculate-kv-scales), fp8_e5m2 and per-head scales are named W2-W5 in .agents/specs/fp8-kv-cache.md. No model can run with an fp8 KV cache yet |
Prefix-cache matching unit (--prefix-match-unit) |
Partial (resolver landed, config/scheduler wiring pending) | 0.26-new knob setting the finest token boundary a prefix-cache hit can land on (the hash_block_size). W1: resolve_kv_cache_block_sizes ported 1:1 (hybrid hash_block_size = prefix_match_unit if set else gcd(group block sizes); single-group inert; back-off on no-consumer / mamba-non-align; throws on non-divisible), CPU unit-gated RED-first (default gcd vs =16). Pending: the config/CLI/ABI field (W2), scheduler threading of a finer-than-block matching unit (W3, needs the KV-block-pool align path), and the benchmark (W4). Default path byte-identical (dense single-group models ignore it). |
CLI chat / complete (SERVE-CLI-CHAT) |
SPIKE (implementation pending) |
The accepted CPU-only spike corrects a stale inventory claim: pinned vLLM 0.26 ships both commands as clients of a running OpenAI-compatible server. The selected design mirrors their URL/auth/model-discovery, quick/interactive, SSE, history, and TTFT/TPS contracts while preserving the existing in-process --model --prompt invocation as a compatibility alias. No remote command code or support claim has landed; next gate is W1 parse/dispatch plus W2 fake-server transcript tests. See the spike. |
| KV offload to CPU / disk | Built, opt-in, off by default; the disk connector is engine-refused | CPU and disk tiers with identity-checked blocks, selected by --kv-transfer-config (or programmatically) over one abstract KVConnector ABI. Worker-side KV store/load is implemented for the LMCache connector only; the CPU/disk connector is scheduler-side only, so the engine now REFUSES it at construction (a loud error, not silently wrong output). Guide: docs/KV-OFFLOAD.md |
LMCache client (lm:// remote KV) |
Built, opt-in, off by default; a working, verified external KV cache | Pure-C++ lm:// client wired as an LMCacheConnector, no lmcache in-process; keys agree byte-for-byte with a real vLLM+LMCache peer, mismatched blocks refused. Proven in a real OPT-125m loop vs a live lmcache.v1.server: connector-ON tokens are BIT-IDENTICAL to the connector-OFF cold run (both after an in-process restart and from a cold second process). See docs/BENCHMARKS.md |
| KV-cache events (for external routers) | Built, off by default; generation and payload gated, live ZMQ transport deferred | Block store/remove/clear events (BlockStored/BlockRemoved/AllBlocksCleared) emitted at the prefix-cache sites when enabled, with a msgpack payload byte-identical to vLLM's msgspec encoding. Behind a publisher seam faithful to --kv-events-config; the live ZMQ transport is not wired yet. Off by default, so the prefix-cache path is byte-identical. |
| Sampling | Supported | Greedy, temperature, top-k/p, min-p, presence/frequency/repetition penalties, seed, stop/stop_token_ids, min_tokens, logit_bias, allowed_token_ids, bad_words, in vLLM's exact order. Custom logits processors are supported through a C-ABI callback (vllm_logits_processor, ABI v8): a per-request host callback invoked each decode step to modify the logits before sampling; absent by default (byte-identical). Sample logprobs are emitted end-to-end for /v1/completions and /v1/chat/completions (logprobs/top_logprobs). Parallel sampling (n>1) is supported: a request fans out into n sequences sharing the prompt, returned as n indexed choices (n==1 byte-identical). Beam search is supported through the BeamSearch driver — an outer engine loop that scores beams by cumulative logprob with a length penalty and returns the top beam_width sequences (deterministic, token-exact vs vLLM's algorithm); it is wired on the OpenAI use_beam_search request field for /v1/completions and /v1/chat/completions over BOTH the synchronous engine AND the production AsyncLLM HTTP server (via an async BeamSearchAsync driver that reuses the same scoring and returns beams token-identical to the sync driver), so beam search runs on the real server (the C-ABI beam params and streaming beam are not exposed yet; per-beam concurrent stepping is a named residual — beams are stepped sequentially, byte-identical to the sync driver). best_of is supported on both endpoints (generate best_of sequences, return the n highest-cumulative-logprob ones; best_of==n is the default no-op) — vLLM 0.26 itself has dropped best_of from its live path, so this follows the classic OpenAI contract. Prompt logprobs and echo are not emitted yet |
| Structured output | Supported (subset), engine-enforced; xgrammar backend W1 (CPU, not yet production-wired) | JSON schema, JSON object, regex, choice, GBNF grammar. Constrained decoding runs in the production engine (native grammar backend, per-step logits bitmask) and is reachable from OpenAI response_format and the C ABI (ABI v2 structured_* fields). A second, xgrammar-faithful backend (XgrammarStructuredOutputBackend, vLLM's default auto) is built behind the same seam: it reuses the native pushdown-FSM/trie matcher (xgrammar's own algorithm) and adds the xgrammar JSON-schema→EBNF converter that preserves property declaration order + any_whitespace + the basic_* grammar, closing the key-order/whitespace/exotic-schema parity gap. CPU-gated (test_backend_xgrammar 6/6, RED-first); production wiring + the auto fallback + GPU oracle parity are the named residuals (see .agents/specs/xgrammar-backend.md) |
| Tool-call parsing | 36 parser families / 40 accepted names, streaming | Every vLLM tool parser at the pin except the three Rust/Harmony-backed ones: pure-text parsers ported 1:1, the six engine-backed families reimplemented from their wire formats, all held to the upstream test suites. Selection via --tool-call-parser (server), tool_parser (C ABI), or template auto-detection; native-syntax forced tool_choice where expressible. Tables: docs/BENCHMARKS.md |
Reasoning parsing (SAMPLE-REASONING, ACTIVE, partial coverage) |
9 parsers, streaming | think_auto (auto-detect default: content unless markers appear), deepseek_r1, deepseek_v3 (passthrough) / holo2 (thinking→R1), mistral ([THINK]), minimax_m2 (+append_think), step3, olmo3 - reasoning split engine-side BEFORE tool parsing, streamed as reasoning deltas in the chat chunks. Coverage: 9 of upstream's ~28 registered names (remaining text families + engine-backed adapters tracked as W2/W3 in specs/reasoning-parsers.md); each ported parser doctest-gated vs its tests/reasoning case |
| Unified streaming parser engine | Core, assembly, serving-SSE dispatch landed, gated; all 10 engine-backed families ported (family parity closed); JSON-schema tool-arg type coercion landed | The vLLM 0.26 declarative parser/engine/ (shared state machine plus all 10 configs: qwen3, seed_oss, kimi_k2, minimax_m2, glm47_moe, deepseek_v4/v32, nemotron_v3, gemma4, inkling) and assembly layer, gated field-for-field vs vLLM 0.26. An engine-backed --tool-call-parser name drives the live chat SSE chunks, off by default. When a request's tools declare typed parameters, the assembled tool-call arguments are coerced to the declared JSON types (int/number/bool/string/array/null) 1:1 with vLLM _fix_arg_types, in both streaming and one-shot; no schema means the arguments pass through as strings unchanged. Details: .agents/specs/parser-assembly-c8.md |
| OpenAI server | Supported (subset) | /v1/completions, /v1/chat/completions, streaming SSE, /v1/models, /health, /version, /ping, /metrics (Prometheus vllm:* names), /tokenize (both the raw-prompt and the chat-messages forms, the latter rendered through the same model chat template as chat-completions), /detokenize, /tokenizer_info (opt-in via --enable-tokenizer-info-endpoint, mirrors vLLM's enable_tokenizer_info_endpoint; surfaces the tokenizer-config fields the BPE tokenizer can genuinely back), /server_info, /reset_prefix_cache, /abort_requests (abort in-flight requests by id; dev-mode via --enable-server-dev-mode, mirrors vLLM's VLLM_SERVER_DEV_MODE). The production vllm-server binary now serves /tokenize, /detokenize by default and /tokenizer_info / /abort_requests behind those flags, wired to the live engine (CLAIM-C8-SERVE-PROD-WIRING); /metrics and /reset_prefix_cache currently have no live backing on the async serving path (handlers exist; production wiring pending an engine stat-logger / reset RPC accessor) |
| Pooling task class (embeddings / classify / score / rerank) | Spiked; pooler op + heads composite + pooling runner path landed (CPU), not yet servable end-to-end (no concrete model / endpoints) | The non-generative task class. W0 spike over the whole vLLM pooling surface (.agents/specs/pooling-task-class.md, CLAIM-POOLING). W1 landed the pooler OP (CLS/LAST/MEAN + Identity/Normalize/MultiLabelClassify/Classify activations, double-precision-gated). W2 landed the pooler HEADS composite — EmbeddingPoolerHead (projector→matryoshka→normalize), ClassifierPoolerHead (classifier→(logit-mean)/sigma→activation), the SequencePooler + PoolerForEmbed/PoolerForClassify factories, the DispatchPooler task routing, and the PoolerConfig/PoolingParams structs (test_pooler_heads 27/27, 240 asserts, RED-first). W3 landed the pooling RUNNER path — PoolingRunner applies the model's Pooler to the last hidden state and returns pooled embeddings instead of sampled tokens, gated by a STRUCTURAL cosine-parity check vs a double-precision LAST+normalize reference (test_pooling_runner 5/5, 14 asserts, RED-first). NOT yet servable / honest residuals: a concrete pooling MODEL forward + the REAL-model oracle cosine gate (vllm.LLM(task="embed").encode) — no cosine-vs-oracle number is fabricated (W3-model); the /v1/embeddings + score/rerank/classify endpoints (W4); tokwise AllPool/StepPool (W5). See docs/BENCHMARKS.md |
| Plugin system (out-of-core registration) | Spiked; first CPU brick landed, not yet wired into any production path | The extensibility-first discovery layer. W0 spike over vLLM's plugin surface (general / platform / io_processor / endpoint groups, the register_model an out-of-tree plugin calls, the invocation seams) is committed (.agents/specs/plugin-system.md, ENG-PLUGIN-SYSTEM ACTIVE, CLAIM-PLUGIN-SYSTEM). W1 landed vllm::plugins::LoadGeneralPlugins() + the out-of-core general-plugin registration seam (RegisterGeneralPlugin / REGISTER_VLLM_GENERAL_PLUGIN) over the existing REGISTER_VLLM_MODEL-style registries: a 1:1 mirror of load_general_plugins (load-once idempotence, the VLLM_PLUGINS allowlist, per-plugin failure isolation). Proven by an out-of-core toy-model plugin that registers a toy architecture through the public RegisterModel seam — unit-gated RED-first (test_plugin_system 1 case / 29 assertions: the toy arch resolves ONLY after LoadGeneralPlugins runs it, and not under VLLM_PLUGINS=""). Python entry points have no C++20 analogue, so discovery is the project's static-init/dlopen registration idiom (recorded porting-inventory §9). NOT yet wired: real shared-object dlopen + the C-ABI vllm_plugin_register entry (W2), the engine/CLI --load-plugins wiring that calls LoadGeneralPlugins from the construction paths (W3), the platform/quant plugin kinds (W4), and the io_processor/stat_logger/endpoint groups (W5) are named residuals. See docs/BENCHMARKS.md |
| Offline Batch API (JSONL file runner) | Spiked; first CPU brick landed, not yet exposed as a CLI | The offline OpenAI Batch API: read a JSONL of OpenAI-format requests, run each through the engine, write a JSONL of responses. W0 spike over vLLM's run_batch.py (schema, endpoint dispatch, run loop, file I/O) is committed (.agents/specs/batch-api.md, SERVE-BATCH-API ACTIVE, CLAIM-BATCH-API). W1 landed RunBatch (RunLine/RunLines/Run) + RunBatchFile — a pure orchestrator over the existing OpenAIServingChat::create_chat_completion (NO reimplemented generation), 1:1 with vLLM's endpoint_registry url→handler map: /v1/chat/completions dispatch, the BatchResponseData/BatchRequestOutput schema (vllm-<uuid> ids, custom_id echoed), the run_request AllResponse/ErrorResponse/stream branches, and the unsupported-endpoint/url error rows. Unit-gated RED-first (test_openai_run_batch 7 cases / 80 assertions over the synthetic serving engine: ordered rows + custom_id echo + per-line BatchRequestOutput schema round-trip, a malformed line isolated into an error row so the batch continues, dispatch + 404 error rows; dropping the custom_id echo fails 9 assertions). Recorded deviation: a malformed line is isolated (batch continues) where upstream aborts the job. NOT yet exposed: the vllm run-batch CLI + BatchFrontendArgs (W2), embeddings/score/rerank dispatch (W3, rides pooling endpoints), audio transcription/translation + media fetch (W4), and http(s)/data-URL file I/O + metrics + overlapped AsyncLLM submission (W5) are named residuals. See docs/BENCHMARKS.md |
| Tokenizers | Supported | Byte-level BPE (Qwen/Llama-3/OPT/GPT-2/DeepSeek/OLMo-2) and SentencePiece BPE (Mistral/Gemma), plus GGUF vocab; added-token lstrip/rstrip whitespace semantics (e.g. Phi-4-mini's special tokens); byte-exact vs the vLLM oracle |
| Multimodal: image to text | Correctness-complete; vision-forward speed beats vLLM; OpenAI-server content-part parse + processor routing + engine mm-request plumbing landed (CPU); end-to-end serving pending | Strict token-exact 32/32 vs vLLM 0.25.0 on Qwen3-VL-4B and Qwen3.6-27B (Qwen3_5ForConditionalGeneration); C++ image processor + vision tower + MRoPE/DeepStack backbone with on-GPU greedy sampling; the 27B decode step is now graph-capturable (routes through the production captured decode, token-exact held). The vision tower now defaults to the flash-tiled non-causal attention (vt::AttentionDenseFlash, byte-identical to the previous warp kernel — token-identical, goldens unchanged): the per-image tower forward is ~142 ms vs vLLM 0.25.0's ~250 ms eager encode = 0.57x (faster). Attribution-first profiling found the tower attention is serial-latency-bound (not K/V-bandwidth-bound), so flash tiling is only a 1.04x same-binary A/B over the warp kernel here — the tower already beats vLLM; batched serving speed still pending. OpenAI-server multimodal wiring — CPU bricks 1+2 landed (CLAIM-MM-SERVING-W1/W2, .agents/specs/mm-serving.md): the chat request parses the OpenAI multimodal content-part array (image_url/input_audio/audio_url), decodes the base64/data: payloads, routes them through the existing single-sequence processors to a placeholder-expanded prompt + mm-feature handles, AND the engine now carries them — additive LLMEngine/AsyncLLM add_request(MultiModalInputs)/generate overloads (via InputProcessor::process_inputs_mm, mm_features onto EngineCoreRequest/Request), chat-template placeholder-string helpers, and a serving_chat MultiModalChatFn seam (default unset ⇒ text path byte-identical). The W3 MultiModalChatFn seam BODY now lands (CLAIM-MM-SERVING-E2E): MakeQwen3VLImageChatFn (chat_mm.{h,cpp}) turns an image chat request into the placeholder-EXPANDED engine input — marker-inject at the mm part position, render the chat template, tokenize (the single `< |
| Multimodal: video to text | Correctness-complete, speed-pending; OpenAI-server content-part parse + engine mm-request plumbing landed (CPU), end-to-end serving pending | End-to-end on Qwen3-VL-4B (near-tie-robust) and Qwen3.6-27B (strict 32/32); reuses the image tower and temporal MRoPE plus video preprocessing (frame sampling, temporal grid) |
| Multimodal: audio to text | Decode beats vLLM; encoder TTFT measured and improved but not yet at parity; OpenAI-server content-part parse + audio processor routing + engine mm-request plumbing landed (CPU), end-to-end serving pending | End-to-end on Voxtral-Mini-3B (near-tie-robust, decoder token-exact 48/48) plus a Whisper-class encoder tower. Decode is graph-captured and runs vLLM's FA2 varlen split-KV kernel (the KV block-size is rounded up to a multiple of 16 to meet its precondition): audio TPOT 39.5 ms/token BEATS vLLM 0.25.0's 40.8 ms (0.97x, non-overlapping bands). The Whisper encoder attention now uses a flash-tiled non-causal head-dim-64 kernel (vt::AttentionDenseFlash): a block of query-warps shares each streamed K/V tile out of shared memory (FlashAttention K/V tiling), with the per-warp online-softmax math copied verbatim from the warp kernel so the output is bit-identical (token-identical). Same-binary A/B: the encoder self-attention drops from 35.11 to 19.29 ms/layer (1.82x, non-overlapping) and the encoder forward from ~1.83 s to ~1.37 s (1.33x), token-identical output (16/16, STRICT golden unchanged, sanitizer 0). The encoder weights are now device-resident: each of the 487 encoder weight tensors is converted to bf16 and uploaded to the GPU once (mirroring the decoder's residency), reused across forwards instead of being re-marshalled every call; same-binary A/B (VT_WHISPER_ENC_REMARSHAL) drops the encoder forward a further ~1.37 s to ~0.73 s (1.89x, non-overlapping), removing ~648 ms of per-call host weight marshalling (nsys: 974 fewer Host-to-Device copies, ~2.5 GB less traffic), byte-identical (16/16, STRICT golden unchanged, sanitizer 0). But encoder TTFT (~0.73 s) is still far above vLLM's 43 ms (~17x, was ~32x): the residual is now GPU-compute-bound (the scalar warp-per-query attention plus the conv GEMMs), so closing it needs a tensor-core MMA head-dim-64 non-causal flash attention. Correctness held under the ratified distributional near-tie gate (teacher-force PASS, 0 divergent, strict prefix exact vs vLLM; STRICT golden unchanged; 16/16). Remaining: a tensor-core MMA encoder-attention kernel and dropping the conv host round-trip (a device im2col kernel), and there is no batched c2+ or audio_url serving ingestion. See docs/BENCHMARKS.md |
| SGLang parity (competitor floor + oracle) | Oracle STOOD UP + first floor MEASURED (cache-neutral, 27B, c8/c16): throughput/TTFT WIN, TPOT/ITL open GAP | SGLang v0.5.15 (f63458b) whole runtime surface inventoried (44 rows: 23 FUSED into our vLLM-derived engine, 8 SGLANG-DISTINCT opt-ins, 5 inventoried, 8 out-of-scope). The v0.5.15-cu130 arm64 image (@sha256:d0a667e) PULLED and RAN the 27B-NVFP4 gate model on GB10 sm_121a with no from-source build. First reproduced SGLang-vs-ours cache-neutral comparison (27B, 3 reps, idle box, one flock, engines sequential): ours beats the SGLang floor on total/output throughput + req/s (2.21×@c16, 1.44×@c8) and TTFT (6–12× lower), but SGLang wins per-token latency (TPOT/ITL 1.18–1.49× below ours) — a reproduced OPEN GAP. SGLang is a competitor, not the mirror source (vLLM stays behavior truth). Residuals: 35B, c1/c2/c4 low-conc sweep, shared-prefix cache-ON arm, token-exact cross-check. Numbers + repro: docs/BENCHMARKS.md, .agents/sglang-matrix.md, .agents/specs/sglang-parity-oracle.md. UPDATE 2026-07-28 (CLAIM-DECODE-LATENCY-EXPLORE, measurement only, no source changed): the TPOT/ITL gap is CONFIRMED batch-composition, NOT a decode-kernel deficiency — on our engine ITL(decode-batch=1)=101.75 ms is already ≤ SGLang's op-point 104–105 ms and rises monotonically with batch (→158.5 ms @ B16); nsys shows every hot decode kernel sub-linear in batch (per-token cost ↓~10×); SGLang's effective decode concurrency is ~4 (not 16) due to its 33 s admission queue, so its low ITL is simply the ITL of a small batch. Our throughput win IS the ITL cost — same lever; knob max_num_seqs/max_num_batched_tokens already exists, latency-oriented point max_num_seqs≈8 = ITL −21% at 1.38× SGLang throughput; default stays throughput-oriented. Full data: .agents/specs/decode-latency-lever.md |
Speculative decoding is available on the Qwen3.5/3.6 checkpoints via
--speculative-config. MTP (k=1) is end-to-end token-exact vs vLLM on both
gate models (the 27B GDN hybrid Qwen3_5MTP and the 35B MoE Qwen3_5MoeMTP):
three-way identical at concurrency 1 (our spec-ON == our spec-OFF == vLLM
--speculative-config mtp greedy) and faster than vLLM there, on par or above
at higher concurrency (mixed-batch), with the draft head alive and acceptance
matched to vLLM.
Block-diffusion DFlash (the z-lab Qwen3.6-27B draft over the 27B NVFP4 target) is correctness-complete and at or above vLLM throughput on GB10 (final concurrency-1 A/B our-on 29.32 tok/s vs vLLM-on 29.24, non-overlapping bands, 1.003x); it is recorded DONE across the engine, model and kernel matrices, built D0 through D14 on the vLLM 0.26.0.dev0 stack (which resolves vllm#40898), and it remains gated behind a spike while its user-facing serving surface is finalized.
DeepSeek-V4 native MTP (DeepSeekV4MTPModel, ACTIVE — W1 self-spec wiring,
2026-07-30) has its nextn draft head wired to the same lossless spec-decode path.
Unlike V3's fused eh_proj, the V4 nextn layer keeps separate e_proj/h_proj,
enorm/hnorm, an MHC-aware mtp_block decoder layer, and its own hc_head
vocab collapse — ported 1:1 from nvidia/mtp.py as a tiny-config host draft
forward (DeepseekV4MtpDraftLogitsHost) reusing the DS4 attention/MoE/MHC
composition, verified by the SHARED greedy RejectionSampler so MTP-on greedy is
token-IDENTICAL to MTP-off (test_deepseek_v4_mtp 5/5, RED-first). The
real-model MTP-on==MTP-off + acceptance/speedup gate is WEIGHT-BLOCKED: both
shipped DeepSeek-V4-Flash GGUFs advertise nextn_predict_layers=1 but the
llama.cpp converter dropped every nextn tensor (DeepseekV4GgufHasMtp returns
false → clean fall-back to MTP-off). The DS4-native propose/verify decode loop +
engine spec-config registration are named residuals. BEAT-ds4 sweep (2026-07-30,
CLAIM-DSV4-BEAT-SWEEP, analysis-only): our raw autoregressive decode is at its
GB10 ceiling (~13.1-13.2 tok/s, Q8_0 front CLOSED after Bricks 3-14), and ds4's
16.5-17.2 is itself RAW decode (its MTP/DSpark spec is opt-in, NOT in that
number). The one path PAST 16.5 is MTP self-spec: 13.19 × (1+p)/1.05, break-even
acceptance p≈0.31, plausible p 0.55-0.85 → ~18-23 tok/s. Lossless (token-identical
to greedy). Gated on unblocking a nextn-carrying GGUF (R1) + the propose/verify
loop (R2) + device draft (R4). fp16 dequant cache refuted net-slower on GB10; MHC
a measured tie; routed-MoE we already win. See
.agents/specs/deepseek-beat-ds4-sweep-2026-07-30.md.
Cross-reference lever scan (2026-07-31, CLAIM-XREF-LEVER-SCAN, .agents/specs/cross-ref-lever-scan-2026-07-31.md). A 5-source parallel scan (vLLM/SGLang/llama.cpp/ds4/ours) for batch-1 raw-decode levers, no-spec, adversarially verified against our code. Corrects the record: ds4/DwarfStar has NO register-prefetch (grep-verified) — the Brick-14 theory below was a misattribution; ds4's 90% comes from full-warp-per-row + lane-strided 34-B Q8_0 blocks (1088-B coalesced burst). Our production Q8_0 ships SUB-warp (8/16-lane/row) for medium-K (LaunchQ8_0Subwarp), so the one non-refuted Path-B lever left is B1: port ds4's exact 32-lane/8-row warp8 for medium-K + measure (DGX A/B, not guaranteed). Verified cross-cutting gap: PDL (Programmatic Dependent Launch) — we have zero, both llama.cpp+SGLang use it on Blackwell. Path A (Laguna beat vLLM): the NVFP4 arm (#230) + the note that vLLM's 18.8 is MARLIN W4A16 (a lower bound; real default FLASHINFER_CUTLASS W4A4 is faster). Full ranked plan + citations in the spec.
DeepSeek-V4-Flash decode ds4-gap — Brick 14 (Q8_0 intra-row register-prefetch —
the ds4 raw mechanism) (2026-07-31, CLAIM-DSV4-DECODE-BRICK14, GB10 sm_121a,
branch ds4-q8-register-prefetch, NOT pushed). Built QuantDotGemmQ8_0PrefetchKernel <OutT,PF> (src/vt/cuda/cuda_quant_dot.cu, env VT_V4_Q8_PREFETCH default-OFF =1,
opt-in =2/=4, wired in MatmulQ8_0Cuda so both eager forward and captured
V4Graph::Step inherit it): same one-row-per-warp map, per-lane block loop
unroll-and-JAMmed by PF so PF weight+act block loads are hoisted into registers
before the dependent __dp4a chains (the ds4 register-resident-MLP mechanism from
ds4-q8-raw-mechanism-2026-07-30). BIT-IDENTICAL to the plain kernel (same block/dp4a
order + f16-scale fold; test_cuda_quant_dot Brick-14 A/B 1/1·1968; real-model
token-identity OFF==PF2==PF4 16/16). sudo-ncu (--replay-mode application +
memory watchdog) grid-256 (n=2048): registers 39→64→96, long-scoreboard 56.1→30.7→28.4
(toward ds4's 17.2), occupancy 72→55→30% (toward ds4's 60%) — the mechanism REPRODUCES —
but per-kernel GPU-active FLAT (46.06→46.27→46.19 µs) and clean decode tok/s FLAT
(13.145→13.140→13.133). MEASURED-NEGATIVE: lowering long-scoreboard ~2× at ds4's
register/occupancy operating point yields ZERO throughput ⇒ long-scoreboard is NOT the
causal gate on the 67%→90% BW gap (the raw-mechanism premise is refuted; the GEMV is
bound by DRAM byte-delivery, not latency-exposure). CORRECTS Brick 13's "intra-row is
the untried right axis" — that axis hits ds4's counter target and still yields nothing.
Kept DEFAULT-OFF like Bricks 4/8/11/12/13; ~13 tok/s stays the honest Q8_0-kernel ceiling
and the raw beat of ds4's 16.5 is not reachable via any Q8_0-GEMV-kernel-structure lever
(the residual is the MoE + glue front). See .agents/specs/deepseek-v4-last-mile.md
Brick 14 + docs/BENCHMARKS.md.
DeepSeek-V4-Flash decode ds4-gap — Step 0 + Lever 1 (2026-07-30,
CLAIM-DSV4-DECODE-LEVER1, GB10 sm_121a). Step 0 (kWarpsPerBlock 4→8 on the
dense Q8_0 GEMV, ds4's 8-rows/block): bit-identical, MEASURED NEUTRAL
(QuantDotGemmQ8_0Kernel 63.2 µs/launch unchanged; the 41 ms GEMV is
memory/latency-bound). Lever 1 (QuantizeQ8_0PreqKernel, VT_V4_Q8_PREQ_QUANT
default-ON): ports ds4's warp-per-32-block quant grid, dropping the dense Q8_0
activation-quant 4.51→0.99 ms/step (6.98→1.53 µs/launch, ds4-parity), in the
captured decode graph body. BIT-IDENTICAL (test_cuda_quant_dot 4/4·106081,
preq==legacy byte-for-byte), token-exact. Decode 11.44→11.92 tok/s (+4.3%,
non-overlapping); SHIPS default-on. Go/no-go for the 11.41→16.5 campaign =
PARTIAL GO: Gate A(ii) 9.1→<1 ms is NOT met — the QuantizeQ8K half (4.60 ms,
grouped MoE) is Brick-8-refuted for fusion and Brick-2-deduped, so the plan
mis-scoped it into Lever 1; realized +0.49 tok/s vs projected +2.3, so the ladder
to 16.5 must be re-based. See docs/BENCHMARKS.md.
DeepSeek-V4-Flash decode ds4-gap — Lever 3 (route → warp-topk) (2026-07-30,
CLAIM-DSV4-DECODE-LEVER3, GB10 sm_121a, base baff8a28). RouteWarpKernel
(VT_V4_ROUTE_WARP_TOPK default-ON) structure-ports ds4's router_select_warp_topk
(ds4_cuda.cu:10113): the E=256/topk=6 expert top-k, previously ONE thread per token
at T=1, now runs one WARP per token (8 experts/lane, per-lane argmax → __shfl_xor
reduce → mask/repeat). BIT-IDENTICAL — argmax under a strict total order
(value, tie-break lower index) is reduction-order-invariant, and the unbiased weight +
our exact renorm run the same float ops in j-order; kept OUR arithmetic, not ds4's.
Wired via a shared RouteDispatch into BOTH the eager forward AND the captured
V4Graph::Step (nsys shows RouteWarpKernel in the graph body, 43 launches/step).
GATE B (bit-exact, RED-first): test_cuda_deepseek_v4 20/20·67072 (new A/B case:
warp == single-thread BYTE-IDENTICAL ids AND weights; deliberate perturbation goes RED),
test_cuda_quant_dot 4/4·106081, test_deepseek_v4_gguf_load 15/15·931; real model
warp-ON == legacy-OFF golden 11111 16 455 6102 294 8760 344 … (TOKEN-EXACT).
Re-profile (nsys 30-step kwin): route RouteKernel 5.9723 → RouteWarpKernel
0.3779 ms/step (138.89 → 8.79 µs/launch, −5.59 ms, 15.8×, 4.7% → 0.5% GPU-active).
Honest realized-vs-projected: the study's route magnitude (6.5 ms → <0.5 ms) was
ACCURATE this time; −5.59 ms of ~81.6 ms/step = −6.9% GPU-active → GPU-active-bound
~+0.7–0.8 tok/s (11.92 → ~12.7). Wall-clock tok/s NOT cleanly resolvable this session
(no sudo drop_caches; churned-box page-cache variance swamps the ~0.7 tok/s delta).
SHIPS default-on (bit-exact, real GPU-work win). See docs/BENCHMARKS.md.
DeepSeek-V4-Flash decode ds4-gap — Lever 2 / Brick 11 (Q8_0 sub-warp GEMV tiling) —
MEASURED NEGATIVE, recorded (2026-07-30, CLAIM-DSV4-DECODE-LEVER2, GB10 sm_121a,
base d1ff7414). The plan's decisive speculative bet: the QuantDotGemmQ8_0Kernel
(41.1 ms/step = 52.7% GPU-active, the whole remaining ds4 gap) is 1-warp-per-output; at
short K a 32-lane warp was theorized to waste ~50% of lanes and starve the device, so a
templated sub-warp kernel (LANES∈{32,16,8}, nb-dispatch, ground: ds4
quarter/half_warp_sum_f32 ds4_cuda.cu:16609, moe sub-warp :17073) should raise
lane-utilization/occupancy. MEASURE-FIRST REFUTED THE PREMISE: the DS4-Flash decode
Q8_0 projections have nb∈{32,64,128,256} (K∈{1024,2048,4096,8192}) — the smallest,
nb=32, fills all 32 lanes (1 block/lane, ZERO idle lanes); there is no nb∈{16,48}
lane-waste regime, and that lone nb=32 projection (n=32768) issues 4096 blocks — NOT
device-starved. nsys 30-step kwin, VT_V4_Q8_SUBWARP A/B, worker down, flock: total
Q8_0 41.11 → 41.06 ms/step (flat, −0.1% = noise); the ONLY sub-warp-eligible
projection (nb=32, LANES=16, grid 4096→2048) 9.13 → 9.19 ms (marginally WORSE) —
well under the study's ~2 ms "worth it" bar. Interleaved warm wall-clock ~12.7 tok/s
both arms (off {12.67,12.80,12.77} / on {12.82,12.71,12.70}, overlapping); token-exact
golden 11111 16 455 6102 294 8760 344 … byte-identical off==on (the near-tie flips zero
tokens). HARD STOP-CONDITION MET → cut losses, no more variants. Gate: A/B unit case
test_cuda_quant_dot 5/5·106483 (sub-warp==plain, LANES=32 byte-identical big-K + NMSE≤5e-4
short-K, RED-first), test_cuda_deepseek_v4 20/20·67072, test_deepseek_v4_gguf_load
15/15·931. Merged default-OFF as a guarded recorded-negative (VT_V4_Q8_SUBWARP=1,
VT_V4_Q8_PROBE for the nb-attribution). Honest reachable ceiling: L1 (+0.49) +
L3 (route −5.59 ms GPU) bank ~12.7 tok/s; the Q8_0 GEMV is the irreducible
DRAM-bandwidth residual for THIS model's nb≥32 dims (occupancy/tiling has nothing to
fix), so ~13 tok/s is the honest ceiling and 16.5 is NOT reachable via this lever —
a valuable measured-negative like Bricks 4/8. See docs/BENCHMARKS.md.
DeepSeek-V4-Flash decode ds4-gap — ds4's Q8_0 kernel DIRECTLY PROFILED → the
"irreducible roofline" is REFUTED, lever RE-OPENED (2026-07-30,
CLAIM-DSV4-Q8-KERNEL-PROFILE, GB10 sm_121a, measurement-only). Nobody had ever
profiled ds4's own Q8_0 GEMV; done now (nsys cuda_gpu_kern_sum + GGUF byte-accounting
— ncu/nsys --gpu-metrics both ERR_NVGPUCTRPERM, no admin, so L2-hit is unmeasurable
and achieved-BW is the counter-immune proxy). Both engines read the identical 6.15 GiB
Q8_0 tower once/step (GGUF --inspect: 345 q8_0 tensors; ds4's fp16-dequant cache was
budget-exhausted this run → it reads the same raw 34-B blocks). Diff (pure decode/step,
(-n60 − -n4)/56): ours QuantDotGemmQ8_0Kernel = 40.93 ms over 646 launches = 161 GB/s
= 67% of the 240 roofline; ds4 = 30.38 ms over 259 launches = 217 GB/s = 90% of roofline.
Same bytes (reading >6.6 GB in 40.9 ms would exceed 240 GB/s — impossible — so our 646 are
row-splits summing to the same tower, not re-reads). Root cause = launch consolidation,
NOT the inner dot / bytes / alignment / lane-occupancy (Bricks 3/4/8/11 all correctly flat):
ds4 packs the tower into fewer, bigger launches — matmul_q8_0_pair_preq_warp8 (2 weights
- 1 shared activation/launch, 38.8 vs our 63.4 µs/matmul = 1.63×),
hc_expand_preq(MHC-expand fused into the matmul epilogue),grouped_q8_0_a— keeping DRAM saturated, where our 646 tiny warp-per-row launches under-subscribe it. Brick 11 built sub-warp lane-splitting (within-launch) and correctly found it flat; it never changed the launch COUNT/SIZE — the one axis that moves 67%→90%. VERDICT (A) PORTABLE LEVER: build Q8_0 projection-pairing (mirrormatmul_q8_0_pair_preq_warp8_kernel, bit-exact) + fold the MHC/residual epilogue in (near-tie); target 40.93→~30.4 ms (−10.5) → ~13.9 t/s (from 12.1 GPU-active), does not reach 16.5 alone (MoE+glue residual ~15 ms is a separate front) but re-opens the ladder. Premise correction: the direct Q8_0 gap is 10.5 ms/step, not the older whole-step-derived "19.8 ms". No code changed. See.agents/specs/ds4-q8-kernel-profile-2026-07-30.md+ docs/BENCHMARKS.md.
DeepSeek-V4-Flash decode ds4-gap — Brick 12 (IMPL): Q8_0 LAUNCH CONSOLIDATION built —
hits ds4's EXACT 259-launch structure BIT-EXACTLY, but the projected bandwidth win does
NOT materialize; the "launch count/size → ds4's 90%" hypothesis is REFUTED (2026-07-30,
branch ds4-q8-launch-consolidation off 1ddcc42c, GB10 sm_121a, VT_V4_Q8_PAIR default-ON,
rollback =0, NOT pushed). Three bit-exact consolidations wired into BOTH the resident-decode
Step and the captured V4Graph RunChain: projection-pairing (QuantDotGemmQ8_0PairKernel,
port of ds4 matmul_q8_0_pair_preq_warp8_kernel :4485) fuses the two A-projections sharing the
layer hidden — MLA wq_a+wkv and shared-expert gate+up — into one launch each (activation
quantized once); block-diagonal o-LoRA (QuantDotGemmQ8_0GroupDiagKernel, ds4
grouped_q8_0_a_preq_warp8_kernel :5509) collapses the ng=8 per-group wo_a GEMVs (344/step,
53% of the 646) into ONE launch; the MHC-expand epilogue fold was DEFERRED. MEASURED (nsys
cuda_gpu_kern_sum, diff (-n60 − -n4)/56): launches/step 646 → 259 (EXACTLY ds4's 259: 130
plain + 86 pair + 43 group-diag); the pair kernel is 37.3 µs/matmul-equiv, at/under ds4's 38.8.
Q8_0 GPU-active/step 40.96 → 39.86 ms (−1.10, −2.7%); achieved BW 161 → 166 GB/s (67.1% → 69.0%
of 240 — NOT the projected ~90%/30.4 ms); wall decode 12.24 → 12.27 tok/s (flat, within
page-noise). HONEST PARTIAL — the launch-count target is fully + bit-exactly met yet the
DRAM-saturation time win is absent: under VT_V4_DECODE_GRAPH=1 the 646 launches were ALREADY one
cudaGraphLaunch, so collapsing graph nodes 646→259 removes host overhead the graph had eliminated
and does not raise on-device BW. ds4's 67%→90% advantage is therefore NOT launch consolidation
(this build reproduces ds4's launch structure and stays at 69%) — it must lie in ds4's fp16-dequant
weight cache or L2 residency (counter-blocked). Token stream byte-identical ON==OFF (golden
11111 16 455 6102 294 8760 344 …, 16/16). GATES (RED-first): new test_cuda_quant_dot A/B cases
(pair==two-separate + group-diag==per-group-loop byte-identical) 7/7·106496,
test_cuda_deepseek_v4 20/20·67072, test_deepseek_v4_gguf_load 15/15·931. Bit-exact,
ds4-faithful, DEFAULT-ON (marginal-positive, not a regression). NOT pushed. See
.agents/specs/deepseek-v4-last-mile.md (Brick 12 IMPL) + docs/BENCHMARKS.md.
DeepSeek-V4-Flash decode ds4-gap — Brick 13 (IMPL): Q8_0 ILP lever (N output-rows per
warp) — MEASURED NEGATIVE via sudo-ncu; ~13 tok/s is the honest Q8_0 ceiling, Q8_0 front
CLOSED (2026-07-31, CLAIM-DSV4-DECODE-BRICK13, branch brick13-q8-ilp off 137c739f,
GB10 sm_121a, VT_V4_Q8_ILP default-OFF, opt-in =2/=4, NOT pushed). The one untried,
measurement-pointed lever from ds4-q8-ncu-2026-07-30 (our QuantDotGemmQ8_0Kernel is
memory-LATENCY-bound: long-scoreboard 54.4 @ 71.9% occupancy, L1-hit 96.7%): raise
memory-level parallelism per thread with N INDEPENDENT weight-load streams. Built
QuantDotGemmQ8_0MultiRowKernel<OutT,NROWS> — each warp computes NROWS consecutive output
columns of the same activation row (activation read once, __dp4a'd against NROWS independent
weight rows), BYTE-IDENTICAL to NROWS plain outputs — wired in MatmulQ8_0Cuda (eager +
captured V4Graph::Step). CAUSAL METRIC (sudo ncu, regex QuantDotGemmQ8_0, decode
steady, n-keyed): the target n=2048 kernel long-scoreboard 57.8 → 87.0 (ILP2) → 51.7
(ILP4) (did NOT drop materially), achieved occupancy 71.4% → 42.5% → 21.3% (COLLAPSED),
per-launch GPU-active 46.7 → 44.9 → 51.0 µs (flat-to-WORSE); n=4096 LS 56.4 → 122.2 →
91.3. Clean wall-clock (sudo drop_caches, 3 reps, --max-tokens 64): OFF 13.19
median, ILP2 13.12 (−0.5%), ILP4 13.05 (−1.0%). ROOT (measured): the load latency was
already hidden by INTER-warp parallelism (71% occupancy), not starved for intra-thread ILP;
folding N rows/warp divides the warp count by N and adds registers ⇒ occupancy collapses and
removes more latency-hiding than the extra streams add. Token stream byte-identical
OFF==ILP2==ILP4 (golden 11111 16 455 6102 294 8760 344 …, 16/16). GATE (RED-first,
byte-identical): test_cuda_quant_dot 8/8·108464 (new Brick-13 A/B 1/1·1968, ILP2==ILP4==plain
over nb∈{16,48,224}, n∈{1,7,16,17} incl. tail n∤N). HARD STOP-CONDITION met on BOTH N-values →
RECORDED-NEGATIVE, kept DEFAULT-OFF (bit-exact-safe, like Bricks 4/8/11/12). This was the LAST
measurement-pointed Q8_0-GEMV lever: ~13 tok/s is the honest Q8_0-kernel ceiling on GB10 and
the campaign's Q8_0 front is CLOSED; ds4's 67%→90% edge stays attributable only to the
unobservable fp16-dequant weight cache / L2 residency (ERR_NVGPUCTRPERM-blocked). NOT pushed.
See .agents/specs/deepseek-v4-last-mile.md (Brick 13) + docs/BENCHMARKS.md.
ngram (method ngram, draft-FREE) proposes the next tokens by matching the
sequence's own suffix n-gram, so it needs no draft model and works on any model;
on the 27B it is token-exact vs vLLM's own --speculative-config ngram on
repetitive workloads (5/5 prompts, every draft accepted).
Generic separate draft-model (method draft_model, SPEC-DRAFT-MODEL
ACTIVE) is spiked with its first CPU brick landed, not yet wired into any
production path. The classic model-agnostic path runs a full smaller standalone
draft LM K autoregressive greedy steps to propose K tokens, which the target
verifies in one forward (longest-accepted-prefix), reusing the landed rejection
sampler unchanged (only the proposer is net-new). W1 landed the greedy propose
(DraftModelProposeGreedy over a next-token-logits oracle) + the draft_model
--speculative-config accept, unit-gated RED-first (test_draft_model_proposer
6 cases / 41 assertions: the accepted tokens equal the target's own greedy run
for every draft/target agreement pattern, a target-matching draft is fully
accepted, and full acceptance depends on the autoregressive feed-back). NOT yet
wired: the real GPU draft-model forward behind the oracle (paged KV, CUDA-graph)
- the e2e greedy our-ON==vLLM-ON token-exact gate + the throughput speed gate are
DGX-offline residuals (W3). Medusa (method
medusa,SPEC-MEDUSASPIKE) is spiked only; its N-head single-pass proposer is deferred to W2 (needs the target's Medusa heads). See docs/BENCHMARKS.md.
The correctness form and the full D0-D14 measured chronology live in docs/BENCHMARKS.md, docs/SPECULATIVE-DECODING.md and .agents/specs/dflash-spec-decode.md.
LoRA (W1 CPU runtime brick landed — see the capability table; not yet usable
end-to-end), multi-GPU, and the full tool-calling template surface. Scale-out /
distributed execution is scoped but unbuilt (spike, 2026-07-28): the engine is
single-GPU today (verified — no NCCL / tensor-parallel / process-group code in
src/). A single-dimension design is on record covering all three legs —
multi-GPU tensor+pipeline parallel, multiple DGX Sparks over the ConnectX-7
200GbE RoCE/RDMA cable (the path that lets DeepSeek-V4-Flash fp8 ~167 GiB run
across 2×119 GiB Sparks), and MLX multi-node over Thunderbolt — all expressed
ONCE against one vt:: collective / process-group abstraction with
backend-specific transports (NCCL / RDMA / MLX-ring), mirroring vLLM's
device_communicators. world_size==1 stays byte-identical. W1 landed
(2026-07-28) the collective ABSTRACTION leg: vt::Communicator
(include/vt/communicator.h + src/vt/communicator.cpp) with
AllReduce/AllGather/Send/Recv, proven by a CPU in-process multi-rank gate
(tests/vt/test_communicator.cpp, a real cross-rank sum, no GPU) and a
byte-identical world_size==1 no-op. W2 landed (2026-07-28) same-host
multi-GPU tensor parallel (CPU-gated): the multi-device backend registry
(per-Device{type,index}, byte-neutral device 0), collective OpId routing, the
NCCL transport (mirrors pynccl.py, built only under -DVLLM_CPP_NCCL=ON), and
TensorParallel wired into the Qwen3-dense forward — the sharded-matmul +
all-reduce is proven equal to the unsharded tp=1 forward (test_tp_forward,
60/60, RED-verified, no GPU); tp_size==1 byte-identical. The real TP-2
multi-GPU RUN + the NCCL build-verify remain HW-blocked (no ≥2-GPU box), as do
pipeline parallel, multi-Spark and MLX. Full scope + seam map + the 2× RTX-6000-Ada
recipe:
.agents/specs/scale-out-distributed.md.
Every parallelism MODE vLLM has (tensor / pipeline / data / expert / sequence /
context parallel) is now enumerated and grounded in upstream source, mapped onto
that one abstraction and priority-ranked, in
.agents/specs/parallelism-modes.md
(2026-07-28) — with the honest note that vLLM's "sequence parallel" is a
tensor-parallel compilation pass, not a separate parallel axis.
Multimodal
(image/video/audio) is correctness-complete; its OpenAI-server wiring has landed
all three CPU bricks (content-part parse + processor routing; the engine
mm-request plumbing — add_request(MultiModalInputs) on both engines, default-inert
to the text path; and the W3 MultiModalChatFn seam BODY that renders an image
chat request into the placeholder-EXPANDED engine input, wired into the production
server). It is not yet servable end-to-end: the closing GPU gate MM-SERVE-E2E is
architecturally blocked on the engine model runner having no multimodal forward
(the vision tower + merge lives in the standalone M2c driver, not in
ModelRegistry::Forward) — folding that forward into the registered engine path is
the named residual (.agents/specs/mm-serving.md).
Open, not root-caused (observed 2026-07-28): the C-ABI custom logits
processor case (tests/capi/test_capi.cpp:410, ABI v8) SIGSEGVs in a CUDA build
on the GB10 box, while the same suite is 232/232 green on a CPU build. Confirmed
present on pristine main ee3d5960 with no local changes, so it is not a
regression from any in-flight work, but it does mean the "supported" claim for
custom logits processors in the capability table above is verified on CPU only.
One build directory, not yet bisected.
Gemma 3 (Gemma3ForCausalLM, google/gemma-3-1b-it) is the first landed
member (correctness-complete, speed-pending): STRICT token-exact 48/48 greedy
vs the vLLM 0.25.0 oracle. It reuses infrastructure already in the tree
(gemma-RMSNorm (1+w), the GLM sandwich norms, the shared dense-attention path,
sliding-window attention, tied embeddings) plus one genuinely-new compute
kernel, GeGLU (gelu_pytorch_tanh), a bf16 embedding-scale multiply, dual
per-layer RoPE theta, and query_pre_attn_scalar scaling. Gemma 2
(Gemma2ForCausalLM, gemma-2-2b-it) and Gemma 1 (GemmaForCausalLM,
gemma-2b) have since landed too, both correctness-complete and speed-pending
at 48/48 greedy vs the same oracle
(spike): Gemma 2 proves the attention + final
logit soft-cap primitives, Gemma 1 the original two-fused-norms block.
Gemma 4, the newest registered variant, now has its text path
correctness-complete and gated (G1b, 2026-07-28): the
Gemma4ForConditionalGeneration language_model stack of unsloth/gemma-4-E4B-it
loads through our engine and greedily emits the exact 32 golden token ids —
STRICT 32/32 token-exact vs the vLLM 0.25.0 golden (gate
tests/parity/test_gemma4_paged_engine.cpp, dgx CUDA). It brings up the large new
primitive stack (per-layer embeddings, YOCO KV-sharing, plain RMSNorm, proportional
partial-RoPE, heterogeneous 256/512 head dims, GeGLU, a per-layer scalar; the
Gemma-4 MoE / k_eq_v / double-wide MLP are off for the E4B checkpoint and stay the
larger-variant follow-on). The G1-named blocker is resolved: the runner now
allocates a per-layer KV head dimension (KVCacheConfig::per_layer_attn_specs),
so Gemma-4's per-layer 256/512 heads each get a correctly-strided cache — and it is
byte-neutral for every uniform-KV model (the field is empty ⇒ the previous single
uniform allocation; the full CPU runner/KV suite plus the OLMo-2 SACRED GPU gate
16/16 are unchanged). Three additive loader gaps were fixed en route to the
first-ever Gemma-4 forward: nested per-layer rope_parameters (config loader), the
Gemma Replace(" "->"U+2581") metaspace normalizer (tokenizer), and reading
Gemma-4's per-arch scalars from raw["text_config"] rather than the full config
(the real 256/512 head-dim source). Speed vs vLLM is pending. Its
multimodal path (image + video + audio, the only
audio-capable model in the pin) has been assessed
(spec) and is now oracle-gateable —
run-verified (W0, 2026-07-28): the pinned 0.25.0 oracle (transformers 5.13.1)
loads, runs, and greedily generates the ungated unsloth/gemma-4-E4B-it
(Gemma4ForConditionalGeneration, 15.99 GB) on GB10, K=5 deterministic, and the
32-token greedy golden is captured
(tests/parity/goldens/gemma4_e4b_text/). The earlier "oracle-blocked at gate
time" concern (transformers lacking gemma4) is refuted — the module is present
and the model runs. The IMAGE modality oracle is now captured too (G2,
2026-07-28): the same E4B model on the pinned 0.25.0 oracle greedily describes a
fixed image (K=5 deterministic ⇒ STRICT gate form, 18 tokens → a coherent
gradient-image caption, 256 soft tokens), committed with four staged vision-tower
reference tensors for per-stage unit-gating
(tests/parity/goldens/gemma4_e4b_image/). Grounding the tower corrected an
earlier assumption: the Gemma-4 vision tower is a custom NaFlex SigLIP2 with
multidimensional vision-RoPE, q/k/v RMSNorm, Gemma-2 sandwich norms, a learned 2-D
position embedding, and a √hidden average-pool-by-position pooler — it does not
drop-in reuse the Qwen3-VL ViT (the block GEMMs/attention are reusable; the
patch-embed, RoPE, norms, pooler and the Gemma-4 NaFlex image processor are new).
The C++ NaFlex SigLIP2 vision tower is now IMPLEMENTED (additive
gemma4_vision.{h,cpp}) and PASSES its four per-stage gates vs the
transformers-eager references on the dgx CUDA build (patch-embed rel-L2 2.15e-3,
encoder 3.14e-2, pooled 1.36e-2, projected 1.85e-2; 220/220; compute-sanitizer 0) —
the tower-in-isolation milestone (mirrors Qwen3-VL M2a before its M2c e2e); grounding
it also revealed the vision Gemma4ClippableLinear layers carry FINITE trained QAT
activation clamps (not the no-ops the port-map assumed), now implemented. So Gemma-4
multimodal image is blocked only on the remaining engine wiring: the C++ Gemma-4
NaFlex image processor + the mm merge-plumbing (register SupportsMultiModal, the
masked-scatter merge of the 256 projected soft tokens at the image placeholder rows,
the tower→merge→decode fork) for the image→text e2e. The USM-Conformer AUDIO
tower is now IMPLEMENTED too (G3, 2026-07-28): the additive
gemma4_audio.{h,cpp} forward (Gemma4AudioModel + the audio projector) PASSES all
seven per-stage gates f32-exact vs the transformers-eager reference (host f32,
1256/1256: subsample rel-L2 5.4e-7, position-embeddings 8.9e-8, block0 4.2e-7,
block_mid 3.8e-7, block_last 4.4e-6, output_proj 5.9e-6, projected 6.3e-6), ported
1:1 from modeling_gemma4.py — a 2×Conv2d subsample, a relative positional
encoding, and 12 Conformer layers (half-step feed-forwards, a chunked-local
attention with Transformer-XL relative-position bias + a tanh soft-cap +
per-dim-scale softplus, and a GLU + depthwise-causal-conv light-conv module), then
output_proj and the audio embedder; the FINITE QAT clamps and the exact sliding
window (dist ∈ [0,12) — a RED-first-caught off-by-one) are grounded in the source.
So Gemma-4 now has all three modalities tower-proven (text STRICT 32/32, vision and
audio per-stage), and the remaining audio work is the same engine wiring as image —
the Gemma-4 audio feature extractor (mel frontend) + the mm merge-plumbing — plus
a device-resident bf16 forward for speed. The Gemma-4 MoE / k_eq_v / double-MLP
backbone stays the larger-variant follow-on. Audio as a standalone modality is also
proven end-to-end on the smallest oracle-runnable audio models (Whisper encoder,
then Voxtral-Mini-3B on the already-landed Mistral backbone).
The OLMo-2 family (Olmo2ForCausalLM / Olmo3ForCausalLM) is the first
landed OLMo member (correctness-complete, speed-pending): token-exact 16/16
greedy vs the vLLM 0.25.0 oracle on allenai/OLMo-2-0425-1B (STRICT 13/16 +
near-tie-band 3/16, max gap 0.094 nats, 0 forward-divergent). It is the cleanest
dense bring-up yet, needing no new compute kernel: its two distinctive traits
both reuse existing infrastructure. The reordered post-norm placement
(norm_after) is a subset of the GLM/Gemma sandwich norms (the same standalone
output-norm plus a plain residual add, without the pre-norms), and its QK-norm
is a full-width RMSNorm reusing the existing norm op at a new shape. Its GPT-NeoX
ByteLevel tokenizer (which prepends no BOS) makes it a real tokenizer-inclusive
gate. Olmo3ForCausalLM rides the same class (the 0.25.0 oracle constructs it);
the Olmo-3 interleaved sliding-window path has since landed and runs, but is
oracle-blocked for a gate (see the capability table above).
Larger DeepSeek / GLM / MiniMax / Gemma-4 variants are recorded as
hardware-blocked (they do not fit 119 GiB of unified memory on this box) or
spiked-only, per the model matrix.
DeepSeek-V4-Flash is scoped with W1/W2 CPU scaffolding landed (not yet
runnable) in deepseek-v4-flash.md: a
~167B/256-expert NEW architecture (DeepSeek Sparse Attention MLA + Manifold
Hyper-Connections + NVFP4/MegaMoE + sqrtsoftplus/hash MoE). As of 2026-07-28 the
additive registry stub + config parse + checkpoint loader name-map are landed and
VERIFIED against the real nvidia/DeepSeek-V4-Flash-NVFP4 safetensors header
(HTTP-range, no download). The forward host-composition (W7) + the W7-device CUDA
kernels are now landed: the four NEW V4 op families (MHC / DSA indexer+seams /
compressor+fp8_ds_mla / sqrtsoftplus-hash MoE) are ported to CUDA, registered
through the OpProvider seam, dispatched by a real DeepseekV4Model::ForwardDevice,
and RUNTIME-VERIFIED on the DGX GB10 at small shape vs the host-ref oracle
(test_cuda_deepseek_v4 11/11·153, compute-sanitizer memcheck 0 errors, RED-first
proven; 2026-07-29, CLAIM-DEEPSEEK-V4-W7-DEVICE). The 512-wide MLA attn + expert
grouped-GEMM REUSE the existing NVFP4/FP8 kernels. The W2b GGUF keep-quant tower
materialization is now landed (2026-07-29, CLAIM-DEEPSEEK-V4-W2B):
LoadDeepseekV4FromGguf (+ DeepseekV4ParamsFromGguf) wires the landed blk.N.*
name-map + keep-quant blocks into the DeepseekV4 weight towers (the MLA linears,
router gate, 256 routed + shared experts, and lm_head KEEP their ~2-3-bit blocks
COMPRESSED via OwnGgufQuantBlocks, the ~91 GiB-vs-~316 GiB OOM enabler; norms /
MHC / DSA / embed / tid2eid dequant), accounts for every tensor, lifts the registry
GGUF reject, and dequants the tiny CPU composition tower so a loaded model forwards
(test_deepseek_v4_gguf_load 5/5·149 — accounting 126/126, keep-quant residency,
load→forward, RED-first) — DERIVED + BUILD-VERIFIED at tiny synthetic shape. The
keep-quant expert/MLA GEMMs now run ON THE GPU (2026-07-29,
CLAIM-CUDA-KEEPQUANT-GEMM): the FIRST CUDA keep-quant GGUF k-quant GEMM
KERNEL-QUANT-CIQ-GEMM-CUDA (cuda_quant_dot.cu, MMVQ-style dequant-in-kernel over
the compressed Q8_K-family blocks) registers a kCUDA kMatmulBTQuant provider,
which flips GgufQuantComputeAvailable TRUE on kCUDA — so on a CUDA runner
vt::MatmulBT dispatches these GEMMs to the GPU instead of the 20 ARM cores
(the biggest DeepSeek-V4 speed lever), weights staying COMPRESSED in the unified
pool. GB10-gated test_cuda_quant_dot 2/2·92401 vs the CPU oracle + f64 dequant,
memcheck 0, RED-first proven. The experts-on-GPU tok/s is the follow-on benchmark.
Still NOT runnable end-to-end on the REAL checkpoint: the 91 GB UD-IQ2_XXS single-Spark
run (keep-quant load + ForwardDevice greedy gen + self-consistency/coherence gate +
benchmark) is the W8-final operational residual (download + DGX), and the NVFP4/fp8
paged-engine strict gate (W8) is multi-Spark-blocked (156.7 GiB does not fit ONE GB10).
HW-fit correction: that NVFP4 checkpoint is 156.7 GiB, NOT the ~83 GiB the
scoping spike estimated (only the 256 routed experts are W4; the MLA and shared
linears are FP8 plus NVFP4 double-scale overhead), so it does not fit ONE
GB10's 119 GiB unified pool. Single-Spark IS viable via a ~2-bit GGUF (user
correction 2026-07-28): unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_XXS = 90.9 GB
(3 shards, ungated) FITS one GB10 with ~28 GiB headroom (also UD-IQ1_S/IQ1_M/IQ2_M/
Q2_K_XL). So two vehicles: single-Spark ~2-bit GGUF, or 2x-Spark NVFP4/fp8 over the
interconnect. The GGUF vehicle's correctness reference is llama.cpp-on-box (the
pinned vLLM cannot load V4-from-GGUF) and needs a V4-GGUF loader + IQ i-quant dequant
(IQ1_S/IQ2_XXS — not in our C4 K-quant set). The NVFP4/fp8 vehicle stays multi-node.
A source-level spike (2026-07-28) confirmed a same-quant IQ2_XXS GGUF benchmark of
our engine vs vLLM is NOT viable today, blocked on both sides: vLLM 0.26 moved GGUF
out-of-tree to the uninstalled vllm-gguf-plugin (which does dequant IQ2_XXS) but
DeepseekV4ForCausalLM has no packed_modules_mapping/GGUF wiring; and our engine
hard-rejects GGUF for DeepSeek-V4/V2 and lacks IQ2_XXS/Q2_K dequant — so even the
UD-Q2_K_XL k-quant fallback does not rescue it. Apples-to-apples for DeepSeek-V4 is
the NVFP4 vehicle; a true same-GGUF cross-engine number is only available on a
Qwen3/dense k-quant both engines already load. W8 KEEP-QUANT MEMORY ENABLER + NAME-MAP
landed (2026-07-29, CLAIM-DEEPSEEK-V4-W8): the ~2-3-bit routed-expert encodings now
carry a keep-quant vec_dot — IQ2_XXS, IQ3_XXS and Q2_K (1:1 ports of ggml
vec_dot_{iq2_xxs,iq3_xxs,q2_K}_q8_K_generic), so HasQuantDotKernel is TRUE and the
loader keeps their blocks COMPRESSED (the memory enabler: ~91 GiB vs the ~316 GiB bf16
OOM). HONEST CORRECTION: the real UD-IQ2_XXS experts are IQ2_XXS (gate/up) +
IQ3_XXS (down) — NOT Q2_K (that is the UD-Q2_K_XL sibling); all three landed.
HONEST finding: there is NO CUDA keep-quant vec_dot for ANY k-quant (kMatmulBTQuant
is CPU-only) — on GB10 it runs on the 20 ARM cores against the unified pool. Gate
test_ops_quant_dot 19 cases / 130444 assertions, RED-first proven. The full 1328-tensor
blk.N.*→V4 name map now has EXACT coverage (0 unmapped, 0 leftover) against the real
manifest (shard GGUF headers read via HTTP-range — no 91 GB download; scripts/check-dsv4-gguf-namemap.py
rc=0). The V4 registry GGUF reject is now a PRECISE message (keep-quant + name-map landed;
W2b tower materialization pending). The RUN did NOT execute and was NOT faked: it is
blocked on the unimplemented W2b (materialize the keep-quant blocks into the DeepseekV4
towers via the name map) — a genuine code brick, then the 91 GB download + GB10 greedy gen.
W8-final entrypoint wiring landed + gated (2026-07-29, CLAIM-DEEPSEEK-V4-W8, base
376e186b): with W2b landed, the top-level GGUF dispatch now recognizes a deepseek4 file —
DeepseekV4HfConfigFromGguf maps general.architecture=deepseek4→the registered
DeepseekV4ForCausalLM (republishing geometry into config.raw for the parse hook) and
LoadedEngine::FromModelDir routes it via HfConfigFromGgufDispatch (was qwen-only). Gate
test_deepseek_v4_gguf_load 6/6·168 (CPU -Werror-clean), qwen path byte-neutral. But the real
single-Spark RUN is still BLOCKED — now on a CODE residual, not download/box, and was NOT attempted
(it would OOM-reboot the box): the DeepSeek-V4 forward composes off the FULLY-DEQUANTIZED f32
weights.host tower, and LoadDeepseekV4FromGguf builds that host tower unconditionally (every
routed expert HostVec→f32) ≈ ~24 GiB/layer × 43 ≈ ~1.0 TiB f32, past the 119 GiB pool by
~layer 5; the keep-quant weights.gguf tower (~91 GiB) is built but never read by the forward. A
real run needs the forward rewired onto the CIQ kMatmulBTQuant keep-quant blocks + the host dequant
gated off (was named residual W2c). No tokens generated (not faked); DGX left as found.
W2c landed (2026-07-29, CLAIM-DEEPSEEK-V4-W2C): the OOM-infeasibility is FIXED.
LoadDeepseekV4FromGguf no longer f32-expands the big MLA/MoE/lm_head weights (only the
small norms/embed/MHC/DSA/hash tensors dequant); a new DeepseekV4ForwardGguf runs the
SAME composition with the 512-wide MLA linears + the 256 routed/shared expert GEMMs +
lm_head consuming the COMPRESSED weights.gguf blocks in place via vt::MatmulBT→the CPU
kMatmulBTQuant CIQ GEMM (no per-layer f32 expansion), and DeepseekV4Model::Forward
gates on has_gguf_weights (the safetensors/NVFP4 + tiny-synthetic host path stays
byte-identical). Gate test_deepseek_v4_gguf_load 7/7·185 (CPU Release -Werror-clean):
the keep-quant forward RUNS finite+deterministic, keep-quant(Q8_0)==dequant(bf16) RelL2
0.0116 (< 0.05 near-tie), RED-first (a no-sink miswire diverges, and a load that rebuilds
the f32 tower fails a load-time VT_CHECK + the host<gguf-bytes assertion). Memory-bound
asserted: at tiny shape host 23,980 B vs keep-quant 141,676 B; projected to the real config
the 256 routed experts alone are ~1032 GiB f32 (OOM-reboots the 119 GiB pool) vs the
keep-quant UD-IQ2_XXS ~91 GiB + small host < 3 GiB = memory-FEASIBLE on ONE GB10. The
real 91 GB run (download + GB10 generate + benchmark) stays the operational W8-run, now
memory-feasible. SACRED-inert: only the V4 forward/loader + its test changed; the W3-W6
primitive tests (the correctness oracle) are untouched and still pass.
W8-run ATTEMPTED 2026-07-29 — HARDWARE-BLOCKED (not run, nothing faked): with the code
complete + memory-feasible @ 2936ff70, the operational run could not start because the
DGX GB10 (dgx.casa) is OFFLINE (ping 5/5 loss, ARP FAILED, SSH no-route; gateway/Thor/
mac-mini up ⇒ DGX-specific, down or crashed). No out-of-band wake path from the dev box.
The DeepSeek-V4 row stays SPIKE (no run, no benchmark). Resume: bring the DGX back online and
re-run the W8-run recipe (free box → download ~91 GB UD-IQ2_XXS → keep-quant load [assert
resident ≈ 91-94 GiB, the W2c VT_CHECK must not fire] → greedy generate → self-consistency +
coherence gate → TPOT/throughput/peak-resident benchmark). No code change needed.
W8-run apples-to-apples oracle SELECTED + ds4-file LOADABILITY PROVEN (2026-07-29, base
edf68c91, NOT pushed): the cross-engine oracle is now antirez/ds4 (DwarfStar,
make cuda-spark on GB10), and the shared apples-to-apples vehicle is ds4's q2-imatrix
GGUF (antirez/deepseek-v4-gguf, a single 80.7 GB file — routed experts IQ2_XXS gate/up +
Q2_K down, Q8_0 attn/shared/out, F16 embed — all keep-quant types we already support).
Verified via HF HTTP-range (NO 80 GB download) that our blk.N.* name-map covers ds4's file
EXACTLY (1328/1328, 0 unmapped, 0 leftover — byte-identical tensor-name set to unsloth's).
Found + FIXED three real-file loader gaps that would throw or zero the model before a token:
(1) compress_ratios length 44 (43 layers + a trailing MTP entry) vs our strict
==block_count — relaxed to >= + prefix-truncate; (2) the clamp limit was read only from
the scalar swiglu_clamp, but ds4 ships the per-layer array swiglu_clamp_exp (all 10.0) and
a 0 limit ZEROES every expert (ClampedSwiGLU min(gate,0)·clamp(up,0,0)) — array fallback added;
(3) the hash ffn_gate_tid2eid is ggml I32 (type 26), unhandled in DequantGgufRowToF32 — I32
case added. Gate test_deepseek_v4_gguf_load 8/8·288 (CPU Release -Werror-clean) with a
new ds4-flavor case reproducing all three quirks (load + keep-quant forward finite +
deterministic); RED-first proven (disabling the I32 case throws "unsupported ggml type 26").
The GB10 run itself did NOT execute (nothing faked): it is blocked on the GPU flock held by a
concurrent agent, two long builds (ds4 cuda-spark + our aarch64 CUDA), the 80 GB download,
and a greedy driver (the engine's incremental paged decode is incompatible with V4's stateless
keep-quant recompute — a manual greedy loop over DeepseekV4ForwardGguf is required; note
Tokenizer::FromGguf also rejects ds4's pre='joyai-llm', so the token cross-check must inject
ds4's token ids). Row stays SPIKE.
W8-run EXECUTED on the DGX GB10 (2026-07-29, base 858b0b15, NOT pushed) — our engine LOADS
the real 158 B model but the forward HARD-FAILS at the real geometry; ds4 reference captured:
the 80.7 GB file was downloaded + integrity-verified (86,720,111,488 B == HTTP Content-Length),
ds4 was built (make cuda-spark), our CPU engine + a new greedy driver (examples/deepseek_v4_gen)
were built, and both were run on the SAME file. ds4 (GB10 GPU) loads 80.76 GiB in ~20–29 s and
greedily generates coherent on-topic text ("We need to answer: "The capital of France is". This is
a straightforward") at prefill 358 tok/s, decode 16.5 tok/s (ctx=1024, ds4-bench). Our engine
LOADS the keep-quant tower on the real model (43 layers, 256 experts, vocab 129280, all 1328 tensors
accounted, has_gguf_weights=1, ~78 s) at PEAK RESIDENT 116.2 GiB — it FITS under the 119 GiB
pool but is ~32 GiB ABOVE the projected ~84 GiB, because OwnGgufQuantBlocks is called WITHOUT the
mmap_src arg so it COPIES ~81 GiB of blocks into owned vectors while the mmap pages stay resident
(fix: mmap-VIEW the blocks like qwen3_5_gguf_weights.cpp). The forward then HARD-FAILS with
deepseek-v4 keep-quant GEMM: weight shape mismatch: want [N=512,K=4096] got [1024,4096]
(deepseek_v4.cpp:234) at layer 2 (the first DSA compressor layer): the W3–W7 forward, gated only
at tiny synthetic shape, conflates head_dim=512 with the DSA compressor's distinct 1024-dim output
(the real attn_compressor_gate/attn_compressor_kv are [1024,4096]). Layers 0–1 pass; layer 2
crashes. NO coherent tokens generated → row STAYS SPIKE. Remaining brick: rework the MLA/DSA
forward to the real DeepSeek-V4 geometry (compressor 1024-dim; audit indexer + grouped output-LoRA at
o_groups=8/nh=64/q_lora_rank=1024) against a reference; plus the mmap-view load fix. The
download + keep-quant load + tokenizer + greedy driver + ds4 oracle are DONE; the forward geometry is
the one blocker before a real single-Spark generation.
W8-run.9 FIXED — COHERENT single-Spark generation; DeepSeek-V4 ADVANCES SPIKE → ACTIVE (2026-07-29,
base eeef1695, NOT pushed). The isolated L00 experiment (aligned dump: ours' post-input-norm attn
input vs ds4's attn_norm) forked H1/H2: attention INPUT bit-exact (rel-L2 0.0000), KV correct
(0.018), but the q operand rel-L2 0.9646 → the MLA q projection. Root cause: our forward OMITTED the
per-head query RMS-norm (ds4 head_rms_norm_inplace after wq_b; the KV correctly gets only its
attn_kv_a_norm). Fix DeepseekV4QHeadRmsNormInplace (+ spec-anchored RED-first doctest) → L00 q
0.9646→0.0013, attn_out 0.5956→0.0175, and the FULL 43-layer curve COLLAPSED to the keep-quant
floor (L33 0.53→0.003, L34 6.02→0.003, MAX 0.0334). Plus the second confirmed fix, the
inverse-RoPE on the attention output (ds4 rope(heads,inverse=true), identity at pos 0, load-bearing
for pos>0). Greedy generation is now "The capital of France is Paris.<｜end▁of▁sentence｜>" (correct +
EOS), self-consistent; test_deepseek_v4_gguf_load 10/10·403. CPU-tier decode ~3.3 s/tok / peak 85.8 GiB
(ds4 GPU: 16.5 tok/s). Gated by coherence + self-consistency + the ds4-oracle per-layer diff (no vLLM V4
GGUF plugin ⇒ not vLLM-token-exact). Named residuals: GPU-expert dispatch (CPU-tier), DSA-sparse ctx>512
(dense-fallback, exact for short gen), paged-engine integration. Row ACTIVE.
W8-run.10 — GPU-expert wiring (--gpu) + apples-to-apples vs ds4 (2026-07-29, base 20d8ccfa). Experts
now run on the GB10: a CUDA queue routes every keep-quant expert/MLA/lm_head GEMM to the kMatmulBTQuant
kCUDA provider (#195), reading unified-memory weight blocks in place (memory-safe; no device copy). GPU
genuinely used (nvidia-smi 38–50% util, compute-app deepseek-v4-gen during the run); GPU path
byte-identical to CPU; CPU default unchanged (test_deepseek_v4_gguf_load 10/10·403). Real table (same 80.7
GB ds4 file): ds4 GPU prefill 325.9 / decode 16.3–16.6; ours --gpu prefill 6.26 / decode 0.68 (1.3–1.4×
over our CPU path). The residual gap is NOT expert placement (fixed) but our stateless full-recompute
driver (no KV cache) + route-(a) host-orchestration/per-GEMM sync. Real speed path = ForwardDevice
(#183): resident on-device activations + KV cache + no per-GEMM sync — the only route to approach ds4's 16.5
tok/s. Route (a) not dressed as a win. Row stays ACTIVE; see docs/BENCHMARKS.md W8-run.10.
ForwardDevice campaign Stage 1 — MLA latent KV cache + incremental decode (2026-07-29, base fd9e191c,
NOT pushed). The single biggest decode lever, equivalence-gated + benchmarked. For the real dense-MLA run
the only cross-token state is the per-layer deck latent [head_dim] (num_key_value_heads=1; MHC manifold
per-token), so DeepseekV4KvCache (mirror of ds4 raw_kv) makes each decode step process ONE new token
against cached KV instead of re-running the whole forward over the growing context. DeepseekV4ForwardGgufCached
- driver
--kv-cache(rollback-able; default +--gpuunchanged). Token-IDENTICAL to full-recompute ("…Paris.", ids11111 16 455 6102 294 8760 344 11111); decode 0.68 → 4.79–5.35 tok/s (7.9×), ~1/3 of ds4's 16.5; GPU used (nvidia-smi 29–46%, compute-appdeepseek-v4-gen); peak 86.33 GiB. Gatetest_deepseek_v4_gguf_load11/11·430 (prefill bit-identical, incremental token-identical, RED-first). Stage 0 spec:.agents/specs/deepseek-v4-forward-device.md. Residual to ds4 = host-orchestration + per-GEMM sync (Stages 2–3: device-resident activations + graphs). Row staysACTIVE; see docs/BENCHMARKS.md. ForwardDevice campaign Stage 2 — decode profiler + deferred-sync drain batching: HONEST NEGATIVE RESULT (2026-07-29, base17ec33ea, NOT pushed). The scoped lever ("drop the per-GEMM host sync") is MEASURED to NOT move decode. AddedVT_V4_PROF, which split the 0.18s decode step into sync 0.073s / gemm-dispatch 0.074s / host-glue 0.034s. Batched the stream drains (21 routed+shared expert GEMMs/layer → 2; grouped o-LoRA → 2;defer_sync+DrainDevice), byte-identical (test_deepseek_v4_gguf_load11/11·430). Real DGX: decode 0.20 s/tok (5.11 tok/s), token-identical ("…Paris."), util ~49%, NO speedup (sync 0.073→0.071s). CORRECTION: the "sync" bucket is the host waiting for GPU compute of ~1100 tiny T=1 GEMMs, NOT sync-call overhead; both buckets scale with GEMM count, not sync-call count. Real lever (named, corroborated by the Qwen3-Coder W7 decode-graph row): a grouped MoE GEMM (18 tiny expert matvecs → ~3 kernels) + a decode CUDA graph to kill launch tax. Recommend re-scoping Stage 2 to the grouped GEMM. Rollback-able; NO speedup claimed. Row staysACTIVE; see docs/BENCHMARKS.md. ForwardDevice campaign Stage 2 (re-scoped) — GROUPED keep-quant MoE GEMM: REAL SPEEDUP (2026-07-29, base05f6318c, NOT pushed). The lever the profiler proved. New vt opkMatmulBTQuantGrouped(mirrorskMoeGroupedGemmBf16, extends the #195kMatmulBTQuantkernel): collapses the 6 routed experts × {gate,up,down} = 18 tiny T=1 matvecs/layer into 3 grouped kernels (CUDAQuantDotGemmGroupedKernel= same dot core, per-group weight-row index; CPU provider loopsMatmulBTQuantper group = byte-identical reference).MoeBlockroutes the routed experts through it (shared stays per-expert);VT_V4_GROUPED_MOE=0rolls back. TOKEN-IDENTICAL (grouped ON == OFF == "…Paris.",11111 16 455 6102 294 8760 344 11111); gatetest_deepseek_v4_gguf_load12/12·531 (+grouped==per-expert byte-identical, RED-first permuted map). DGX--gpu --kv-cache: decode 5.83 tok/s (23-tok) / 6.60 (11-tok) — ~22% over Stage 1's 4.79 (equal length); prefill 7.6–8.1; peak 86.33 GiB; ~35% of ds4's 16.5. nvidia-smi util did NOT rise (36%) — fewer/ larger kernels leave the GB10 idle on host launches, so decode is now clearly host-launch-bound → motivates Stage 3 (decode CUDA graph, mirroring Qwen3-Coder W7). Row staysACTIVE; see docs/BENCHMARKS.md. ForwardDevice campaign Stage 3 — DECODE CUDA GRAPH: BLOCKED (architectural finding; honest last residual) (2026-07-29, base3eb018df, NOT pushed). A decode CUDA graph CANNOT be built on the host-orchestrated forward — no code change. The vt capture contract (cuda_backend.cu:173-182) demands pure async stream work (no Synchronize / host<->device copies / malloc; fixed pointers). The dense/DFlash graphs capture a device-residentForwardLayersover persistentDBufs (qwen3_5.cpp:5820/5747). DeepSeek-V4's GGUF forward (ForwardComposeImpl,device=false) instead runs MHC/Sinkhorn, the attention QK/softmax/AV Dot loop, rope, RMSNorms, the MoE router (which decidesexpert_ids), SwiGLU and combine ON THE HOST overstd::vector, between GPU GEMMs, reading each output on the host — no contiguous async device sequence to capture; the host syncs abort capture; the #183 glue kernels Upload/Download+sync per call and there is no device attention kernel. No host-node fallback exists. Prerequisite: a device-resident decode forward (all glue as real parallel device kernels on persistentDBufs, incl. a device MLA attention kernel), mirroringQwen3_5DenseDecodeGraph— a multi-brick port, NOT the #183<<<1,1>>>correctness kernels. FINAL campaign result: decode 0.68 → 5.83 tok/s (~8.6×), token-identical, ~35% of ds4's 16.5 (Stage 1 KV cache 7.9× + Stage 2 grouped MoE GEMM +22%; both merged). Named residual = host-orchestration launch overhead (util ~36%), removable only by the device-resident forward. RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick 0 (scope) + Brick A (device MLA attention kernel) (2026-07-29, base59260579, NOT pushed). User commissioned the device-resident decode forward (unblocks the Stage-3 graph). Brick 0 =.agents/specs/deepseek-v4-device-decode.md(bricks A→B→C→D, DBuf layout, gates, honest ceiling: a graph likely reaches ~10-13 tok/s but SHORT of ds4's 16.5 — fp8 KV + tuned MMQ are a named follow-on). Brick A = first real device V4 forward kernel:DecodeAttnKernel(cuda_deepseek_v4.cu) replaces the host QK/softmax-sink/AV Dot loop, reading the unified KV latent in place (num KV heads = 1), preserving hostSoftmaxWithSinkorder (only expf vs std::exp differs). FlagVT_V4_DEVICE_ATTN(default OFF). Gate = correctness: CUDA unittest_cuda_deepseek_v412/12·412 (+device-vs-host RelL2<1e-5, RED-first no_sink);test_deepseek_v4_gguf_load12/12·531; real-model TOKEN-IDENTICAL (ON==OFF=="…Paris."). Speed (not Brick A's gate): decode 6.23 tok/s (24-tok) vs 5.83 baseline (~7%, grows with ctx); util still ~35% (payoff at Brick D's graph); peak unchanged. Rollback-able (flag OFF default). Bricks B→C→D await review. RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick B increment 1: device clamped-SwiGLU (in place) (2026-07-29, base21191ce2, NOT pushed). Brick B reimplements the decode glue as real in-place device kernels (no Upload/Download/Sync — the Brick-A pattern), landed in gated increments. Increment 1 = MoE clamped-SwiGLU:ClampedSwiGLUInPlaceLaunch(reuses the tested #183ClampedSwiGLUKernelon unifiedgate_up/out);DispClampedSwiGLUroutes to it underVT_V4_DEVICE_GLUE=1+ CUDA (default OFF). CHARACTERIZED NEAR-TIE (RelL2<1e-5, deviceexpfvs hoststd::expin the SiLU — stated, not bit-identical). Gate = correctness: CUDA unittest_cuda_deepseek_v413/13·671 (+in-place==host RelL2<1e-5, RED-first clamp-limit);test_deepseek_v4_gguf_load12/12·531; real-model TOKEN-IDENTICAL despite the near-tie (VT_V4_DEVICE_GLUE=1, and combined withVT_V4_DEVICE_ATTN=1, both = "…Paris."). Speed flat (6.12 tok/s attn+glue 24-tok vs 5.83; util ~38% — payoff at Brick D). Remaining Brick-B glue (router, MHC pre/post/head/Sinkhorn, RMSNorm, RoPE, combine) = next increments. RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick B increment 2: MHC-post/head + router device glue; MhcPre FINDING (2026-07-29, base21191ce2, commit3046f087, NOT pushed). In-place launchers reusing the tested #183 kernels for MHC pre/post/head + router;DispMhcPost/DispHcHead/DispRouteroute to them underVT_V4_DEVICE_GLUE. KEY FINDING (measured): the #183MhcPreKernel/HcHeadKernelare<<<1,1>>>single-thread stubs — routing MhcPre (86 calls/step over 16K, one GPU thread) REGRESSED decode ~10× (0.59 vs 6.5 tok/s), so MhcPre stays HOST until a real PARALLEL MhcPre kernel exists (folded RMSNorm + 20-iter Sinkhorn + gates + collapse — the hardest remaining piece). The Grid-launched glue (SwiGLU, router, MHC-post, hc_head) IS on device + FLAT. Gates: CUDAtest_cuda_deepseek_v414/14·749 (in-place == round-trip #183 bit-identical, RED-first);test_deepseek_v4_gguf_load12/12·531; real model TOKEN-IDENTICAL all-device-on "…Paris.", decode 6.38 tok/s (flat). BRICK B NOT COMPLETE: parallel glue done + gated; remaining = parallel MhcPre kernel + RMSNorm + RoPE + MoE combine. RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick B increment 3: PARALLEL MhcPre kernel + per-op-sync-tax finding (2026-07-29, base21191ce2, commit9bd51523, NOT pushed). The crux:MhcPreParallelKernel(one block, 256 threads over the hc·H=16K width; block-tree reductions in double; Sinkhorn+gates on thread 0 in host order) replaces the<<<1,1>>>stub. CHARACTERIZED NEAR-TIE (width reductions reorder → RelL2<1e-3).DispMhcPrere-enabled. Gates: CUDAtest_cuda_deepseek_v414/14·736 (MhcPre parallel==round-trip RelL2<1e-3, RED-first);test_deepseek_v4_gguf_load12/12·531; real model all device on (incl parallel MhcPre): TOKEN-IDENTICAL "…Paris.", decode 5.43 tok/s (24-tok). FINDING — per-op SYNC tax, not poison: parallel MhcPre is fine (10× poison gone); the ~21% dip vs host-glue (5.43 vs 6.84) is the ~560 per-op device-glue drains/step — inherent to Brick B's per-op-synced structure, removed by Brick C (drop syncs) + Brick D (graph). No single-thread on the hot path. Remaining for full device-residency: device RMSNorm, RoPE, MoE combine (can fold into Brick C). RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick C part 1: folded-in device glue kernels (RMSNorm/RoPE/combine) (2026-07-30, basef2fff1ee, commitd68768b4, NOT pushed). The device kernels the resident assembly needs:RmsNormKernel(parallel block reduction; has_w=false → per-head q-RMS),RopeKernel(per-row sequential YaRN recurrence, inverse flips sin),MoeCombineKernel(Σ_a w_a·eo_a). Added to the seam (rms_norm/ropein DsaDeviceKernels,moe_combinein MoeDeviceKernels); NOT wired into the forward yet (they land in the resident assembly). ALL THREE CHARACTERIZED NEAR-TIES (stated): RMSNorm (reduction reorder), RoPE (cos/sin lib), and combine (a new FMA finding — device fused multiply-add vs host separate mul+add; my bit-identical assertion correctly failed at last-ULP → fixed to a near-tie). GATE: CUDAtest_cuda_deepseek_v415/15·1106 (RMSNorm/RoPE/combine == host RelL2<1e-5, RED-first each);test_deepseek_v4_gguf_load12/12·531. Brick C part 2 (the SPEED part) REMAINS: thread the T=1 decode through persistent DBufs across the 43-layer stack, consume these + the Brick-A/B kernels resident, and drop the ~560 per-op drains (sync only at the step-boundary logits read) — that recovers + exceeds 5.83. RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick C part 2: the resident T=1 decode assembly — CORRECT + token- identical, but eager-SLOWER (honest launch-bound finding) (2026-07-30, base24bddf15, commit3a635624, branchdeepseek-v4-resident-assembly, NOT pushed).ForwardResidentDecodeGgufruns the whole 43-layer T=1 step as ONE async device chain over the unified buffers — every GEMM defers (GemmIntoKq/RowSliceInto/ GroupedInto), the small host primitives run on the Brick-C device kernels IN PLACE (q/kv/final RMSNorm + the batched per-head q-RMSrms_norm_rows, dual+inverse RoPE, MHC pre/post/head, router, clamped-SwiGLU, combine, decode_attn — none draining), routing stays RESIDENT (device router → i32 topk_ids the grouped GEMM consumes - topk_weights the combine consumes — no host gather), the KV append writes the new deck row into the cache
slot, async cudaMemcpyAsync for the grouped-GEMM broadcast + [gate|up] pairing. Flag
VT_V4_RESIDENT_DECODE=1(default OFF). New tiny glue:rms_norm_rows(batched nh=64 q-RMS in ONE launch). Fixed a guard bug (rejected on config compress_ratios; the keep-quant run is densedsa_denseregardless → resident never engaged on the first run). GATES (DGX GB10, deepseek-v4-gen --gpu --kv-cache, real 80.7 GB, 24 decode): CUDA unit 16/16·33877 (+rms_norm_rows case);test_deepseek_v4_gguf_load12/12·531; resident ON == host OFF TOKEN-IDENTICAL ("…Paris.", exact 25 ids). SPEED (honest): resident-eager 5.17 tok/s vs host 6.44 (−20%); util 38%→55%. Attribution: NOT a correctness bug (tokens identical), NOT the router-gate drains (bf16 gate → 43 CPU-GEMM drains ≈ 1.4%/step) — launch/host-gap bound (~1700 small device-kernel launches/step replace fast host-ARM glue →45% GPU idle). The §4 "may improve" precondition, NOT the payoff — Brick D (collapse the ~1700 launches into ONE cudaGraphLaunch) is the payoff (55%→100% busy ≈ 1.8× → ~9 tok/s, would exceed 6.44), PROVIDED the bf16 router gate is first device-ified (the one remaining non-capturable host op). Rollback- able (flag OFF). STOPPED for review before Brick D. RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — Brick D: decode CUDA graph — step 1 gated; step 2 infra gated, capture BLOCKED on the model's Q8_0 weights (grounded finding; REVISES the Brick C attribution) (2026-07-30, base24bddf15, branchdeepseek-v4-brick-d, commits3c262cea+f62de523, NOT pushed). STEP 1 = device router gate (RouterGateKernel: [ne,H] BF16 gate × f32 x, sequential f32 + exact bf16 upcast, bit-identical to the host CPU MatmulBT) → resident routing is device. Gate: CUDA unit 17/17·33919; real model resident==host TOKEN-IDENTICAL. STEP 2 = the graph:decode_attn_g(KV length from a DEVICE buffer, attends cache[0..len)+deck_new, fixed shmem → one graph serves a growing context) +V4Graph(persistent buffers, fixed-cap KV, cold→warm→capture→replay, deck_new→cache appended between replays on-stream; every captured input a persistent buffer) + opaque holder onDeepseekV4KvCache; flagVT_V4_DECODE_GRAPH=1(default OFF). Unit 18/18·34176 (decode_attn_g == eager decode_attn BIT-IDENTICAL + RED-first);test_deepseek_v4_gguf_load12/12·531. BUT capture ABORTS:matmul_bt_quant: keepquant CPU-fallback drain: operation not permitted when stream is capturing. ROOT CAUSE (gguf dump): 345 Q8_0 tensors — the MLA projections (attn_q_a/attn_kv/wq_b), o-LoRA (wo_a/wo_b), shared experts,output.weight(the AProjQ8/SExpQ8/ OutQ8 in the filename).IsCudaKeepQuantSupportedis FALSE for Q8_0 (only IQ2/IQ3/Q2_K…Q6_K have a CUDA keep-quant GEMM), so each of ~15 Q8_0 GEMMs/layer CPU-fallbacks with acudaStreamSynchronize— legal eager, ILLEGAL during capture. REVISES Brick C: the eager-resident ~45% GPU-idle is dominated by these ~646 Q8_0 CPU-fallback GEMM drains/step (MLA/shared/lm_head run on the ARM CPU), NOT device-launch overhead. UNBLOCK = a CUDA Q8_0 keep-quant GEMM = the named "tuned expert GEMM" follow-on (a USER decision — not started). Graph infra is COMPLETE + unit-gated + capture-hazard-safe behind a default-OFF flag; replays once Q8_0 lands. FINAL honest table (real 80.7 GB, 24 decode): host 6.44 · eager-resident ~5.0-5.2 · graph BLOCKED (Q8_0) · ds4 16.5. Rollback-able (both flags OFF). RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — the Q8_0 CUDA keep-quant GEMM (graph unblocker + dominant lever): CAMPAIGN COMPLETE with an honest speed finding (2026-07-30, base93e72203, branchdeepseek-v4-q8-cuda-gemm, commitfcf3003c, NOT pushed). Added an on-GPU Q8_0 path tocuda_quant_dot.cu(both providers), moving the ~646 Q8_0 GEMMs/step (MLA/o-LoRA/shared/lm_head — AProjQ8/SExpQ8/OutQ8) off the ARM CPU onto the GB10:QuantizeQ8_0Kernel(bit-exact port of the CPUQuantizeRowQ8_0, 32-block Q8_0 activation) +QuantDotGemmQ8_0Kernel(+grouped) (warp/output, int core bit-identical toVecDotQ8_0Q8_0, float scale reassociated — NMSE 5e-4 band); wired BEFORE the CPU-fallback ⇒ NO stream sync (the capture unblocker). GATE:test_cuda_quant_dot2/2·105601 (Q8_0 CUDA==CPU nmse≤1e-6 + f64≤5e-4);test_cuda_deepseek_v418/18;test_deepseek_v4_gguf_load12/12; real 80.7 GB model host + resident + GRAPH all TOKEN-IDENTICAL ("…Paris."); THE DECODE GRAPH NOW CAPTURES + REPLAYS (0 errors, token-verified 24 steps). FINAL honest table (2 stable runs, 24 decode): host 7.22 tok/s · eager-resident 5.64 · GRAPH 5.57 · ds4 16.5. Q8_0-on-GPU lifted the HOST (shipped-default) path 6.44→7.22 (+12%) — the real win. The resident+graph track does NOT beat host (5.6 vs 7.22): Q8_0 on GPU makes the resident path GPU-compute-bound (util 94-95%), the correctness-grade device glue does more GPU work than the ARM-CPU glue, and the graph recovers nothing (already 94% busy). REVISES the thesis: eager-resident is glue-kernel-efficiency bound, not launch-bound; the graph works + is token-identical (infra proven) but is a speed dead-end on GB10 vs the fast ARM CPU. NAMED last mile to ds4's 16.5 (not started, user decision): tune the device glue kernels, fp8 KV, and/or an async host path overlapping CPU glue with GPU GEMMs. Rollback- able (device flags OFF; the Q8_0 GEMM is always-on + benefits every path). RowACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — GLUE-KERNEL TUNE (top-3 offenders): THE GRAPHED DECODE NOW BEATS HOST (2026-07-30, basef83c8065, branchdeepseek-v4-glue-tune, commite232146f, NOT pushed). After the profile-first tunable-win verdict (glue 44% of GPU time), tuned the top-3: (1)RouterGateKernelone-thread/expert → ONE WARP/expert — 18× (7.5%→0.6%); (2)MhcPreParallelKernel(the #1 kernel,<<<1,256>>>= 24 sequential block reductions in ONE SM) → splitMhcPreDotsKernel(one block per mix dot → concurrent, BIT-IDENTICAL) +MhcPreFinishKernel— 4.8× (25.5%→7.5%); (3)HcHeadKernel<<<1,1>>>→ one-block parallel — 42× (3.75 ms→88 µs/instance). GATE:test_cuda_deepseek_v418/18·34176 (tuned kernels' equivalence cases stayed green);test_deepseek_v4_gguf_load12/12·531; real 80.7 GB host + resident - graph all TOKEN-IDENTICAL ("…Paris."), graph 0 errors. NEW TABLE (2 runs, 24 decode): host 7.20 ·
eager-resident 7.96 · GRAPH 7.92 · ds4 16.5 — THE GRAPHED DECODE BEATS HOST (+10%). New nsys split: GEMM
~78% · GLUE ~21% (RoPE 7.3%, MhcPreFinish 6.0% incl. Sinkhorn floor, Route 4.7% — partly irreducible) · ATTN
~1.3% — the step is now GEMM-bound (full glue elimination caps ~10 tok/s; the true last mile to ds4 16.5 is
GEMM/quant microarch — fp8 KV, tuned MMQ). The H2H-memcpy lever is NOT on the critical path (util 95%,
beats host). RECOMMEND flipping the device-resident/graph default ON (fastest path; flags still default OFF
pending that call). Rollback-able. Row
ACTIVE; see docs/BENCHMARKS.md. Device-resident decode campaign — DEFAULT FLIPPED: device-resident decode is the SHIPPED DEFAULT (2026-07-30, base73d4799c, branchdeepseek-v4-resident-default, commit75c3697a, NOT pushed). Broadened validation (real 80.7 GB, resident vs host, 4 prompts × 256 tok): P0 (factual) TOKEN-IDENTICAL (258 toks); P1/P2/P3 (open-ended) diverge only at genuine near-tie positions into COHERENT, deterministic continuations (bounded kernel noise can only flip host's ~tied top-2 → by construction near-ties, not bugs), 0 errors — the ratified coherent-near-tie gate is met. Resident is ~1.8× the host path at a grown 256-tok context (7.8 vs 4.3). FLIP:ResidentDecodeEnabledOFF→ON (VT_V4_RESIDENT_DECODE=0= rollback off-switch → host forward); the guard is UNCHANGED so CPU / non-dense / T>1 (prefill) / no-KV still fall back to host (verified: prefill→host every run; gguf_load 12/12). The CUDA graph stays OPT-IN default OFF (graph 7.92 ≈ eager-resident 7.96 — GPU-bound at 95%, the graph adds nothing on GB10 + carries the capture-hazard surface). GATE:test_cuda_deepseek_v418/18·34176;test_deepseek_v4_gguf_load12/12·531; real model new DEFAULT (resident) 8.01 tok/s vs=0(host) 7.24, both "…Paris." token-identical. Shipped DeepSeek-V4 decode = device-resident ~7.96 tok/s (~48% of ds4, GEMM-bound); last mile = fp8 KV + tuned MMQ (named residual). Rollback viaVT_V4_RESIDENT_DECODE=0. RowACTIVE; see docs/BENCHMARKS.md. Last-mile campaign — Brick 0 (PROFILE-ONLY): the keep-quant GEMM roofline (2026-07-30, baseaed4a498, branchdeepseek-v4-last-mile, commit42a99471, NOT pushed). Profile-only (holding for the coordinator's synthesis with parallel source research — no MMQ implementation). MEASURED the T=1 decode keep-quant GEMM efficiency vs roofline (real 80.7 GB, nsys per-kernel time ÷ exact weight bytes; GB10 peak = 240 GB/s via a float4 copy microbench):QuantDotGemmQ8_0(35.6%, 6.60 GB/step) = ~150 GB/s = 63% of peak → MEMORY-bound (~1.6× headroom);QuantDotGemmGrouped<IQ2_XXS>(20.2%) = ~45 GB/s = 19% → DEQUANT/LATENCY-bound (~5×);<Q2_K>(10.0%) = ~57 GB/s = 24% (~2.5×);QuantizeQ8K/QuantizeQ8_0(12%) = LAUNCH-bound (~795 tiny launches/step). KEY: every T=1 GEMM is <1% of int8 compute peak → memory-bound; tensor cores do NOT help a decode matvec (mmq.cu is prefill; the reference is llama.cppmmvq.cu). Ranked levers: (1) grouped-MoE dequant (the 5× gap, mmvq.cu vectorized dequant + dp4a), (2) fuse activation-quant, (3) Q8_0 coalescing, (4) fp8 KV (parity/long-ctx, NOT the short-ctx lever). Honest projection: ~14-16 tok/s target (ds4 16.5 = 58% of the BW roofline, fully fused); ~28 tok/s is the hard ceiling. Full table:.agents/specs/deepseek-v4-last-mile.md. RowACTIVE; see docs/BENCHMARKS.md. Last-mile campaign — Brick 1: __dp4a vectorized-dequant matvec for the grouped keep-quant GEMMs — SPLIT result (2026-07-30, baseaed4a498, branchdeepseek-v4-last-mile, commitc1f92d24, NOT pushed). Ported llama.cppmmvq.cu/vecdotq.cuhSIMD dequant intoDotIQ2XXS(vecdotq.cuh:920-928) +DotQ2K(vecdotq.cuh:329-354), keeping warp-per-output + Q8_K activation. BIT-IDENTICAL (__dp4a= exact int32); thetest_cuda_quant_dotnmse≤1e-6 gate CAUGHT a signed-overflow UB in the IQ2 sign broadcast (RED-first worked) → fixed to unsigned → 2/2·105601 zero drift;test_cuda_deepseek_v418/18;test_deepseek_v4_gguf_load12/12; real model resident-default TOKEN-IDENTICAL "…Paris.". SPLIT: Q2_K grouped 2.35× (24%→56% of peak, MAC was the bottleneck → now memory-bound; 10.0%→4.4% of step); IQ2_XXS grouped FLAT (~17-19% of peak) — GROUNDED:d_iq2xxs_gridis__constant__, our 32 lanes look up different indices → divergent constant-memory reads ~32-way serialized/warp = the bottleneck, not the MAC. Decode 8.01 → 8.51 tok/s (+6%, 4 stable warm runs) vs ds4 16.5. NEXT (Brick 1b): the IQ2 grid-lookup (move to GLOBAL / mmvq.cu handling) — the bigger grouped kernel's win. STOPPED for review. RowACTIVE; see docs/BENCHMARKS.md. Last-mile campaign — Brick 1b: IQ2_XXS grid-lookup fix (__constant__→ GLOBAL) — the flat kernel unblocked, +12.5% (2026-07-30, basef6c34252, branchdeepseek-v4-last-mile, commite4d8845b, NOT pushed). Brick 1's finding: IQ2_XXS grouped stayed at ~19% of peak becaused_iq2xxs_grid[256]is__constant__and the 32 warp lanes read DIFFERENT indices → divergent constant reads serialize ~32×/warp. FIX: movedd_iq2xxs_grid+d_ksigns_iq2xsto__device__GLOBAL (L2-cached, cross-lane parallel; llama.cpp's mmvq does the same). BIT-IDENTICAL (same literals):test_cuda_quant_dot2/2·105601 nmse≤1e-6 ZERO drift;test_cuda_deepseek_v418/18;test_deepseek_v4_gguf_load12/12; real model resident-default TOKEN-IDENTICAL "…Paris.". RESULT: IQ2_XXS grouped 2.45× (median 265→108 µs; 19% → 46% of peak — now memory-bound-ish like Q2_K 56% / Q8_0 63%; 22.8% → 10.0% of step). Decode 8.51 → 9.58 tok/s (+12.5%, 5 stable warm) vs ds4 16.5 (~58% of ds4; campaign-cumulative host 6.44 → 9.58 = +49%). The grouped-MoE dequant lever (Bricks 1+1b) is DONE. NEXT: Brick 2 (activation-quant fusion, ~14% launch-bound) → Brick 3 (Q8_0 coalescing, 43% at 63% of peak). STOPPED for review. RowACTIVE; see docs/BENCHMARKS.md. Last-mile campaign — Brick 2: routed gate/up activation preq-reuse (broadcast) — bit-exact +4.6% (2026-07-30, base67bb8d1cafter rebase, branchdeepseek-v4-last-mile, commit99d2b282→amended, NOT pushed). The resident-decode routed experts fedxrep— topk-identical copies of the shared hidden x — into the gate + up grouped GEMMs, re-quantizing an IDENTICAL row per expert (6× redundant) + a per-layer topkAsyncCopyF. FIX (ds4's preq pattern): the grouped providers (Q8_K + Q8_0) detect a 1-row activation (act.shape[0]==1 && P>1) → quantize ONE row, kernels read block-set 0 for all p (bcast); the resident forward (eager + graph) passes x withact_rows=1, dropping xrep. BIT-IDENTICAL (identical input → identical block-quant → identical dot):test_cuda_quant_dot2/2·105601 nmse≤1e-6;test_cuda_deepseek_v418/18·34176;test_deepseek_v4_gguf_load13/13·631 (+1 broadcast==replicated BYTE-IDENTICAL + RED-first case); real 80.7 GB model resident-default TOKEN-IDENTICAL "…Paris." (ids byte-equal to host=0). RESULT (nsys re-profile, 50 tok):QuantizeQ8K9.8%→7.3%,QuantizeQ8_04.5%→4.7% (bucket14.3%→12.0%). Decode 9.58 → 10.02 tok/s (+4.6%, 6 stable warm runs 9.99–10.04; nvidia-smi 92% util) (host=08.47→8.50, MoE host path unchanged) vs ds4 16.5 (~61% of ds4; campaign-cumulative host 6.44 → 10.02 = +56%). HONEST: the quant bucket is LAUNCH-bound (fixed per-launch overhead dominates → cutting per-launch rows 6→1 barely moves the %); the +4.6% comes mostly from eliminating the 6×/layer xrep host-copies. Remaining quant lever = LAUNCH-COUNT reduction (fuse quant into GEMM / dedup gate+up). NEXT: Brick 3 (Q8_0 coalescing — now 45% of step at 63% of peak, the dominant lever). STOPPED for review. RollbackVT_V4_RESIDENT_DECODE=0. RowACTIVE; see docs/BENCHMARKS.md. Last-mile campaign — Brick 3: Q8_0 GEMM vectorized dp4a dot — BIT-EXACT but NEAR-FLAT +0.5% (the Q8_0 matvec is at the 34-byte-block MEMORY WALL) (2026-07-30, basedd6cc93c, branchdeepseek-v4-last-mile, commit91e51aea→amended, NOT pushed).QuantDotGemmQ8_0Kernel(+grouped,45% of step at 63% of BW peak) read the 32 int8/block as 32 scattered int8 loads + a scalar-MAC loop. FIX (mirrors llama.cpp64% of peak. Decode 10.02 → 10.07 tok/s (+0.5%, 6 warm runs 10.05–10.09, non-overlapping vs Brick 2's 9.99–10.04; hostvecdotq.cuh:vec_dot_q8_0_q8_1_impl): read 8 int32 + 8__dp4a. The Q8_0qsis at offset 2 in the 34-byte block → mis-aligned for a naked int32 load, soGetIntB2(bit-exact port ofggml-cuda/common.cuh:get_int_b2, two uint16 loads) reconstructs the bytes →__dp4agets the same signed int8 lanes as(int)qs[p]. BIT-IDENTICAL:test_cuda_quant_dotq8_0 2/2·105601 nmse≤1e-6 ZERO drift;test_cuda_deepseek_v418/18·34176;test_deepseek_v4_gguf_load13/13·631; real 80.7 GB model resident-default TOKEN-IDENTICAL "…Paris." (ids byte-equal to host); SASS shows exactly 8 IDP in the Q8_0 float kernel (Iron-Law compile-proof). **RESULT (nsys, 50 tok) — HONEST NEAR-FLAT: the Q8_0 kernel moved only 2.182→2.149 s (−1.5%; med 30,112→29,472 ns) = 45.0%→44.6% ≈ 63%→=08.50→8.56) vs ds4 16.5 (~61% of ds4; campaign host 6.44 → 10.07 = +56%).** FINDING: dp4a cut instructions but not the bottleneck — the kernel is DRAM-BW-bound and the 34-byte-misaligned block precludes true 128-bit coalesced loads; reaching ~85% needs an ALIGNED WEIGHT REPACK at load (numerics-neutral, but trades the in-place unified-memory mmap design). THE BIT-EXACT IN-KERNEL GEMM LEVERS ARE NOW LARGELY EXHAUSTED (both grouped dequant kernels memory-bound after B1/B1b; Q8_0 at its block wall after B3). Remaining: (a) aligned Q8_0 repack (numerics-neutral, design-changing); (b) glue (Rope+MhcPreFinish+Route ≈ 23%); (c) numerics-delicate (dequant-cache, fp8-KV). At 10.07 ≈ 61% of ds4 vs the ~13–15 bit-exact ceiling. STOPPED for review. RollbackVT_V4_RESIDENT_DECODE=0. RowACTIVE; see docs/BENCHMARKS.md. Last-mile campaign — Brick 4: aligned Q8_0 weight repack for coalesced CUDA loads — a MEASURED NEGATIVE (bit-exact, NO speedup, +4.7 GiB peak RSS); the coalesced-load hypothesis is REFUTED (2026-07-30, based1c63dc7, branchdeepseek-v4-last-mile, commitb8a3f691→amended, NOT pushed). Repacked the Q8_0 tower (6.146 GiB) at load into[all qs (16B-aligned) | all scales](RepackQ8_0Cuda) so a warp lane reads via alignedint4loads (QuantDotGemmQ8_0AlignedKernel); memory-disciplined (copy off mmap +DropSpanResidency); opt-inVT_V4_Q8_0_ALIGN(default OFF);wo_aexcluded (row-sliced). BIT-EXACT:test_cuda_quant_dot3/3·105841 (new aligned==plain BYTE-IDENTICAL + RED-first);test_cuda_deepseek_v418/18;test_deepseek_v4_gguf_load13/13; real 80.7 GB ALIGNED resident ids BYTE-EQUAL to plain, "…Paris.". RESULT — NEGATIVE: decode 10.07 → 10.03 tok/s (FLAT). nsys: aligned kernel 1.671 s + still-plain wo_a 0.506 s = 2.177 s vs Brick 3's 2.149 s (+1.3% WORSE). PEAK RSS: aligned 91.05 GiB vs plain 86.33 GiB = +4.7 GiB (safe, 28 GiB headroom; not net-flat). WHY: the aligned int4 loads didn't help + separating the scales added a scattered per-block gather. With Brick 3 (dp4a also no help), the Q8_0 matvec is NOT load/ALU/alignment-bound but LATENCY/OCCUPANCY-bound (1-warp/output, ~33k tiny launches/step). THE BIT-EXACT IN-KERNEL GEMM LEVERS ARE EXHAUSTED. RECOMMEND: do not ship — revert (keep record) or merge default-OFF. Residual = a kernel-structure change (risky) OR fp8-KV/dequant-cache (numerics-delicate) — a USER CALL. At 10.07 ≈ 61% of ds4 vs the ~13–15 bit-exact ceiling. STOPPED for review. RollbackVT_V4_Q8_0_ALIGNdefault-OFF. RowACTIVE; see docs/BENCHMARKS.md. W8-run (2): geometry FIXED — the forward now RUNS the real 158 B model end-to-end; generation still INCOHERENT (2026-07-29, basefba56f9b, NOT pushed). The layer-2 hard-fail is fixed: the DSA compressor projects to2*head_dim(ds4coff=2), nothead_dim, so the real keep-quant run uses DENSE MLA (EXACT forseq ≤ index_topk=512; the sparse compressor/indexer at real geometry is a named residual), guarded by a new real-ish gate (test_deepseek_v4_gguf_load9/9·388, RED-first). Also landed: mmap-VIEW keep-quant load (PEAK RESIDENT 116.2 → 86.1 GiB, byte-identical, load 78 s → 6 s warm) and per-layer YaRN RoPE (compressed layers use base 160000 + freq_scale 1/16, ds4rope_tail_ext_inplace). The forward now runs all 43 layers on the real model at ~1.4–1.7 s/tok / peak 86 GiB, but greedy output is a degenerate 2-cycle loop (201 7249 465 7249…) — INCOHERENT. A different prompt gives different ids, so the body is input-responsive; the failure is a numerical/scaling fidelity error (suspects: MoE routing/expert-scale/shared-expert, MHC Sinkhorn mixing, keep-quant IQ2_XXS/Q2_K numerics), needing instrumented per-layer comparison vs the ds4 oracle. Ruled out: geometry, attention scale/structure (match ds4), fp8 KV roundtrip, memory path, rope. ds4 reference (same file, GB10 GPU): coherent, prefill 358 tok/s, decode 16.5 tok/s. Row stays SPIKE (no coherent generation). W3 attention primitives landed (2026-07-28): the genuinely-new-vs-V2/V3 math is ported as portable host references and unit-gated — the DSA "Lightning Indexer" sparse top-k SELECTION (a weighted multi-query logit with a load-bearing per-head ReLU, then a causal top-512 token select), plus the two 512-wide-MLA output pieces V2/V3 lack: per-head attention-sink softmax and grouped output-LoRA.test_deepseek_v4_dsapasses 13/13 (hand-derived literal cases plus double-precision references, clean CPU-Wall -Werror -Wextra); the gate is honest hand-case + structural review, not a dumped-oracle comparison (the fixed-config 167B model cannot be built at a tiny shape). It is additive and byte-neutral for DeepSeek-V2. SGLang v0.5.15 registers and implementsDeepseekV4ForCausalLM(the full DSA/MHC stack), so it is a viable second reference, subject to the same single-GB10 memory limit. W4 attention primitives landed (2026-07-29): the second half of the DSA stack is ported as portable host references and unit-gated — the DSA COMPRESSOR forward (the softmax-weighted window pool that compresses the KV-state window into one latent, computed per head-dim column, plus the fused save-time position-embedding add and the RMSNorm) and the fp8_ds_mla KV-cache state read/write layout (the compressed latent split into a 448-wide part quantized to FP8 with per-64 UE8M0 power-of-two block scales and a 64-wide RoPE part stored bf16, at a 576-byte token stride, plus the dequant read).test_deepseek_v4_compressorpasses 12/12 (hand-derived literal cases plus double-precision references plus an independent scale-byte recompute, RED-first proven); ported 1:1 from vLLM and cross-checked against SGLang, additive and byte-neutral for DeepSeek-V2. W5 Manifold Hyper-Connections landed (2026-07-29): the hardest V4 brick — the residual stream becomes a[tokens, hc_mult, hidden]manifold mixed at every attn/ffn boundary by a 20-iteration Sinkhorn-normalized doubly-stochastic matrix, with the layer RMSNorms folded in and a learned head that collapses the streams back to one hidden. Ported as portable host references (the Sinkhorn, the mHC pre/post mixes, the hc_head collapse) and unit-gatedtest_deepseek_v4_mhc14/14 (hand-derived literals plus from-first-principles double-precision references, RED-first proven on both the iteration count and a normalization axis). The W0 scope's "no eager reference upstream" premise is corrected: vLLM ships an eager PyTorch reference (mhc/torch.py) and four upstream implementations agree on the Sinkhorn, so the port is grounded 1:1 and cross-checked against an independent derivation. Additive and byte-neutral. W6 sqrtsoftplus + hash-routed MoE landed (2026-07-29): the three genuinely-new-vs-V2/V3 MoE pieces are ported as portable host references and unit-gated — the router score functionsqrt(softplus(x))(distinct from V2/V3's sigmoid/softmax), the router that adds the correction bias for SELECTION only then either picks top-k or looks the experts up directly in thetid2eidtoken-id→expert hash table (bypassing top-k), gathering the weights from the UNBIASED scores, and the asymmetric clamped SwiGLU expert activation (gate clamped max-only, up clamped both-sided).test_deepseek_v4_moepasses 12/12 (hand-derived literals plus from-first-principles double-precision references, RED-first proven on all three load-bearing levers: dropping the sqrt, gathering weights from the biased scores, and symmetric-clamping the gate). The shared grouped-GEMM / expert / shared-expert / NVFP4 machinery is REUSED, not re-ported (only scoring + hash + clamp are net-new); MegaMoE is SM100-only so GB10 mirrors the FusedMoE-fallback router. Additive and byte-neutral. W7 forward assembly landed (2026-07-29): theVT_CHECK(false, "W3-W8 pending")stub is replaced by a realDeepseekV4Model::Forwardthat composes the four landed host primitives (W3 DSA/MLA seams, W4 compressor + fp8_ds_mla KV, W5 MHC + Sinkhorn, W6 sqrtsoftplus/hash MoE) into an end-to-end logits producer on the CPU path at a small synthetic config — the brick that finally makes V4 structurally runnable. It is a STRUCTURAL/composition gate (test_deepseek_v4_forward6/6·26: the layer interleave runs, the residual is a[T, hc_mult, H]stream, the hash layers route by token id, the DSA sparse path selects and the compressor pools, logits are finite/deterministic, and a deliberately-miswired interleave changes the output — all-gated hash, skip-final-fold, no-sink), NOT a real-checkpoint token gate. Honest 3-state: the CPU forward assembly at tiny shape is DERIVED + BUILD-VERIFIED (structural); it does not claim V4 "runs" a real model. Additive and byte-neutral (shared MLA/MoE + the W3-W6 primitive TUs untouched). The remaining bricks — the device kernels (which reuse the existing grouped-GEMM) +ForwardDevice, the real-checkpoint tower materialization, and the full strict model gate (W8) — stay multi-Spark-blocked; the single-Spark IQ2_XXS-GGUF vehicle also needs the V4-GGUFblk.N.*name map (download-blocked). The frontier families Kimi / MiniMax / GLM-latest are scoped for mechanical porting in a dedicated spike: Kimi-Linear-48B is the one that fits GB10 (91.5 GiB) and is e2e-gateable, while MiniMax-M2 (214.3 GiB fp8), Kimi-K2, MiniMax-M3 and GLM-5 remain hardware- or dependency-blocked (honesty-pass only). Kimi K3 (released 2026-07-27, after the pin) is scoped derive-and-ship in kimi-k3.md: a 2.8T MoE whose text backbone is the Kimi-Linear KDA+MLA+MoE hybrid massively scaled (H=7168, 93 layers = 69 KDA + 24 MLA, 896 experts, MXFP4 + MoonViT-V2 vision). It reuses heavily (our GDN, DeepSeek MLA, DeepSeek MoE, the Qwen3.6-35B GDN-hybrid skeleton); net-new = the KDA kernel delta, MXFP4 (group-32/e8m0), AttnRes (report-only, unconfirmed) and the MoonViT-V2 tower. It does not fit one GB10 (~1.56 TB MXFP4, ~12× the 119 GiB pool) and is not in the pinned oracle, so there is no on-box golden — like the beyond-vLLM CUDA bricks; the real signal is a primitive gate on the fitting Kimi-Linear-48B proxy plus build-verify, with full K3 verification left to a multi-Spark / 16×H200-class box. The W2/W5 CPU scaffolding is now build-only (2026-07-28): an additive registry stub, nested text/vision/quant config descent, the 93-layer KDA/MLA + 896-expert MoE text-backbone structural name-map (grounded inkimi_linear.py), a REFUSE-by-name forward and an MXFP4-refuse loader, green on a CPU build with a 6/6 scaffold gate; MXFP4 and the MoonViT-V2 tower stay not-yet-buildable (deferred to the shared DeepSeek-V4 MXFP4 row and W7). The KDA kernel delta itself now has a first buildable brick (2026-07-28,KERNEL-KDA-DELTA): the four gated-linear-attention pieces plain GDN lacks — the per-channel[H,D]low-rank decay (f_a/f_bbottleneck), the sigmoid-gated output norm (FusedRMSNormGated), the three q/k/v short convs and the q/k L2-norm — are ported as portable host references and unit-gated (test_kimi_kda14/14, clean CPU-Wall -Werror -Wextra) against hand-derived literal cases plus double-precision references. KDA subclasses GDN, so its recurrence is reused and the brick is additive (the Qwen3.6 GDN gate is byte-untouched); the gate is host-reference + structural review, not a dumped-oracle comparison — the real e2e gate stays the DGX-blocked Kimi-Linear-48B proxy, and the CUDA device kernel is a named residual. Registry-metadata correction (2026-07-29): the outerKimiK3ForConditionalGenerationwrapper's_ModelInfonow setshas_inner_state=false(mirroring the Qwen3.5 hybrid-mm sibling wrappers, whose inner language-model class carries the recurrent state), fixing thetest_model_registryinvariant that every registered outer wrapper is inner-state-free;test_model_registry24/24. Test-golden sync (2026-07-29):test_model_loader_gguf's hardcoded supported-architectures string was likewise stale (predated the breadth-sweep + frontier registrations); synced to include DeepseekV4/Gemma4/KimiK3/Qwen3VL,test_model_loader_gguf3/3 (was RED on main). Env-doc hygiene (2026-07-29): two production env vars recent lanes added were undocumented (leftcheck-env-docRED) —VLLM_PLUGINS(user-facing plugin allowlist) documented indocs/ENVIRONMENT.md,VLLM_GEMMA4_MM_DEBUG(a diagnostic print switch) classified onscripts/env-doc-allowlist.txt;check-env-docrc=0. And thecheck-device-leakageDSR ratchet (RED from the W7-device lane): the 8kCUDAop-lookups indeepseek_v4_device.cpp(the DeepSeek-V4 CUDA device-forward resolver TU) are allowlisted with a reason + a recorded FOLLOW-UP to thread the runnerDeviceTypethrough them (deferred, needs a GB10 re-gate — DGX offline);check-device-leakagerc=0. All six record checkers now green on main. The matrix opens with an architecture-support checklist (a per-architecture status roll-up covering every engaged model) that a CI checker keeps in lockstep with the detailed rows.
From a next-tier batch of recent dense text families
(spike), the top three are now
implemented (correctness-complete or gating): Granite-3
(GraniteForCausalLM, granite-3.3-2b-instruct) is correctness-complete,
token-exact 16/16 vs the vLLM 0.25.0 oracle (Llama plus four scalar multipliers,
no new compute kernel); Phi-3 / Phi-4 (Phi3ForCausalLM,
Phi-4-mini-instruct) is implemented and runs but is not recorded as a clean
pass: it scores 15/16, and although its LongRoPE cos/sin cache is bit-identical
to the oracle's, two positions on one prompt fall outside the near-tie band as a
cascade after an exact-tie divergence, so closing it needs a cascade-aware gate
or a strict check on the larger phi-4 (14B); OLMo-3 (Olmo3ForCausalLM,
dual rope plus interleaved sliding window) is implemented and runs in our engine
but has no SACRED gate because the gate-time vLLM 0.25.0 oracle cannot run the
OLMo-3-1025-7B checkpoint (its transformers predates OLMo-3's per-layer-type
rope schema).
Command-R / Cohere (CohereForCausalLM, parallel-residual with a weight-only
LayerNorm and a logit scale) is now implemented (a faithful zero-new-kernel
port that compiles, links and self-registers) but has no SACRED gate yet:
every real small CohereForCausalLM checkpoint is HF-gated with no token on the
GPU box, the only ungated checkpoints are tiny-random (head_dim 8/2, outside the
validated attention path), and the GPU box is disk-full; the oracle was
run-verified at W0 (arch confirmed CohereForCausalLM, not
Cohere2ForCausalLM).
Phi-1 / Phi-2 (PhiForCausalLM, microsoft/phi-2) is now
correctness-complete (token-exact 16/16 vs the vLLM 0.25.0 oracle: 9 strict +
7 bf16 near-ties, max gap 0.25 nats, 0 forward-divergent), the OLDER Microsoft
Phi architecture, DISTINCT from Phi-3/Phi-4: a GPT-J parallel-residual decoder
(one nn.LayerNorm with bias feeds both attention and MLP), biased q/k/v/dense
projections, partial NeoX RoPE, a non-gated NewGELU MLP and an untied biased
lm_head, all reusing landed ops (its gelu_new maps to the landed
vt::GeluTanh, so no new compute kernel), with an F16-checkpoint dtype-aware
loader mirroring vLLM's f16->bf16 cast.
MiniCPM (MiniCPMForCausalLM, openbmb/MiniCPM-2B-sft-bf16) is now
correctness-complete (token-exact 16/16 vs the vLLM 0.25.0 oracle: 10 strict
- 6 bf16 near-ties, max gap 0.0 nats, 0 forward-divergent): the landed
Llama/Granite dense forward plus three scalars (a scale_emb embedding scale, a
scale_depth/sqrt(layers) scaled residual add on each sublayer, and a
dim_model_base logit scaling), tied embeddings, no new compute kernel. Its
checkpoint ships only as a pickle
.binwith no safetensors and no native transformers config, so the vehicle was prepared with a trusted torch.bin->safetensors conversion of the official weights (both the oracle and our engine read the identical bf16), and the gate feeds the oracle's exact prompt ids (its SentencePiece tokenizer is a follow-up).
MiniCPM3 (MiniCPM3ForCausalLM, openbmb/MiniCPM3-4B) is now
correctness-complete (token-exact 16/16 vs the vLLM 0.25.0 oracle: 13 strict
- 3 bf16 near-ties, max gap 0.0 nats, 0 divergent): the first MLA-attention
MiniCPM, and the model that closes the non-trivial part of this tier. It is the
MiniCPM three-scalar skeleton with its attention swapped from GQA to
DeepSeek-style MLA, reusing the landed DeepSeek-V2 MLA block (weight absorption,
absorbed-MQA decode, materialized-MHA prefill) with only three faithful deltas: a
neox-style rope flag threaded through a new default-false field on the shared MLA
dims (so DeepSeek-V2 stays byte-identical, re-gated 8/8), a LongRoPE cos/sin
cache instead of YaRN, and a reuse-only zero-pad of the FlashAttention-2 MLA
prefill from head_dim 96 up to the already-compiled 128. No new compute kernel;
its
.bin-only weights were converted via trusted torch like MiniCPM.
The remaining recent-dense families are the trivial tail only: Yi
(YiForCausalLM, a Llama alias) and InternLM3 (InternLM3ForCausalLM,
InternLM2 plus a sliding window).
Measured against an oracle built from source at the ACTUAL parity pin, the gap is 0.9970x total throughput, not the 0.9819x published against the older pip-installed 0.24.0 release: that release is 1.25% faster than the pin, so the old denominator was understating us. TPOT (+12.4%) is the one real gap. Re-validated 2026-07-29 after rebasing onto 139 upstream commits: 0.9972x / TPOT 1.1247x, every axis inside noise and token-identical to the prior series.
One open lead is on record from the same profiling pass: cuBLASLt resolves Ampere-class GEMM kernels on this Blackwell device. It is unmeasured and may be a non-issue, since it lands on the axis we already pass.
Device-specific references in the device-agnostic layer stay at their ratchet floor: capability questions ("is memory unified?") rather than device-type tests.
Benchmark provenance is gated as well as measured: both the comparison and the same-binary A/B harness refuse to start a leg on a GPU that is not idle, after a cooldown rather than before it.
Alongside the default build, -DVLLM_CPP_SANITIZE=address,undefined and
-DVLLM_CPP_SANITIZE=thread build the CPU tier under the dynamic detectors, and
CI runs both as separate jobs. Verified end to end: the ASan+UBSan lane builds
and passes test_input_batch, test_combine_tokens and test_arena with leak
detection on. The lanes keep the warnings but drop -Werror, because sanitizer
instrumentation makes GCC's range and initialization analyses fire inside
libstdc++ on correct code; the plain build is the one that enforces -Werror. The lane refuses to configure with the CUDA
backend on, because a host sanitizer runtime does not instrument nvcc device
translation units and reports false positives against the CUDA driver; the CUDA
tier's equivalent is compute-sanitizer, and VT_POOL_BYPASS=1 makes the device
scratch pool hand out exact-size, really-freed allocations so that tool can see
tensor boundaries and use-after-free the caching pool otherwise hides.
CI concurrency is split by GATE SCOPE, not applied workflow-wide. The two
DIFF-scoped gates, documentation-checkpoint and commit-protocol-tag,
evaluate github.event.before..github.sha, so each owns a commit range that no
later run re-covers (the next run's before is this run's sha). They
deliberately carry no concurrency group and are never cancelled. The six
TREE-scoped jobs (agent-record, cuda-arch-features, device-leakage,
build-test-cpu, and both sanitize-cpu lanes) validate HEAD, so only the
newest push to a ref is meaningful; each carries a per-job group keyed on
github.ref with cancel-in-progress. Net effect: a superseded push to main
keeps its 2 per-commit gates and drops the other 6 jobs. The sanitize-cpu
group includes matrix.lane, without which the two mutually exclusive
sanitizer legs of the SAME run would share a group and cancel each other. At
the workflow level the group is keyed on the SHA for a push, so two pushes to
main never share a group; only pull-request pushes are deduplicated there.
Known failing on discrete sm_120 (RTX 5070 Ti), pre-existing and not introduced
by the lanes: test_cuda_ops "CUDA matmul (cuBLASLt) matches CPU on odd sizes"
fails 11 of 442 elements on the bf16-in/bf16-out 17x31x13 case. Verified present
on a pristine build of the same commit. The project's development GPU is GB10 /
sm_121a, so this is an unattributed per-architecture numerics difference, not a
regression.
Local Qwen3.5-4B plain BF16 direct loader, speed-pending: revalidated on
current main (7f620e74) on an RTX 5070 Ti at 6600.66 total tok/s versus vLLM
0.24.0 at 6722.24 tok/s (0.9819x) on the identical 128-request workload. Direct
loading reduces peak PSS from 8.59 to 2.28 GiB, stable PSS from 8.59 to 0.76
GiB, and mean TTFT from 835.0 to 729.2 ms. Mean TPOT/ITL is 38.22 ms versus
vLLM's 33.53 ms. Every token matches the previous series exactly (128/128 per
repetition, both arms), so the 109-commit upstream advance moved no output here.
The failing TPOT axis has a corrected diagnosis as of the W4 work below: the per-step synchronization it was blamed on is the synchronous engine loop waiting for its own sampling, about one per step, not a removable defect in the async sampler. Closing it means running the async engine loop, which is what the (opt-in) device-resident sampled-token mirror now makes legal on a discrete GPU; the measurement that would bind it is a serving A/B and is pending a harness this host can run.
Two things to know about the numbers. First, the earlier 5769.99 tok/s figure is VOID, not superseded by an improvement: that whole series ran against a GPU held at 11-13% utilization by a graphics consumer the harness's idle check could not see, and both arms gained ~14% once measured on a genuinely idle box. The check now fails on non-idle utilization. Second, the residual is a specific missing mechanism rather than a diffuse gap: this GPU is discrete, so the async-scheduling device combine/scatter falls back to a host path that must synchronize the main stream for the sampled ids, and the depth-2 scheduler therefore has nothing to overlap. Upstream keeps that state GPU-resident unconditionally. Scoped as ENG-ASYNC-SCHED W4.
This local 4B diagnostic does not establish 27B/35B support. Exact evidence and reproduction: Qwen3.5-4B post-pull revalidation.
There is no front-page race clip yet; when one is produced it will follow the LocalAI house style (side-by-side, identical output, honest measured ratios).
CUDA architectures. The runtime-gated production arch is GB10 sm_121a
(every gate model, every benchmark). A build-supported cross-family fan-out
(sm_80/86/87/89, sm_90a, sm_100a/103a, sm_110) compiles single-arch,
portable-kernels-only (all fp8/fp4/CUTLASS/Marlin/FA2 fast paths resolve EMPTY).
As of 2026-07-27, sm_110 is RUNTIME-VERIFIED (portable bf16 path) on real
silicon — the FIRST non-GB10 runtime proof. vllm.cpp was built natively for
sm_110 on an NVIDIA Jetson Thor board (aarch64, JetPack R38, nvcc 13.0,
on-box compute_cap=11.0; cutlass absent and not needed), Release -Werror
0 warnings, 16 TUs of real sm_110 SASS. It then ran the Llama-3.2-1B
paged-engine greedy gate and was STRICT token-exact 12/16 prompts (192/192
tokens) vs the committed vLLM oracle golden (every vLLM-deterministic prompt),
15/16 bit-identical to the GB10 sm_121a anchor; the remaining 4/16 are the
ratified bf16 near-tie prompts (committed teacher-forced gap 0.000 nats). Scope
is precise: RUNTIME-VERIFIED covers only the portable bf16 path that actually ran;
the fp8/fp4/CUTLASS fast paths on sm_110 remain DERIVED/NOT-YET (a cutlass-backed
kernel campaign). The other fan-out boards remain build-supported only (no board
here). Repro and evidence: docs/BENCHMARKS.md,
.agents/backend-matrix.md BACKEND-CUDA-SM110.
GDN Triton-AOT cubins are now vendored per-arch (2026-07-28). The vendored
Triton-AOT GDN fast-path cubins (the measured codegen-win packed decode plus the
delta_h/chunk_o FLA kernels for the Qwen3.6 GDN-hybrid models) previously existed
for GB10 sm_121a ONLY. Because a cubin loads only on the SM it was compiled for
and the cross-family arch builds ship -DVLLM_CPP_TRITON=OFF, GDN decode on the
other arches ran the slower spilling hand kernel. The full GDN AOT set is now
regenerated and vendored for sm_80/86/89/90a/100a (cuobjdump shows real
per-target SASS), so a -DVLLM_CPP_TRITON=ON single-arch build on those arches
selects the non-spilling path too — DERIVED+BUILD-VERIFIED (testing-welcome):
no non-sm_121 board runs a GDN model here, so this is not a runtime
GDN-decode-parity claim on any arch. sm_121a is byte-unchanged (SACRED gate
intact). Evidence: .agents/specs/triton-aot-per-arch.md.
As of 2026-07-28, sm_87 is also RUNTIME-VERIFIED (portable bf16 SYNC path) on
real silicon — the SECOND non-GB10 runtime proof. vllm.cpp was built
portable-only for sm_87 on an NVIDIA Jetson AGX Orin (Tegra R36.4.3 / JetPack 6,
aarch64; the integrated GPU self-reports sm_87, unified memory). Because the
Orin driver advertises only CUDA 12.6, the CUDA-13 container is refused by the
NVIDIA container runtime, so the build used the JetPack-6 l4t-jetpack:r36.4.0
image (nvcc 12.6) with g++-13 (the tree needs GCC ≥ 13). All fp8/fp4/CUTLASS/FA2
fast paths resolve EMPTY on sm_87 (Ampere: bf16 + int8, no fp8/fp4; cutlass not
present). It ran the Llama-3.2-1B paged-engine greedy gate and was 13/16 prompts
STRICT token-exact vs the committed vLLM 0.25.0 oracle golden, 16/16 under the
near-tie distributional gate, 0 forward-divergent (exceeding Thor's 12/16), plus
test_cuda_backend/test_cuda_ops (461 assertions of real on-device kernel
execution). One honest sm_87 bug surfaced: the DEFAULT asynchronous runner path
crashes on the first forward with an illegal memory access, so RUNTIME-VERIFIED is
scoped to the portable bf16 synchronous path (VT_ASYNC_RUNNER=0) — the async
runner on sm_87 is a tracked unblock item. Repro and evidence:
docs/BENCHMARKS.md,
.agents/backend-matrix.md BACKEND-CUDA-SM087.
As of 2026-07-28, the datacenter-Blackwell sm_100a fan-out gained its first
FAST-PATH body: the NVFP4 tcgen05 block-scaled GEMM is BUILD-VERIFIED
(DERIVED+BUILD-VERIFIED, testing-welcome). A faithful 1:1 type-port of vLLM's
Fp4GemmSm100 (ArchTag=Sm100 + KernelScheduleAuto — CUTLASS 4.5.0 selects the
5th-gen tcgen05 collective) compiles single-arch 100a on nvcc 13.0 -Werror-equiv
0 warnings and cuobjdump shows a real sm_100a cubin; it is gated by its own
cutlass-nvfp4-sm100 feature cell (enabled only for 100a, so the GB10 sm_121a
gate build is byte-unchanged). The native consumer mma.sync kind::mxf4nvf4 fp4
path stays sm_12x-only (it does not port to sm_100's tcgen05). No B200/sm_100
board ran it — a green compile + SASS is not execution evidence, not runtime
support, and not vLLM-competitive; the other sm_100 fast paths (CUTLASS C3x FP8,
MoE, MXFP4, MLA) remain scoped.
As of 2026-07-28, the Hopper sm_90a fan-out gained its first FAST-PATH body: the
CUTLASS C3x FP8 (W8A8) scaled-mm GEMM is BUILD-VERIFIED
(DERIVED+BUILD-VERIFIED, testing-welcome). A faithful 1:1 type-port of vLLM's
cutlass_3x_gemm_sm90_fp8 (ArchTag=Sm90 + KernelTmaWarpSpecialized*FP8FastAccum
— CUTLASS 4.5.0 emits the 4th-gen wgmma/TMA warp-specialized collective) compiles
single-arch 90a on nvcc 13.0 -Werror 0 warnings and cuobjdump shows a real
sm_90a cubin (ptxas even names wgmma.mma_async in the emitted kernels); it is
gated by its own scaledmm-c3x-sm90 feature cell (enabled only for 90a, so the
GB10 sm_121a gate build is byte-unchanged). The consumer sm_12x FP8 scaled-mm
body (ArchTag=Sm120) is untouched. No H100/H200/sm_90 board ran it — a green
compile + SASS is not execution evidence, not runtime support, and not
vLLM-competitive; the sm90 int8/blockwise C3x legs and the other Hopper fast paths
(FA3, Machete, CUTLASS MoE) remain scoped.
As of 2026-07-28, the datacenter-Blackwell sm_100a fan-out gained its second
FAST-PATH body (after the DC1 NVFP4 tcgen05 GEMM): the CUTLASS C3x FP8 (W8A8)
scaled-mm tcgen05 GEMM is BUILD-VERIFIED (DERIVED+BUILD-VERIFIED, testing-welcome) — the intersection of DC1's tcgen05 arch and the Hopper DC2 C3x
fp8 kernel. A faithful 1:1 type-port of vLLM's cutlass_3x_gemm_sm100_fp8
(ArchTag=Sm100 + KernelScheduleAuto/EpilogueScheduleAuto — CUTLASS 4.5.0
selects the 5th-gen tcgen05 warp-specialized collective, a 2SM
ClusterShape<_2,_2,_1> default) compiles single-arch 100a on nvcc 13.0
-Werror all-warnings 0 warnings and cuobjdump shows a real sm_100a cubin (the
emitted kernels are Sm100TmaUmmaWarpSpecialized naming SM100_MMA_F8F6F4,
SM100_TMA_2SM_LOAD and SM100_TMEM_LOAD; SASS carries LDTM/tmem
tensor-memory ops); it is gated by its own scaledmm-c3x-sm100 feature cell
(enabled only for 100a, so the GB10 sm_121a gate build is byte-unchanged). The
consumer sm_12x FP8 body (ArchTag=Sm120) and the DC1 sm100 NVFP4 leg are
untouched. No B200/sm_100 board ran it — a green compile + SASS is not
execution evidence, not runtime support, and not vLLM-competitive; the sm100
int8/blockwise C3x legs and the other datacenter fast paths (CUTLASS MoE Sm100,
MXFP4, CUTLASS MLA) remain scoped.
A separate beyond-vLLM breadth lane targets the older NVIDIA arches vLLM
DROPS but llama.cpp still runs (Pascal/Volta/Turing). Its first brick landed
2026-07-28: Turing sm_75 is BUILD-VERIFIED. The bf16-WMMA prefill kernels
(bf16 tensor-core fragments are Ampere+ only) are now guarded
#if __CUDA_ARCH__ >= 800, so a single-arch 75 build compiles -Werror 0
warnings on nvcc 13.0 and cuobjdump shows real sm_75 cubins — the TU falls
back to the existing scalar CUDA-core attention path (the llama.cpp fp16 fast
body is a later brick). On sm_80+ the guard is a no-op (preprocessor identity),
so the GB10 gate build is byte-identical. This is
DERIVED+BUILD-VERIFIED (testing-welcome), NOT runtime — no Turing board ran
it; there is no vLLM oracle on Turing (vLLM dropped it), so correctness is
referenced against llama.cpp-on-card plus a portable cross-check. Volta/Pascal
need a CUDA <13 toolkit (nvcc 13 rejects sm_70/sm_60/sm_61). See
.agents/backend-matrix.md BACKEND-CUDA-SM075.
Metal (Apple Silicon), indicative, not binding. Two models run end to end
and pass correctness (OPT-125m, Qwen3-0.6B); 18 of 75 ops are native, the rest
fall back to CPU on unified memory. Kernel work including mma prefill attention
(4.3x), a vectorised decode V accumulation (+1.66%) and vectorised prefill
attention staging (+0.50%) puts warm b=1 throughput at about 95.9% of MLX-LM by default, and 97.6% with the optional MLX GEMM provider
shape-gated to prefill (-DVLLM_CPP_MLX=ON), which is a numerics-deviating
configuration. Both figures are against an interleaved 6-run MLX-LM baseline;
an earlier 99.1% claim used a 2-run baseline containing an outlier. With MLX
gated, prefill is 1.5% AHEAD of MLX-LM and the entire remaining deficit is
decode. The GEMV runs at 83% of memory peak; decode attention runs 2.6x slower
per byte than the GEMV because at b=1 it gets 8x fewer threads in flight, and
three separate attempts to exploit that (flash-decoding, GQA grouping, wider
threadgroups) have each measured worse. Bisecting the GEMM puts it at 97% of MLX's own
(mma issue rate 3.91 TFLOP/s), so the residual is NOT GEMM-led: it is spread
across prefill attention, small prefill kernels and decode, none dominant. The
decode share is 93% weight streaming at 83% of the part's memory peak, so the
remaining items sit against roofs rather than against defects: GEMV at 83% of
memory peak, GEMM at 97% of MLX's, and a re-test of GEMV memory-level
parallelism under the paired harness confirmed the earlier exclusion at -1.03%.
The decode residual is non-GEMV overhead (2.7 ms/token against MLX's implied
2.0), so the next lever is a fused qk-norm-RoPE attention
preamble. The dense recipe (kAttnQkNormRope) is composite-only on every backend
(fast_op = kNoFastOp), so this meant writing the kernel rather than porting one.
That kernel now exists and is validated on Metal (worst element error 1.4e-06 vs
the composite); the default bf16 path is now routed through it
(gated on the rope cache), worth a measured +0.35% and taking warm b=1 throughput
to about 96.0-96.4% of MLX-LM, with decode at 98.0%. The largest remaining lead
is prefill attention, which at 547 GFLOP/s is still 5.2x slower per FLOP than
this device's own GEMM; its online softmax now runs one simdgroup per query
row rather than one thread, worth -16% on the kernel and +0.19% end to end. Its
remaining 4.4x FLOP-rate gap to this device's own GEMM has no open structural
idea left: BK deepening is blocked by a bf16-P precision constraint and eliding
the identity softmax rescale measured no gain. Prefill's
small-kernel bucket is ~36 ms against MLX-LM's implied ~25, so improving our own
architecture there caps out near 97% rather than parity. A spike comparing MLX's
lazy-graph execution model against ours found our matmul chain already +1.5%
ahead of MLX's own ceiling, so adopting MLX's execution model would not help
decode; the residual is small-op fusion, which is ours to do. The V-accumulation win came from BISECTING the attention
kernel once a paired ABBA harness with cooldown made 0.2% differences
measurable: the V loop moved 2 bytes per lane per load where the score loop
moved 8, giving 29 GB/s against 64 on identical traffic. That also explains why
five earlier decode hypotheses (fusion, split-K, KV layout, threadgroup width,
barrier depth) were each refuted — the limiter was bytes per instruction in one
loop, not parallelism or layout. An optional MLX GEMM provider is available via
-DVLLM_CPP_MLX=ON and currently measures net slower than our own kernels. Full
per-lever chronology: docs/BENCHMARKS.md and
.agents/specs/metal-dispatch-attribution.md.
CUDA architectures. The production target is GB10/sm_121a (runtime-gated,
both gate models token-exact + at/above vLLM speed). The arch-additivity
framework makes adding a CUDA architecture a data edit; sm_120a, sm_90a,
Ampere/Ada sm_80/86/87/89, datacenter Blackwell sm_100a/103a and sm_110 are
BUILD-supported, portable-kernels-only (they compile -Werror-clean and emit
real per-arch SASS on dgx, but no such board runs here, so none is runtime
support). The Ampere major-8 fast-path bring-up is now SPIKED as a
derive-and-ship campaign
(.agents/specs/cuda-arch-ampere-fastpath.md):
the fast-path kernel bodies vLLM builds for Ampere (FA2, Marlin int4 W4A16 +
fp8-input, AllSpark W8A16, CUTLASS scaled-mm C2x) will be ported 1:1, build-verified,
and shipped labeled DERIVED+BUILD-VERIFIED (testing-welcome) — a build is never
a runtime claim. AGX Orin (sm_87) is the one reachable non-GB10 board and the
sole target that reaches RUNTIME-VERIFIED, gating the first non-GB10 runtime proof
token-exact vs the vLLM 0.25.0 oracle. FA2 Ampere (WA-1) has LANDED
DERIVED+BUILD-VERIFIED (testing-welcome) — the ROAD-V1-D1-CUDA first brick: the
fa2 FEATURE-TABLE cell was widened to 8.0,8.6,8.7,8.9,12.0a,12.1a (the vendored
kernels are already the _sm80.cu FlashAttention-2 bodies, __CUDA_ARCH__>=800)
and the FA2 build gate decoupled from the sm_12x NVFP4 flag onto arch-independent
CUTLASS-header availability; single-arch 87 and 80 dgx builds are -Werror
0-warn with cuobjdump-proven real sm_87/sm_80 FA2 cubins (86/89 inherit
same-major), and sm_121a is unchanged (still resolves 121a, NVFP4 GEMM still
ON) — verified by a re-run gate. NO Ampere board ran it here. AllSpark and
scaled-mm C2x are new bodies; Marlin needs its int4/fp8 instantiations. The whole CUDA-arch
fast-path + beyond-vLLM breadth effort is tracked on the roadmap as
ROAD-V1-D1-CUDA (derive-and-ship + a public tested/untested signal matrix), with
AGX Orin (sm_87) and NVIDIA Thor (Blackwell) as the reachable runtime-gate boards.
Datacenter CUDA arches (Hopper sm_90a, Blackwell sm_100/103/110) — build-only, fast-path SPIKED. These arches compile today as portable-kernels-only (build-only; every wgmma/tcgen05 fast-path FEATURE-TABLE cell resolves EMPTY, no board here) — see .agents/backend-matrix.md. The FRAMEWORK (arch-additivity seams) is done; the FAST-PATH kernel bodies are now scoped for derive-and-ship (the llama.cpp model) in .agents/specs/cuda-arch-datacenter-fastpath.md. The CUTLASS C3x / NVFP4-tcgen05 / grouped-MoE / MLA kernels are CUTLASS template instantiations (the wgmma/tcgen05 MMA lives in cutlass 4.5.0, selected by ArchTag), so they are 1:1-portable and BUILD-VERIFIABLE here via a single-arch 90a/100a cross-compile + cuobjdump SASS proof, with NO Hopper/B200 board. Such kernels will ship LABELED DERIVED+BUILD-VERIFIED (testing-welcome) — a faithful port with a compile+SASS proof, HONESTLY untested (a green link is not execution evidence); cloud-GPU token-exact + every-axis runs upgrade the label to runtime-verified. DeepGEMM (runtime JIT/autotune) is the one path that is not-yet-buildable / needs-real-port and ships no fake.
CUDA arch breadth beyond vLLM (Pascal / Volta / Turing). These
are older NVIDIA arches vLLM DROPS but llama.cpp still supports, so this is a
"support more than vLLM" lane with llama.cpp as both the kernel source and the
competitor floor. Our portable attention uses bf16 tensor-core WMMA, which does
not exist before Ampere, so these arches need the bf16-WMMA path compile-guarded
(and, for a fast path, a new fp16/non-tensor-core kernel body ported 1:1 from
llama.cpp's fattn-tile/fattn-vec). The lane is derive-and-ship, and its first
brick has LANDED: Turing (sm_75, e.g. the cloud-common T4) is BUILD-VERIFIED
(2026-07-28, CLAIM-CUDA-TURING-SM75). The bf16-WMMA prefill kernels are now
guarded #if __CUDA_ARCH__ >= 800, so a single-arch 75 build compiles -Werror
0-warn on our current CUDA 13 toolkit and cuobjdump shows real sm_75 cubins —
the TU falls back to the existing scalar CUDA-core attention path (the fp16 fast
body is a separate later brick). On sm_80+ the guard is a no-op (byte-identical
GB10 build). It ships labeled "derived, build-verified, not hardware-tested here,
community testing welcome" — no Turing board ran it; a green compile + SASS is not
execution evidence. Volta (sm_70, V100) and Pascal (sm_60/sm_61, P100/P40)
are not-yet-buildable because CUDA 13 dropped their code generation and no
CUDA 12.x toolkit is provisioned here. There is no vLLM oracle on these cards
(vLLM will not run there), so a real correctness test uses llama.cpp on the same
card plus a newer-card/CPU cross-check; nothing is runtime-verified yet.
- Automatic prefix caching (APC) is on by default for dense models (hybrid /
GDN and attention-free default off, mirroring vLLM), and it now has an
end-to-end cache-ON gate on
Qwen/Qwen3-4B(a shared common prefix reused across requests). A cache hit never changes the output: cache-ON and cache-OFF produce token-identical greedy decodes (differing only where the bf16 model itself is a genuine near-tie, exactly as vLLM's own greedy does), and both match vLLM's own greedy under teacher-forcing. The prefix cache measurably hits (0.807 hit rate on that workload) and pays off: on a cached request the time to first token drops about 1.76x versus the cold first request (70 ms to 40 ms on the GB10), because the shared prefix's attention is not recomputed. Hit-rate statistics are counted per vLLM's own counters. Block-hash extra keys (multimodal hash, LoRA adapter name,cache_salt) are folded into the block hash exactly as vLLM does, so requests differing only in those never share a cache block (proven no-false-share). Thecache_saltOpenAI request field is not surfaced yet. A Prometheus/metricsendpoint is available (opt-in): it exposes vLLM'svllm:*metric names, help text, types, buckets and{model_name, engine}label schema in text format 0.0.4, gated against vLLM's own scrape spec. It now serves live per-step values, not just the schema: each engine step folds the scheduler snapshot (running/waiting counts, KV-cache usage, prefix-cache queries and hits) and the iteration stats (prompt and generation token counts, request-success counts, and the time-to-first-token, inter-token-latency and end-to-end latency histograms) into the registry, matching vLLM's own mapping. The scheduler also records per-request QUEUED/SCHEDULED/PREEMPTED engine-core events (1:1 with vLLM, behind the same stats gate), which the frontend folds into the per-request queue, prefill and inference timing histograms and the preemption counter, so those series now carry real durations rather than staying at zero. A behavioural CPU gate drives the reference engine for several steps and checks the values track the run. The remaining work is the async production-serving path wiring, the chat/completion response-body timing surface, and the config-gated metric families (speculative decoding, KV connector, multimodal cache, LoRA). matching vLLM's own mapping. A behavioural CPU gate drives the reference engine for several steps and checks the values track the run. The remaining work is the async production-serving path wiring and the config-gated metric families (speculative decoding, KV connector, multimodal cache, LoRA). - SGLang RadixAttention behavior parity — scoped 2026-07-27, W1+W2 flags now
IMPLEMENTED (CPU-gated). SGLang's radix-tree prefix cache is functionally
equivalent to our block-hash APC (both do automatic longest-prefix KV sharing
with LRU eviction; the only delta is token/page vs block sharing granularity,
which is bounded and never changes the output), so RadixAttention is fused into
APC and
--enable-radix-attention/--disable-radix-attentionare aliases for the existing prefix-cache toggle (also reachable via the C-ABI tri-statevllm_model_params.enable_prefix_caching, ABI v7), not a second cache. SGLang's overlap scheduler is likewise already our async scheduler. The one genuinely-distinct SGLang behavior — cache-aware admission that reorders the waiting queue by longest prefix match — is implemented behind--schedule-policy=lpm(SchedulerPolicy::kLPM, output-neutral: it changes only which request is admitted first, never any request's tokens; default staysfcfs). SGLang's in-batch prefix-collision de-prioritization (SW2) is now also implemented inside thelpmpath (a later request that would collide on the same not-yet-cached prefix as an earlier waiting request is sorted behind the non-colliding requests), still output-neutral. Note: this de-prioritization changes only admission order here — our prefix cache already caches a request's blocks at allocation time, so a second same-step request sharing an uncached prefix already reuses the first's blocks (the redundant-prefill this avoids in SGLang, whose cache updates only after the forward pass, does not occur for us). Jump-forward constrained decoding (SW3) now has its SAFE subset implemented: an opt-in driver (envVT_ENABLE_JUMP_FORWARD, default off) emits a grammar-forced token WITHOUT a model step ONLY where the grammar leaves exactly one valid token at a non-accepting state — provably byte-identical to normal per-token constrained decode (the constrained sampler has a single finite-logit token there), so no re-tokenization is needed. The general case (a forced byte run with several possible tokenizations, which SGLang handles by re-tokenizing and rolling back the boundary token) is deliberately not jumped, and wiring the driver into the live decode loop (jumped tokens need their KV recomputed) is a named residual — so the flag stays off by default. The radix eviction-policy knob (SW4) stays a named/deferred opt-in. See.agents/specs/sglang-radixattention.md. - SGLang-compatible knobs are now first-class, documented toggles (2026-07-28,
CLAIM-SGLANG-ABI-DOCS, reconciled to ABI v10). LPM scheduling and jump-forward are exposed on the C++ library API AND the C ABI (not env-var/internal-only), joining the already-exposed prefix-caching/RadixAttention and custom-logits-processor knobs. LPM is selected through the ABI v9 string fieldvllm_model_params.scheduling_policy("fcfs"/"priority"/"lpm"); jump-forward through the new ABI v10 tri-state fieldvllm_model_params.enable_jump_forward(0=default/1=on/2=off). The server has--scheduling-policy lpmand--enable-jump-forward/--disable-jump-forward. All default to today's behavior — an all-zero config is byte-identical to before ABI v10, andvllm_abi_version()returns 10.VT_ENABLE_JUMP_FORWARDis retained as an env override (wins when set). The user-facing enablement guide for all four behaviors isdocs/SGLANG-COMPAT.md(engineering record:.agents/specs/sglang-enablement.md). - KV persistence to disk / CPU offload is built (CPU and disk tiers,
identity-checked blocks, a size-budgeted disk tier) and wired opt-in into the
scheduler through an abstract
KVConnectorABI selected by aKVTransferConfig(a disk-offload connector and the LMCachelm://client, over the same seam), off by default. The LMCachelm://client interoperates with a real vLLM+LMCache peer (its cache keys are byte-for-byte identical to LMCache's ownChunkedTokenDatabase, chunk 256, vLLM'ssha256_cbor, and a KV prefix a real peer wrote is found and loaded back byte-identically over the wire), and it is now proven END TO END inside a real generation loop: with the connector ON against a livelmcache.v1.server, an OPT-125m run loads a previously stored prefix's KV and shortcuts prefill, generating tokens BIT-IDENTICAL to the connector-OFF cold run (verified both after an in-process restart and from a genuinely cold second process). The worker-side KV store/load is implemented for the LMCache connector only: the CPU/diskOffloadingConnectorimplements the scheduler half alone (it shortcuts prefill for matched blocks, but nothing moves their bytes into the KV pages), so the engine now REFUSES to wire it at construction, on every device, naming the connector and the device, rather than serving output computed over KV that was never written. Both connectors are selectable from the server with--kv-transfer-config '<json>', mirroring vLLM's own flag and JSON shape. The disk connector's worker half remains unimplemented and is the named remaining step. Usage guide: docs/KV-OFFLOAD.md. - KV-cache events (the stream external KV routers and prefix-cache-aware load
balancers subscribe to) are built and off by default. When enabled, the
prefix-cache block pool emits
BlockStored,BlockRemoved, andAllBlocksClearedevents at the store, evict, and reset sites, and they serialize to amsgpackpayload that is byte-for-byte identical to vLLM'smsgspecencoding (checked against captures for both the default int-truncated block-hash form and the raw-bytes form). Event generation and the wire payload go through an abstract publisher seam that mirrors vLLM's--kv-events-config. The live ZMQ socket transport (the actual publish to an external router, with its replay buffer and per-rank ports) is not wired yet, so today the events can be generated and encoded but not yet streamed over a socket; the seam is built so a real ZMQ publisher drops in without touching the emission sites or the encoder. Because it is off by default, the prefix-cache path is byte-identical to before. - Tool calling covers the dialect set in the capability table above,
including the Qwen3-Coder XML parser (
qwen3_coder/qwen3_xml/mimo). Selection over HTTP is--tool-call-parserand--reasoning-parser, mirroring vLLM's own flag names, alongside the C ABI'svllm_model_params.tool_parser/.reasoning_parser. Both flags keep the server's previous behaviour as their default (hermes, reasoning off), takeautoto opt into the same chat-template detection the C ABI uses, takenoneto disable, and abort startup listing every registered name if given an unknown one. - Engine configuration through the C ABI reached parity with the bundled
server's flags at ABI v9 (2026-07-28).
vllm_model_paramsnow carriesmax_num_batched_tokens(chunked-prefill token budget),scheduling_policy(fcfs/priority/lpm), andkv_transfer_config(the external KV connector / LMCache JSON, connector name validated againstKVConnectorFactoryat load), and honourstokenizer_config_path, which had been declared since v1 and never read. Every addition is inert at its default, so a v8 caller that zero-fills the struct growth gets the byte-identical pre-v9 engine. Malformedspeculative_config/kv_transfer_configdocuments reportVLLM_ERR_INVALID_ARGUMENTrather thanVLLM_ERR_MODEL_LOAD(the contractvllm.hhas documented since v6). Driver: an embedder (the LocalAI vllm-cpp backend) could not expose LMCache or the prefill budget in a model config. - The stable C ABI contract is now spiked and current at ABI v10
(
SERVE-C-ABI,CLAIM-SERVE-C-ABI-SPIKE). The accepted contract inventories ownership, no-throw/error mapping, concurrency, versioning and packaging gates. The row remainsANCHOR-BACKFILL: the next CPU-only bricks are all-19-symboldlsymcoverage, historical-layout compatibility fixtures, allocation-failure no-throw proof, and lifetime sanitizer stress. - DFlash speculative decoding from a GGUF DRAFT WORKS end to end on the GB10
release target (
SPEC-DFLASH-GGUF,DONEfor correctness on both axes, speedPENDINGby design, spike). A DFlash draft packaged as a singledflash-arch GGUF now loads AND generates against a safetensors target (axis A): config from thedflash.*KVs, weights through a resolver that reuses the existing safetensorsLoadQwen3DFlashbody, a.ggufbranch in the draft-path resolution, and the target-shared embedding and lm_head still taken from the target, as the GGUF contract intends. Binding result, MEASURED 2026-07-29 on a production-configured build (CUTLASS + FlashAttention-2 + vendored Triton-AOT, sm_121a, proven three ways; earlier numbers came from a build carrying none of those and are void): on the Qwen3.6-27B NVFP4 safetensors target, a DFlash draft read from an UNQUANTIZEDBF16GGUF is the same draft as the z-lab safetensors one down to the count - token-for-token identical greedy continuations AND exactly equal accept counts, 47/96 at 48 tokens (reproduced) and 27/64 and 15/144 at 24 tokens on two prompts, k=16, concurrency 1. The publishedQ4_K_Mdraft emits the same tokens and costs at most ONE acceptance (46/112 against 47/96 at 48 tokens, zero delta at 24). That one acceptance is ordinary 4-bit cost, established rather than assumed. It stood RED against an exact bar until 2026-07-29, and what closed it was finding out what the bar was measuring, not relaxing it. Loaded through the SAME code path theBF16GGUF is BYTE-IDENTICAL to the safetensors draft on all 58 tensors (a CPU gate,tests/vllm/models/test_qwen3_dflash_gguf.cpp, 302/302, functionally RED against theQ4_K_Mfile); our Q4_K/Q6_K dequant is bit-equal togguf-py's reference on the real tensors; the published ladder's weight error is monotone and uniform (BF16 0, Q8_0 5.6e-3, Q6_K 1.85e-2, Q5_K 3.85e-2, Q4_K_M 7.6e-2 mean relative, no outlier tensor); the only numeric config difference isrms_norm_epsat 2.5e-9 relative. The end-to-end run then closed it: restoring the draft's precision and nothing else moves the count back to exactly 47/96. So the accept-count bar is now EXACT across formats and BANDED across quantization levels (|d_accepted| <= 2,|d_proposed| <= k*2, both derived from the measurement and mutation-proved non-vacuous at band 0), with the arm selected by reading the draft file's ggml types rather than by a flag. Three conventions this had to undo, all invisible to shape and name checks:dflash.target_layersis stored+1-offset; the draft's RMSNorm weights are RAW (its converter class does not inherit the Qwen3Next+1shift, so unlike the trunk and the MTP head they must NOT be un-shifted, which points the opposite way); andvocab_size, which the GGUF contract deliberately omits, must be back-filled from the target or the draft's shared embedding table is empty and the first proposal throws. Only the last survived the load-level tests, and only generating found it. Axis B (the TARGET is a GGUF too) now WORKS as well.SharedHeadSourcereplaces the safetensors-typed parameter that used to make the draft's shared embedding and lm_head unavailable from a GGUF - which is why the GGUF path refused dflash outright - and serves both tensors out of either container. On the Qwen3.6-27B NVFP4 GGUF target with the sameQ4_K_MGGUF draft, the DFlash-ON continuation is token-for-token identical to that same target's spec-OFF and the drafter is alive at 14/160 accepted. Re-measured 2026-07-29 on the production build and BROADENED from one prompt to three, strict-form green on all of them: "The capital of France is" IDENTICAL with acceptance 14/160, "Write a Python function that reverses a string:" IDENTICAL with 24/64, "Photosynthesis is the process by which" IDENTICAL with 15/128, all exit 0. Two things are worth stating plainly rather than rounding off. The spike's highest-ranked risk - the draft scoring with a head dequantized out of a quantized target - does not exist on this asset: the GGUF storestoken_embd/outputas BF16 beside its NVFP4 body, byte-identical to the safetensors sibling of the same quantization run, so the read is verbatim. And acceptance against the safetensors target is PROMPT-DEPENDENT rather than uniformly lower: 14/160 against 15/144 on the first prompt, but EQUAL at 24/64 with IDENTICAL DFlash-ON streams on the second. That belongs to the GGUF target's compute path, not to the shared head, because there is no NVFP4 GGUF GEMM yet, so the GGUF target expands to bf16 while the safetensors target runs the true W4A4 kernels; the two containers' own spec-OFF streams diverge at index 4 on the first prompt and index 16 on the second, with no speculation anywhere, and the DFlash-ON pair tracks that. Operational note for anyone re-running the axis-B arm: it peaks at ~81 GiB RSS against GB10's 119 GiB unified pool, so a run launched on top of a large page cache can fail withvt cuda: embedding: out of memory. That is contention, not a result; on a settled box the arm reproduces its numbers exactly. Both axes are now gated and closed (gates 1-5 and 7 of the spike), so the row movedPARTIAL->DONEfor correctness in closing commitc62f2fa3. What remains open is deliberately NOT this row's: no speed A/B is owed until a native NVFP4 GGUF GEMM exists to make the two target containers comparable, anddocs/BENCHMARKS.mdcarries that as an explicitPENDINGwith its reproduction command. - MTP speculative decoding from a GGUF target WORKS, on CPU and on the GB10
release target (
SPEC-MTP-GGUF,DONE, spike). The head loads from a head-carrying GGUF -HfConfigFromGgufrepublishes the depth it already read from<arch>.nextn_predict_layers, andLoadQwen3_5MTPFromGgufreads thenextnblock with the trunk loader's own helpers, so it inherits the GGUF (w+1) norm storage, the quantization routing and the torch [N, K] shape order. A GGUF exported without the head is refused, naming that as the reason. Both correctness gates pass: spec-ON output is token-identical to spec-OFF on CPU against a real llama.cpp-converted Qwen3.5-2B, and on the GPU against the 35B A3B NVFP4 GGUF with 13 drafts proposed / 11 accepted. RE-ANCHORED 2026-07-29 to a production-configured build (CUTLASS + FlashAttention-2 + vendored Triton-AOT, sm_121a; the original close-out ran without any of them, so its evidence was not trustworthy): dgx GB10 underflock, 3/3 cases, 10/10 assertions, exit 0, spec-ON token-identical to spec-OFF, 13 proposed / 11 accepted, unchanged, 90.26 GiB peak RSS, 7m13.59s. Acceptance parity is measured against the safetensors sibling of the same quantization run: 12 proposed / 11 accepted there vs 13 / 11 here. Cross-format token agreement is NOT APPLICABLE - it needs an F16/F32 GGUF sibling and every head-carrying export on hand is quantized. Throughput is the one open item and is deliberately PENDING, see docs/BENCHMARKS.md. Closing commitedf91449. Two things had held this row open and both are now settled by measurement. The build was believed to beCMAKE_CUDA_ARCHITECTURES=75on an sm_121 GB10: that came fromCMakeCache.txt, which records the CUDA compiler's own probe default and is shadowed by the normal variable the project sets, so it never described the compiled code. The real flags and every emitted cubin aresm_121a. And the GPU arm's 24 tokens differed from the CPU arm's: that was read as a near-tie, and the 2026-07-29 re-measurement shows it was an artifact of the defective build. There is no CPU-versus-GPU token delta on a production build: both devices emit the same 24 tokens, because with CUTLASS and the Triton-AOT GDN kernels in, the GPU takes the CPU's side of the 0.06-nat near-tie at position 1 (GPU rank111-0.763897 over rank213-0.824083; CPU rank111-0.765499 over rank213-0.830374, bit-identical to before, as expected for CUDA-only kernels). Zero exact ties in either arm, minimum margins 0.060 and 0.065 nats. The near-tie is real and unchanged in size; only which side the GPU landed on moved, and the earlier cascade narrative is retracted. - The 35B A3B NVFP4 SAFETENSORS target has EXACT logit ties, and its
speculative token identity does not hold there (open, surfaced by
SPEC-MTP-GGUF's gate, owned bySPEC-MTP). Pointed at the safetensors sibling of the same quantization run, the same concurrency-1 gate shows spec-ON inserting one token relative to spec-OFF, and two runs of spec-OFF itself disagreeing. Both divergences land exactly on positions where the top candidates carry BIT-IDENTICAL logprobs: that arm has three such exact ties in 24 tokens, because keeping NVFP4 packed puts its logits on a 1/16 grid where distinct tokens collide. The GGUF arm, which expands to bf16, has ZERO exact ties over the same 24 tokens (minimum margin 0.048 nats) and reproduces its sequence every run - so the GGUF result stands and this is a quantized-GEMM determinism question, not a speculative-decoding logic one. Not root-caused beyond the tie measurement (one prompt, two runs). - Speculative decoding on CPU corrupted the target's own state, and is FIXED
(
CPU-SPEC-DIVERGENCE).qwen3_5.cpp:3616sized the GDN state gather/scatter row by(Kw-1)while the speculative persistent row is(Kw-1)+num_spec, so the contiguous row helpers mis-strode the cache slot and every channel past the first, corrupting post-prefill recurrent state; speculation therefore changed the target's own greedy output even when no draft was accepted. Fixed with a stride-aware per-channel state copy used only when the two widths differ, so every non-speculative path stays byte-identical. CPU-only in effect - the fp16/bf16 cache arm routes through theGdnStateGather/Scatterops, so CUDA was never exposed and no GPU result is affected or retracted. It survived because there was no CPU spec-decode gate anywhere in the tree; there is one now. /healthreports process liveness rather than a full engine-health probe.- Speculative decoding ships user-facing via
--speculative-config(OpenAI server, example CLI, and C ABI v6), unset by default and byte-identical to the non-speculative engine when unset. MTP (k=1) on the Qwen3.5/3.6 gate checkpoints (those shipping anmtp.*head) is token-exact to both our own spec-off output and vLLM's own MTP greedy at concurrency 1, and about 1.04x faster than vLLM's spec-on decode; the concurrent (multi-request) path shares a scheduler step with ordinary prefills over a bit-exact GDN split-merge. Honest caveat: above concurrency 1 the 27B greedy output is not bit-stable across batch shapes, so exact spec-on/spec-off token agreement is a concurrency-1 property. Block-diffusion DFlash (methoddflash, the z-lab Qwen3.6-27B draft over the 27B NVFP4 target, wired through the same--speculative-configwith an external draftmodelpath) is correctness-complete (ratified near-tie) and at or above vLLM throughput at concurrency 1 (29.32 vs 29.24 tok/s, 1.003x), built D0 through D14 on the vLLM 0.26.0.dev0 stack (which resolves vllm#40898) and recorded DONE across the engine, model and kernel matrices. ngram (methodngram) is the draft-free suffix-match proposer, reusing the same verify/reject loop with no draft model; it is token-exact vs vLLM's own ngram on repetitive workloads (27B, 5/5 prompts, every draft accepted) and off by default (byte-identical). Measured A/Bs and the D0-D14 chronology: docs/BENCHMARKS.md; usage guide: docs/SPECULATIVE-DECODING.md. - LoRA and multi-GPU are not supported yet. Multimodal is
correctness-complete on a single-sequence eager path and the top roadmap
priority; the OpenAI-server wiring has advanced through its CPU bricks —
brick 1 (
CLAIM-MM-SERVING-W1) parses the OpenAI multimodal content-part array and routes the decoded media through the existing processors, and brick 2 (CLAIM-MM-SERVING-W2) carries the result into the engine (additiveadd_request(MultiModalInputs)/generateoverloads on both engines viaInputProcessor::process_inputs_mm, chat-template placeholder-string helpers, and a default-inert serving_chat multimodal seam), but the closing end-to-end GPU gate (MM-SERVE-E2E— the model tokenizer/processor seam body plus the mm forward) is still pending, so it is not yet servable end-to-end. Image-to-text is strict token-exact 32/32 vs the vLLM 0.25.0 oracle on both Qwen3-VL-4B and our own gate model Qwen3.6-27B (Qwen3_5ForConditionalGeneration); video-to-text works end-to-end on both (Qwen3-VL-4B near-tie-robust, Qwen3.6-27B strict 32/32); and audio-to-text works end-to-end on Voxtral-Mini-3B (near-tie-robust, decoder proven token-exact 48/48), the first audio understanding in the tree. The pipeline is a pure C++ port end-to-end: the Qwen3-VL image/video processors and the Whisper-class audio processor are bit/byte-identical to the oracle, the Qwen3-VL vision tower and the Whisper encoder tower are proven faithful within the bf16 envelope, and the MRoPE/DeepStack text-backbone contracts are unit-exact. Concurrency-1 speed is measured against vLLM 0.25.0's graphed config: on Qwen3.6-27B image the decode is at parity and the vision tower, after a warp-scoped online-softmax attention kernel plus a resident-weight load, is 148 ms/image (0.59x of vLLM's eager encode, faster) with the strict image/video gates still 32/32. On Voxtral audio the decode is now graph-captured and measured at 60.9 ms per token (vs 61.7 ms eager, a small real non-overlapping win), against vLLM's graphed 40.8 ms, so the audio gap is about 1.49x: the removable per-step launch overhead turned out to be only about 1.25% of the per-token time, so the residual is per-step compute (kernel) efficiency, not launch overhead. That residual is now fully attributed: nsys shows all of it is the naive scalar paged-attention decode kernel (723 us per call over 30 layers, about 120x the KV memory floor). vLLM runs decode attention on flash-attention varlen, which our tree already ships and enables by default for this head shape; Voxtral only missed it because the decode driver allocates one KV block of size 444, not a multiple of 16, which disqualifies the fast path. Rounding that to a multiple of 16 routes decode through the flash kernel and reaches 38.2 ms per token, a 36% win that beats vLLM's 40.8 ms, and the resulting sequence teacher-forces vLLM with zero divergences (a valid greedy branch). It is a bf16 near-tie ceiling: it changes which side of an exact tie is taken, so it no longer reproduces the committed near-tie golden byte-for-byte and is not adopted here; the byte-exact scalar path is kept at 14/14 with the golden unchanged, and the win is one line plus a golden regeneration away. On-GPU greedy sampling is in place on both mm decode loops. The open work is batched/graphed mm serving, adopting the audio decode win, and higher-concurrency throughput. Design and status: .agents/specs/multimodal-track.md, .agents/specs/audio-track.md, and .agents/specs/multimodal-speed.md. - Environment variables. A handful of behavior-changing knobs (CPU thread
count, prefix-cache hash seed, LMCache host/port, the default-on async and CUDA
fast paths and their rollbacks, the portable-reference bisect switch, GGUF
loading behavior, the Vulkan device selector) are documented in
docs/ENVIRONMENT.md. A CI check
(
scripts/check-env-doc.py) keeps that reference honest: any new production env var must be documented or explicitly classified kernel-internal.
Every model is gated token-for-token against the vLLM 0.25.0 oracle (the
gate-time pin, re-validated bit-identical on the current 0.26.0.dev0 pin) on the
same workload, and every change that could affect correctness or performance is
compared apples-to-apples against vLLM with both numbers and the ratio recorded.
Behavioral CPU tests run under CTest; CUDA correctness, sanitizer, trace, and
performance evidence is recorded per feature rather than inferred from source.
The protocol is in .agents/gates.md and
.agents/benchmark-protocol.md. A CI check
(scripts/check-fusion-consistency.py) additionally keeps model forwards routing
their fusable add+RMSNorm glue through the portable fusion catalog rather than
hand-fusing it. The Gemma family (Gemma-1/2/3) was MIGRATED off the drift
allowlist 2026-07-30 (Tier-B2 of the cross-arch fold plan): their
residual-carrying gemma-(1+w) input/pre-ffn/final norms now route through
vt::FusedChain(kFusedAddRmsNorm) behind FusedChainAdoptEnabled() (byte-exact
else fallback), and the sandwich post-norms (no residual add) correctly stay
standalone — so the allowlist shrinks to glm4/phi3. In the SAME wave (Tier-C1)
all four text Gemma MLPs fold onto a new shared GeGLU gate-up method
(layers::UnquantizedMlpGateUpGeluMethod, the GeluAndMul sibling of the SwiGLU
UnquantizedMlpGateUpMethod A1 landed) — one merged [2I,H] descriptor family now
serves BOTH activation families. Gemma is the strongest reuse proof because it has
REAL committed goldens: see docs/BENCHMARKS.md for the DGX SACRED-gate results.
The 27B SACRED gate needs the production kernel build, and now says so. The
27B W4A4 acceptance gate (tests/parity/test_qwen27_paged_engine.cpp) reproduces
the oracle's production continuation only when the binary carries BOTH the CUTLASS
sm120a NVFP4 fp4xfp4 GEMM (-DVLLM_CPP_CUTLASS_DIR=<cutlass 4.5.0+>) and the
vendored Triton-AOT GDN kernels (-DVLLM_CPP_TRITON=ON). Drop either and the
engine falls back to the emulation-grade fp4 GEMM / hand GDN recurrence, which
deterministically takes the other side of the documented tok6 whitespace near-tie
and emits the greedy_ids_emulation.npy stream. Measured 2026-07-29 on GB10
sm_121a from ONE source tree at main d4492c03: CUTLASS off + Triton off 234/235,
CUTLASS on + Triton off 233/235, both on 235/235. The gate now refuses to run
in a build missing either flag and names the configuration, instead of reporting a
token regression that is not there. The gate's own status is UNCHANGED and GREEN.
A whole-surface sweep of pinned vLLM 555967922 (0.26.0.dev0) against our
matrices ranked what vllm.cpp is MISSING: 8 HIGH, ~19 MED, ~16 LOW gaps, each
grounded in vLLM file:line. The material HIGH-priority misses (common,
single-box, user-facing) are LoRA / multi-LoRA runtime, the
pooling/embedding/classify/rerank task class, AWQ + GPTQ native quantized
compute, the xgrammar structured-output backend, fp8 KV cache, and reasoning
parsers. One medium gap has no tracked row yet (generic draft-model + Medusa
speculative decode); the offline Batch API (SERVE-BATCH-API) and the plugin
system (ENG-PLUGIN-SYSTEM) were picked up 2026-07-29. vLLM has
removed prompt adapters, so that is not a gap. The full prioritized list lives
in .agents/specs/vllm-feature-gap-analysis.md;
this is analysis only, no capability state changed.
The CPU CTest suite is green (0 real regressions). A 2026-07-27 hygiene pass
triaged the previously-red tests as stale assertions or -j contention, not
regressions: the model-registry count assertion tracks the 24 registered
architectures, the GGUF loader rejection case uses a genuinely-unregistered arch,
the DFlash proposer test supplies the now-required "model" key, and the two
HTTP-server tests run under CTest RUN_SERIAL so a CPU-hog concurrent test can no
longer starve their accept thread.
env-doc hygiene (2026-07-29): allowlisted VT_SPEC_TRACE (a per-block
spec-decode acceptance trace, off by default, output-neutral — runner.cpp
spec_trace) on scripts/env-doc-allowlist.txt; introduced by the concurrent
DFlash-GGUF session and left check-env-doc RED on main. No behavior change.
fusion-consistency gate-test repair (2026-07-30): VOID, no behavior change and
no teeth weakened. tests/scripts/test_check_fusion_consistency.py's mutation test
hardcoded gemma2 as its in-tree known-drift example, so the Tier-B2 fold broke it
by doing what the gate wants (Gemma adopted vt::FusedChain, its allowlist entries
were retired, and the assertion went looking for a model that no longer drifts). It
was RED on clean main before this change. The example is now DERIVED from the
allowlist: emptying the allowlist must expose every in-tree hand-fusing model, and
the allowlist must suppress at least one model the detector really sees, so the next
migration closes the gate instead of breaking the test.
scripts/check-fusion-consistency.py itself is unchanged; tests/scripts 87/87.
Landing-page restructure + project logo (2026-07-30): DOCUMENTATION-ONLY, no
capability state moved and no number changed. The thesis is now stated as the
project's actual ambition (smallest, fastest, every feature people want, with vLLM
as the PROOF rather than the identity) instead of "a C++ port of vLLM", the
llama.cpp and MLX-LM comparisons carry their real per-axis numbers alongside the
vLLM grid, and a Credits section attributes vLLM / SGLang / llama.cpp / MLX / the
kernel upstreams and records that this tree does NOT use ggml (own portable vt::
runtime, GGUF read natively) without disparaging it. Two accuracy fixes landed with
it. (1) The footprint claim was STALE: the README said "9.4 GiB venv vs one 10 MiB
binary" while the committed spec (benchmarks/demo/footprint_gb10.json, measured
2026-07-28 on the same GB10) and the rendered figure both say 9318.6 MiB (9.1 GiB)
vs a 65.78 MiB stripped CUDA server binary, i.e. ~140x, not ~960x; every instance is
now the measured number. (2) The concurrency ratios are now read against the gated
0.5% non-tail run-to-run noise band: c2-c32 (0.7%-1.7%) are presented as a TIE
with the c1 4.5% as the one win outside noise, llama.cpp decode 0.97x as parity inside
its own spread, and the Metal 2.4% as REAL and outside both spreads (ours 0.12%,
MLX-LM's 0.34% over 6 interleaved runs) rather than waved off as noise. README.md becomes a landing page
(logo lockup, scoped badges, the six-axis concurrency grid, the footprint figure,
and a features block); its reference tables move to the new
docs/USAGE.md (CLI, endpoints, server flags, C ABI, C++) and
docs/BUILD.md (per-backend recipes, every CMake option, hardware and
quantization state), and the 26-row architecture matrix plus the full feature list
become <details>. The sections scripts/check-readme-structure.py requires are
retained as compact linking sections, so the AGENTS.md landing-page contract holds
with the tables elsewhere (README 23k -> 19.1k chars against the 30k budget). The
logo ("Steps": the V1 scheduler timeline, bars crossing step boundaries) is
rendered by scripts/make-logo.py from one source-of-truth geometry, reusing the
benchmarks/demo/concurrency_race.py palette. Honest-numbers rule held: the Metal
row states prefill 1.5% ahead of MLX-LM with warm total 97.6% MLX-gated / 95.9%
default and the 18-of-75-ops caveat, and the performance badge is scoped to the
Qwen3.6-27B gate model rather than implying every architecture beats vLLM.
Bounded mechanical sync cycle. Enumerated + tier-classified the 198-commit delta
(ranked queue: .agents/specs/upstream-sync-2026-07-30.md;
report: .agents/sync/2026-07-30-e04a30a.md).
Ported (T1, CPU-gated): vllm#49754 per-request stream_interval sampling
param, mirrored into SamplingParams, the output-processor clamp, and both
OpenAI request paths (only raises the engine --stream-interval, clamped up;
first/final chunks always emitted). Gates: test_output_processor 11/11,
test_sampling_params 9/9, test_openai_protocol 28/28 (CPU). Parity pin
UNCHANGED at 555967922 (subset landed); the T1 correctness queue (48245
preemption-underflow, 50297 P/D-race, 49343 eagle-draft) and the T0 RMSNorm
perf/determinism pair (49750, 48391 — GPU runtime gate) are the ranked
carry-over. vllm#49030 (video temporal padding) was found ALREADY-CORRECT in
qwen3vl_processor.cpp (our ceil formula equals the fix); no port.
Increment 2 (2026-07-30, T1-correctness triage). Ported the portable core of
rank 3, vllm#49343 (EAGLE draft max_position_embeddings clamp): the pure
config helper SpeculativeConfig::MaybeOverrideDraftMaxPositionEmbeddings
(include/vllm/config/speculative.h) — raise an eagle/eagle3 draft's
max_position_embeddings up to the target's max_model_len (never lower;
prevents the draft rotary cos_sin_cache OOB gather our RotaryEmbeddingBase
would take). Gate: test_speculative_draft_max_position_embeddings 4/4 (CPU,
RED-first verified — the two raise/override-flag assertions fail with the clamp
neutered). We do not yet LOAD eagle/eagle3 drafts (draft-config resolution is
deferred, EAGLE3 is T2), so there is no live call site yet — this is the tested
clamp the eagle loader wires in when it lands; the 2 upstream model-integration
tests (real ModelConfig(EAGLE3_DRAFT/AR_MODEL)) are SKIPPED (need HF-download +
draft ModelConfig machinery), recorded in porting-inventory.
Ranks 1-2 SKIPPED as unported-code-path, not force-fitted: vllm#48245
(preemption num_output_placeholders underflow) and vllm#50297 (P/D
preemption race) both rebuild the async stale-output machinery around
num_in_flight_tokens / num_stale_output_tokens / drop_stale_output + the
reset_prefix_cache(reset_running) scheduler force-preempt + a skipped_waiting
queue + (50297) the KV-connector producer/consumer requires_kv_delivery role —
NONE of which this tree carries (our async_tokens_to_discard is never set, the
force-preempt path is unported, the async engine has no PP concurrent-batch
window, and KV P/D transport is N-A by sync policy). The underflow they fix
cannot occur here; porting the prerequisite machinery is a scoped feature port,
requeued in the sync doc. Parity pin still UNCHANGED at 555967922.
Increment 3 (2026-07-30, quant/config triage — verified SKIP-all). Worked the
quant-loader + config-validation slice of the ranked queue (ranks 6-8). All three
are SKIPPED as unported-code-path after source verification against the real
upstream diffs — not force-fitted, no invented code. vllm#49483 (compressed-
tensors: prioritize fused-name match over class match in find_matched_target)
reorders two matching strategies inside find_matched_target
(vllm/model_executor/layers/quantization/compressed_tensors/utils.py:145-150); our
tree has NO such generalized matcher — compressed-tensors schemes are resolved
per-projection by direct tensor-name probing (IsCtNvfp4Projection /
LoadCtNvfp4W4A16 in include/vllm/model_executor/models/dense_weight_loaders.h
check has(proj + ".weight_packed") and honor the config ignore list inline),
so there is no targets-list walk, no class-name substring match, and no
fused-vs-class ordering to fix. The bug is structurally impossible here.
vllm#48589 (INC extra_config layer-name suffix match) adds an endswith
fallback to the Intel-Neural-Compressor config parser
(vllm/model_executor/layers/quantization/inc/config_parser.py); we carry no INC
quant path (Intel/XPU, N-A bucket) — our only extra_config is the KV-connector's,
an unrelated map. vllm#49134 (reject contradictory custom-op directives)
validates CompilationConfig.custom_ops (+op/-op/all/none), a
torch.compile/Inductor plumbing config we do not carry (our fusion catalog is a
portable vt::FusedChain applied by declaration, not a runtime custom-op
enable/disable list). The T0 RMSNorm pair (49750 uncontiguous-residual perf, 48391
batch-invariance) stays QUEUED, GPU-gated: our RmsNormRowKernel
(src/vt/cuda/cuda_ops.cu:62) indexes contiguously (row*h, no input_stride),
so 49750's residual-stride relaxation is a multi-backend structural change with no
current strided-residual caller, and 48391 depends on the vllm_is_batch_invariant
determinism subsystem we do not carry — neither is a clean 1:1 mirror. No code
landed; the queue + records were refreshed. Parity pin UNCHANGED at 555967922.
Best-GEMM-path research (2026-07-30): .agents/specs/best-gemm-path-2026-07-30.md — best-in-class GEMM path per regime (fp8/full-tensor/low-bit × decode/prefill) across vLLM+SGLang+ds4+llama.cpp + FusedChain-mapping verdict. Two actionable findings: (1) fp16 dense inputs THROW (cuda_matmul.cu:220-227) — a breadth gap, one-branch cuBLASLt fix; (2) nvfp4 W4A4 M=1 decode is mis-routed to the prefill tensor-core tactic on our production 27B/35B gate models — a real decode perf hole, dp4a-matvec fix. Fp4GemmSm120 confirmed LIVE (not dormant). No behavior change (research/spec only).
nvfp4 W4A4 M=1 decode "mis-route" — MEASURED NON-ISSUE (2026-07-30, best-GEMM-path rank-2 CLOSED).
DGX-measured (Iron Law): at M=1 the cutlass autotuner ALREADY selects swap-AB small-M tactics, and
the weight-dominant 27B W4A4 projections run FLAT across M at the HBM copy roofline (gate_up 258
GB/s, down 283 GB/s) — decode is memory-bound, so a dp4a matvec reads the same bytes at the same
bandwidth and cannot raise tok/s. NO production code changed; SACRED test_ops_nvfp4_fp4 30/30·27006.
The gate models' W4A4 decode is at the achievable roofline. See .agents/specs/nvfp4-w4a4-m1-decode-measured-2026-07-30.md.
DeepSeek-V4-Flash Brick 6 — MoE fusion, +5.2% decode (2026-07-30, 6795717f, MERGED). Corrects the Brick-4 "bit-exact GEMM levers exhausted" framing: that was the GEMM mainloop; the CROSS-kernel fusion lever (ds4's real edge) was still open. Fused routed-MoE gate+up+silu into ONE kernel (mirrors ds4 moe_gate_up_mid) — BIT-EXACT (byte-identical 49-token sequence, real 80.7 GB model), decode 10.27 → 10.80 tok/s (~65% of ds4 16.5), kernel instances −12.2%. VT_V4_FUSED_MOE=0 rollback. NEXT (Brick 7): fuse the un-fused per-step glue (Rope 119.9ms + MhcPreFinish 98.7ms + Route 77.4ms) — the LARGE remaining ds4 gap. Row stays ACTIVE.
DeepSeek-V4-Flash Brick 7 — fused per-head RMSNorm+RoPE + parallelized RoPE tail, +5.5% decode (2026-07-30, branch brick7-fused-rope, commit <this>, NOT pushed). Attacks Brick 6's top glue (RopeKernel). New NormRopeRowsKernel (ds4 head_rms_norm_rope_tail_kernel/dsv4_qkv_rms_norm_rows_kv_rope_kernel) folds the {rms_norm_rows;rope} launch pair into ONE block-per-row kernel and splits the RoPE pairs across threads — BIT-EXACT (the theta *= theta_scale recurrence is a left-fold, so pair p via p sequential mults reproduces the exact double product; unit case norm_rope_rows == split BYTE-IDENTICAL, real 80.7 GB model VT_V4_FUSED_ROPE=1 vs =0 byte-identical 49-token sequence). Wired into BOTH the eager forward AND the captured V4Graph::Step (the graph path was a separate code path — the first eager-only build measured FLAT, an RED-first catch). Decode 10.82 → 11.41 tok/s (+5.5%, non-overlapping bands, ~69% of ds4 16.5); nsys: RoPE glue GPU time 129.1 → 63.9 ms (−50.5%) + ~1204 rms_norm launches folded; PEAK RSS 86.68 GiB unchanged. VT_V4_FUSED_ROPE=0 rollback. NEXT: activation-quant fusion into the GEMM prologue (QuantizeQ8K launch-bound, now the top glue at 7.6%) + MhcPre reduction glue. Row stays ACTIVE.
Cross-arch merged-GEMM fold plan (2026-07-30, research-only, .agents/specs/arch-fusion-fold-plan-2026-07-30.md). Read-only survey of ALL implemented forwards (~60 fold sites) for the merged-GEMM + fused-glue patterns, ranking how to fold every arch onto ONE shared descriptor family + the FusedChain glue catalog — the "use it consistently" rollout so DeepSeek's fused-kernel wins stop being DeepSeek-only. Headline: the merged-GEMM weight layouts already exist almost everywhere; the gap is launch/epilogue granularity + seam routing, not weight marshaling — most high-value folds inherit an already-tuned kernel for ZERO new kernel code. Tier A (bit-exact, free-kernel, wide reuse) first: A1 bf16-SwiGLU-MLP seam (5 archs: OLMo-2/Granite/StableLM/dflash/deepseek_v2), A2 MLA A-proj 3→1 merge, A3 keep-quant grouped-MoE → kMatmulBTQuantGrouped, A4 bf16 grouped-MoE → promote ds4 moe_gate_up_swiglu to shared kMoeGateUpSwiGLU, A5 MLA norm-rope → kFusedNormRope. Feeds the generalization front (shared ops + descriptor). End state: a new dense/MoE arch is "born fused" (kimi_k3 the proof — inherits A2/A4/A5 with zero model-specific work). No code changed (research/spec only).
Keep-quant fused MoE → SHARED op + declarative merged-GEMM descriptor (2026-07-30, bb864b0d, MERGED). The extensibility answer to "do other archs gain from what we do?" — the DeepSeek-V4 Brick-6 fused routed-MoE gate+up+SwiGLU kernel, previously an arch-private MoeDeviceKernels::moe_gate_up_swiglu function pointer, is PROMOTED to a first-class shared vt:: op OpId::kMoeGateUpSwiGLUGrouped, MIRRORING kMatmulBTQuantGrouped (enum + typedef + vt::MoeGateUpSwiGLUGrouped host shim + CUDA registration of the SAME kernel via an append-only registrar [BIT-IDENTICAL move] + a CPU composite golden = two MatmulBTQuantGrouped + clamped-SwiGLU). Plus the DECLARATIVE DESCRIPTOR vt::MergedGemmGroup (include/vt/merged_gemm.h) — the contraction-tier analog of FusedRecipe: N GEMMs SHARING operand A + a fused epilogue opcode, selected by (dtype, arity, epilogue), realized end-to-end by kKeepQuantGateUpSwiGLU with the byte-exact Tier-0 composite (arity× MatmulBTQuantGrouped + epilogue) as the portable fallback — same tiering as fused_recipe.h. Any keep-quant MoE arch now inherits the tuned single-launch kernel by calling ONE op; the inheritable win is the MEASURED Brick-6 +5.2% decode (6795717f). CPU gate test_deepseek_v4_gguf_load 15/15·931 (2 new RED-first cases: fused == composite byte-identical; descriptor fast==Tier-0 byte-identical). DGX: DeepSeek CUDA regressions GREEN + bit-exact (promotion is a bit-identical move). deepseek_v4.cpp UNTOUCHED (Front B owns it for Brick 8). NAMED RESIDUAL: a fresh second-model e2e decode-tok/s A/B is BLOCKED — DeepSeek is the only in-tree keep-quant grouped-MoE model (qwen3_5/Coder MoE are fp4/bf16), there is no K-quant MoE checkpoint on the box (root 96% full) and no host K-quant encoder to synthesize one, so the win is proven at the kernel level (shared op + Brick-6 measurement) + framework level (descriptor byte-exact across tiers). Spec .agents/specs/keepquant-shared-ops-2026-07-30.md.
DeepSeek-V4-Flash Brick 8 — FUSE QuantizeQ8K into the grouped-GEMM prologue: BIT-EXACT but a MEASURED NEGATIVE (−22%); hypothesis REFUTED, code REVERTED (2026-07-30, branch brick8-fused-quant, records-only, NOT pushed). Brick 7 named QuantizeQ8K (7.6%) the top glue and called it launch-bound (lever-2). Brick 8 folded the activation-quant into the two decode grouped kernels' prologue (fused gate+up+SwiGLU + down), quantizing one row/block cooperatively into a shared tile (VT_V4_FUSED_QUANT, both eager + graph). BIT-EXACT (test_cuda_quant_dot 4/4·106909 incl. a new grouped fused-quant nmse≤1e-6 case; test_cuda_deepseek_v4 19/19·66949; test_deepseek_v4_gguf_load 13/13·631; real 80.7 GB =1 vs =0 byte-identical 50-token "…Paris…"). RESULT: decode 11.46 → 8.91 tok/s (−22.3%, 4 interleaved warm runs, non-overlapping). nsys (grounded): the fused path removes QuantizeQ8K (125M ns) but the grouped kernels balloon 2.5–4.9× (+630M ns net) — folding the quant in forces THOUSANDS of GEMM blocks (down ≈6144/layer) to re-quantize their row behind a __syncthreads barrier, vs the ONCE-per-row standalone quant whose L2-cached Q8_K output all blocks read; "launch-bound" doesn't hold under the cudagraph. REFUTED, no salvage (per-block re-quant is intrinsic to the fusion); code reverted, production default unchanged (11.41). Same mechanism refutes fusing QuantizeQ8_0. CORRECTED next lever: the Q8_0 matvec QuantDotGemmQ8_0Kernel is now 49.3% of the step (Brick-3/4 latency/occupancy-bound) → a kernel-STRUCTURE change (multi-output-per-warp/persistent), NOT a fusion; then MhcPre/Route reduction glue (near-tie). Row stays ACTIVE.
DeepSeek-V4 ds4-gap STUDY — 16.5 IS reachable, ranked lever plan (2026-07-30, research/profiling only, .agents/specs/deepseek-ds4-gap-lever-plan-2026-07-30.md). Fresh nsys-diff (ours 85.35 ms/step = 11.2 t/s vs ds4 54.86 ms = 17.2 t/s; gap 30.5 ms, 98.5% GPU-bound; roofline ours 42–45% vs ds4 65–70% of 240 GB/s). No irreducible gap — the 30.5 ms is fully attributed to specific kernels, each with a ds4 file:line precedent. RANKED LEVERS (measured-%-anchored): L1 dense Q8_0 preq-prologue fusion (~9 ms; a MEASURED still-present 9.1 ms activation-quant DRAM round-trip — DISTINCT from the refuted Brick 8, which fused into the grouped-MoE kernel; near-tie; → ~13.5 t/s, the go/no-go), L3 route→warp-topk (6.5 ms, BIT-EXACT, → ~14.5), L4 QKV+norm+RoPE+KV fusion (4.0 ms, near-tie, → ~15.3), L2 Q8_0 sub-warp GEMV tiling (speculative occupancy lever, short-K only, 5–8 ms → ~15.8–16.5). Measurement REJECTED as wishful: MHC (measured TIE), routed-MoE (we WIN +2.6 ms via Brick 6), attention/fp8-KV (0.6 ms short-ctx), stream-overlap (98.5% GPU-bound). Honest band: 15.3 t/s sure (L1+L3+L4 bit-exact/near-tie), 16.5–17.2 if L2 tiling lands. Implementation follows the gated sequence with explicit stop conditions. Row stays ACTIVE.
Tier-A1 fold — 5 bf16 SwiGLU-MLP archs onto the SHARED gate-up seam layers::UnquantizedMlpGateUpMethod (2026-07-30, branch worktree-wf_9054a036-47d-1, commit <this>, NOT pushed). First EXECUTED tier of the cross-arch merged-GEMM fold plan (arch-fusion-fold-plan-2026-07-30.md §A1) AND the empirical REUSE PROOF the keep-quant generalization front (bb864b0d) could not run in keep-quant (DeepSeek is the only keep-quant MoE model). Each arch's hand-rolled {ResidentWeight; MatmulBT[2I,H]; SiluAndMul} gate-up sequence is REPLACED by the already-shared bf16 gate-up MLP method (exemplar qwen3.cpp:91), so multiple real gated archs now route their MLP through ONE tuned method and inherit the nvfp4 GateUpFusedMarlinD fused arm FOR FREE the moment a quantized checkpoint ships. BYTE-FOR-BYTE identical op sequence — no loader concat needed (every arch already packs a merged [2I,H] gate_up_proj OwnedTensor). Folded: OLMo-2 (olmo2.cpp), StableLM (stablelm.cpp), qwen3_dflash (qwen3_dflash.cpp, was hand-merged but off-seam), deepseek_v2 (deepseek_v2.cpp DenseMlp — one fold covers BOTH the dense-layer MLP and the shared-expert epilogue, its two call sites) via the direct UnquantizedMlpGateUpMethod arm; Granite (granite.cpp) via the full layers::MakeMlpGateUpMethod factory (its MLP weights ARE Qwen3DenseMlpWeights, so it is the exact exemplar and is nvfp4-ready today). The shared method's TU (linear.h/nvfp4.h) is UNTOUCHED — only new #includes in the 5 model TUs — so 27B/35B/qwen3_dense stay strict without re-gate. EMPIRICAL GATES (DGX GB10 sm_121a, RelWithDebInfo + CUTLASS 4.5.0 + Triton-AOT — the PRODUCTION stack; configure log proves CUTLASS … enabling sm120a NVFP4 cutlass GEMM + Triton AOT … vendored … sm_121a). (1) CPU composite-golden RED-first unit case in test_linear_method — the fused seam UnquantizedMlpGateUpMethod::Apply is BYTE-IDENTICAL to the standalone {MatmulBT[2I,H]; SiluAndMul} (RED-first PROVEN on the dev box: a perturbed reference weight fails all 15 byte checks; GREEN after revert) → on the DGX production CUDA stack test_linear_method 4/4 cases · 31 assertions GREEN. (2) OLMo-2 SACRED paged-engine gate RAN + PASSED: test_olmo2_paged_engine 1/1 · 92 assertions, 16/16 prompts (13/16 STRICT token-exact vs vLLM per-prompt greedy + 3/16 near-tie-band, max gap 0.094 nats, 0 forward-divergent) — IDENTICAL to the pre-fold bring-up result (5c00fc4), and the test's HARD anchor-drift REQUIRE against the committed pre-fold our_ids.npy PASSED, i.e. the folded engine produced the EXACT SAME token sequence → the fold is a proven byte-exact NO-OP on the real GPU forward through the shared seam. Golden md5 UNCHANGED before/after (no golden overwrite). HONEST SCOPE: only OLMo-2 has a committed golden in-tree — Granite and StableLM goldens were NEVER committed (the granite_greedy_2b/stablelm_greedy_1_6b dirs the test comments name do not exist in the repo or on the box), so their paged-engine tests LINK + LOAD-verify + skip-on-absent-golden (rc=0, build-verified on the production stack) but do NOT run a vs-vLLM token gate THIS pass; capturing their oracle goldens is available follow-up. Granite/StableLM route through the IDENTICAL shared method OLMo-2 just proved token-exact on GPU, and are covered by the CPU composite golden. dflash + deepseek_v2: build-verified on the production stack + covered by the CPU composite golden. Net: 1 arch empirically token-exact-gated end-to-end on GPU (OLMo-2, the shared-seam anchor) + the reuse proof (CPU composite RED-first) + 5/5 build-verified on the production stack. No new env flag (the fold is unconditional and bit-exact; the shared method carries the existing nvfp4 arm selection). Next tiers: A2 (MLA A-proj 3→1 merge) / A4 (bf16 grouped-MoE kMoeGateUpSwiGLU).
Q8_0 GEMV sudo-ncu root cause (2026-07-30, measurement-only, .agents/specs/ds4-q8-ncu-2026-07-30.md). sudo ncu (now unblocked) on our QuantDotGemmQ8_0Kernel: LATENCY-bound (long-scoreboard stall 54.4) at 71.9% occupancy, L1 hit 96.7% (34-byte over-fetch L1-absorbed → not the DRAM bottleneck), L2 10.7%. Explains why Bricks 4/11/12 (alignment/occupancy/launch-count) were all flat — the axis is memory-LATENCY, not those. The one untried lever = ILP (multiple output-rows/thread → independent load streams). ds4 ncu side not captured (graph-decode needs --graph-profiling). No production code changed.
ds4 Q8_0 GEMV — ds4 SIDE ncu captured + DIFFed (2026-07-31, measurement-only, .agents/specs/ds4-q8-raw-mechanism-2026-07-30.md). Closes the residual of ds4-q8-ncu-2026-07-30 (ds4 side never captured). sudo ncu --replay-mode application --graph-profiling node on ds4's matmul_q8_0_preq_warp8_kernel (the 1:1 kernel) on the byte-identical 6.15 GiB Q8_0 tower. DIFF (ds4 vs ours): long-scoreboard 17.2 vs 54.4, regs 56 vs 39, occupancy 60.3% vs 71.9%, L2-hit 2.6% vs 10.7%, bytes/sector 1.70 vs 2.0. Mechanism = register-resident memory-level parallelism (longSB inversely tracks regs across 3 kernels; occupancy non-monotone). Task's L2-residency + coalescing hypotheses both REFUTED (ds4 is LOWER/worse on both). VERDICT (A): a portable, bit-exact lever exists — deepen our single-row load pipeline (intra-row multi-block prefetch at 56 regs), DISTINCT from the failed Brick-13 multi-ROW ILP; but bounded (+1.5-2 t/s), does NOT overtake ds4's raw 16.5 alone. No production code. Box OOM-rebooted once on --replay-mode kernel (fixed → application replay). Decision surfaced to the user.
Hygiene (2026-07-31): VT_MLA_FUSED_NORM_ROPE allowlisted (88ae8fdf) — the A5 MLA kFusedNormRope fold rollback A/B switch (kernel-internal, no behavior change). Recorded here to pair with the env-doc allowlist entry.
Tier-A4 fold — bf16 grouped-MoE gate+up+SwiGLU onto a SHARED fused op vt::MoeGroupedGemmBf16GateUpSilu (kMoeGroupedGemmBf16GateUpSilu) (2026-07-31, branch fold/a4-bf16-grouped-moe off d05fff2d, NOT pushed). Last A-tier of the cross-arch merged-GEMM fold plan (arch-fusion-fold-plan-2026-07-30.md §A4). OPEN-QUESTION RESOLVED (§3 gap 3): qwen3_coder ALREADY routes its bf16 grouped MoE through task #90's fast vt::MoeGroupedGemmBf16 (qwen3_moe.cpp RunMoeBlock → qwen3_5.cpp MoeBlock → MoeBlockBf16Cuda, default-ON), NOT the qwen3_5.cpp:4256 reference loop — so A4 for coder is a launch-reduction/consistency fold, not a perf-gap fix. Both bf16 consumers (qwen3_coder MoeBlockBf16Cuda, deepseek_v2 MoeBlock) now replace their {gate grouped-GEMM (f32); up grouped-GEMM (f32); MoeSiluMul} triplet with ONE vt::MoeGroupedGemmBf16GateUpSilu call. The new op is BIT-IDENTICAL to that composite in EVERY regime by construction: decode/non-WMMA reuses the exact MoeGroupedGemmBf16NaiveSplitK partials + a fused reduce+SwiGLU launch (5→3 launches, drops the two f32 [P,I] reduce round-trips); prefill/WMMA reuses LaunchGroupedBf16 twice + the identical silu-mul (byte-for-byte the old three ops). HONEST DEVIATION: the bf16 grouped path uses per-expert weight-POINTER ARRAYS [E] i64 + a pair→token row_map, NOT the contiguous [E*N,K] + broadcast-act convention the keep-quant kMoeGateUpSwiGLUGrouped takes, so the bf16 arm is a SIBLING op (bf16 twin, documented in merged_gemm.h) rather than literally the same OpId. Route-weight stays in moe_combine post-down (invariant preserved); no loader change (both archs already resident-pack the per-expert bf16 gate/up towers). Default-ON behind VT_MOE_BF16_FUSED_GATEUP (=0 = composite everywhere, same-binary rollback; allowlisted). EMPIRICAL GATES (DGX GB10 sm_121a, RelWithDebInfo + cutlass-4.5.0 + triton). A/B unit test_ops_moe_grouped_bf16_gate_up_silu PASS — fused == {2× vt::MoeGroupedGemmBf16 (f32) + vt::MoeSiluMul} BYTE-IDENTICAL (RED-first memcmp) across the naive(splits==1), split-K-decode, and WMMA-decode/prefill regimes; sibling regression test_ops_moe_grouped_bf16 PASS. test_deepseek_v2_paged_engine SACRED 8/8 UNCHANGED (223 assertions; 5/8 STRICT + 3/8 near-tie, 92/128 tokens strict, max 0.25 nats @ prompt[3] tok=9, 0 forward-divergent — the near-tie prompts/tokens/gaps are IDENTICAL to the A2/A5 baseline ⇒ end-to-end BIT-EXACT), plus test_deepseek_v2_load/_forward PASS. test_qwen3coder_paged_engine vLLM-0.25.0 golden 6/6 PASS (138 assertions; 5/6 STRICT + 1/6 near-tie, 0 forward-divergent). Shared-TU (qwen3_5.cpp) inertness canary test_qwen3_32b_nvfp4a16_paged_engine 6/6 (142 assertions, 4/6 STRICT = 67/96 tok, max 0.062 nats) + test_qwen27_paged_forward PASS (27B/35B MoE take the fp4 path, never MoeBlockBf16Cuda; ops.h change is additive-before-kCount). coder speed non-regressed by construction (decode fuses to fewer launches with the SAME grouped-GEMM kernels/accumulation; prefill/c4+ decode is the byte-identical composite) — a launch-count/consistency fold, no isolated tok/s A/B claimed (mirrors the A2/A5 disposition; the same-binary rollback flag is available for an A/B). See docs/BENCHMARKS.md. A-tier of the fold plan now COMPLETE (A1/A2/A3-carry/A4/A5); remaining rollout: Gemma B1 (deferred not-bit-exact), C2/C3 vision towers, D/E tiers.
Laguna-S-2.1 W4 CUDA-verify (2026-07-31, .agents/specs/laguna-s21-w4-cuda-verify-2026-07-31.md). The full CUDA fast-path build of the W1-W3 tree is CLEAN (1059/1059, 0 warn; all 4 laguna TUs). BUT the real-GGUF forward is NOT yet built: LoadLagunaFromGguf is a VT_CHECK(false) stub, LagunaModel::Forward is an f32 reference fed only synthetic weights, no run entrypoint — so the engine cannot yet run Laguna on real bytes (this is W5, not W4). Reference established: llama.cpp Poolside-fork runs the UD-Q4_K GGUF COHERENTLY ("…Paris." @ 27.8 tok/s). Two W5 blockers scoped: a multi-shard GGUF reader (3 shards) + disk (build tree ~137GiB alongside the 69GiB model). Row stays ACTIVE (NOT runnable). W5 fully scoped + de-risked in the spec.
D3 GPU gate CLOSED (2026-07-31, GB10). OLMo-2 full-width qk-norm fold empirically confirmed: olmo2 16/16 + fused_chain 10/10 + qwen3/27B canaries unchanged. See BENCHMARKS.
A3 fold SCOPING VERIFIED (2026-07-31, pending impl). qwen3_5 GGUF keep-quant grouped MoE fold: premise confirmed bit-exact (ResidentWeight keeps weights quantized → per-expert loop already runs the ds4 vec_dot). Remaining: loader stacking + a keep-quant grouped MoE block (ds4 MoeFusedResident analog) + RED-first byte test. DGX SACRED gate OWED. See .agents/specs/arch-fusion-fold-plan-2026-07-30.md A3.
Laguna W8 embed-fix LANDED + GATED (2026-07-31). LagunaEmbed gathers only the T needed rows instead of converting the whole 1.23 GB embed table to f32 every token. GATED on real UD-Q4_K_XL GGUF (GB10, --gpu): TOKEN-IDENTICAL to the W5/W6 golden + decode 0.66 → 0.17 s/tok (3.9×; 1.5 → 5.9 tok/s). Bit-exact by construction. See .agents/specs/laguna-s21-w7-speed-2026-07-31.md §W8.
Laguna W9 grouped-expert MoE LANDED + GATED (2026-07-31). The 30 un-grouped per-expert GEMV launches/step fold onto the shared vt::MatmulBTQuantGrouped (3×top_k → 3 launches/token, no loader change). Same-binary A/B on real UD-Q4_K_XL (GB10, --gpu): grouped==per-expert BYTE-IDENTICAL (md5 754728c6, == W6 golden) + decode 0.18 → 0.13 s/tok (1.38×). Cumulative with W8: 0.66 → 0.13 s/tok (5.1×, 1.5 → 7.7 tok/s). See spec §W9.
(Laguna W8+W9 decode-speed campaign handoff recorded in .agents/state.md; cumulative 0.66 → 0.13 s/tok, next lever #1 device-resident decode.)
(Laguna W10 re-profile: current 0.13 s/tok decode is 89.7% host-sync (14,034 cudaStreamSynchronize/24tok ≈ 610/step); lever #1 device-resident decode confirmed top. See spec §W10.)
(Laguna W11 go/no-go: decode is GPU-COMPUTE-bound (GPU-busy 2.56s ≈ sync 2.59s); device-residency DEMOTED, real levers = QuantizeQ8K-dedup + keep-quant GEMV BW-tuning. See spec §W11.)
qwen3_5 A3 keep-quant grouped-MoE fold — W1 spike + scaffolding (2026-07-31). Spike spec committed (.agents/specs/qwen35-a3-grouped-moe-2026-07-31.md) + additive MoeBlockWeights.expert_*_kq stacked fields (dead until W2 loader). Applies the landed Laguna W9 grouped pattern to qwen3_5 GGUF; memory-neutral (stacked replaces per-expert copies). W2 loader / W3 forward / W4 gate (test_qwen36_gguf_engine) to follow.
qwen3_5 A3 W3 byte-exactness hazard recorded (2026-07-31). The grouped down GEMM needs bf16-output byte-exactness with ExpertMlp MatmulBf16 on the STRICT-gated 35B; shared MatmulBTQuantGrouped is f32-out only. Provable-by-construction but A/B-confirm-before-strict; may need a bf16-output grouped-down variant. See spec §W3 hazard.
qwen3_5 A3 W2/W3a attempt (2026-07-31): CPU gate caught an incomplete multi-consumer change. Emptying the per-expert expert_gate vector for keep-quant breaks 3 consumers (main MoE forward + MTP VT_CHECK + the split-lossless test); all must route through the stacked tower together (both-copies = 35B OOM). Nothing broken pushed. See spec §W2/W3a attempt.
qwen3_5 A3 W2/W3a attempt #2 (2026-07-31): DGX-gated, device staging bug ROOT-CAUSED, 1/2 passing (not landed). Fixed the primary bug (keep-quant slice must go via ResidentWeight — CUDA needs_weight_staging=TRUE); case 1 (APEX-Balanced) now golden-correct, case 2 (APEX-Compact) residual all-zeros. CPU gate 37/37. Verify-before-push held (nothing broken shipped). See spec §W2/W3a attempt #2.
qwen3_5 A3 W2+W3a LANDED (2026-07-31). Keep-quant experts load as one stacked [E*out,in] tower + byte-exact ExpertMlpKq slice forward (via ResidentWeight-resident slice — CUDA needs_weight_staging=TRUE). ATTRIBUTED byte-identical to pre-A3: APEX-Compact keep-quant golden-correct + CPU 37/37; the APEX-Balanced 2nd-case all-zeros is PRE-EXISTING (reproduced on pre-A3, tasks #50/#65). Unblocks W3b grouping. See spec.
qwen3_5 A3 W3b LANDED (2026-07-31) — A3 COMPLETE. Per-expert keep-quant MoE loop → 3 grouped vt::MatmulBTQuantGrouped launches (VT_QWEN35_GROUPED_MOE default-ON). BYTE-EXACT: DGX APEX-Compact grouped(=1) vs per-expert(=0) continuations byte-identical + strict-passes golden. Shared descriptor with Laguna W9. See spec §W3b.
Fold-plan A3 tier CLOSED (2026-07-31). Both keep-quant models (Laguna W9 + qwen3_5 GGUF W2/W3a/W3b) route their routed-expert MoE through the shared vt::MatmulBTQuantGrouped descriptor, byte-exact gated. See .agents/specs/arch-fusion-fold-plan-2026-07-30.md A3.
Canonical record-owner repair (2026-07-31). No lifecycle or support status changed. Six existing DONE rows now point to their reachable binding closure commits; this is records-only metadata maintenance.