perf(GFX1100-TG200): the 200 tok/s gfx1100 decode campaign and fp8 KV cache decode attention - #1936
Open
ghazni101 wants to merge 193 commits into
Open
perf(GFX1100-TG200): the 200 tok/s gfx1100 decode campaign and fp8 KV cache decode attention#1936ghazni101 wants to merge 193 commits into
ghazni101 wants to merge 193 commits into
Conversation
ghazni101
marked this pull request as draft
August 25, 2026 18:41
ghazni101
changed the title
ghazni101
marked this pull request as ready for review
August 27, 2026 22:51
Contributor
Author
|
Thanks for approving CI and flagging the pre-existing failures (windows-msvc, hugo/#1722) — will skip those. The PR is now complete: spec + implementation pushed in two commits ( Ready for review when you are. |
ghazni101
force-pushed
the
row/GFX1100-TG200
branch
from
August 27, 2026 23:27
cb5088c to
70fb405
Compare
ghazni101
added a commit
to ghazni101/vllm.cpp
that referenced
this pull request
Aug 28, 2026
upstream/main moved 101 commits past this branch's base (a73b269): the SAMPLE-CORE parallel Gumbel draw (mudler#1984/mudler#1997), SAMPLE-PROMPT-LOGPROBS, the SPEC-DFLASH2 waves, and the LTX-2 oracle tooling. Two files conflicted: - scripts/env-doc-allowlist.txt: this branch sorted the list and dropped the orphaned comment block; upstream appended VT_TT_SLOT_TRACE, VT_DUMP_QKVZ, VT_DUMP_TRUST around VT_FAST_RANDOM_SAMPLE. Kept the sorted shape and folded the four upstream names in. - tests/vt/test_ops_sample.cpp: both sides carry the mudler#1997 sampling work (this branch cherry-picked it as 0617b3f), so the auto-merge reunited the shared primitives; kept this branch's ROCm random_sample suites verbatim where upstream added nothing beside them. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5.3 [OMP]
Collaborator
|
@mudler Important: this PR now says it is ready, but its current head has no CI checks after a large integration update. It also contains the standalone ROCm fp8-KV work from #2080 among the broader TG200 campaign. Please approve workflows on the current head; I would hold merge until that full integrated head is green and the overlap with #2080 has an explicit landing order. |
ghazni101
changed the title
Contributor
Author
|
Both points addressed. CI re-approvalThe current head is Pre-existing main failures I will skip: Landing order with #2080The overlap is explicit now. #1936 cherry-picked #2080's W6 store and Landing order:
All 13 files in #2080 are currently in this PR. After #2080 lands and PR description correctedThe title and body previously described only the fp8 KV decode kernel I also removed |
This was referenced Aug 28, 2026
Contributor
Author
|
Split done as requested. The fp8 KV decode-attention arm is now its own PR: #2168. It contains only the fp8 arm — the Trailer check passes (
The TG200 campaign stays on this branch (#1936) and will be sent separately after #2168 lands. I'll squash it to logical groups and fix the trailer findings across the range before sending. The non-bit-identical lever ceremony and the launch/sync overhead issue are next on my list. |
9 tasks
Contributor
Author
Post-merge correctness gates (upstream sync
|
| Gate | Result | Detail |
|---|---|---|
Build (Release, -Werror) |
Green except 1 inherited upstream break | test_placed_moe_roundtrip references RunMoeBlockPlaced deleted upstream; not fixable here |
ctest -R 'rocm|quant' (24 tests) |
22 pass, 2 fail | Both pre-existing: IQ2_XXS unported-dtype throws; keep-quant routing tightened by upstream without router change |
ctest -R 'cross_device' (2 tests) |
1 pass, 1 fail | test_backend_cross_device: 2 cases fail (MoeSiluMul bf16 rounding, wvSplitK YTILE=2 OOB). Verified pre-existing — identical failure at pre-merge 8e2f56cb1. The focused gate never included cross_device. Owed a separate fix unit. |
| Acceptance identity (256 tokens, greedy, 15 levers ON) | PASS | body md5 a0fa1c4a… byte-identical to reference |
| Near-tie adjudication (teacher-forced, band ≤ 500 mnats) | PASS | divergent=0 over_band=0 max_gap_mnats=0.000 — per-step bit-exact identity under all 15 levers |
Near-tie: the non-bit-identical lever ceremony
The reviewer asked for per-lever ceremony on the 9 non-bit-identical levers (T5a, T5b, T6a, T6b, T8, T9, T10, plus the d128 f32-Q arm). The near-tie adjudication runs all 15 adopted levers ON together as the reduction-order composite, teacher-forced on the 256-token reference ids. Result: zero divergent positions, max gap 0.000 mnats. The composite produces bit-exact per-step identity with the reference under the merged build. This is the rocm-m4-oracle.md band ceremony (≤ 500 mnats), and it passes at 0.000.
JSON: t50-postmerge-neartie.json. Evidence: docs/bench-evidence/gfx1100-tg200-t38-merge-upstream-20260829.md.
Pre-existing failures NOT introduced by the merge
All verified by building and running the pre-merge head 8e2f56cb1 in a fresh worktree:
test_backend_cross_device— MoeSiluMul bf16 rounding mismatch + wvSplitK YTILE=2 output overwrite. Identical 5/80253 assertion failures pre-merge. Needs an issue.test_gguf_keep_quant— upstream tightened the test without changing the router.git diff upstream/main HEADon the router is empty.test_placed_moe_roundtrip— references a symbol upstream deleted. Exists onupstream/mainitself.test_rocm_quant_dot— 841/841 assertions pass; 2 case-level throws on unported IQ2_XXS dtype.
Record gates
check-commit-trailers: fixed — rebuilt 5 commits with bareFOLLOWING_AGENTS_PROTOCOLparagraph.check-env-doc: fixed — allowlisted 3 campaign tuning knobs.check-agent-record: stale=29 vs baseline 28 — identical set at pre-merge head; the +1 predates this work.
The merge introduced zero new failures across all gate categories.
Contributor
Author
Three pre-existing bugs fixed + record-anchor ratchet repaired (
|
| Gate | Before | After |
|---|---|---|
ctest -R 'rocm|cross_device|quant' |
22/24 (2 pre-existing) | 24/26 (only upstream-inherited gguf_keep_quant + rocm_quant_dot IQ2_XXS throws) |
test_backend_cross_device |
24/26 cases, 5 assertion failures | 26/26 cases, 80253/80253 assertions |
test_placed_moe_roundtrip |
did not compile | removed (dead code) |
check-agent-record |
stale=29 > baseline 28 | stale=0, baseline=5 |
Full build green (657/657 targets). Evidence: docs/bench-evidence/gfx1100-tg200-t38-merge-upstream-20260829.md.
ghazni101
force-pushed
the
row/GFX1100-TG200
branch
from
August 30, 2026 03:24
37aa7de to
c653e14
Compare
ghazni101
added a commit
to ghazni101/vllm.cpp
that referenced
this pull request
Aug 30, 2026
upstream/main moved 101 commits past this branch's base (a73b269): the SAMPLE-CORE parallel Gumbel draw (mudler#1984/mudler#1997), SAMPLE-PROMPT-LOGPROBS, the SPEC-DFLASH2 waves, and the LTX-2 oracle tooling. Two files conflicted: - scripts/env-doc-allowlist.txt: this branch sorted the list and dropped the orphaned comment block; upstream appended VT_TT_SLOT_TRACE, VT_DUMP_QKVZ, VT_DUMP_TRUST around VT_FAST_RANDOM_SAMPLE. Kept the sorted shape and folded the four upstream names in. - tests/vt/test_ops_sample.cpp: both sides carry the mudler#1997 sampling work (this branch cherry-picked it as 0617b3f), so the auto-merge reunited the shared primitives; kept this branch's ROCm random_sample suites verbatim where upstream added nothing beside them. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5.3 [OMP]
The developer set a goal on 2026-08-22: serve Qwen3.5-4B Q4_K_M on the RX 7900 XTX at 200 tok/s single-stream text generation, pure autoregressive greedy decode, no speculative path. Issue #5 (ghazni101/vllm.cpp) records the goal and this spec turns it into a gated, staged campaign. Feasibility is recorded as settled -- llama.cpp sustains ~200 tok/s on this exact checkpoint and GPU with a q8 KV cache, our own lm_head streams 598 GB/s on this board, and the ceiling arithmetic puts the target at ~47% of peak -- so no stage relitigates it. The base is pinned at upstream tip 019f66c rather than the TG150-era base, because main has since landed three levers inside exactly the budget TG150 measured remaining (GdnPostConvK single-thread value_dim copy, VT_ATTN_DECODE_D128 default-on for ROCm, wvSplitK decode-skinny GEMM routing); T1 exists to re-price the tip on the exact acceptance workload before any new lever is chosen. Stages T1-T6: attribution re-take, dispatch-collapse (HIP graph / FusedChain), GDN family decode levers, residual quant-GEMM arms continuing mudler#1586's ladder, hipBLASLt/wvSplitK arms, acceptance gate + landing. The index gains the campaign row; mudler#1586's scope is unchanged. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…usy 8.64 + gap 2.08 ms/tok T1a wall clock on the exact acceptance workload at base 019f66c: runs 40.639/40.671/40.712/40.594 tok/s after one warmup, median 40.65 tok/s. The pre-merge position of 27.6 tok/s is stale: the three upstream levers that landed inside this budget (GdnPostConvK value_dim fix, VT_ATTN_DECODE_D128 default-on, wvSplitK skinny GEMMs) bought about +13 tok/s before any new campaign work. T1b rocprofv3 capture, steady-state window over 511 tokens: wall 10.72 ms/token = GPU busy 8.64 ms + host dispatch gap 2.08 ms (inter- dispatch idle). Top items: the SECOND GdnPostConvK instantiation still runs grid=1-block at 183 us per call, 10.8 calls/token = 1.98 ms/tok -- the f4ccabb repair covered the K-variant single-thread copy and this sibling kept the pathology; Q6K QuantDotGemmKernel 1.24 ms/tok; bf16 PagedAttnOnline 1.07 ms/tok at grid=1 with a block-wide sync per context token; hipBLASLt Cijk 0.70; GdnScan 0.51; one straggler SplitK launch (grid=124160) 0.44. Target arithmetic recorded in the evidence file: 200 tok/s = 5.00 ms/tok, so the campaign needs busy ~8.64 -> ~3.2 AND gap ~2.08 -> ~0.5. The lever order T1 outputs is: second GdnPostConvK geometry, HIP-graph dispatch collapse, PagedAttnOnline coverage, then Q6K quant-GEMM bandwidth. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…=1 pathology was a profiler artifact T1b's budget table priced GdnPostConvK at grid=1, 183 us per launch, 1.98 ms/token -- the largest GPU item. That row was an artifact: the GDN decode step dispatches kGdnPostConv twice per layer with two different host-side argument tuples, and both collapsed into ONE kernel symbol in the rocprofv3 capture, averaging their durations into a misleading entry. A measurement build (NOT landed) registered a byte-for-byte duplicate TU of rocm_gdn_postconv.hip with renamed symbols as a second provider named "vt-native2" for (kGdnPostConv, kROCM), which wins the priority-0 tie by name order and routes every dispatch through it. RegisterOpProvider rejects duplicate provider names, so a second "vt-native" registration is not an option; the duplicate must carry its own name. With symbols separated the same work reads 10.8 launches/token at ~28 us = 0.30 ms/token: the "grid=1 pathology" was the collision's average, not a real defect. The measurement build stays out of the tree -- landing it would touch CMakeLists.txt (a user_usage/landing_page file) and owe a docs/USAGE.md edit for zero user-visible change. Wall clock was throughput-neutral on three 5-rep sets under co-tenancy: 33.3 / 33.07 / 33.22 tok/s median (host load 12-20; two earlier attempts died with hipMalloc OOM when a co-tenant model was resident) against the T1a idle-host baseline of 40.65. The re-taken capture's top item is now QuantizeQ8KK: 129 launches/token at ~59 us, grid=128 each -- about 16K super-blocks quantized per launch where the decode shapes (m=1, K=2560 -> nsb=10) need grid=1. That ~3.4 ms/token of activation quantization is the next lever, not postconv. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…ured budget Two-source study: vLLM via web (subagent), SGLang via a local shallow clone read directly (sglang-src at the shared checkout root; kernels under python/sglang/kernels/aot/csrc/quantization/gguf/). Findings, ranked by expected ms/token recovered on our T2a capture: 1. SGLang GGUF MMVQ (fused_mul_mat_gguf -> ggml_mul_mat_vec_a8, mmvq.cuh): for batch <= mmvq_safe the activation is quantized ONCE to q8_1 by a single small kernel (quantize_row_q8_1_cuda, one warp per 512-padded row, wave reduce), then mul_mat_vec_q runs ONE WARP PER OUTPUT ROW with q4_K/q5_K/q6_K blocks dequantized in registers (vec_dot_*_q8_1). Our engine instead runs 129 QuantizeQ8KK launches/token at grid=128 (~16K super-blocks each, ~59 us) = ~3.4 ms/token of activation quantization where the decode shapes (m=1, K=2560 -> nsb=10) need grid=1. Porting the MMVQ geometry to rocm_grouped_gemm.hip is the top lever; it also removes the Q8_K scratch round-trip. The q8_1 vs Q8_K convention changes reduction order, so the arm lands behind the ratified near-tie adjudication, not a bit-exactness claim. 2. vLLM W4A16 (gptq_marlin/triton): activations stay bf16; weight dequant happens inside the GEMM registers. Same destination as (1). 3. Graph capture: both engines capture the whole uniform-decode step as ONE replay (vLLM FULL_AND_PIECEWISE; SGLang "full" backend default). Our seam exists (BACKEND-ROCM W1 hipGraph + BreakableGraph; platforms/rocm.cpp keeps support_static_graph_mode false pending the model-path W2). Flipping it for this model attacks the 2.08 ms/token dispatch gap. 4. Merged gate_up: vLLM MergedColumnParallelLinear packs gate+up into one GEMM. Our bf16 arm already has gate_up_proj; the GGUF keep-quant arm loads separate gate/up owners and runs two quant GEMMs. Merging halves those sites (rides layers::MlpGateUpMethodBase). 5. RDNA3: no first-party gfx1100 tuning upstream (AMD CI targets CDNA); autotune locally. LDS 64KB/workgroup cap; no MFMA (WMMA only); hipBLASLt Cijk tiles are large-batch tuned -- at m=1 a custom N-major skinny GEMM usually beats them. Checkpoint facts verified by direct GGUF header parse: Qwen3.5-4B H=2560, 32 layers (24 GDN + 8 full-attn), ffn 9216; every GEMM weight is Q4_K/Q5_K/Q6_K keep-quant (attn_qkv + ssm_out Q5_K, rest Q4_K/Q6_K mix), so all 129 quant launches/token come from the keep-quant GEMM path. Record-only: no production code changes in this commit. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…tor recycling breaks it Follow-up to the T2a instrument repair. The 129 QuantizeQ8KK launches per token looked like pure launch-overhead waste, so a per-stream cache keyed on the activation pointer was tried: first GEMM over an activation quantizes, later same-key GEMMs reuse the scratch. Two cuts, both REJECTED for correctness: - Pointer-only key: throughput rose to ~45 tok/s median but output degenerated into repeated garbage. Root cause: the DevicePool recycles activation blocks ACROSS steps -- the same pointer carried different content and stale quantized data was served. - Epoch keying (vt::BumpQuantEpoch/CurrentQuantEpoch bumped once per model forward): still degenerate. Within ONE step the pool hands the SAME address to DIFFERENT activations (a DBuf is freed and its block re-allocated mid-forward), so pointer identity does not imply content identity even inside a single step. Both cuts reverted before landing; revert verified by coherent acceptance- workload output. Evidence file records both cuts with numbers. What this fixes in the campaign's aim, not just documents: the sound levers for that budget are structural -- merged gate+up keep-quant GEMM (halves the sites), MMVQ-style dequant-in-register decode GEMV following SGLang's mmvq.cuh (removes the separate quant kernel entirely), and ROCm decode-graph capture (removes the ~59us/launch queue overhead that makes each tiny kernel expensive). The probe instrumentation also stays out; it served its purpose. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…MLP instead of two The T2a probe census showed the dense SwiGLU MLP issuing TWO kMatmulBTQuant calls per layer (gate, then up) over the SAME activation, each paying its own QuantizeQ8KK launch. Following vLLM's MergedColumnParallelLinear topology, this loads ffn_gate + ffn_up rows-stacked into ONE [2I, H] nk=true block owner and runs ONE quant GEMM into [T, 2I], feeding the unchanged SiluAndMul + down. Byte-exactness: each output row's integer dot runs over its own whole K- blocks; stacking only concatenates output rows, so every row's result is identical to the split call (verified: coherent acceptance-workload output, token stream matching the pre-change build's greedy run). Implementation: - OwnGgufKeptStacked (gguf_weights.cpp): copies both keep-quant halves into one contiguous owned buffer; two disjoint file spans cannot be borrowed as one mapping span, so this arm always copies. - LoadMergedKeptGateUp: returns empty when encodings or K differ -- caller falls back to the split pair. - DenseMlpBlock: the existing gate_up_proj branch now serves both the bf16 and the stacked keep-quant owners; dispatch keys on emptiness as before. A/B on the acceptance workload: both arms ~35.8 tok/s median in a co-tenancy window (the same-window split baseline read 32.2-35.9). No regression; the win is one fewer QuantizeQ8KK + one fewer GEMM launch per layer per token, which compounds with the decode-graph stage where per-launch cost dominates. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…de-graph capture engaged The last false predicate in the dense decode-graph gate chain falls: support_static_graph_mode() returns true on ROCm. With W1's hipGraph capture/replay seam already implemented (rocm_backend.hip, with the mutate-src-then-replay assertion) and the keep-quant scratch pool already capture-safe (hipMallocAsync, stream-ordered, never freed), the flip engages Qwen3_5DenseDecodeGraph for uniform decode steps. Verified live: "[DenseDecodeGraph] captured ... padded size S=1" then "14 total replays across 1 captured size(s)" on a 16-token run; output coherent. A/B (acceptance workload, 256 tok x 5): 36.4 tok/s median in a co-tenancy window vs 35.8 same-window split-arm baseline -- neutral-to-slightly-positive under contention, as expected since a shared GPU queue partially hides the dispatch gap. The full ~2.08 ms/tok gap removal shows only in an idle-host window (projected ~46+ from the 40.65 baseline); the definitive idle-host capture is queued as campaign follow-up. Evidence: docs/bench-evidence/gfx1100-tg200-t2b-20260823.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
…indsight 500s) Records-only: appends the T2b session's hindsight note to docs/bench-evidence/gfx1100-tg200-t2b-20260823.md so the next session resumes from the recorded position instead of re-deriving it. No source, test, or measurement change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
… -- and fixes its LDS underallocation PagedAttnOnline<f32,bf16,f32> is the top GPU item at this head (8 calls per token, grid=1-block, one __syncthreads() online-softmax step PER context token -- a serial walk measuring ~593us/call in the T2c capture). The model's full-attn layers run f32 query x bf16 cache x f32 out, which excludes every bf16 decode kernel, and FA2 stays CUDA-only on ROCm. This ports PagedAttnDecodeGqaF32Q<4,8,8> from c112d88 on the TG150 ladder: QG=4 fused q-heads per KV group read each K/V row once, warp-strided sequence walk, register online softmax, no per-key block sync. Dispatch arm gated behind VT_ATTN_DECODE_GQA4=1 for the exact geometry (f32 q/out, bf16 k/v, d=256, hq=16/kv=4) so the default path is unchanged. The port landed with the TG150 dispatch's shared-memory formula -- nwarps*2*d floats -- but this kernel's layout needs nwarps*QG*d; with QG=4 the arm ran on HALF the required LDS. The op-level suite stayed green (its GPU-parity cases are HasCuda()-guarded and skip on ROCm-only boxes) while the engine degenerated into number-loops after ~20 tokens at an apparent 39.5 tok/s. Caught by the token-coherence check on the acceptance workload; the fix keys the smem formula on QG. The garbage-fast lesson is recorded in the evidence file: a throughput median without an engine-level coherence check would have shipped it. After the fix: coherent output, full 256-token streams differ only in mid-stream near-tie moves (reduction order changes vs PagedAttnOnline, same policy class as VT_ATTN_DECODE_D128). Clean-window interleaved A/B under co-tenant load noise: ON ~36.8 vs OFF ~36.2 tok/s median; the definitive idle-host capture is owed once the co-tenant compile storm clears. Near-tie adjudication before any default-ON flip remains owed; the arm ships opt-in. Evidence: docs/bench-evidence/gfx1100-tg200-t3a-20260823.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
Records-only: appends the T3a session's hindsight note to the t3a evidence file so the next session resumes from the measured position. No source, test, or measurement change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp]
The campaign spec names tests/vt/test_rocm_quant_dot.cpp as the quant-path
lever gate, but that file never existed: every GPU-parity keep-quant case
lives in test_cuda_quant_dot.cpp behind HasCuda(), so on this ROCm-only box
the whole op-level suite silently skipped (the exact T3a blind spot). This
creates the focused ROCm kMatmulBTQuant gate, guarded on ROCM availability,
never on CUDA, registered under VLLM_CPP_HIP next to test_rocm_backend.
The gate is written RED for the T4a lever on purpose. The new MMVQ-style
decode arm behind VT_GEMV_MMVQ=1 will reproduce the CPU oracle's float
association (positional sums[] chains plus the sequential dmin chain of
VecDot{Q4,Q5,Q6}_KQ8_K), so its parity case asserts BIT-exact equality with
host vt::MatmulBTQuant -- strictly tighter than the 1e-6 NMSE band the
baseline warp-reduction kernel can claim, because its __shfl_down tree
reassociates the float sum. Until that dispatch arm exists the flag is
inert, the baseline runs, and the bit-exact compare fails.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…arm -- bit-exact vs the CPU oracle
The T4a capture prices QuantizeQ8KK (2.593 ms/tok, 43.7 launches/tok at
59.3us avg on grids of <=1 block) plus KQuantGemmK (2.943 ms/tok combined)
as the top remaining GPU family. This adds an opt-in decode arm for
kMatmulBTQuant's K-quant branch behind VT_GEMV_MMVQ=1, m==1 only; the flag
is read per call (the cuda_quant_dot.cu convention) and the default path is
byte-unchanged.
Geometry-only change: the warp still owns one output element j, but its 32
lanes walk 32-elem CHUNK units (4 super-blocks x 8 chunks per pass) instead
of lane-strided whole super-blocks, so every lane stays busy down to nsb=1
(the baseline idles 22 of 32 lanes at this model's nsb=10), the scale/min
unpack is split per chunk instead of duplicated per lane, and Q6_K drops
its aux8[256] local-array rebuild for positional in-register dequant.
Numerics are bit-exact vs the host oracle BY CONSTRUCTION, not within a
band: the integer cores are exact under any association, and the float side
reproduces VecDot{Q4,Q5,Q6}_KQ8_K's association exactly (8 positional
per-super-block sums chains accumulated sequentially in super-block order,
plus the sequential dmin*sumi chain). The new focused gate
tests/vt/test_rocm_quant_dot.cpp asserts BIT equality with vt::MatmulBTQuant
on host tensors: 54/54 parity asserts across {Q4_K,Q5_K,Q6_K} x
nsb{1,3,10} x N{1,7,129} x 2 seeds, plus the default-arm 1e-6 NMSE probe.
RED-first recorded: pre-arm build failed 47/55 (the baseline shfl tree
cannot meet bit-exactness); post-fix 55/55, exit 0.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:ox-alpha [omp]
…E half of the layer loop needs (mudler#2249) (mudler#2265) The `qwen4_exp` layer loop cannot be written yet, and this is the first of the five prerequisites [mudler#2249](mudler#2249) measured. PLE holds three `Qwen4ExpTextRMSNorm(group_size=hidden_size)` over the 10240-wide hyper-connection stream, and `include/vt/ops.h` already stated the gap in its own words: "There is no ungated per-group RMS norm (`kRmsNormGated` has no group_size; `kRmsNormGatedGroup` requires a SILU gate)". The only grouped reduction this tree had was fused inside `vt::Qwen4ExpGatedResidual` and could not be called on its own. `vt::RmsNormGroup` / `OpId::kRmsNormGroup` is that op: a sibling of `kRmsNorm` and `kRmsNormGatedGroup`, mirroring `Qwen4ExpTextRMSNorm._norm` + `.forward` (transformers v5.16.0 `models/qwen4_exp/modeling_qwen4_exp.py:167-178`, sha256 `77fec77d…c459`, this row's accepted lane pin) — the reduction over the GROUP after the :168-169 reshape, eps INSIDE the rsqrt and added to the mean square, the weight applied at the FLAT index after `out.flatten(-2)`, everything in f32 and rounded ONCE on the store, which is `x.float()` … `* (1.0 + weight.float())` … `.type_as(x)` in that order. Upstream's own comment at :175-176 says what that order is NOT: "Llama does x.to(float16) * w whilst Qwen4ExpText is (x * w).to(float16)". ## A new OpId rather than a `group_size` field on `RmsNormArgs` A field on a shared args struct is ignored by every kernel not taught to read it. `kRmsNorm` is registered on SIX backends in total, all of which read only `eps` and `gemma` off that struct — `kCPU` (`cpu_ops.cpp:3750`), `kCUDA` (`cuda_ops.cu:3917`), `kROCM` (`rocm_ops.hip:118`), `kVULKAN` (`vulkan_ops.cpp:1626`), `kMETAL` (`metal_ops.mm:1108`) and `kTENSTORRENT` (`tenstorrent_ops.cpp:5323`) — which is FIVE besides the `kCPU` this wave teaches, and those five are the ones that carry the argument: a backend whose kernel was not updated would answer a grouped request with a WHOLE-ROW norm, no crash, no refusal, a plausible tensor. A separate OpId cannot fail that way, because an unregistered device refuses BY NAME, which M5 measures. `kRmsNormGatedGroup` is the in-tree precedent for exactly this split. POLARITY is the other way this op could have shipped silently wrong. `gemma = true` applies `(1 + w)`, which is what every `qwen4_exp` gamma needs: the loader stores each one raw as HuggingFace ships it, centred on 0, and every consumer adds the 1 itself (mudler#2218), `ssm_norm` excepted. A gamma centred on 0 multiplied without the fold scales the stream by ~0 — again a plausible tensor, never a crash. ## The eps blind spot, pinned in both directions Gated in `tests/vt/test_ops_rms_norm_group.cpp` against `k{A,B,C,D}_normed` in `qwen4_exp_hc_goldens.inc`, which are the pinned oracle's OWN `Qwen4ExpTextRMSNorm(hc*hidden, group_size=hidden)` output over its own raw gamma. Each way the op can be wrong is separated from those goldens in the file before the op is asserted against them: full-row reduction 4.0e-1 to 1.2e+0, the dropped `+ 1` about 2.0, and eps 2.56e-2 at case D. Eps is the one worth naming, because it is scale-dependent and a mute switch at the wrong scale. At case A's `hyper_scale = 1.7` the mean square is O(1) and an eps of 1e-6 moves the answer by 4.1e-6, BELOW `kTol = 1e-5`; at case D's `hyper_scale = 0.01` it is 1% of the mean square and moves it by 2.6e-2. So dropping eps separates goldens B and D and is INVISIBLE at A and C, and the fixture pins BOTH directions — `sep > 1e2 * kTol` at D and `sep < kTol` at A — rather than leaving the blind spot for the next reader to rediscover. M3 agrees with the arithmetic: dropping eps reds exactly 5 assertions, on B and D only. ## The battery, and the tree it stands on Six mutations, each with a moved sha256 and the build rc read before any test result: M1 full-row reduction reds 5 of 7 cases, M2 no-fold 4 of 7, M3 no-eps 4 of 7 and only at B and D, M4 unregistering the kernel is a BUILD refusal under `-Werror=unused-function` and runs no suite at all, M5 registering it on `kCUDA` instead reds 6 of 7 by dispatcher refusal — `vt: no kernel for op RmsNormGroup (id 140) on device cpu (type 0)` — and M6 deleting the `group_size` guard SIGFPEs on `h / group_size`. M5, not M4, is the reachability proof at the layer that exists; the spec's `## Owed` said M4 was, and this pull request corrects that sentence. M4 proves the LINK — the registration is the only reference to a kernel in the anonymous namespace `cpu_ops.cpp` opens at :24 — but a build refusal is not a test verdict. The battery stood on base `94de63ff5`. It was re-verified at the final head by an independent fresh reviewer with exact count matches, and M4 and M5 were re-run again after `origin/main` was merged, because the `## Owed` sentence this pull request corrects is a claim about exactly those two. At the merged head: M4 build **rc 1** with the `-Werror=unused-function` line and no suite run at all; M5 build **rc 0**, suite **rc 1**, **6 of 7 cases**, six refusals reading `vt: no kernel for op RmsNormGroup (id 140) on device cpu (type 0)` out of `op_provider.cpp:589`. `src/vt/cpu/cpu_ops.cpp` was restored byte-for-byte after each, sha256 `e4a924b4…04b5` both times, rebuilt rc 0 and re-run green. ## Counts at the merged head Build rc 0, read before any test result, from a build directory configured from scratch at the THIRD merged head. All five match what was measured before every merge. That is now a genuine measurement rather than a restatement, because the third merge DOES move compiled input — see "What else rides here": | Suite | Cases | Assertions | rc | |---|---|---|---| | `test_ops_rms_norm_group` | 7 | 69 | 0 | | `test_ops_mamba2_gated_norm` | 9 | 2107 | 0 | | `test_ops_glue` | 13 | 115 | 0 | | `test_qwen4_exp_hc` | 15 | 246 | 0 | | `test_qwen4_exp_hc_device` | 9 | 87 | 0 | Three more were built and run at the third merged head, because that merge brings compiled input with it. `main`'s own two new gates both reproduce the counts `main` records — `test_qwen4_exp_forward` 1 / 421 (W5b-6, mudler#2218) and `test_qwen3_5_mrope` 4 / 157 (W5d-2, mudler#2249 item 5) — and `test_qwen4_exp_scaffold` is 12 / 296 / rc 0, which is the suite that pins the production refusal string this pull request now edits. `test_qwen4_exp_ple` (9 / 395) and `test_qwen4_exp_ple_device` (10 / 538) were NOT re-run after any merge and are not restated as if they were; they stand at `94de63ff5` in the spec. Neither was the mutation battery re-run at the third merged head: M4 and M5 stand at the second, and what was re-checked at the third is the throw site `op_provider.cpp:589` and the printed `id 140`. ## Three limits this change does not clear - **No CUDA arm and no CUDA gate.** Nothing registers `kRmsNormGroup` for any device but `kCPU`. An arm written on this CPU-only host could not be gated on it, and an ungated kernel is worse than an absent one. - **Toy width only.** The group widths exercised are 4, 5 and 6, not the 2560 the released config uses, so the f32 sum-of-squares accumulator is gated at toy width and the accumulator-width decision is left explicitly to the CUDA arm. - **The battery's tree.** M1, M2 and M3 were measured at `94de63ff5`, not at the head that lands. They were re-verified there by an independent fresh reviewer with exact count matches, M6 was re-measured on the final head when its refusal string changed, and M4 and M5 were re-run again after the merge. So the three value mutations rest on a reviewer's re-run rather than on a fresh one here, and that is stated rather than implied. ## This lands UNREACHED `vt::RmsNormGroup` is reached at this merge commit only by its own suite. Its four sites are the declaration in `include/vt/ops.h`, the dispatcher in `src/vt/ops.cpp`, the CPU kernel `RmsNormGroupKernel` in `src/vt/cpu/cpu_ops.cpp` and the name entry in `src/vt/op_provider.cpp`. No production entry point calls it: `ModelRegistry::Forward` is the only one this architecture has and `ForwardQwen4ExpForConditionalGeneration` still refuses by name, so the PLE block that will hold the three norms does not exist to call it from. The wiring is owned by row `MODEL-MM-QWEN4-EXP` under [mudler#2031](mudler#2031), tracked by campaign [mudler#1978](mudler#1978) and gated by [mudler#2249](mudler#2249). `.agents/specs/qwen4-exp-flash-next.md` lists it under `## Owed`, together with the CUDA arm. ## Landing order: after W5d-2 NO `.agents/issue-index.md` ROW IS APPENDED HERE, deliberately. mudler#2249 covers five prerequisites and this change is item 1; the sibling wave W5d-2 carries the SINGLE row for that issue. The file carries `merge=union`, so a second row for the same number would merge cleanly into a DUPLICATE that `check-agent-record.py` refuses while `check-issue-index-append-only.py` refuses its removal. **W5d-2 HAS NOW MERGED, AS `3ed2378a3`, AND THE DEPENDENCY IS SATISFIED.** This paragraph previously read "W5d-2 IS NOT MERGED — `origin/main` carries ZERO mudler#2249 rows as this is written", and both halves are stale: `origin/main` carries EXACTLY ONE mudler#2249 row, W5d-2's own, and it is the only one this issue may ever have. The ordering constraint it stated is met rather than dropped — this pull request lands after W5d-2, which is what it asked for. It still appends no row, so the index in this branch is byte-identical to `origin/main`'s: 865 rows, zero duplicate row ids matched on `^| \[#NNNN\]`, and exactly one `mudler#2249` row. Per [mudler#2266](mudler#2266), that set is resolved against the MERGE BASE and not from a row count, because a count cannot tell "union dropped my row" from "the other side added one". ## What else rides here `origin/main` kept moving under this branch, so it was merged back FIVE times: `1bc16ca3c` + `37fbccea8` first, then `fa9903b86`, `3ed2378a3`, `2e9f4d88d` (`QUANT-GGUF-IQ-VECDOT`, mudler#2247) and `7d41de27d` (`PERF-LAGUNA-GROUPED-GEMV` W1, mudler#2250), the last of which landed while the previous merge's gate was still running. `scripts/agent-preflight.sh` skips both trailer gates whenever `origin/main` is not an ancestor of HEAD — "this branch is behind it and the trailer gates did NOT run" — and a skip is not a pass. Every merge commit carries a written message rather than `git merge --no-edit`, because `check-commit-trailers.py` walks `rev-list` without `--no-merges` and a default merge subject has no trailer block and no non-force repair. MERGES ONE, TWO AND FIVE carried no compiled input: `git diff --name-only` over each delta returns nothing under `src/`, `include/`, `tests/` or `third_party/`, so nothing they touch can reach what this wave builds and the suite counts are not restated as if they had been retaken. **THE THIRD AND FOURTH ARE DIFFERENT**, which is why this section is rewritten rather than extended. The fourth, `2e9f4d88d`, moves `include/vt/dtype.h` and three `src/vt/cpu/cpu_quant_*` files, which are inputs to the translation units these suites link, so the build directory was reconfigured and every named target rebuilt: configure rc 0, build rc 0, and all eight counts byte-for-byte unchanged. The third: `fa9903b86..3ed2378` is fourteen commits and it moves both files this wave edits — `include/vt/ops.h`, where W5b-6 rewrote the `vt::Qwen4ExpGatedResidual` contract comment in place (mudler#2218 via `a6f933b81`), and `tests/CMakeLists.txt`, where W5d-2 registered `test_qwen3_5_mrope`. Git merged both without a conflict and both are purely additive against `main`. The one consequence a conflict-free merge could still have hidden is an id shift, because `kRmsNormGroup` is appended before `kCount` and so is every other new op: nothing on `main` appended an enumerator, and the merged header compiles to `kRmsNormGroup == 140`, `kCount == 141`, so the `id 140` M5's refusal prints is still this op's id. `.agents/specs/qwen4-exp-flash-next.md` conflicted on that merge and was resolved as a SET against the merge base `fa9903b86`, not by taking either side, because AGENTS.md `## Records` refuses an automatic three-way merge of a keyed record. Both sides added at the SAME two anchors and BOTH additions are kept: `main`'s `## Mutation record — W5b-6` and `— W5d-2` sections stand beside this branch's `— W5d-1`, in landing order, and `main`'s W5d-2 `## Owed` entry stands beside this branch's two. **A FIGURE IN AN EARLIER REVISION OF THIS BODY WAS MEASURED AT THE WRONG POINT AND IS CORRECTED HERE.** It said `git diff origin/main` over the resolved spec is "two hunks and ZERO deletions". That was true when it was taken — after the SET resolution, which is purely additive — and false by the time it was written, because the `## Now` and `## Owed` survey repairs below edit `main`'s own text. The true figure is **FIVE hunks and TWELVE deletions**, and all twelve are enumerated and intended: the `## Owed` item 1 heading, the two `## Owed` item 5 heading lines, and the nine-line `## Now` "five things are absent from `main`" paragraph. Each is replaced by text carrying the original fact TOGETHER WITH the correction, so no key `main` authored is dropped — which is the property that mattered, and it still holds. The count is stated at the HEAD that lands rather than at the point it was convenient to take. ## THE MERGE FALSIFIED PRODUCT OUTPUT, AND THAT IS WHY THIS TOUCHES A FIFTH FILE Tracked as [mudler#2288](mudler#2288), filed for traceability once the defect was found and **fixed in this same flow** rather than deferred, per AGENTS.md "Every change starts from an issue". `src/vllm/model_executor/models/qwen4_exp_registry.cpp` was not in this pull request before the third merge. The refusal thrown by `ForwardQwen4ExpForConditionalGeneration` enumerated FIVE missing prerequisites, and two of them are now present. Item (1) was "a standalone grouped RMSNorm op — the PLE block needs three and the only grouped reduction in this tree is fused inside vt::Qwen4ExpGatedResidual". That is `vt::RmsNormGroup`, which is what this pull request adds. Leaving the clause would have landed a commit whose own product output denies what the commit adds. Item (5) was "a mRoPE cos/sin builder with external linkage — qwen3_5.cpp's BuildMropeCosSinHost is static". It stopped being true at `3ed2378a3`: `BuildMropeCosSinHost` is declared at `qwen3_5_mrope.h:55` and defined without `static` at `qwen3_5.cpp:9475`. W5d-2 corrected the row spec's prose for this and did not correct the string, so the refusal has been naming a finished seam since it merged. Both clauses are removed rather than reworded, because a refusal enumerates what is MISSING and a present item is not missing; the message now names the three that remain. WHAT PINS THAT STRING was checked rather than assumed. `tests/vllm/models/test_qwen4_exp_scaffold.cpp:767` drives the hook with a foreign handle and asserts five substrings — "Qwen4ExpForConditionalGeneration", "forward is not ported", "W2", "W4", "mudler#1978" — plus the ABSENCE of "was not produced by". All five survive and the absence still holds; the suite is unchanged at 12 / 296. The two removed clauses are prose no assertion reads, which is the point: the gate holds the refusal REACHABLE and names the owing waves, it cannot hold the enumeration TRUE. The same two corrections ride into the spec's `## Now` and into its five-item survey under `## Owed`, where items 1 and 5 are marked CLOSED with the landing wave named and the survey argument kept, and a one-line reference sits at the defect site itself so the next reader of `qwen4_exp_registry.cpp` meets the issue number where the mistake was. **WHAT mudler#2288 DOES NOT CLOSE IS RECORDED UNDER `## Owed`**, because repairing two instances of a recurring class and leaving the class open is visible debt, not success. `test_qwen4_exp_scaffold.cpp:767` pins that five substrings are PRESENT and can never pin that any of them is still TRUE, so a refusal enumerating finished work satisfies every assertion — a spelling gate, not a truth gate. This is the third instance on this row in one day (mudler#2276 for the paged QSA consumer, mudler#2254 for the opposite polarity, a refusal that UNDERSTATED what refuses), and nothing mechanical prevents a fourth. The durable fix is a truth-linked check, or a convention that the refusal enumerates issue numbers rather than prose; both are larger than this flow and belong to whoever owns the reachability convention, so the residual is given a named home in the spec rather than living only in the issue. **NO `.agents/issue-index.md` ROW FOR mudler#2288 IS APPENDED HERE**, and that is a coordination decision rather than an omission — the same one this pull request already makes for mudler#2249. Three other row branches are appending to that `merge=union` file right now, and a fourth concurrent writer is the shape that produces the duplicate [mudler#2266](mudler#2266) documents. The row is being appended separately. The spec lists mudler#2288 under `## Owed`, which is what `check-agent-record.py` needs to resolve that row's owner when it lands. Fixes mudler#2288 A FALSE ALARM IS RECORDED IN THE SPEC RATHER THAN BURIED, because it came within one commit of landing a duplicate index row. `origin/main` is a shared ref in a shared checkout and another session fetched it mid-repair. Comparing the merged index against the ref after that fetch read exactly like the `merge=union` driver's known tail-row drop — 855 rows against the ref's 856, the missing one at the tail — and a commit was written to "restore" it. The merge had dropped nothing: the row arrived with `fa9903b86`, not yet merged, and appending it by hand would have put a SECOND mudler#2220 row on `main`. That commit was dropped before it left the worktree. `check-issue-index-append-only.py` returns rc 0 on the pre-restore AND post-restore heads, so the gate would not have caught it either way; what caught it was diffing the row-number list against the ref and then asking which commit AUTHORED the row. The index is byte-identical to `origin/main`'s copy at every merge, including this one: one mudler#2220 row, ONE mudler#2249 row (W5d-2's, which arrived with `3ed2378a3` — this paragraph used to say zero), and no duplicated issue number anywhere in the file. The `kernel-matrix.md` anchor repair rides here because this change is what made it stale: 52 new lines in `cpu_ops.cpp` moved `DFlashBlockAttentionKernel` from :2929 to :2979, re-verified at the merged head. The `## Owed` entry's "four backends that already register `kRmsNorm`" is corrected in the same paragraph the review sent back: six register it, enumerated with anchors, which is five besides `kCPU`. `include/vt/ops.h:634` says "five backends" and is deliberately NOT corrected — its next clause is "so a CUDA or Metal caller would get a whole-row norm back", so it counts the five OTHERS and agrees with the enumeration. That reconciliation is a spec edit and touches no product file; the separate product edit this pull request now carries is the refusal string above, which is a different correction for a different reason. One further clause in `include/vt/ops.h` is annotated rather than corrected: the `kQwen4ExpGatedResidual` rationale opens "There is no ungated per-group RMS norm", which is the sentence that produced this op and which this op makes historical. The words stay, because this wave's own comment quotes them, and a following sentence records that `kRmsNormGroup` now exists while the other four primitives that rationale names still do not — so the fused-family-op argument is unchanged. `test_cpu_x86_llamacpp_floor::test_a_contended_leg_is_discarded_and_never_summarised` is [mudler#618](mudler#618): load-dependent, red on a busy box, and untouched by this diff. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…d tensor-core tiling does not apply to decode (mudler#2250) (mudler#2297) Closes the Laguna decode-speed campaign that came out of the W7/W11 attribution, and closes it with an answer rather than an absence of one. Tensor-core tiling was the obvious reopening candidate: it is the one mechanism that shortens the unpack dependency chain rather than feeding more warps to it, and `cuda-keepquant-gemm.md` defers it explicitly. Reading the upstream reference closes it for this shape. llama.cpp's MoE dispatch at `b10451` takes MMVQ when `ne2 <= get_mmvq_mmid_max_batch`, and for Q4_K and Q5_K — exactly Laguna's expert dtypes — the Turing+ table falls through to `default: MMVQ_MAX_BATCH_SIZE`, which is 8. Laguna decodes one token, so upstream takes MMVQ: warp-per-output, the structure we already have. It reaches for tensor-core MMQ only above batch 8, which is prefill. There is therefore no upstream existence proof that MMQ wins at decode on this dtype, and the shape agrees — MMQ tiles want at least 16 rows while the decode grouped GEMM has P = 10 top-k experts, so six of sixteen rows would be padding. Tensor-core tiling stays a real deferral for PREFILL, which would need its own attribution first because Laguna's measured gap was decode. So the ranked list is spent: lever #1 measured and shipped default-OFF because it moves a token on 6 of 6 prompts; lever #2 refuted at the counter, latency-bound at 101.5% occupancy with memory SOL 28%, which refutes the bandwidth, occupancy and dp4a levers before any is attempted; lever #3 demoted by W11 itself at ~0.02 s/tok. With the Q8_0 kernel's five refuted structural bricks and its recorded MEASURED FLOOR beside them, the plan is complete. The conclusion is specific rather than a shrug: Laguna's remaining decode cost is a memory-latency dependency chain in the keep-quant unpack, and nothing on the enumerated list moves it. What that does NOT say is that Laguna is at its floor for all time — a new lever needs a new mechanism and a fresh attribution, not another pass at this list. Gates: `check-agent-record` rc=0. No product code; this is a record. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
… its own refusal asks for (mudler#2286) (mudler#2287) `deepseek_v4.cpp`'s own refusal ends with "The DSA port itself is OWED and has no owning row". This creates it. Closes mudler#2286. The composition is the blocker between a DeepSeek-V4-Flash artifact that **loads** (mudler#2186/mudler#2283) and one that **runs**: the forward refuses by name on the 21 `compress_ratio == 4` layers, and dense MLA is not a substitute at any sequence length (mudler#1964). **What had no owner is narrower than "the DSA port."** Both kernel primitives already have rows — `KERNEL-ATTN-DSA-SPARSE-INDEX` and `KERNEL-ATTN-DSA-COMPRESSOR`, both at `SPIKE`. Missing was the code that ASSEMBLES them into `AttentionBlock`, which upstream writes as three layer shapes selected by `compress_ratio` (`attention.py:454-533`), every one ending at `forward_mqa` then `_o_proj`. ## Three findings that change the shape of the work **1. The 3-way stream overlap is PERFORMANCE, not correctness.** `attention_impl` dispatches through `execute_in_parallel` with `enable=aux_streams is not None`, and ROCm runs the same work sequentially. So a sequential first wave MIRRORS upstream rather than diverging from it. Recorded explicitly so a later reader does not "restore" the overlap believing correctness depended on it — this is the largest scope reduction available on the row. **2. `coff = 1 + (compress_ratio == 4)` is a per-token ROLE chosen by offset within the gathering window**, not a property of the token. The state cache holds two head-sized rows per token and the gather picks a half by `head_offset = (tokens >= COMPRESS_RATIO) * HEAD_SIZE`, with emission at boundary tokens only, so a token in the overlap belongs to two windows **with a different role in each**. That is the whole of what the refusal means by "never recoverable from the tensor alone" — and since our loader already materializes the width correctly (mudler#1970), this is a **forward** change, not a loader one. **3. The compressor is two stages with the second boundary-gated**, and its RoPE is exactly specified: GPT-J style, `is_neox_style=False` (interleaved, not split-half), on the last `rope_head_dim` elements, at position `(positions // compress_ratio) * compress_ratio`. ## Two constraints recorded because getting either wrong is silent - **Hard ordering.** The row CONSUMES `KV-DSV4-MULTICACHE` (mudler#1925); its W1/W2 landed, and this row's W1 cannot start before that row's W3 hands the forward the cache. - **Not gateable at or below 512 tokens.** The one arm that caches today runs `dsa_dense` with indexer and compressor forced OFF and is exact only while `seq_len <= index_topk` (=512), so every gate on this row must exceed it. The spec also records that `config.json`'s `compress_ratios` has **46** entries `{0:5, 4:21, 128:20}` while older records say "43 layers" — 43 is the trellis shard count — and makes W1 reconcile which number each claim means before writing code. ## Why spec-only `AGENTS.md` requires the spec before implementation, and a capability this size benefits from agreement on scope before implementation waves start. W1 is blocked on the multicache row regardless. ## Gates `scripts/agent-preflight.sh --staged`: **110 checks ok**, rc=0. Two checker interactions worth naming, both resolved in this change rather than worked around: - `check-agent-record` required the structured `Upstream chain` / `Our baseline` / `Port map` / `Tests to port` sections; they are written with real content, and `Tests to port` says explicitly that a green port of the three upstream tests does **not** by itself gate the composition. - The new row's `## Gates` section carries runnable commands, so the row joined the gated-runnable population and `RUNNABLE_BASELINE` in `scripts/check-gate-commands.py` is re-pinned in this same change, as that ratchet requires. - That edit made `check-pr-size` demand semantic mutation evidence for a `governance_checker` path, which is the right demand even though this particular edit is a data re-pin the checker itself instructed. **The evidence is a test that pins the REASON rather than the membership**: `MODEL-DSV4-DSA-COMPOSE` must be runnable *because its spec's Gates section names commands that can fail*, and it is scoping-only, so a reader finding a spec with no implementation in the runnable set does not conclude the ratchet is broken. MUTATION-PROVEN, and it proves the test is not redundant with the exact-equality pin beside it: removing the runnable command from the spec's Gates section AND the baseline entry together leaves `test_the_baseline_matches_the_shipped_record` **green** (both sides moved consistently) and takes **only** the new test red — 1 failure of 57, restored byte-for-byte after. That is the case where a row silently leaves the gated population, which nothing else here could see. `tests/scripts/test_check_gate_commands.py`: **57 tests, OK**. `scripts/check-pr-size.py --base origin/main --head HEAD`: rc=0. No `src/`, `include/` or `tests/` change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: CLAUDE:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…s the CPU guards cannot see (mudler#2274) (mudler#2304) fix(SPEC-DFLASH2): read the DEVICE bounds back, which is the one class the CPU guards cannot see (mudler#2274) `DflashBlockPagedAttention` has two checks that read `seq_ext` and `slot_map`, and both are guarded on `device.type == kCPU` because they dereference the tensors. On CUDA neither runs — and CUDA is where mudler#2274's illegal access is. The comment on the slot-map check names the variant that therefore goes unchecked: the host values were right and the UPLOAD did not land on the tensor this call reads (a stale graph buffer, a copy that went elsewhere) That is now the only class left. `ddd527f3f` added four shape-only checks that run on every backend, and on the failing configuration they are SILENT: the extended bound fits the pool, the last write slot fits it, and the block table addresses it. So the host accounting this call derives is correct, and a host/device divergence is what remains. `VT_DFLASH_BOUNDS_DEVICE=1` downloads the device `seq_ext` and the slot-map endpoints and asserts they equal the host derivation. OFF by default and read once: the download synchronizes, and this call sits on the no-sync path the whole paged route exists to keep, so a diagnostic that silently degraded production would be the wrong trade. It is a DETECTOR. mudler#2274 is not root-caused, and seven candidates are now excluded by measurement rather than by argument — the CUDA graph, the FA2 block lane, merged QKV, the entire seam adoption of mudler#2207 (built `c9b2049bc~1`, still faults), FA2 being compiled out, the `max_seq_len` replay staleness, and the host bounds. `VT_DFLASH_PAGED=0` remains the only configuration that completes. Also adds the index row mudler#2274 has been missing since it was filed. The issue was opened and a fix for it landed without one; `check-agent-record` tolerates a recorded number of unowned rows, so nothing failed, but the obligation is real and the row carries the exclusion table so the next reader does not re-test what is already excluded. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…am, not in each caller (mudler#2313) `RunMoeBlockPlaced` refused to place a layer whose routed experts are fp4-resident. W3c moved every architecture onto the shared `RunMoePlaced` / `RunMoePlacedPair` seam and did not carry that refusal across, so the old helper became dead code and the live path accepted the arm. Placing an fp4-resident arm uploads every expert at load and then computes on the host across the bus, which is slower than not placing at all — the opposite of what the user asked for. A token gate cannot see it, because the tokens stay correct and only the placement is wrong, so nothing in this tree would have reported it. The refusal now lives on the seam as a `placeable` / `unplaceable_reason` contract rather than in each caller, so a newly wired architecture inherits it instead of having to remember it. It fires only when a placement is actually in force, leaving an ordinary unplaced load untouched; a guard that fired there would break every load, which is the opposite failure and just as bad. Proved by a COMPILING mutation: with the guard rewritten never to fire, `test_device_placement` goes red at 1 case and 2 assertions. The first mutation attempt failed to compile under `-Werror` on the now-unused parameters, and the stale binary reported 19/19 SUCCESS — a passing mutant that proved nothing, so the mutant build's rc=0 is part of the evidence. Also documents `VT_DFLASH_BOUNDS_DEVICE`, which mudler#2304 landed undocumented and which made `check-env-doc` red on `main` itself. That is a base failure every later branch inherits, found while gating this change and fixed in the same flow. FOLLOWING_AGENTS_PROTOCOL Closes mudler#2309 Closes mudler#2312 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…pec, and the wave splits because W3 left no attention block to call (mudler#2238) feat(MODEL-MM-GLM53-FLASH): W5 lands the 288+1 MoE and the KV-cache spec, and the wave splits because W3 left no attention block to call `MakeGlm5NextKVCache` returned a refusal. It now returns three groups, reached through the production `make_kv_cache` factory hook, and that is the first thing on this row a user can arrive at. The MoE lands beside it and is not reached. **Review found five defects and all five are repaired; see "The review round" below.** One was real -- the KV spec's convolution dtype came from a resolver its own comment named and the code never called, and no assertion read it -- and four were prose this branch's diff had falsified. The branch is merged up to `origin/main` at `866075b2f`. **Rebased onto `origin/main` at `e541be98a`, and the merge falsified six of this branch's own claims.** W5c (mudler#2242) landed the weight tower while this was in review, so the model LOADS: the GGUF refusal this branch corrected no longer exists, O24's "unreachable BY CONSTRUCTION" premise is retired, O18/O19 are taken on `main` and this branch's entries are renumbered O23/O24, and the artifact reading below is history rather than state. Each repair is named where it lands rather than left for a reader to find. The `sanitize-cpu` red was a `main` defect fixed by `6e805abcf`, not this branch's -- see below. ## The wave split, and the reason is a gap rather than a scope decision mudler#2223 named four deliverables. Two landed. The other two -- the per-layer control flow and the assembled `Glm5NextTextModel::Forward` -- need an assembled `Glm5NextTextAttention` for the DSA arm to call, and there is none. W3 landed `SelectIndexerTopk`, the indexer's SELECTION, and relaxed `MlaBlockDims::Validate`; it landed no block over either. `q_a_proj` / `q_a_layernorm` / `q_b_proj`, `kv_a_proj_with_mqa`, `kv_b_proj`, `expand_kv` (`modeling_glm5_next.py:1136-1153`), the attention itself and `build_attention_mask_from_topk` (`:1218-1257`) are all unwritten. A decoder layer whose sparse arm throws is a control-flow shell with one live branch, so the remainder is W5b (the layer and the forward) and W5c (the weight tower and `load_weights`), each with its own gate row in the spec. ## The MoE binds; it does not reimplement The router is `vt::MoeRouterTopK`'s grouped `noaux_tc` arm and the epilogue is `deepseek_v4::ClampedSwiGLU` at `alpha=1, beta=0`, which is `_apply_gate`'s "Simple swiglu instead of alpha" line for line (`modeling_glm5_next.py:137-142`). Gated at the PUBLISHED 288 experts / top-8 against goldens RUN out of an unmodified `Glm5NextTextTopkRouter.forward` at transformers v5.16.1. **The selection is asserted as a SET and the margin is printed.** Top-k error is bimodal: the routing is either the oracle's set or a different one, and a different set can carry values that are numerically close, so a tolerance on `topk_weights` passes a wrong routing. The gate asserts set equality, asserts the weight AT each id, and prints the gap between the last accepted and the best rejected selection score -- 1.84e-3, 2.76e-2, 8.58e-3, 1.98e-2 over four tokens. It compares by id rather than by position on purpose: our seam emits descending with the lowest index winning ties, upstream calls `torch.topk(sorted=False)`, and positional comparison would gate this tree's determinism convention against torch's implementation detail. Five silent-failure axes are each a killed mutation rather than a comment: sigmoid scoring swapped for softmax, `routed_scaling_factor` dropped, `norm_topk_prob` dropped, the `e_score_correction_bias` dropped (it SELECTS while the unbiased score WEIGHTS), and the factor applied TWICE by also handing it to `vt::MoeCombine`. The clamp case exercises a row that reaches past the limit in BOTH halves, because a fixture that never leaves `[-limit, limit]` passes a port with no clamp and one that only exceeds `+limit` passes a port that clamps the gate symmetrically. ## Three KV groups, and three numbers a plausible port gets wrong | # | layers | spec | geometry | |---|---|---|---| | 0 | 11 DSA | `MLAAttentionSpec` | head **512** | | 1 | 34 KDA | `MambaSpec`, 2 states | conv `[24576, 4]` at the CACHE dtype (bf16 default) + recurrent `[64,128,128]` f32 | | 2 | 11 DSA | `MLAAttentionSpec` | head **257**, `compress_ratio` **1** | **512, not 576.** Every DeepSeek variant and Kimi-Linear publish `kv_lora_rank + qk_rope_head_dim` = 576. Here `qk_rope_head_dim` is ZERO and upstream requires it to be, so the latent row is 512 and reusing 576 over-allocates by 12.5% with nothing downstream reading the difference. **`conv_kernel_dim`, not `conv_kernel_dim - 1`.** The reference allocates the conv state at the full kernel width -- `LinearAttentionLayer.lazy_initialization` builds `torch.zeros((*shape[:-1], conv_kernel_size))` (`cache_utils.py:1015-1024`) and the layer passes `conv_kernel_size=self.conv_kernel_size` (`modeling_glm5_next.py:669-671`) -- and `causal_conv1d_update` reads `state_len = conv_state.shape[-1]` back off it, so the slack column is part of the contract. `kimi_linear_registry.cpp:156` publishes `K-1` for ITS model; copying that across hands the runner a cache one column short of what the layer reads. **257, not 128.** `PackIndexerStates` stores `concat[k(128), gate_scores(128), valid(1)]` per token (`:798-801`). The DeepSeek-V4 parent stores the key alone. And `compress_ratio` is 1, not `index_kpool`: the k-pool compresses at READ time inside `GetPooledStates`, which is the opposite of MODEL-MM-QWEN4-EXP's QSA side cache where the compression is in the store. ONE conv state, not three: the checkpoint's `self_attn.{q,k,v}_conv1d` concatenate into one grouped depthwise conv (`:620-628`). Group 2 must be an `MLAAttentionSpec` -- a `FullAttentionSpec` there is absorbed by the runner's leftover scan as the single `fa_draft` slot, `multi_cache_topology` stays false, and the side cache is published and never allocated, in silence. ## Reachability, proven twice The cases enter through `ModelRegistry::Resolve` and the `make_kv_cache` factory hook -- the pair `LoadedEngine::FromModelDir` uses -- and nothing constructs `MakeGlm5NextKVCache` by name. Unwiring `.make_kv_cache` reds `test_glm5_next_scaffold` at three `REQUIRE` sites. DELETING the row does not compile: `-Werror=unused-function` fires, so the toolchain proves the factory row is the only reference to the function. The MoE is NOT reached and this body says so: nothing calls it until W5b assembles the layer. That is **O23** in the spec's `## Owed`, with the row and the issue named, per AGENTS.md "Nothing lands dead". One more claim the split falsified, in `main`'s own prose rather than this branch's: W4 recorded "that code is not reached from any production entry point (O16); W5 owns the wiring". W5 no longer owns it -- this wave lands the MoE and the KV-cache spec and explicitly not the decoder layer -- so O16's wiring owner is corrected to W5b (mudler#2241). Left alone the spec would state both halves of a contradiction, because O23 already says W5b for the equivalent MoE wiring. **O18 and O19 were taken on `main` while this branch was in review**, for the per-layer config arrays and the CUDA keep-quant disclosure, so this branch's two entries are renumbered O23 and O24 rather than colliding. O24 is retired in place rather than deleted: as written it said that while `load_weights` refuses, `ModelRegistry::Forward` is unreachable BY CONSTRUCTION and no wiring alone can discharge O15, O16, O17 or O23. W5c (mudler#2242) ended that -- `load_weights` now returns a real `LoadedModel` -- and the entry is kept because it is the reason the numbering skips and because it was live when the tests under it were written. ## What the published artifact did on 2026-08-29, and why that reading is now history **This section is a dated measurement, not a current state, and the merge is what moved it.** mudler#2245 landed the IQ2_XS and IQ4_XS decoders, mudler#2247 made both keep their blocks, and W5c (mudler#2242) now resolves all 1383 backbone tensors of this artifact. The reading is kept because the census under it is still the census and because three of the facts it established are what the MoE was built from; only the "decodable here" column has moved. The spec, the claim file and the index row each say so in place rather than carrying the old sentence forward. `unsloth/GLM-5.3-Flash-GGUF` rev `d425e572f`, arm `UD-Q2_K_XL`, staged read-only, run through `LoadedEngine::FromModelDir` on 2026-08-29: ``` REFUSED: gguf: tensor "blk.3.ffn_gate_exps.weight" has unknown ggml type id 17 ``` It opens the file, resolves `general.architecture = glm5next` against our own registration, walks the 4-way split into shard 2, and stops on a TENSOR TYPE. Type 17 is IQ2_XS. A census of every tensor header across all four shards reads F32 638, Q8_0 346, Q2_K 2, Q3_K 1, Q4_K 1, Q5_K 181, Q6_K 117, IQ2_XS 82, IQ3_XXS 41, IQ4_XS 3. **"Q2_K" in the arm name is a floor, not a format.** Unsloth Dynamic mixes eight encodings and only two tensors in the whole file are Q2_K; the 288 experts are IQ2_XS and IQ3_XXS. Six of the eight were undecodable here when this was read -- Q3_K/Q4_K/Q5_K per O8 and IQ2_XS/IQ3_XXS/IQ4_XS per O5 -- and that is the half of the reading which has since changed. O8 has been corrected on `main` to say what it always meant: it is about the ENCODERS, and the matching decoders have been present since the k-quant port. Three facts from the same read confirm ports made blind. `expert_shared_feed_forward_length` is 2048, so the shared expert is `moe_intermediate_size` and not `intermediate_size` (12288). `ffn_gate_inp` is F32 on all 43 sparse layers, which is the fp32 router GEMM as an on-disk fact. And `ssm_conv1d_{q,k,v}` are three separate tensors per KDA layer, which is why the cache is one grouped conv state. `glm5next.layer_types` is absent and the schedule is `attention.head_count_kv`, exactly as mudler#2177 measured -- but that question is not yet reachable on this file, because the type refusal preempts the config read. mudler#2177 keeps it: every case here DECLARES its schedule, so nothing in this wave rests on the `idx % 4 != 3` coincidence. ## Three refusal messages that named landed waves as owing (mudler#2230) The forward refusal read "W3 the NoPE MLA block -- `MlaBlockDims::Validate` still refuses `qk_rope_head_dim == 0`", which W3 made false by relaxing exactly that validator; it named W2's forget gate and W4's collapse as owed too, both landed. The GGUF refusal said no `.gguf` of this model exists anywhere. The KV refusal said the KDA layers carry three separate conv states. **Two of those three surfaces have since been overtaken, and the third is the reason this merge needed a decision rather than a side.** W5c deleted the GGUF refusal outright, because the GGUF arm now loads, so that correction survives as a test negative rather than as a message. The forward refusal is the interesting one: `main`'s copy of it still named W2, W3 and W4 as owing primitives that landed in `199c44578`, `e511a614b` and `6c715de00` -- the file had only ever been touched by W1 and W5c, so no wave in between updated it -- while this branch's copy still said W5c owed the tower. **Taking either side whole would have RED the suite**, because this branch's own test asserts the message does NOT contain "MlaBlockDims::Validate still refuses" and `main`'s side does contain it. The merged message names the four primitives as LANDED and the load as DONE, and it names `load_weights` positively because the suite asserts that string and both sides had dropped it. `test_glm5_next_scaffold.cpp` asserted all three sentences, so the gate passed BECAUSE nothing had corrected them -- it was pinning stale text rather than checking it. The repair adds the negatives, so a revision that reintroduces either claim reds. `docs/FEATURES.md` and `docs/USAGE.md` carried the same two stale claims and are corrected in the same pull request. ## The sanitizer red was never this branch's defect `sanitize-cpu (address,undefined)` failed one case out of 661, `test_exl3_native_loader`, on a UBSan misaligned load: ``` include/vllm/model_executor/models/dense_weight_loaders.h:640:80: runtime error: load of misaligned address 0x7f8744d87a89 for type 'const short unsigned int', which requires 2 byte alignment #0 vllm::dense_loaders::LoadF16AsBf16Direct(...) dense_weight_loaders.h:640 #1 vllm::LoadLlamaForCausalLMWeights(...) llama_weights.cpp:143 #2 DOCTEST_ANON_FUNC_15 tests/.../test_exl3_native_loader.cpp:242 ``` A header, a model and a test this branch does not touch, in a code path it does not reach. `LoadF16AsBf16Direct` read the safetensors payload through a `reinterpret_cast<const uint16_t*>`, and a safetensors tensor offset is the running byte total of everything before it and carries no alignment guarantee, so an odd offset makes the widened load undefined. **`6e805abcf` (mudler#2229) fixed exactly this on `main`**, replacing the cast with `vt::LoadUnaligned<uint16_t>` and quoting the same UBSan sentence in its own comment. That commit is on `main` and was not on this branch, which is the whole of the difference: the same job on `main` at `a36add6a8` ran the same test as mudler#482 and passed 661/661. **Merging is the fix.** No sanitizer setting was weakened and no assertion was deleted. ## The review round: five findings, five repairs **The conv dtype came from a resolver the comment named and the code never called.** `MakeGlm5NextKVCache` said its two recurrent dtypes come from "the SAME resolver every other hybrid in this tree uses". Both were literals. The sentence was copied from `qwen4_exp_registry.cpp:378-382` without the call, and this file does not include `qwen3_5_internal.h` at all. Calling that resolver would have been the wrong repair. `ResolveMambaSsmCacheDType` mirrors `_mamba_state_dtype` (`mamba_utils.py:96-108`), the Mamba/GDN calculator, and this model's linear layers are KDA. The mirror is `kda_state_dtype` (`:130-137`), which returns `(get_kv_cache_torch_dtype(mamba_cache_dtype, model_dtype), torch.float32)` and takes no `mamba_ssm_cache_dtype` argument at all: the conv half follows the paged-KV storage dtype, the recurrent half is f32 unconditionally, and `mamba_ssm_dtype` is not read. `kimi_linear_registry.cpp:161` already publishes that pair for the other KDA model here. So `conv_dtype` is `v1::ResolveKvCacheDType()`, which also ends this model being the only cache in the tree that would not move under a `VT_KV_CACHE_F32` A/B. Ignoring `mamba_ssm_dtype` is kept and is now argued rather than denied: honouring a key the KDA calculator does not read would be an invention, and a `bfloat16` value in some future config would silently halve a state upstream keeps in f32. The default does not move, so no golden shifts. **`dtypes[0]` was asserted nowhere.** The gate checked only `dtypes[1] == kF32`, so mutating the conv dtype passed green, which is how the defect shipped. Asserting the bf16 default alone would gate nothing either, because a hardcoded `kBF16` passes it. The case now builds BOTH arms explicitly, setting and restoring the environment around two `make_kv_cache` calls rather than reading it ambiently, through the portable `tests/support/test_env.h` shim. Both mutations are killed on a real build: hardcoding `kBF16`, the code that shipped, reds `f32_mamba->dtypes[0]` at `:815`, and hardcoding `kF32` reds `bf16_mamba->dtypes[0]` at `:814`, each 2659/2660 at rc=1. The file was restored to sha256 `e4368554a31e1` and reruns 2660/2660. **Four claims this branch's own diff had falsified.** O10's heading, its body and `## Now` all assigned the KV-cache spec to W5b (mudler#2241) alongside the forward, while this branch lands `MakeGlm5NextKVCache` and wires it into `kGlm5NextFactory`. `glm5_next_loader.h`, `test_glm5_next_gguf_load.cpp` and the `model-matrix.md` row repeated it. **That is mudler#2230's failure, a refusal naming a landed wave sending the next reader to redo finished work, reintroduced inside the change that closes mudler#2230.** Each now names only the forward. `glm5_next_moe.h` said the MoE is not reached "because `load_weights` does (O10)"; `load_weights` has returned a real `Glm5NextLoadedModel` since W5c. The conclusion held, the reason did not: what refuses is `ForwardGlm5NextForConditionalGeneration` itself, because W5b's decoder layer does not exist. The claim file's earlier blanket exclusion of `model-matrix.md` is narrowed rather than quietly broken: the row's lifecycle state still does not move, and only the sentence this wave falsified was repaired. **Two nits.** `kimi_linear_registry.cpp:157` is the recurrent-state shape; the `K-1` conv publish both comments cite is `:156`. And `kOracle` was the only one of 44 golden symbols no assertion read, so the fixture was not pinned to its lane revision. It is asserted now the way `test_glm5_next_mhc.cpp:92` asserts its own, and moving it to `5.17.0` reds `test_glm5_next_moe` at 1613/1614. **The merge was taken twice.** `origin/main` at `21ef6f053` was RED on `check-env-doc`: mudler#2304 added `VT_DFLASH_BOUNDS_DEVICE` and documented it nowhere. `main` moved while the first merge was open and `866075b2f` documents it, so the second merge takes that repair instead of this branch filing a fix for another row's defect. `.agents/issue-index.md` auto-merged and was verified from the DIFF against the merge base rather than from its tail: 874 rows, 868 base plus this branch's mudler#2223 and mudler#2230 plus `main`'s mudler#2274, mudler#2286, mudler#2309 and mudler#2312, no duplicate id and no row dropped from either side. Nothing `main` brought touches this row (`git log e029857..origin/main -- '*glm5*'` is empty), so the merge falsified no claim here. ## Evidence, and which tree each number was taken on **Taken on the pre-merge tree, and not re-measured here.** RED first, on the same tree: a declared-surface stub plus the unmodified refusal gives `test_glm5_next_moe` 7/8 cases and 22/30 assertions failed, and `test_glm5_next_scaffold` 6 failed -- four on the stale forward message, two on the KV refusal. Green is 1613/1613 and 2530/2530; full `ctest` 649/649 passed with 6 pre-existing skips; SACRED inertness 35/35 across DeepSeek-V2/V4, Kimi-Linear, Kimi-K2/K3 and GLM-4/4.7 with non-zero assertion counts, checked rather than assumed because a zero-assertion suite is a skip wearing a pass. Sixteen negative mutations, each sha256-proved applied, built and restored byte-for-byte; all sixteen kill their gate. **Taken on this merge head (`866075b2f` merged in).** `scripts/agent-preflight.sh` is rc=0 with **zero skips** after the two commits land, and every record and script gate is green including `issue-index append-only`, `check-gate-commands`, `check-env-doc`, `commit-trailers` and `commit-style`; `check-commit-trailers.py --range origin/main..HEAD` and `check-commit-style.py` over the same range are both rc=0. The earlier run on `e541be98a`, quoted for its one failure below, had the same shape. On the `e541be98a` head one gate failed, `test_cpu_x86_llamacpp_floor`, and it was a box-load flake rather than a result -- it does not reproduce on this head, where preflight is rc=0: the file is byte-identical to `origin/main`, this branch does not touch it, and the harness itself names the cause in its own vocabulary -- it has a quiet gate, and the gate never opened: ``` waiting for quiet: 15s busy=158% builders=0 load=35.94 NO_QUIET_WINDOW after 30s (busy=114% builders=0 load=35.55 36.86 28.97) ``` Re-run in isolation on the same head it fails 2 of 10, both on `NO_QUIET_WINDOW` at load average 31-36, which is the same other-session build load that took the disk. The cases expect returncode 0 and 2 and get 4, the harness's own "I could not measure" code. This is the failure mode `.agents/` already records for this suite: it reds under box load. It is not a verdict on this tree. **Two targets were built and run on this head; the SUITE was not, and CI is the first thing that builds it.** Disk recovered enough in the repair round for a cold CPU-only configure (`VLLM_CPP_CUDA=OFF`, no `nvcc` on the box) and a `-j 2` build of exactly `test_glm5_next_scaffold` and `test_glm5_next_moe`. On the merge head they read 38/38 cases with 2660/2660 assertions and 8/8 with 1614/1614, both rc=0, and the two mutations below were killed on that same build. **That is two executables. Nothing else in `tests/` was compiled, no sanitizer or Windows arm was run, and the whole-suite counts quoted above stand only for the pre-merge tree.** The last green CI on this branch was `96e5c018e`, which predates both the W5c integration in `glm5_next_registry.cpp` and 849 changed lines of the scaffold test, so **CI is the first build of this head as a whole and this body claims no green it does not have.** The static check the earlier round ran in place of a build still stands and is worth keeping, because it covers the exact class of break a merge causes: every `find(...) == std::string::npos` negative assertion in `test_glm5_next_scaffold.cpp` was checked against the concatenated string literals of the merged refusal messages. All 14 hold, including the two the merge put at risk -- `"MlaBlockDims::Validate still refuses"` and `"the weight loader is not ported"` are both absent from every message. The positive assertions resolve against the merged literals too, apart from 19 that are composed at runtime (`"block_count is " + std::to_string(...)`) and belong to `main`'s own passing per-layer-config cases. Goldens are the RUN output of unmodified reference modules at transformers v5.16.1; the installed `modeling_glm5_next.py` hashes `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b`, the value W3 recorded. Every golden the fixture emits is now consumed by an assertion, including the selection margin and, since this review round, `kOracle` -- W3's review found a captured `kIndexScores` golden that nothing read, which let two scale defects pass 1602 of 1602 assertions, and `kOracle` was this wave's last unread one. GPU: PENDING. `dgx:gpu0` was held by other sessions throughout the window and `strix:gpu0` can hold neither the artifact nor a CUDA kernel. No result invented, and W3's committed CUDA arm stays unmeasured under O17. The artifact probe above needed no GPU and is a CPU measurement. Closes mudler#2223 Closes mudler#2230 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
mudler#2283) (mudler#2285) W1d's pull request body said "Closes mudler#2186", so merging `c9ad53fee` closed that issue. Two `## Owed` entries in `.agents/specs/model-dsv4-exl3.md` were written in the same change and name mudler#2186 as what tracks them — and both describe work W1d explicitly did NOT do: 1. **The ~97.7 GiB residency figure has never been observed.** It is arithmetic on the measured 108.59 / 26.64 split. The last real measurement (2026-08-28, `dgx:gpu0`, worker `rc-worker-4b8lj`, tree `525d2b991`) REFUSED, and nothing has re-run since W1d landed. 2. **The carried tower's BF16-sourced half is still widened to f32**, worth a further ~2.62 GiB. W1d took the FP8-sourced half only. A reader following either entry landed on a closed issue and would reasonably conclude the work was done. mudler#2283 is filed to own both, and the entries now name it — the first says explicitly that mudler#2186 closed with W1d, so the redirection is not mistaken for a second issue about the same thing. **Why this is record-only and still a legitimate unit of work.** AGENTS.md: a record-only pull request "is still correct when the record is the work: a stale row, a corrected pin, a newly filed gap." This is a newly filed gap. It is not a restatement of what the landing change already knew — W1d could not have cited mudler#2283, which did not exist until mudler#2186 closed. **What is NOT claimed.** Nothing is measured or fixed here. mudler#2283 stays open with both items owed, and the DeepSeek-V4 forward still refuses by name on the 21 `compress_ratio == 4` layers (mudler#1961, mudler#1970, mudler#1976) regardless of residency. ## Gates `scripts/agent-preflight.sh --staged`: **110 checks ok**, rc=0 — including `check-issue-index-append-only.py`. The `mudler#2283` row appears exactly once (counted explicitly: an append-only index duplicate is invisible to that checker, since two new rows are both "new"). No `src/`, `include/` or `tests/` change, so no suite is implicated. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: CLAUDE:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
Resolving the ops-seam conflict with @both concatenated both sides, but each side's function-closing brace had been unified by git into the single shared closer, so PermuteVHeadsKernel and its vt wrapper were left unclosed and the merge did not compile. Adds the two missing closers; no other content changes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
Records the post-merge correctness gates in the T38 evidence file and moves the spec's Now to the synced position. Both red gates are traced to upstream main rather than assumed: the keep-quant routing test tightened without the router changing, and test_placed_moe_roundtrip does not compile on upstream main itself. The acceptance body md5 reproduces the re-minted reference bit-for-bit under all 15 levers. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
VT_GDN_MERGED_BA_ROCM, VT_GEMV_SPLITK and VT_GEMV_WARPS are kernel-internal, default-OFF tuning switches whose behavior contracts live in the row spec and evidence files; they belong on the env-doc allowlist, not in docs/ENVIRONMENT.md. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
The four sync commits were rebuilt with commit-tree to add the bare FOLLOWING_AGENTS_PROTOCOL paragraph their first versions lacked; the merge kept both parents (upstream tip 3015aad unchanged) and no tree changed. This pins the new SHAs in the T38 record and spec, and records the two gate dispositions: check-env-doc repaired via the allowlist, while the +1 stale record anchor pre-dates the merge (identical set at 8e2f56c) and is owed separately. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
Adds two gate results the initial T38 evidence omitted: - ctest -R cross_device: test_backend_cross_device fails 2 cases (MoeSiluMul bf16 rounding, wvSplitK YTILE=2 OOB). Verified pre-existing by building and running the pre-merge head 8e2f56c in a fresh worktree — identical 5/80253 assertion failures. The campaign's focused gate regex never included cross_device. These are owed a separate fix unit, not a merge regression. - Near-tie adjudication: teacher-forced logprob-band ceremony (rocm-m4-oracle, band <= 500 mnats) under all 15 adopted levers. verdict=PASS, divergent=0, over_band=0, max_gap_mnats=0.000, body_md5 matches. The 9 non-bit-identical levers are covered by this all-levers-ON composite: zero per-step divergence from the reference under the merged build. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
Three pre-existing bugs surfaced when the focused gate was widened to
include cross_device, plus the record-anchor ratchet is repaired.
1. MoeSiluMul bf16 rounding (rocm_moe_router.hip): the GPU kernel
skipped the intermediate RoundThrough to the gate dtype that the CPU
oracle does (cpu_ops.cpp:545). For bf16 gate the silu value was not
narrowed to bf16 before the multiply, so the exact-equality check
in test_backend_cross_device failed. Add the intermediate
__float2bfloat16/__bfloat162float round-trip when Tg is bf16,
matching the CPU oracle and upstream's silu_and_mul.
2. wvSplitK output overwrite (rocm_skinny_gemm.hip): the kernel's store
loop unconditionally wrote YT output rows starting at offset m, but
when M is not a multiple of YT the last wave writes past the output
buffer. Shape {tok=2, k=256, feat=254} with YT=4 wrote C[254] and
C[255], corrupting the guard band. Add a bounds check
(m + y < M) before each store.
3. test_placed_moe_roundtrip (tests/): upstream commit 866075b
deleted RunMoeBlockPlaced as dead code (W3c moved all architectures
onto the shared RunMoePlaced seam) but left the W3b test behind.
The test references a symbol that exists nowhere on upstream/main.
Remove the dead test and its CMake entry.
4. Record-anchor ratchet: 29 stale citations across 4 matrix files
pointed to symbols that moved during upstream growth. All 28
relocated symbols updated to their new file:line; one renamed
symbol (LoadFromDir -> FromModelDir) updated in prose. Baseline
lowered from 28 to 5 (the 5 BROKEN anchors are pre-existing and
not in scope).
All fixes verified: ctest -R 'rocm|cross_device|quant' 24/26 passed
(only the 2 documented upstream-inherited failures remain);
test_backend_cross_device 2/2 passed; check-agent-record OK.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:GLM [omp]
The three static dispatch helpers (DispatchCombineSampledAndDraftTokens, DispatchScatterLastSampled, DispatchApplyLastSampledOps) had empty bodies on CPU-only and Vulkan builds (no VLLM_CPP_CUDA or VLLM_CPP_HIP), so every parameter was unused and the functions themselves were unused. Under -Werror=unused-parameter and -Werror=unused-function the CPU and Vulkan CI lanes failed to build. All call sites are already inside #if defined(VLLM_CPP_CUDA) || defined(VLLM_CPP_HIP) blocks, so wrapping the definitions in the same guard is the minimal fix. No behavior change on any build. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
…r differences a shape comparison cannot see (mudler#2249, mudler#2275) (mudler#2281) `LoadMoe` produced `Qwen4ExpMoeWeights` and nothing read it; `RunMoeBlock` consumes `MoeBlockWeights`. This adds the adapter between them, plus the one production composition that runs `Qwen4ExpTextSparseMoeBlock` through the shared sparse-MoE seam rather than a second MoE path. It is mudler#2249 item 4, under wave issue mudler#2031 and campaign issue mudler#1978. A fresh review returned FAIL on records and prose, not on engineering: it reproduced the headline finding, re-ran all seven mutations with a byte-for-byte restore, verified the fixture's four properties, confirmed the borrowed views cannot dangle, and found the neighbours untouched. The second commit is that repair. `qwen4_exp_moe.cpp` is byte-identical to the reviewed head, so the mutation evidence below still describes this tree. ## "The `_kq` arm's shapes match" is measured FALSE on the default route mudler#2249 item 4 says the keep-quant arm is a reinterpretation, so the work "looks small". It is false on the arm every shipped checkpoint takes. `LoadStackedExperts` records the tower as RANK 3 `[E, N, K]` (`qwen4_exp_weights.cpp:160-164`); `MoeBlockWeights::expert_*_kq` is RANK 2 `[E*N, K]`, and the default keep-quant route (`Qwen35GroupedMoeEnabled`, ON) hands that tensor straight to `vt::MatmulBTQuantGrouped`, whose first check is "matmul_bt_quant_grouped: rank-2 out/act/weight required" (`ops.cpp:223`). That throw is a permanent case in the suite, so the claim cannot be re-derived from the shapes. THE ROUTE QUALIFIER IS LOAD-BEARING, and the first draft stated the claim flatly. With `VT_QWEN35_GROUPED_MOE=0` the seam takes the per-expert `ExpertMlpKq` path, which reaches `KqResidentSlice` (`qwen3_5.cpp:5665-5678`); that helper rebuilds a rank-2 view from its `N`/`K` ARGUMENTS by pointer arithmetic, sets `wt.rank = 2` itself and never reads the declared rank. A rank-3 tower does not throw there, and the tower being contiguous `[E, N, K]`, `row_off = e * N` is the right slice, so it answers correctly — measured at max|diff| `0.00865547`, the keep-quant arm's own number. mudler#2249 item 4's sentence is literally true of that route and false of the default one. Three more the same comparison cannot see: * `LoadMoe` keeps the router and the shared gate f32 on purpose. The seam reads both against a BF16 activation and `MatmulBTKernelCuda` accepts only (bf16,bf16) or (f32,f32) (`cuda_matmul.cu:397-403`), so passing them through runs on CPU and dies on every GPU. They are converted here, which is also what the oracle has: upstream's router `weight` and `shared_expert_gate` are ordinary model-dtype parameters (`modeling_qwen4_exp.py:905`, `:925`) and `F.linear` returns bf16 logits that `softmax(..., dtype=torch.float)` upcasts (`:909-910`) — the seam's own bf16-logits / f32-softmax split. * `MoeBlock` picks the whole expert path from `expert_gate_kq` ALONE (`qwen3_5.cpp:7257`). `GgufLoadPolicy::Route` is per tensor, so a policy that keeps `gate` quantized and expands `down` yields a set the seam reads as keep-quant and then dereferences an empty tower. Refused by name. * A bf16 tower cannot use the stacked fields at all (`ops.cpp:231`), on the same route condition as above. The bf16 arm therefore fills the per-expert vectors, with zero-copy borrowed views: at the released 512 x 640 x 2560 geometry three per-expert copies per layer is 240 GB across the stack, so zero-copy is the only representable arm, not a tuning. ## The oracle, and the gate vLLM registers `qwen4_exp` at no revision, so the algorithm oracle is transformers 5.16.0, this row's accepted lane pin: `Qwen4ExpTextSparseMoeBlock.forward` (:927-938), `Qwen4ExpTextTopKRouter.forward` (:907-916), `Qwen4ExpTextExperts.forward` (:869-894) and `Qwen4ExpTextMLP` (:842-855). `tests/vllm/models/test_qwen4_exp_moe.cpp` reimplements those four in double precision from the SOURCE stacked towers, so a mis-sliced expert, a swapped gate/up or a transposed slab has nothing on the other side agreeing with it. Routing is gated non-trivial rather than assumed: five distinct selected sets across five tokens — gated `== kT`, not `>= 2`, which is what the prose claims — a set outside the {0,1,2} prefix, one expert never chosen, and a smallest top-k probability margin of 0.0740. A discrete selection has bimodal error, so the margin is what bounds it, not a tolerance. Seven mutations, and they red in TWO different ways rather than one. The index, orientation and slice mutations (expert off-by-one on both arms, gate/up swap on both arms, transposed slice on both arms) red on VALUE, at max|diff| 0.757 to 1.59 against |reference| 0.799. The rank-3-tower and transposed-keep-quant mutations red as SHAPE REFUSALS: they throw before any value exists, so no max|diff| can be reported for them and a single range covering all seven would be arithmetic nobody ran. The keep-quant tolerance is MEASURED, not derived, and the derivation the first draft gave was wrong. All three of its GEMMs quantize their activation to q8_0 and none is exact: `QuantizeRowQ8_0` (`cpu_quant_act.cpp:52-81`, per-block loop `:58-69`) takes `amax` per 32-ELEMENT BLOCK, `HiddenCodes()` forces `|code| = 127` at element 0 of each ROW only, and `kH = 64` is TWO blocks, so block 1 takes an arbitrary amax and gate and up carry quantization error as well as down. Forcing the code in every block would restore the derivation and is deliberately not done: it moves the router logits, hence the routing, hence the seven margins an independent review has already reproduced against this fixture. The bf16 arm's exactness argument was rechecked and holds. ## Both routes are gated, and that is a result `Qwen35GroupedMoeEnabled()` caches in a function-local `static const` (`qwen3_5.cpp:6299-6302`), which blocks a MID-PROCESS flip, not one before launch. The suite therefore runs on both routes and names the route beside every number it prints: | route | cases / assertions | bf16 max\|diff\| | keep-quant max\|diff\| | |---|---|---|---| | grouped (default) | 5 / 112 | 0.00218359 | 0.00865547 | | `VT_QWEN35_GROUPED_MOE=0` | 5 / 114 | 0.00218359 | 0.00865547 | BIT-IDENTICAL, which turns `qwen3_5.cpp:7261`'s byte-identity claim into a reading rather than an inherited assertion. Before the repair the suite was 111 assertions and RED on the alternate route, because its `CHECK_THROWS_WITH_AS` asserted the default route's refusal unconditionally. Neighbours, re-run unchanged on this tree after the merge below: `test_qwen4_exp_forward` 1/421, `test_qwen4_exp_gguf_weights` 11/2975, `test_qwen4_exp_qsa_block` 8/2831 — 8 and not 11, because mudler#2276 has not landed and the paged QSA arm is still owed. `test_qwen4_exp_scaffold` is 12/296, unchanged by the refusal edit below. `scripts/agent-preflight.sh --fail-on-skip` exits 0 with zero skips. ## What lands unreached, and what is owed This lands UNREACHED, by AGENTS.md "Nothing lands dead": `Qwen4ExpMoeBlockWeights`, `Qwen4ExpMoeHfConfig` and `RunQwen4ExpMoeBlock` are reached at this merge commit only by their own suite. `ForwardQwen4ExpForConditionalGeneration` EXISTS (`qwen4_exp_registry.cpp:142`) and IS registered as the model's `.forward` hook (`:456`); its entire body is one `VT_CHECK(false, ...)` refusal-by-name (`:228`) placed ahead of any downcast, so `ModelRegistry::Forward` reaches a real hook and that hook refuses `Qwen4ExpForConditionalGeneration` before a layer runs. An earlier draft said the function "does not exist" — the conclusion was right, the mechanism was not, and under `squash_merge_commit_message = PR_BODY` a mechanism nobody can check would have become permanent. The reachability mutation is therefore VACUOUS rather than passing: there is no production call site to delete, and a gate that stays green without one measures nothing. What the suite does prove is that it enters through the adapter, which is what the substitute mutation — deleting the keep-quant wiring — measures. The wiring is owned by row `MODEL-MM-QWEN4-EXP` and tracked by mudler#2031 under campaign issue mudler#1978. The spec's `## Owed` records it with these debts: * `norm_topk_prob` is not representable through this seam (the seam hardcodes `renormalize = true` and `HfConfig` has no field for it; upstream's default is `True`). * The bf16 arm is ineligible for the CUDA fast grouped-bf16 MoE, because `MoeBf16FastLayoutOk` wants the other orientation and a transposing copy is 240 GB across the stack. * The NVFP4 expert arm is refused by absence rather than by name. * The routed top-k weights reach the experts f32, which is WIDER than the oracle (`router_top_value.to(router_logits.dtype)`, `:914`) by one bf16 rounding. An earlier draft defended that in a test comment as "the seam is the more precise of the two"; AGENTS.md §"Inherit vLLM defaults" exists to refuse exactly that argument, since a token gate cannot see a dtype that is too wide. The width is INHERITED from `vt::MoeRouterTopK` and narrowing it is a seam-level change. * The keep-quant arm is value-proven at Q8_0 ONLY — every released checkpoint is a k-quant and none is executed here — and the fixture's `nk = false` source towers are not the loader's `nk = true` output. * The keep-quant tolerance's derivation, per the block-amax finding above. ## A filed bug that nothing could find mudler#2275 — `LoadStackedExperts` implements one of the three residencies `GgufLoadPolicy::Route` can return for a stacked expert tensor and silently expands the other two to bf16 — appeared in NEITHER `.agents/issue-index.md` NOR the spec by number. The spec's `## Owed` described the fall-through in prose without citing it, and the issue's own body claimed the opposite. AGENTS.md requires the link in three places that must agree, and requires an unfixed issue to name an owner. It now has an index row naming `MODEL-MM-QWEN4-EXP` and the number sits on the spec bullet that describes it. No `mudler#2249` index row is appended: mudler#2264 landed that one and it is on `main` already, so appending a second is the duplicate mudler#2266 says no gate would catch in either direction. ## One more gate The fourth refusal the header contracts — a stacked dtype that is neither arm (`qwen4_exp_moe.cpp:200-205`) — had no case. It has one, and the case is ARMED: neutralising the `Refuse(... == kBF16 ...)` predicate reds exactly it and nothing else (`CHECK_THROWS_WITH_AS ... did NOT throw at all!`, 111/112 passed), and the tree was restored by sha256 afterwards. It is also the refusal that catches mudler#2275's fall-through once the loader stops expanding silently. ## Brought current, and the three things that made it stale This branch sat 18 commits behind. A MERGE, not a rebase: it is published and was reviewed at its head, so a rewrite would need a force push and would discard the reviewed identity of `qwen4_exp_moe.cpp`. That file is still byte-identical to the reviewed head, so every mutation above still describes this tree. `.agents/issue-index.md` was resolved as a SET against the merge-base rather than left to the `merge=union` driver, which had interleaved the new row into main's tail and so left main's file not a prefix of the result. Rebuilt as main's complete file plus this branch's added rows: 876 rows, every base row present byte-for-byte, all of main's added rows and this branch's mudler#2275 present, and zero duplicate ids matched on the leading `^| [#NNNN]`. mudler#2249's row appears exactly once and it came from main. `.agents/specs/qwen4-exp-flash-next.md` is a keyed record, so the three-way result was discarded and main's complete version taken instead, with this branch's two scoped edits re-applied onto it. Main added three `## Owed` bullets at the same anchor this branch adds one; both sides are kept. Six lines are deleted and all six are the sentence W5d-4 measured false. **THE MERGE FALSIFIED THE PRODUCTION REFUSAL, WHICH IS THE FOURTH INSTANCE OF mudler#2288 AND EXACTLY THE ONE ITS OWN `## Owed` BULLET PREDICTED NOTHING WOULD CATCH.** Item (3) of the `ForwardQwen4ExpForConditionalGeneration` refusal was "an adapter from the stacked [E, I, H] qwen4_exp MoE tensors onto MoeBlockWeights" — which is what this PR adds. Landing the merge unedited would have shipped a commit whose own product output denies the commit. The clause is removed, `W5d-4` joins the landed-seams list, and the message now enumerates two. Verified on the EMITTED BYTES read out of the linked binary, not by grepping the source, because `test_qwen4_exp_scaffold.cpp:767` pins that five substrings are PRESENT and never that any is TRUE — it passes a wrong message. All five survive and the asserted absence still holds. The two remaining items were re-checked against the tree rather than inherited: no paged arm of `RunQwen4ExpQsaBlock` exists, and `gather_block_table` still has exactly three call sites reaching only `full_attn_group_id_` and `gdn_group_id_`. The same recount reaches the spec, where "TWO OF THE FIVE ARE NOW CLOSED" and "THREE remain" become three and two, keeping main's W5d-1 and W5d-2 facts beside the new one, and "what has no production shape yet" stops listing the MoE weight adapter. Closed as a SEAM and not as a call: this still lands unreached. Finally, the merge shifted `qwen3_5.cpp` by exactly one line, so fourteen `file:line` anchors in the spec, this header, this suite and this description pointed one construct off. Each was re-pointed by READING THE LINE BACK. The anchors in `ops.cpp`, `qwen4_exp_weights.cpp`, `gguf_keep_quant.cpp`, `cpu_quant_act.cpp` and `cuda_matmul.cu` were checked the same way, did not move, and were left alone rather than shifted on the assumption that one file's offset applies everywhere. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…n3-VL server is unservable (mudler#2300) (mudler#2303) File [mudler#2300](mudler#2300) and correct three records that state the opposite of what the tree does. `ForwardQwen3VLForConditionalGeneration` makes `ModelForwardInput.mm` mandatory (`qwen3_vl_registry.cpp:127`), and the GPU runner never sets it: `runner.cpp:2234` builds `ModelForwardInput` without `.mm` and calls `ModelRegistry::Forward` at `:2340`, and no line of that 4443-line file names `mm_features`, `MultiModalForwardInput` or `.mm`. A server started on a Qwen3-VL checkpoint therefore throws on the first forward step of every request, text or image. The records read the other way: `README.md` told a user that image chat requests reach the registered forward, and `mm-serving.md` headed a section "the block is RESOLVED". Both are true up to the engine and false one hop later, and this lands before the engine change so the record is true when that work starts. No product code moves. The fix is either a text arm in the registered forward or the runner building `mm` from staged encoder outputs, and the choice between them changes what an image request does, so it takes its own spec, gate and independent review. ## What changed `README.md`'s multimodal bullet now says what is gated (the CPU serving seam, the token-correct single-sequence drivers) and what is missing (the codec residual and the runner hop), in place of "correctness-complete" plus a wiring claim that ends in a `VT_CHECK`. The usage section restricts the parse claim to `image_url`, and the top-of-README "Everything." bullet no longer lists image, video and audio input without a qualification. `.agents/specs/mm-serving.md` gains a `## Owed` section and a pointer under the `MM-ENGINE-FORWARD` heading. Nothing is deleted. It corrects two statements: `MM-SERVE-PARSE`, `MM-SERVE-ENGINE` and `MM-SERVE-E2E` are wave names in that document rather than roadmap rows, and the registered mm forward is reached by `tests/vllm/multimodal/test_qwen3vl_registry_e2e.cpp:167` and by nothing else. `.agents/specs/multimodal-track.md` lists mudler#2300 under `## Owed`, beside the mudler#607 L3 bullet that already recorded the text half of the same refusal as a benchmark condition. `.agents/issue-index.md` gains one appended row for mudler#2300. `docs/FEATURES.md` gets a corrected seam anchor and the mudler#2300 link, so it and `README.md` point at the same issue and name the same two residuals. ## The review repair (`8539d1b39`, rewrapped in `9d672e3b3`) A fresh review of `bfb4f87f3` returned FINDINGS with one MEDIUM, and this branch carries the repair. The rewritten `README.md` line replaced one false claim with another. It said the server parses `image_url`, `video_url` and `input_audio` and carries them into the engine. That is true of `image_url` only. `video_url` is parsed as its bare `type` with an empty payload (`include/vllm/entrypoints/openai/protocol.h:384`), and both `video_url` and `input_audio` are refused at the seam with HTTP 400, because `Qwen3VLChatSupportedMmLimits()` allows one image and nothing else (`chat_mm.cpp:295`), `MakeQwen3VLImageChatFn` runs `ValidateChatMmLimits` first (`chat_mm.cpp:311`), and `ValidateNumItems` defaults an absent modality to 0 (`context.cpp:41`). The claim also contradicted `docs/FEATURES.md`, which says video, audio and multi-image over HTTP are not started. Both changed `README.md` passages named mudler#2300 as the single blocker. The residual a real client meets first is the codec: the production `ImageCodecFn` accepts `image/x-raw-rgb` and throws on any container format (`server_main.cpp:1538`), so a PNG or JPEG data URI never reaches the runner. Both passages now name the codec beside mudler#2300 and give the order in which a request meets them. The new `## Owed` section said the CPU seam body is wired in `examples/server/main.cpp`. That file is a 23-line ABI shim with no multimodal reference. `ARCH-ONE-SURFACE` moved the construction into the library, and the wiring is `src/vllm/entrypoints/openai/server_main.cpp:1545`. The reachability conclusion is unchanged, because the shim reaches that wiring through `vllm_server_main`. `docs/FEATURES.md` cited `server_main.cpp:826` for the seam attach. That line is the `--enable-auto-tool-choice` validation at this head, so the anchor is corrected to `:1545`. ## How to verify The claim is statically derived and was not run. No binary was built and no server was started, which the issue body states in those words. At `e541be98ab070d7961b8e43ce3afb9327dc6f0ba`: ```sh grep -n "VT_CHECK(input.mm.has_value" src/vllm/model_executor/models/qwen3_vl_registry.cpp grep -n "input.mm.has_value" src/vllm/model_executor/models/gemma4_registry.cpp grep -n "input.mm.has_value" src/vllm/model_executor/models/muse_glimmer_registry.cpp grep -cE "mm_features|MultiModalForwardInput" src/vllm/v1/worker/gpu/runner.cpp grep -rn "Qwen3VLGenerateGreedyViaRegistry" src include tests examples tools ``` They return line 127 (a mandatory check), lines 145 and 113 (guards with a text fallback), 0, and one caller under `tests/`. `scripts/now.py --offline` lists no `MM-SERVE-E2E` row, and `.agents/engine-matrix.md` carries `ENG-MM-INPUT-PIPELINE` as `READY` and `ENG-MM-QWEN36-VL-FORWARD` as `ACTIVE`. For the repair anchors, at `bfb4f87f3`: ```sh grep -n "image_embeds / video_url" include/vllm/entrypoints/openai/protocol.h grep -n "Qwen3VLChatSupportedMmLimits\|ValidateChatMmLimits(info, messages)" src/vllm/entrypoints/openai/chat_mm.cpp sed -n '1538,1546p' src/vllm/entrypoints/openai/server_main.cpp sed -n '826p' src/vllm/entrypoints/openai/server_main.cpp grep -c "multimodal\|mm_features" examples/server/main.cpp ``` ## Out of scope `runner.cpp`, the scheduler, `NewRequestData`, `EncoderCacheManager`, `MultiModalForwardInput` and every registry forward are untouched. The runtime confirmation is owed and needs a Qwen3-VL safetensors checkpoint, because `LoadQwen3VLForConditionalGeneration` (`qwen3_vl_registry.cpp:95`) refuses any other source. ## Correction to the dispatch brief mudler#2257's index row was NOT missing. It landed in `3ed2378a3` between the brief being written and this branch's base, so appending it again would duplicate a row that `scripts/check-agent-record.py` gates. This pull request appends one row, not two. ## Known state before merge: resolved by `bf82f5e2d` The branch was `mergeable: CONFLICTING` and had **zero** check-runs on head `9d672e3b3`. That is not a slow queue. GitHub schedules `pull_request` workflows against `refs/pull/N/merge`, and it stops recomputing that ref once the merge fails, so no `pull_request` event fired for the last two pushes and no check-run was ever created. The ref still existed, which is the trap: it resolved to `7b84f5cb2`, whose parents are `a4b333329` (a pre-conflict `main`) and `bfb4f87f3`, the first of this branch's three commits. The only two workflow runs on the branch are pinned to that stale head and had sat at `queued` and `in_progress` since `2026-08-29T20:25:47Z`. The conflict was the appended `.agents/issue-index.md` tail row for mudler#2300 against the rows added on `main` after this branch's base. `git merge-tree --write-tree` reports it clean **locally only**, because `.gitattributes:7` sets `merge=union` on that path. GitHub does not apply the union driver: over the identical three blobs, `git merge-file -p ours base theirs` exits 1 with a conflict hunk at `:890-896` while `merge-tree` exits 0. One variable, and the forge is on the `merge-file` side of it. `bf82f5e2d` merges `origin/main` `3015aad08` into the branch with an authored message. A merge commit is the only remedy AGENTS.md allows here; a rebase would also clear the ref and needs a force-push it withholds. The merge commit is inside the `e541be98a..HEAD` walk that `scripts/check-commit-trailers.py` performs with no `--no-merges`, and both commit gates were run explicitly rather than through preflight, which SKIPs them on a behind branch: ``` python3 scripts/check-commit-trailers.py --range e541be9..HEAD -> rc 0 python3 scripts/check-commit-style.py --range e541be9..HEAD -> rc 0 python3 scripts/check-agent-record.py -> rc 0 python3 scripts/check-issue-index-append-only.py --base origin/main -> rc 0 ``` The index resolution keeps both sides' appended rows and edits none, resolved from the diff rather than from the tail. Counted by row key `^| [#N]` rather than by a bare grep, since a prose cross-reference inside a row body is not a duplicate: 868 base + 1 ours (mudler#2300) + 7 theirs (mudler#2223 mudler#2230 mudler#2283 mudler#2286 mudler#2274 mudler#2309 mudler#2312) + 2 appended in the merge = **878**, which is what the merged file holds. Zero rows dropped, zero duplicated, zero row keys twice, all 868 base rows present byte-for-byte. ## The two rows the merge appends They ride in the merge commit deliberately, so that filing them does not immediately re-conflict this branch. [mudler#2317](mudler#2317) records the systemic defect this branch is a symptom of. 19 of 23 open pull requests touch `.agents/issue-index.md`; GitHub has computed mergeability for 9 of them; 8 are unmergeable; 7 of the 8 conflict on the index, and for 6 of those the index is the only conflicting path. Under AGENTS.md `## Records` the second admitted record shape, "a genuinely append-only file that can union-merge", does not exist on this forge, so the surface degrades into the lock the same section forbids. Filed, not fixed: the repair changes `check-agent-record.py` and `check-issue-index-append-only.py` semantics and needs its own spec, a red-before test and a fresh review. Owned by `ENG-RECORD-CONFLICT-SURFACES`. [mudler#883](mudler#883) first observed the behaviour and left the mechanism unestablished. It had **no row** in the index at all, so the issue that named this class was untracked by the surface it is about. Its row is appended in the same commit. No prose and no code from the reviewed content changes in `bf82f5e2d`. This pull request uses no closing keyword. mudler#2300 records a defect that stays open until the engine change lands, and mudler#2317 stays open for the row that owns it. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…vice when the box can hold them (mudler#2327) (mudler#2328) perf(PERF-QWEN35-STAGE-WEIGHTS): stage dense decode weights to the device when the box can hold them (mudler#2327) The target decode step is weight-bandwidth-bound — ~50 GB swept per forward against GB10's 273 GB/s roof, a ~184 ms floor, activations at 8 query rows ~100 KB — so where the weights LIVE is the step time. `ResidentWeight` handed every decode GEMM a HOST pointer tagged with a device wherever `host_memory_is_device_addressable()` answers true. vLLM's parameters are constructed inside the torch CUDA allocator and are real device memory. MEASURED on `dgx:gpu0`, one boot, one binary at `21ef6f053`. Qwen3.8-27B bf16 + DFlash2 k=7, concurrency 1, `max_num_seqs=1`, 64 tokens, greedy. `VT_QWEN35_ALIAS_HOST_WEIGHTS` is the only variable. Four warm repeats per leg, interleaved A,B,A,B,A: alias ON 11.677 11.693 11.690 <- the third is the TERMINAL CONTROL staged 14.288 14.337 The three ON arms agree to 0.14% and the control matches the opening arm, so drift cannot masquerade as effect and the run is admissible. **+22.4%.** vLLM on this prompt is 14.305; staged, we are 14.31 — parity, from one variable. This is not a new idea. `laguna.cpp:130-132` records the same mechanism and the fix shipped default-ON for two other models (Laguna to 1.03x vLLM, DeepSeek-V4 to 1.144x). Qwen3.5 never received it, and `docs/ENVIRONMENT.md:259` names the decode risk in the alias's own row and leaves it open. IT CANNOT BE A BLANKET FLIP, and that is the whole design. This file also serves `Qwen3.8-2.4T-A95B`, and mudler#1299 measured that checkpoint exhausting a 119.631 GiB box precisely BECAUSE the CUDA arm paid for its weights twice. Staging is right for a model that fits and fatal for one that does not, so the policy asks the BOX: stage only while `VT_QWEN35_STAGE_MIN_FREE_FRAC` (default 0.55) of total device memory is still free AFTER the copy. A 50 GiB model on a 119.6 GiB box satisfies that; the 2.4T model is already past the floor when its first dense weight arrives and never stages, so mudler#1299's invariant is preserved exactly. `=1` pins the retag, `=0` forces staging, and an unanswerable `DeviceMemoryInfo` keeps today's behaviour — an unknown budget is not a licence to double a model's residency. The arithmetic is a PURE function so it gates without a device: a fake `vt::Backend` would need every pure virtual stubbed and would gate less than five direct cases do. Red-first by mutation — deleting the reserve floor fails exactly the mudler#1299 case and the floor case, BUILD rc=0 / TEST rc=1, 18/18 and 84 assertions with it restored. Two things went wrong writing this and are recorded rather than smoothed over. One case asserted the EXACT floor boundary, which gates the binary representation of 0.55 rather than the policy; it is replaced by clearly-above/clearly-below plus a case where moving the floor moves the answer. And the declarations sit at the END of the header because the first draft put them mid-file, shifting `Fp8Weight` from `:628` to `:658` and breaking an anchor three records cite — `check-agent-record` red at 29 stale against a baseline of 28. Relocating cost nothing and left every citation true. NOT a full gate result. This is one prompt of the four the committed gate runs. The gate re-run on the staged default is the next step and is what settles the overall ratio. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…the bridge answers O22 by refusing to decode the tower (mudler#2325) W5b as scoped was two changes with two oracles. The attention block and the `OwnedTensor` bridge answer to `transformers` v5.16.1 and to the llama.cpp [#27752](ggml-org/llama.cpp#27752) container, and both can be gated with no KV cache and no decoder layer over them; the decoder layer, the mHC threading and the forward answer additionally to `MakeGlm5NextKVCache` and to the `[T, hc_mult, hidden]` manifold. This is the first half. [mudler#2241](mudler#2241) stays open for the second. ## The attention block `Glm5NextTextAttention` (`modeling_glm5_next.py:1064-1257`) as a host f32 reference, the same shape `glm5_next_dsa.cpp`, `glm5_next_mhc.cpp` and `glm5_next_moe.cpp` already have. Three things a fluent wrong port gets wrong, each with a case that separates it rather than a tolerance that does not. **The converter SPLITS `kv_b_proj` and transposes only the K half.** The file carries `attn_k_b` at `[H, kv_lora, qk_nope]` and `attn_v_b` at `[H, v_head, kv_lora]`, so K contracts over its FIRST inner axis and V over its SECOND. At the published geometry a swap is a shape error, so the gate also runs a SQUARE case where the untransposed reading is perfectly shape-valid and merely wrong; it separates by **2.9469 over all 900 values**, where 900 is `kBatch(2) * kNumHeads(3) * kSeqLen(25) * kSqDim(6)`. The case prints the figure rather than leaving it to prose: ``` test_glm5_next_attn.cpp:373: MESSAGE: square k_b: transposed vs untransposed separation = 2.9469 over 900 of 900 values ``` **Cross-layer top-k sharing.** A `shared` layer builds no indexer and reuses the previous full layer's selection (`:1130-1134`, `:1181-1191`). A layer that recomputes runs, selects a plausible key set and emits plausible tokens, and nothing about the output's shape, finiteness or scale says otherwise. So the fixture carries BOTH the correct output and what a recomputing port produces from a decoy indexer, both captured from the same oracle run, and asserts ours is the first: 320 of 800 values differ, max separation 1.52, over 20 of 50 query rows. A separate case proves the decoy golden really is a recomputation, so the inequality is between two reference values and not between a reference and an arbitrary number. **Read that one with its caveat attached.** The `shared` arm is CONFIG-KEYED, and the published `GLM-5.3-Flash` `config.json` selects it on **zero of its 45 layers** — the suite measures that and prints `published schedule: 0 shared layers of 45`. So the sharing gated here is correct against `transformers` v5.16.1 on a schedule the released checkpoint does not contain. It is the same "unselected branch" shape as the rope half below, with one difference: the rope branch is REFUSED and this one is IMPLEMENTED and gated. It is also the shape that SURVIVES W5b-2 — once O25's reachability half is discharged the two files become reached and the `shared` arm still is not. O25 carries this. **The all-masked row is `finfo.min` and not `-inf`** (`:1253`). A left-padded query row has every key masked; `finfo.min` gives it a uniform softmax and a finite output, `-inf` gives it NaN through `o_proj` and into the residual stream for the rest of the stack. The `-inf` mutation reds 49 of 160 assertions. There is no rope branch, because upstream can reach none. `validate_architecture` (`configuration_glm5_next.py:225-228`) raises for any positive `qk_rope_head_dim` — measured by constructing one in the golden generator rather than described — so `expand_kv`'s concat has a zero-width second half and `key_states` IS `k_nope`. `MlaDims::Validate` mirrors the refusal in upstream's own words instead of half-implementing a branch no released config selects. ## The bridge, and O22 O22 left the residency choice open on purpose: "Whoever writes the forward decides whether to decode per layer or to go device-native." The decision is **decode ONE DSA layer at a time, on demand, and never retain the tower in float.** | what | GiB | |---|---:| | the published `UD-Q2_K_XL` artifact, block-resident as loaded | **101.14** | | the same tower with every tensor expanded | **426.72** | | all-bf16 | 597.46 | | usable on `dgx:gpu0`, the largest device this project reaches | **~119.63** | | ONE bridged DSA layer, f32 | **0.4654** | | all ELEVEN DSA layers held at once | 5.12 | A decoded tower is 3.57x over the box, and that is the figure [mudler#2245](mudler#2245) and [mudler#2247](mudler#2247) spent six pull requests removing. A float tower is not expensive; it does not exist on any hardware this project can reach. One layer is 499,657,728 bytes, 0.39% of the box, and the caller's peak is one layer because the mirror is a value it can drop. There is deliberately no `BridgeTower`, no cache and no map keyed by layer index, because each of those turns "one layer" into "every layer visited so far", which is the tower again with a slower ramp. Device-native was not chosen, for a stated reason rather than a preference: there is nothing to be device-native against while every glm5_next primitive on this row is a host reference and W3's CUDA arm is committed and unmeasured. That would be the "unpassed parameter" shape. W5b-2 revisits it. O19 / [mudler#2260](mudler#2260) stays live and this bridge cannot make it reachable. Structurally there is no overload taking `Glm5NextMoeWeights`, `Glm5NextMlpWeights` or any expert bank. Numerically the 1 GiB per-tensor ceiling sits EXACTLY 4x above the largest legitimate tensor (`o_proj`, 0.25 GiB) and EXACTLY 9x below the smallest expert bank (`up_exps`, 9.0 GiB); both sides are asserted, because a ceiling above everything is a mute switch and one below the real population fires on ordinary work. The check runs from the shape before any allocation, proved by handing the bridge a published-size bank carrying no bytes at all. `byte_ceiling` is a DEFAULT ARGUMENT, so the structural claim binds unconditionally and the numeric one binds every call that takes the default, which is every call in this tree; O25 says so. ## The bridge's four advertised refusals are now gates `glm5_next_bridge.h` lists four cases `DecodeOwnedTensorToF32` refuses by name. Review found that the block element-count check, both byte-span checks and the `default:` dtype arm could each be deleted with the suite staying 8/8 and 56/56. Two of those are not cosmetic: * without the elementwise byte-span check, `std::memcpy(out.data(), src, need)` reads `need` bytes out of a shorter `t.bytes` and the bridge **serves the heap as weight values** — finite, plausible, wrong, and invisible to a token gate; * without the `default:` arm, an encoding the bridge cannot widen falls off the end returning the **zero-filled buffer it allocated**, which is the failure the `host_released` refusal already exists to stop, reached by another door. Five cases pin them, each proved by disabling the refusal in a scratch copy, with the mutant's BUILD rc recorded beside its TEST rc because a mutant that does not build reads as a passing test, and the file restored byte-for-byte (sha256 `bee9a0f66d88914193e3c0f1d1d89d9840f7ca91bef79b0538f08945b54f99a7`) after each: | refusal | mutation | result | |---|---|---| | block element count is a whole number of blocks | `if (false)` | BUILD 0 / TEST 1, 1 assertion | | block byte span equals `RowSizeBytes` | `if (false)` | BUILD 0 / TEST 1, 3 assertions | | a block dtype has a `BlockToFloat` decoder | `if (false)` | BUILD 0 / TEST 0 — **SURVIVES** | | elementwise byte span equals `numel * SizeOf` | `if (false)` | BUILD 0 / TEST 1, 3 assertions | | `default:` refuses a non-float encoding | `return out;` | BUILD 0 / TEST 1, 4 assertions | The survivor is disclosed rather than chased, and it refines the review's finding. `vt::IsBlockQuant` is true for exactly the 18 dtypes `vt::cpu::BlockToFloat` answers for, so **no input can reach that arm in this build**: it is the unselected-branch shape, a guard for the encoding that lands next without a CPU decoder, which is the state IQ2_XS and IQ4_XS were in before [mudler#2245](mudler#2245). What the suite gates instead is the PREMISE, and that gate is ARMED rather than assumed. Rewriting `BlockToFloat`'s `kQ8_0` case to `return nullptr` (BUILD rc=0, after referencing the now-unused function so `-Werror=unused-function` does not turn the mutant into a build failure) reds the premise case at `CHECK(vt::cpu::BlockToFloat(d) != nullptr)` AND makes the refusal fire by name in two more, with the bridge's own message: ``glm5_next bridge: `moe.gate_exps` is q8_0, which this build has no `BlockToFloat` decoder for``. ## Seventy-four upstream citations did not resolve at v5.16.1 The ported behaviour is right everywhere it was checked; the citations were not. Re-resolving every anchor in this wave against `modeling_glm5_next.py` sha256 `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b` at `refs/tags/v5.16.1` found **74 citation sites carrying 38 distinct wrong values**, off by 1 to 6 lines. AGENTS.md requires citing the `file:line` that was ported, and an anchor that lands on a blank line or excludes its own symbol sends the next reader somewhere useless. The worst: | cited | actual | what the citation missed | |---|---|---| | `:1126-1131` for `skip_topk` / `next_skip_topk` | 1130 and 1132-1134 | the range **excluded** `next_skip_topk` | | `:1165` for `q_resid` | 1167 | 1165 is **blank** | | `:1167-1171` for `CompressKv` | 1170-1172 | **excluded** `kv_a_layernorm` | | `:1157-1216` for `Attention` | 1155-1216 | `def forward` is at 1155 | | `:1180-1186` / `:1188-1192` for the selection and mask | 1181-1191 / 1193-1197 | the mask block sat **wholly outside** its anchor | | `configuration_glm5_next.py:219-226` for the NoPE refusal | 225-228 | 219-220 is an unrelated `index_topk % index_kpool` clause | One citation is deliberately left alone: the whole-class span `:1064-1257`, whose only slack is the blank line after the class's last statement at 1256. It excludes no symbol, and it is baked into the generated `glm5_next_attn_goldens.inc` banner, which cannot be regenerated without a live oracle install. ## Evidence **Oracle identity asserted, not assumed.** The golden generator hashes the INSTALLED `modeling_glm5_next.py` and refuses unless it is `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b`, the value W3 (mudler#2213) and W5c (mudler#2242) both recorded, alongside the version string. **RED first**, from the plausible wrong port on the same tree in one build: 9 of 14 cases and 63 of 150 assertions in `test_glm5_next_attn`, 3 of 8 and 3 of 55 in `test_glm5_next_bridge`. That red also found two defects in the TESTS rather than the product — the refusal golden carried huggingface_hub's wrapper class name, and the bridge's shape case moved a dim `q_b_proj` also depends on, so it threw on the wrong tensor — both repaired before green. **GREEN**: `test_glm5_next_attn` **14/14 cases and 160/160 assertions**, `test_glm5_next_bridge` **13/13 and 96/96** (up from 8/8 and 56/56 with the five refusal cases), both exit 0, both rerun by hand after merging `origin/main`, since `agent-preflight.sh` runs the record and script gates and not the C++ suites. **Twenty-four negative mutations**, each sha256-proved applied, built and restored byte-for-byte. **Twenty-two kill their gate.** Two do not, and both are recorded with their reason rather than as passes: `host_f32_bytes` taken from the dims instead of from the buffers is an EQUIVALENT mutant while `DecodeShaped` refuses any shape disagreement (the test now pins the sum against the buffers themselves — an earlier version pinned it against the predictor and that mutation passed it), and the `BlockToFloat`-null arm is unreachable in this build, whose premise gate is armed instead and proved by removing the Q8_0 decoder. Two further findings came out of the earlier run and are fixed: the fixture could not tell `min(l+1, n-1)` from a wrapping `(l+1) % n`, so a schedule where they disagree was added, and one mutant failed to BUILD under `-Werror` on an unused parameter, which is a passing mutant proving nothing. **Every emitted golden is read by an assertion**, and a case names each array and the case that reads it. W3 emitted a `kIndexScores` golden that nothing consumed and two real scale defects then passed 1602 assertions; that is why this is a requirement here and not a courtesy. **`scripts/agent-preflight.sh --fail-on-skip`: ZERO skips, and two failures that are BASE rather than branch.** `check-env-doc` and `test_check_env_doc` both report `VT_QWEN35_STAGE_MIN_FREE_FRAC`, added by `207c12932` (mudler#2328). At `origin/main` that variable is read once under `src/` and appears **zero** times in `docs/ENVIRONMENT.md` and **zero** times in `scripts/env-doc-allowlist.txt`; this branch touches none of those three files. It is the [mudler#2312](mudler#2312) shape again, it is already tracked by open issue [mudler#2329](mudler#2329) with a live `row/ENV-DOC-2329` worktree on it, and it is deliberately NOT fixed here: a second fix to the same line is a conflict, not an in-flow repair. ## What is NOT reached, and who owns it Nothing in `glm5_next_attn.{h,cpp}` or `glm5_next_bridge.{h,cpp}` is called from a production entry point at this merge commit. `grep` over `src/`, `include/` and `examples/` for `glm5_next::Attention`, `BridgeDsaLayer`, `DecodeOwnedTensorToF32`, `IndexerRoleFor` and the two headers returns nothing outside the four files of this change. There is therefore no production call site to delete, so `.agents/reachability.md`'s reachability mutation is already answered: the change has no entry-point chain, and saying so is the answer. The wiring belongs to **W5b-2**, on row `MODEL-MM-glm5-next-glm5-next-for-conditional-generation`, tracked by [mudler#2241](mudler#2241) under campaign issue [mudler#1998](mudler#1998), and the spec lists it under `## Owed` as **O25**. O25 also records that the `shared` indexer arm stays unreached even after that wiring, and that `Numel` iterates `i < t.rank` against a fixed `shape[vt::kMaxRank]` — unreachable through the loader, which is the bridge's only producer. ## What is NOT claimed No token, no load and no speed number, and none was observed. No oracle for this model runs on any device this project reaches — the reference needs 305.78 GiB (FP8) or 598.5 GiB (BF16) against ~119.63 GiB — so what is gated is the NUMERICS of one block against a tiny-shape reference and nothing about the MODEL. The staged artifact was not opened at all: W5c already measured that a materialising load on this box stops at 8.09 GiB RSS in uninterruptible CIFS I/O, and this wave needs no artifact. No GPU was used, no `rc` lease was taken and no `ssh` to a fleet device was attempted. Refs mudler#2324, mudler#2241, mudler#1998. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…e stops failing every branch (mudler#2329) (mudler#2332) docs(ENV-DOC): document VT_QWEN35_STAGE_MIN_FREE_FRAC so the base gate stops failing every branch (mudler#2329) `check-env-doc` was red on `origin/main` itself. `207c12932` (mudler#2327 / mudler#2328, `PERF-QWEN35-STAGE-WEIGHTS`) introduced `VT_QWEN35_STAGE_MIN_FREE_FRAC` at `src/vllm/model_executor/models/qwen3_5_weights.cpp` and documented it only in a code comment, so the checker refused it in neither `docs/ENVIRONMENT.md` nor `scripts/env-doc-allowlist.txt`. Because preflight runs that checker, **every branch cut after `207c12932` inherited a failing gate and could not reach a green preflight before push**. Three open PRs on `MODEL-MM-QWEN4-EXP` were blocked behind it at the time of writing. Red to green, measured on bytes rather than inferred: extracting `origin/main` into a clean tree and running `scripts/check-env-doc.py` there exits 1 naming that one variable; the same checker on this branch exits 0. The entry records what the code does rather than restating the knob's name: the predicate is `stage only while free - bytes >= VT_QWEN35_STAGE_MIN_FREE_FRAC * total`, the default is `0.55`, and unset, empty, unparsable, `<= 0` and `>= 1` all fall back to that default rather than refusing — so a typo cannot silently disable the floor and double this model's device residency, which is the failure mudler#1299 recorded when the CUDA arm paid for its weights twice. This is mudler#2312 recurring. That row records the identical failure for `VT_DFLASH_BOUNDS_DEVICE` from mudler#2304, fixed in flow by mudler#2313 with the note that it is "a base failure every later branch inherits". Twice in one day suggests the gap is structural: nothing forces a documentation entry at the point a knob is introduced, so the omission is only caught later, by whoever happens to branch next. That observation is left for `PERF-QWEN35-STAGE-WEIGHTS` and whoever owns the checker; this change does not attempt it. Scope is two files: the documentation entry and one append-only issue-index row naming `PERF-QWEN35-STAGE-WEIGHTS` as the owner. No product code is touched, the knob is not renamed, and neighbouring entries are unchanged. Fixes mudler#2329 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5 [Claude Code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…AGED cache the engine actually allocates (mudler#2249) (mudler#2276) `MakeQwen4ExpKVCache` publishes the QSA layers' K/V as a paged `FullAttentionSpec` group, and the block W5b-5 landed reads a contiguous `[max_kv, num_kv_heads, head_dim]` array. Nothing could serve from the cache a runner hands a forward. That is [mudler#2249](mudler#2249) item 2 in its own words: > `Qwen4ExpQsaCaches` is contiguous `[max_kv, ...]`; `MakeQwen4ExpKVCache` > publishes **paged** specs. The block landed by W5b-5 reads the contiguous form, > so nothing can serve from the cache the engine actually allocates. ## An address MODE, not a second op — and why W5d-1 called it the other way The fork between the two shapes is the resolution of ONE address: the key/value row for logical position `p`. So this is a mode on `vt::Qwen4ExpQsaGatherAttention` (`kv_block_table` + `kv_block_size`), not a new op. It mirrors vLLM's paged read as `vt::PagedAttention` already states it — `page = block_table[p / block_size]`, `offset = p % block_size`. The expansion, the ascending visit order, the two softmax passes and the f32 accumulation stay ONE body, so the arms cannot drift apart the way two kernels would. **This is the opposite call from sibling wave W5d-1 (mudler#2265), and the two are not in conflict; the arm COUNT is what separates them. A later reader must not read the two waves as contradictory.** W5d-1 added a whole new `OpId::kRmsNormGroup` rather than a `group_size` field on the shared `RmsNormArgs`, because a new field on a shared args struct is IGNORED by every kernel not taught to read it — a backend whose `kRmsNorm` was not updated would answer a grouped request with a whole-row norm, no crash and no refusal. `kRmsNorm` has **six** registered arms (`cpu_ops.cpp`, `cuda_ops.cu`, `rocm_ops.hip`, `metal_ops.mm`, `vulkan_ops.cpp`, `tenstorrent_ops.cpp`), so five of them could have silently ignored the field. `kQwen4ExpQsaGatherAttention` has **exactly one**, checked at this head rather than assumed: `grep -rn kQwen4ExpQsaGatherAttention src include` returns the enum in `include/vt/ops.h`, the dispatch in `src/vt/ops.cpp:2688`, the provider case in `src/vt/op_provider.cpp:540`, and one single `RegisterOp` at `src/vt/cpu/cpu_qwen4_exp_qsa.cpp:357`. There is no un-updated arm to silently ignore the new fields, so the hazard that forced mudler#2265's split cannot occur here, and the cost of a split — two bodies that drift — would be paid for nothing. A future second arm inherits the mode in the same body it inherits everything else. `RunQwen4ExpQsaBlockPaged` is the block-level entry, on the `RunGdnBlockPaged` precedent this row already produced, and it shares `QsaBlockCore` with the contiguous one. What it changes is where the new rows are STORED (`dense_attn::WriteKvCache` at a slot mapping, which is `dense_attn::AttnBlock`'s own order and upstream's) and how the consumer ADDRESSES them. `Qwen4ExpQsaIndex` now takes the indexer side cache TENSOR rather than the cache struct, because it reads nothing else from it and handing it a struct with two dead fields is a landmine now that the K/V has two shapes. An fp8 paged cache is refused BY NAME: the QSA consumer has no dequantising read, so reading fp8 bytes as floats would be wrong tokens rather than a crash — the exact failure `kv_cache_route.h` exists to prevent. ## The gate, and why it discriminates Three cases in `tests/vllm/models/test_qwen4_exp_qsa_block.cpp`, over a block table of `{5, 3, 7}` that shares no fixed point with `{0, 1, 2}`, and a last page that is partial. An identity table would make a paged read and a contiguous read agree for every input, so the case would pass over a body that ignored the table. The whole flash cache is bf16-NaN before the block runs, so every row a correct read never addresses is not a number. The paged output is asserted against the lane-pinned oracle's own `Qwen4ExpTextAttention.forward` golden AND bit-for-bit against the contiguous arm — no tolerance, because paging moves where a row lives and nothing else. The control reads with an identity table what was written with the permuted one and must come back NaN on all 23 rows. **RE-MEASURED TWICE**, once after merging `cff257622` (which brings in W5d-4, mudler#2281) and again after merging `207c12932`; the figures below are the second run and they are identical to the first. Every BUILD RC WAS READ BEFORE ANY TEST RESULT, because under ENOSPC a failed build reads as a passing test. `cmake` rc 0, `ninja -C build -j 2 <five named targets>` rc 0 both times, `df -h .` 17 GiB free at the start and 13 GiB at the end — never near the floor where a build failure would masquerade as a passing suite. | Suite | At this head | |---|---| | `test_qwen4_exp_qsa_block` | **11 / 4382 / rc 0** | | `test_qwen4_exp_qsa_block -tce="*PAGED*"` | **8 / 2831 / rc 0, 3 skipped** — the pre-existing population, unmoved | | `test_qwen4_exp_scaffold` | **12 / 296 / rc 0** — pins the rewritten refusal | | `test_ops_rms_norm_group` | **7 / 69 / rc 0** — mudler#2265's gate, run because the merge shares `include/vt/ops.h` | | `test_qwen4_exp_moe` | **5 / 112 / rc 0** — W5d-4's gate, run for the same reason | | `test_qwen3_5_mrope` | **4 / 157 / rc 0** — W5d-2's gate | **THE PARTITION IS THE POINT OF THE FILTERED ROW, AND IT HOLDS: 8 + 3 = 11.** The exclude filter reports THREE skipped, which are exactly this wave's three new paged cases, so the eight that remain are the pre-existing population and no pre-existing case has silently become a new one. A filtered run that merely passed would not have shown that; the skip count is what makes it evidence. ## Mutations Each was sha256-proved applied, the **build rc was read BEFORE any test result** (a failed build reads as a passing test), and the tree was restored byte-for-byte after each one. | # | Mutation | Site | Build rc | Result | |---|---|---|---|---| | RED | the pre-W5d-3 body: `RowBase` always returns the CONTIGUOUS address `(p * HKV + kvh) * DH` | `cpu_qwen4_exp_qsa.cpp` | 0 | RED, 2 of 11 cases, 1537 assertions. `CHECK(std::isfinite(v))` x1472, `differing 1472 of 1472`, 64 more on decode. The capture of the gap mudler#2249 item 2 names | | M1 | off-by-one page-table read: `pages[(p + 1) / page_size]` | `cpu_qwen4_exp_qsa.cpp` | 0 | RED, 2 of 11 cases, 195 assertions. `rel 0.309431 < 0.03`, `differing 703 of 1472`, decode `rel 0.667465`. **DELIBERATELY NOT ALL-NaN**: an off-by-one lands on a WRITTEN page most of the time, which is exactly why the value comparison sits beside the finiteness one rather than replacing it. A battery whose every mutation is all-NaN proves only that the finiteness guard fires | | M2 | partial final page read at FULL length: the ragged tail runs to `ceil(kv_len / page) * page` | `cpu_qwen4_exp_qsa.cpp` | 0 | RED, 2 of 11 cases, 515 assertions. 448 NaN outputs, `rel 1.24777`, `differing 1339 of 1472`, plus the `keys_visited` equality with the contiguous arm | | M4 | the PAGED STORE deleted (`dense_attn::WriteKvCache` never called) | `qwen4_exp_qsa_block.cpp` | 0 | RED, 2 of 11 cases, 1537 assertions, all NaN. The store site is gated, not merely present | Restore proof: `cpu_qwen4_exp_qsa.cpp` back to `d95eea49e1800a25fb0b920a43c92936973c5e51e576651fa74400064a2497cc`, `qwen4_exp_qsa_block.cpp` back to `feb0eccd41d39a1588a9ffd32db9bfec50ac01dc40e44e494ad426d7ca3c43b7`. **M3, the reachability mutation, has NO SITE, and that is the finding rather than an omission.** `grep -rn 'RunQwen4ExpQsaBlock\|Qwen4ExpQsaPagedCaches' src include examples` returns only the block's own header and translation unit. There is no production caller to delete. ## The harness finding: `MaxRelDiff` cannot see a NaN Measured, not feared. `MaxRelDiff` folds with `std::max`, and `std::max(x, NaN)` returns `x`. Under two of the mutations above the run came back **all NaN**, the harness printed `max relative difference vs the oracle 0`, and the tolerance **PASSED**. The two new paged cases therefore assert FINITENESS FIRST and the oracle bound second, and the test says why (`test_qwen4_exp_qsa_block.cpp:915-924`, `:1020-1022`). The tree-wide fix is **NOT** in this pull request; it is owned by [mudler#2272](mudler#2272), which this wave's finding opened — 13 suites carry their own NaN-blind max-abs helper and `MaxRelDiff` has two blind local definitions with no hardened equivalent anywhere. **So this suite is NOT fully hardened, and the next reader should not read it as such.** Four tolerance-only sites remain in it with no finiteness guard beside them: `:467` (`rel_k`), `:485` (`rel_q`), `:576` (decode `rel`), `:594` (`rel_row`). They sit on the CONTIGUOUS arm over zero-filled caches, so a NaN is less reachable there, and they are deliberately left to mudler#2272 rather than fixed here. ## THIS LANDS UNREACHED AGENTS.md "Nothing lands dead" requires this be said in the body. The paged QSA consumer — `Qwen4ExpQsaPagedCaches` and `RunQwen4ExpQsaBlockPaged` (`src/vllm/model_executor/models/qwen4_exp_qsa_block.{h,cpp}`) together with the `kv_block_table`/`kv_block_size` address mode on `vt::Qwen4ExpQsaGatherAttention` — is reached at this merge commit ONLY by `tests/vllm/models/test_qwen4_exp_qsa_block.cpp`. `ForwardQwen4ExpForConditionalGeneration` still refuses `Qwen4ExpForConditionalGeneration` by name because the layer loop is unwritten, so the reachability mutation has no production call site to delete. The wiring is owned by row `MODEL-MM-QWEN4-EXP`, tracked by [mudler#2031](mudler#2031) and [mudler#2249](mudler#2249), and the row spec's `## Owed` records the wave as UNREACHED with those issues — the entry is **W5d-3 (mudler#2249 item 2) lands UNREACHED, by AGENTS.md "Nothing lands dead"**, and it names the same four residuals listed under "What this does NOT close" below. ## The production refusal, repaired for the POST-MERGE state `qwen4_exp_registry.cpp` carries its own rule three lines above its `VT_CHECK`: "A refusal that names finished work sends the next reader to rebuild it." This string has now gone stale five times on this row, and the fifth is this merge. `origin/main` had already removed the MoE-adapter clause when W5d-4 (mudler#2281, `3f9177f7f`) landed mudler#2249 item 4, and still listed the paged QSA consumer that this branch closes. This branch had already removed the paged-consumer clause and still listed the MoE adapter. **Both sides say TWO items remain, both are exactly one item too long, and git merges the prose around them without a conflict**, so taking either side whole puts a refusal naming finished work on `main` — the exact failure [mudler#2288](mudler#2288) records, and the exact failure that already caught the sibling wave. The enumeration is the SET DIFFERENCE of the original five-item survey against EVERY landed wave: | survey item | fate | |---|---| | (1) a standalone grouped RMSNorm op | closed by **W5d-1** (`25ee19464`, mudler#2265) | | (2) a PAGED QSA consumer | closed by **W5d-3**, this branch, K/V half only | | (3) reach for the indexer side cache, group-2 block table | **REMAINS**, owed as W5c-2 | | (4) a MoE weight adapter onto `MoeBlockWeights` | closed by **W5d-4** (`3f9177f7f`, mudler#2281) | | (5) an externally linked mRoPE builder | closed by **W5d-2** (`3ed2378a3`, mudler#2264) | So **FOUR of the five are closed and ONE remains**, and one is a number neither side of this merge reached. That item 3 genuinely remains was checked against the tree and not against either side of the conflict: `git grep -n gather_block_table origin/main -- src` shows it called for `full_attn_group_id_` and `gdn_group_id_` only (`src/vllm/v1/worker/gpu/runner.cpp:2031`, `:2048`, `:4390`, anchors read back at this head rather than offset), and `git log --grep W5c-2` over `origin/main` returns no commit that lands it. **VERIFIED BY READING THE EMITTED BYTES, NOT BY GREPPING THE FILE.** A temporary `MESSAGE(...)` was inserted in the `SUBCASE("the forward")` of `tests/vllm/models/test_qwen4_exp_scaffold.cpp`, the suite rebuilt (build rc 0) and run with `-s`. What the hook actually threw: ``` vt: Qwen4ExpForConditionalGeneration: the forward is not ported yet. The ops and block seams ARE on main (W2/W3/W4/W6a/W5a/W5b-1..6, W5c-1, W5d-1, W5d-2, W5d-3, W5d-4); what the layer loop still lacks is reach for the indexer side cache — W5d-3 gave the QSA layers' PAGED K/V a consumer (RunQwen4ExpQsaBlockPaged), but KV group 2 is still contiguous and GPUModelRunner::gather_block_table never gathers its block table (W5c-2). ModelRegistry::Forward additionally refuses any multi-cache topology by name, and this model publishes one. See .agents/specs/qwen4-exp-flash-next.md and issues mudler#2031 and mudler#1978. ``` The instrumented file was then restored byte-for-byte, proved by sha256 `32ee46d64b3045cbf852c387a2cb106bc46d391ed9fd19b4b6cbc013afac9b07` and by a clean `git status` on that path, and the target rebuilt from the restored source. The suite is unchanged at 12 cases / 296 assertions, because `MESSAGE` adds no assertion and no substring assertion was touched. **No gate could have caught the wrong resolution, and that is the durable finding.** That `SUBCASE` pins five substrings — "Qwen4ExpForConditionalGeneration", "forward is not ported", "W2", "W4", "mudler#1978" — plus the ABSENCE of "was not produced by". All five survive every wrong variant of this message, because the suite pins that those substrings are PRESENT and can never pin that the enumeration is TRUE. It is a spelling gate, not a truth gate. A file-wide grep is misleading here too: "was not produced by" DOES appear in the file, in a comment, and never in the string the test reads. Two comment paragraphs in that file were REMOVED rather than merged, because they closed on "the count is TWO" and that is the sentence this merge falsifies; the history they carried is restated in the paragraph that replaces them. The seam list gains both W5d-3 and W5d-4, which is the convention W5d-4's own landed message set when it listed itself. ## What this does NOT close - The INDEXER side cache is still contiguous — mudler#2249 item 3, KV group 2 is never gathered, owed as W5c-2. Not smuggled into this wave. - The block still serves ONE request per call. `kv_block_table` is `[1, max_pages]` and anything else is refused by name; a ragged multi-request batch needs the `query_start_loc` plumbing `vt::PagedAttention` carries. - There is no CUDA arm of the address mode, so nothing here says a GPU resolves a page the same way. Inherited debt: the op has one registered arm, CPU. - No real checkpoint is decoded through the paged arm. - An fp8 paged cache is refused by name rather than served. ## Records **No `.agents/issue-index.md` row is appended.** mudler#2249's single row is already on `main`. The merged index is byte-identical to `origin/main` at this head, resolved as a SET against the merge base per [mudler#2266](mudler#2266) rather than from a row count, because a count cannot tell "union dropped my row" from "the other side added one". Observed invariants: `main`'s file is a strict prefix of the result (898 lines each, `cmp` clean); every merge-base row present byte-for-byte, 0 missing; **880 rows, 0 duplicate ids** matched on `^| \[#NNNN\]`; exactly **one** `mudler#2249` row. Both merges were checked this way, and `main`'s own appended rows arrive unopposed because this branch appends none. `.agents/specs/qwen4-exp-flash-next.md` is a keyed record, and AGENTS.md `## Records` refuses an automatic three-way merge of one. `origin/main`'s complete file was taken, this branch's five scoped edits re-applied, and the result diffed back: **5 hunks, 21 deleted lines**, each enumerated as intended — 1. `@@ -2261` (+52, −0): adds `## Mutation record — W5d-3`, after `main`'s W5b-6, W5d-2 and W5d-1 records in landing order. 2. `@@ -2779` (+14, −6): the `## Owed` "PAGED cache" bullet becomes HALF DISCHARGED, naming the K/V half as done and the indexer half as W5c-2. 3. `@@ -2854` (+35, −0): adds the `## Owed` entry recording this wave as UNREACHED. 4. `@@ -3605` (+11, −1): survey item 2 marked CLOSED for the K/V half, survey text kept because the layer loop still has to CALL the paged arm. 5. `@@ -3772` (+46, −14): the `## Now` recount, which BOTH sides rewrote. This is the one hunk that conflicted, and it is hand-authored to carry BOTH facts — W5d-4's closure of item 4 and this wave's closure of item 2 — so the section now reads FOUR closed, ONE remaining, and shows the set-difference arithmetic rather than only the result. Every unrelated key in the file is byte-for-byte equal to `origin/main`; the diff has exactly these five hunks and no others. `origin/main` moved FIVE times under this branch and every merge carries a hand-authored message, because `git merge --no-edit` produces a trailer-less merge commit that reds the trailer walk with no non-force repair. `cff257622` is the product merge that forced the recount above: it carries W5d-4 (mudler#2281), whose `qwen4_exp_moe.{h,cpp}` closes survey item 4. `207c12932` (`PERF-QWEN35-STAGE-WEIGHTS`, mudler#2327/mudler#2328) is the last, and its intersection with this wave's eight files is empty, so no suite count is restated as if retaken for it. ### The gate, stated exactly `scripts/agent-preflight.sh --fail-on-skip` was run twice, once per merge. At `469ffb588` (merged with `cff257622`): **rc 0**, `All gates green.`, **zero skips and zero failures**. The three `skip` matches in that log are test NAMES (`test_agent_preflight_skip_report`, `test_tower_skip_rss_report`, `test_tower_skip_rss_arm`) and the one `FAIL` match is prose inside the appended `.agents/NOW.md`. At this head, after merging `207c12932`: **rc 1**, `2 gate(s) failed: check-env-doc test_check_env_doc`. Zero skips; every other gate `ok`, including `now-current range`, `issue-index append-only`, `commit-trailers` and `commit-style`. **THAT FAILURE IS INHERITED FROM `main` AND IS NOT THIS BRANCH'S**, and it is stated here rather than worked around: ``` ERROR: production env var(s) read from src/+include/ are neither documented in docs/ENVIRONMENT.md nor on scripts/env-doc-allowlist.txt: - VT_QWEN35_STAGE_MIN_FREE_FRAC ``` `VT_QWEN35_STAGE_MIN_FREE_FRAC` is read from `src/vllm/model_executor/models/qwen3_5_weights.cpp:218` and documented only in a comment at `include/vllm/model_executor/models/qwen3_5_weights.h:1340`. Both files arrived with `207c12932` (mudler#2327/mudler#2328, row `PERF-QWEN35-STAGE-WEIGHTS`); neither is among the eight files this branch changes, and `git diff origin/main | grep -c VT_QWEN35_STAGE_MIN_FREE_FRAC` is **0**. It is already tracked by open issue [mudler#2329](mudler#2329), which names mudler#2328 as its cause, so it is another row's obligation and documenting it here would be a record edit riding in the wrong pull request. Every branch cut after `207c12932` inherits it, and CI builds the merge commit, so it is inherited whether this branch merges forward or stays behind. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5 [Claude Code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…ock table, so the QSA indexer side cache stops being unaddressable (mudler#2249) (mudler#2330) `GPUModelRunner::gather_block_table` had three call sites and reached exactly two group ids, `full_attn_group_id_` and `gdn_group_id_`. A model that publishes a THIRD group therefore had that group's cache allocated and no map from a logical position into its pages. `qwen4_exp` is such a model: its QSA indexer side cache is an `MLAAttentionSpec` at `compress_ratio` 4, published as group 2 beside the QSA layers' paged K+V and the uniform recurrent group. That is mudler#2249 item 3. **It is the LAST of the five prerequisites mudler#2249 surveyed.** With W5d-3 (`787373626`, mudler#2276) on `main`, this wave takes the survey's remaining set to ZERO — item 1 is `vt::RmsNormGroup` (W5d-1), item 2 is `RunQwen4ExpQsaBlockPaged` (W5d-3), item 3 is `GPUModelRunner::gather_group_block_tables` (here), item 4 is `qwen4_exp_moe.{h,cpp}` (W5d-4), item 5 is `BuildMropeCosSinHost` declared in `include/vllm/model_executor/models/qwen3_5_mrope.h` (W5d-2). Each was checked by resolving the symbol in this merged tree, not by reading either side's prose. The forward still refuses, because prerequisites being BUILT is not the loop being WRITTEN. ## Upstream has no two-named-group shape at all The two named ids were OUR divergence, not a mirror of anything. At the parity pin `5559679229`, vLLM builds per-group attention metadata by iterating `enumerate(kv_cache_groups)` and handing each group its own table — `cm.block_table_tensor = _get_block_table(kv_cache_gid)`, `vllm/v1/worker/gpu_model_runner.py:2551-2567`. `_get_block_table` (`:2318-2334`) is `self.input_batch.block_table[gid].get_device_tensor(...)`, which is byte-for-byte what this tree's `gather_block_table` already does. There is no "the two special groups" anywhere in that path. So the fix is a LOOP over every published group, `GPUModelRunner::gather_group_block_tables`, and not a third named id. Group 0 is re-gathered inside that loop, at the cost of ONE EXTRA COPY of a table the step already built, and the reason is stated rather than assumed. Upstream gathers group 0 once BEFORE its loop (`block_table_gid_0 = _get_block_table(0)`, `:2337`), carries it into every iteration on `cm_base`, and guards the loop body with `if kv_cache_gid > 0:` (`:2565`) — so upstream does NOT re-gather group 0 inside the loop and this tree does. That guard is an OPTIMISATION, not a second convention for "which groups are special": the table it skips rebuilding is the same table. Paying the copy buys a vector with no index the reader has to know is special. ## The carrier, and what was rejected The tables travel on `MultiKvCacheIndex` — the forward channel KV-DSV4-MULTICACHE W3 built in mudler#2068 (`ca3dcda21`), which already carries which group each published cache came from. This wave EXTENDS that channel rather than opening a second one: two new vectors, `group_block_tables` and `group_block_table_cols`, sized to the published group count and INDEXED BY GROUP ID rather than parallel to `attn_kv`, because a block table belongs to a group and every layer in that group shares one. That is upstream's own fan-out ("make layers in the same group share the same metadata", `:2551-2552`). Both stay empty on every uniform topology, which is the same empty-means-unchanged contract mudler#2068's other vectors keep, so no model shipping today pays a copy or can observe the channel. `CommonAttentionMetadata` was rejected as the carrier because it holds exactly one table, the target group's. Widening it would put a multi-cache field on every uniform step. ## What is UNREACHED The gather runs on the production `execute_model` path and its count is READ by the multi-cache refusal in `ModelRegistry::Forward`. Nothing CONSUMES the tables. No forward reads them, because `ForwardQwen4ExpForConditionalGeneration` refuses by name and `ModelRegistry::Forward` refuses every multi-cache topology by name. What is unreached is therefore the per-group block-table channel's VALUE. The row that owns the wiring is `MODEL-MM-QWEN4-EXP`; the issues that track it are mudler#2031 (the layer loop) and mudler#2249 (the prerequisite survey). It is listed under `## Owed` in `.agents/specs/qwen4-exp-flash-next.md`, which is the entry AGENTS.md §"Nothing lands dead" requires. ## The mutation battery Every mutation is a single edit to `src/vllm/v1/worker/gpu/runner.cpp`, applied to a file whose pre-mutation sha256 is `b1fcc71a36ee1ac02b87cbc8786958f301177661f3ce07287911d5785c3ae889`, built to a recorded rc BEFORE any test result was read, and restored to that same sha256. | ID | Mutation | Build rc | Result | |---|---|---|---| | M1 | gather every group with the WRONG id — `gather_block_table(gdn_group_id_, …)` instead of `(g, …)` | 0 | **RED**, 6 of 44 | | M2 | OFF-BY-ONE in the group index — `g == 0 ? 0 : g - 1` | 0 | **RED**, 4 of 44 | | M3 | REACHABILITY — the production call site in `execute_model` DELETED | 0 | **RED**, 2 of 1001 over the whole suite | M3's bound is stated rather than left to be inferred. It is NOT vacuous — there IS a production call site and deleting it reds the gate — but what it proves is that the RUNNER reaches the gather, not that any forward reaches the tables. No forward does. The fixture discriminates because it has no fixed point: `MakeQwen4ExpShapedKvConfig()` gives each of the three groups a DISTINCT two-block list, `{6, 2}`, `{4, 7}` and `{5, 3}`, over a 20-token prompt on a 16-token block. An identity table would have proven nothing, because a body that ignores the table, returns the logical indices, or reads another group's table all agree with `{0, 1}`. ## The refusal string, and why NEITHER side of the merge was true `ForwardQwen4ExpForConditionalGeneration`'s message enumerates the prerequisites the layer loop still lacks. Two waves rewrote that same literal on `main` while this one was under review, and the final merge against `787373626` left both sides false: `main`'s literal enumerated the group-2 block table THIS change closes, and this branch's enumerated the paged QSA consumer W5d-3 closed. Each side was exactly one item too long, both said ONE, and taking either verbatim would have put a refusal naming finished work on `main` — the sixth turn of the staleness class mudler#2288 names on this row. The literal is resolved BY HAND and the count is the SET DIFFERENCE of the five-item survey against every landed wave: ZERO. The enumeration is not deleted, because a survey that falls silent reads as an unfinished one. The message now says the prerequisites are done and names what is left — the LAYER LOOP itself, `Qwen4ExpTextModel::Forward` under mudler#2031, with the QSA indexer side cache's PAGED STORE owed beside it — and it keeps the clause that `ModelRegistry::Forward` refuses any multi-cache topology by name, which this model triggers and which mudler#2249 records as an engine seam rather than this row's. **The emitted bytes were RE-VERIFIED on this merged head**, by a temporary `MESSAGE` in `test_qwen4_exp_scaffold.cpp`'s `SUBCASE("the forward")` run with `-s`, after which the file was restored to its pristine sha256 `32ee46d64b3045cbf852c387a2cb106bc46d391ed9fd19b4b6cbc013afac9b07`. They were not grepped out of the source: `test_qwen4_exp_scaffold.cpp:767` pins five substrings plus one absence, so it pins the message's PRESENCE and never its truth, and all five survive every wrong variant. All five were confirmed present in the captured bytes and the absence still holds; the suite printed 12 cases / 296 assertions on that same run. The captured string is quoted verbatim in the spec's `## Mutation record — W5c-2` section. ## Records, resolved rather than auto-merged `.agents/specs/qwen4-exp-flash-next.md` is a keyed record, so AGENTS.md §"Records" forbids accepting its automatic three-way merge. `origin/main`'s complete version was taken and this branch's eight scoped edits re-applied onto it: **nine hunks, 206 added lines and 43 deleted lines**, every deletion enumerated and intended, every unrelated key byte-for-byte equal to `main`. The recount paragraph is one of them and it is why the rule exists — it merged with NO conflict marker while saying "FOUR OF THE FIVE ARE NOW CLOSED" and "**ONE remains**", both false after this merge. It now reads ALL FIVE and **NONE remain**, with items 1-5 each named by the symbol that closes it, and the `## Owed` survey says the prerequisites are done while the loop is not. `.agents/issue-index.md` is byte-identical to `main` (1373251 bytes, 882 rows, zero duplicate ids, exactly one mudler#2249 row): this branch appends no row, because mudler#2249's row is already on `main`. Two comments `main` carries were falsified by this merge through conflict-free hunks and are repaired in the same flow: `qwen4_exp_qsa_block.h` said twice that the runner does not gather KV group 2's block table. Two line anchors are re-pinned by reading the line back rather than by applying an offset — the refusal at `qwen4_exp_registry.cpp:266` and its `.forward` registration at `:494`, and `glm5_next_registry.cpp`'s citation of `qwen4_exp_registry.cpp:433`, which named `:382` and was already stale on `main` before this merge moved it further. ## Gates `scripts/agent-preflight.sh --fail-on-skip` on this merged head: **rc 0, ZERO skips**, 109 gates green, none red. `check-env-doc` and `test_check_env_doc` were inherited red from `main` on the previous head; mudler#2329 landed as part of `ebc154341` and both are now green, so the branch carries no inherited red. Suites on this head, every build rc read before any test result (both builds rc 0): `test_runner` 34 cases / 1038 assertions, `test_qwen4_exp_scaffold` 12 / 296, `test_qwen4_exp_kv_cache` 4 / 399, `test_qwen4_exp_qsa_block` 11 / 4382 — the figure mudler#2276 recorded when W5d-3 landed its paged cases, measured here rather than assumed. All rc 0. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: CLAUDE:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…hrough the decoder layer and binds the DSA block to its published KV groups (mudler#2334) W5b-2a threads the four-stream manifold through the decoder layer and binds the DSA block to its published KV groups `Glm5NextTextDecoderLayer` (`modeling_glm5_next.py:1279-1329`), `Glm5NextTextModel.forward` (`:1431-1494`) and the KV binding, as `glm5_next_layer.{h,cpp}` plus an additive `DsaCache*` on `glm5_next::Attention` and a `SelectIndexerTopkFromPacked` lifted out of `SelectIndexerTopk`'s body. Both new parameters default to the uncached shape, so W5b-1's two suites are unchanged and still green at 160 and 1934 assertions. The oracle is `transformers` v5.16.1, this row's lane pin; vLLM registers no `glm5_next` at the parity pin `555967922` nor at its `main`, and neither do vllm-omni, SGLang or llama.cpp. `modeling_glm5_next.py` sha256 `2092bbb4efa2a8087b74f4a4da37635c503fe1df9ae73f1e6e8342af8b4b8e8b`, asserted by the generator against the module it imported. ## The manifold is what this exists to get right `:1477` expands the embedding to `hc_mult` residual streams and nothing collapses them until `hc_head` at `:1493`. A port that threads `[T, hidden]` and collapses early runs, is finite, is the right shape, emits fluent tokens, and leaves every sublayer gate on this row green -- and there is no end-to-end token gate for this model on this fleet to catch it downstream. Three non-overlapping gates: * the per-layer `[B, S, 4, H]` streams asserted elementwise; * the oracle's own minimum pairwise stream separation, **6.4703**, so those assertions are shown to discriminate rather than being four copies of one value; * an early-collapse decoy produced by the SAME oracle modules with the manifold collapsed to its mean and re-broadcast after every layer, which the gate asserts we differ from by the oracle's own **2.4032**. The fixture is a five-layer mixed schedule carrying all four control-flow combinations `:1261-1272` selects between plus a `shared` DSA layer, at the published `hc_mult` of 4, `seq_len` 12 against `index_topk` 8, row 1 left-padded by three. ## The cache stores the latent, not what the reference stores Upstream caches the expanded K/V at `:1175-1179` -- 32,768 values per token per layer. `DsaCache` stores the 512-wide `k_pass` and the 257-wide packed indexer row, which is exactly what `MakeGlm5NextKVCache`'s groups 0 and 2 already publish, and a case asserts that against the production `make_kv_cache` hook on the published `config.json`. The spec decided this and said the equivalence was to be proved rather than asserted; the proof is that `ExpandKv` is token-wise under NoPE, and it is a case asserting `ExpandKv(a ++ b) == ExpandKv(a) ++ ExpandKv(b)` exactly. An 8-token prefill plus 4-token continuation reproduces the reference's own `DynamicCache` run and agrees bit-exactly with the 12-token one-shot tail. ## Red first, and the red was in the oracle configuration The first run read 4 of 10 cases and 7 of 1647 assertions failed, layer 0 green and every DSA layer red by 2.7 to 8.1. Bisected against oracle intermediates rather than guessed: the mHC pre plus `input_layernorm` agreed to 4.8e-07 and the attention output did not. `Glm5NextPreTrainedModel` sets `_supports_sdpa = True`, so a default config resolves `_attn_implementation` to `sdpa`, whose `build_attention_mask_from_topk` returns a boolean mask (`:1249-1250`) where the eager arm returns the additive `finfo.min` one (`:1252-1256`). The two disagree on a left-padded query row with every key masked: SDPA emits 0.0 and eager's uniform softmax emits the mean of the values, measured 0.0 against 0.509. The generator pins eager, which is the arm W5b-1 gated and the only interface `:1227-1228` says a 3-D per-(query, key) mask can reach. No token gate could see this, because the rows that differ are padding. ## One mutation found the instrument rather than the port Fourteen negative mutations, each sha256-proved applied, built and restored byte-for-byte, with BUILD and TEST exit status recorded separately. All fourteen kill their gate -- after a fifteenth finding that is the one worth reading. The mutation that truncates the attention's key range under a filled cache SURVIVED at 1647 of 1647. Its output is all-NaN, `NaN - want` is NaN and `NaN > x` is false for every x, so the running maximum in the test's own `MaxGap` helper never left its initial zero: an all-NaN forward read as a PERFECT MATCH on every gap assertion in the file. `MinStreamSeparation` and the cached-tail loop were blind the same way, because `std::max(m, NaN)` returns `m`. All three now treat a non-finite value as an infinite gap and report the count separately, so a failure distinguishes "wrong number" from "not a number"; the mutation then reds 3 assertions and the suite grew from 1647 to 1656. One earlier mutant also failed to BUILD under `-Werror` on an unused variable, which is a passing mutant proving nothing, and was rewritten to keep the variable used and the arithmetic wrong. ## Not reached, and this wave's own scope said it would be `ForwardGlm5NextForConditionalGeneration` still refuses by name, so the only call site of the new files is their own gate's, and **O15, O16, O17, O23 and O25 are NOT discharged**. `.agents/reachability.md` is explicit that an intermediate hop that is itself unreached does not carry, so a KDA arm now called from `glm5_next_layer.cpp` is exactly as unreached as it was. What changed is that five separate dead ends became one gated assembly point. Spec `## Owed` O26 carries the disclosure in that strong form, and **W5b-2b owns the wiring**, which is why this does not close mudler#2241. W5b-2b's two halves are scoped from measurement rather than guess. The weight bridge for the KDA, MoE, dense-MLP and mHC arms is the larger half, because the 42 sparse layers' routed experts are ~1,150 GiB in f32 against a ~119.63 GiB box, so an on-demand decode of the 8 experts a token selects is the only shape that fits -- `kBridgeTensorF32ByteCeiling` correctly refuses a 9.0 GiB bank today, which is O25's gate working as designed. The engine binding is the smaller half and has a house pattern: `NemotronHForCausalLM` (`nemotron_h_registry.cpp:200-213`) and `KimiLinearForCausalLM` (`kimi_linear_forward.cpp:462-476`) both carry a host arm that ignores the paged caches and re-runs the whole prefix, and a survey of every `: public LoadedModel` found no model keeping per-request state on it. ## Gates `scripts/agent-preflight.sh --fail-on-skip` green, all gates. Suites built and run by hand: `test_glm5_next_layer` 10/10 cases and 1656/1656 assertions, plus the eight siblings unchanged and green -- attn 160, dsa 1934, bridge 96, moe 1614, mhc 98, scaffold 2660, gguf_load 8731, kda 342. No token, no load and no speed number is claimed, and none was observed. No oracle for this model runs on any device this project reaches, so what is gated is the control flow and the numerics of a five-layer stack at tiny shapes and nothing about the 321.32B model. CPU only: no `rc` lease, no GPU, no checkpoint download and no materialising load. Refs mudler#2241 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…n on the only server it will ever face, and the two arms were not commensurable (mudler#2060) (mudler#2267) A fresh review returned FAIL on the llama.cpp arm of `MODEL-MM-QWEN4-EXP`. The `gateable = yes` promotion it examined is JUSTIFIED and is not touched here: every load-bearing number was re-checked against the primary artifacts (`props.json`, `completion.json`, `llama-server.log`, `cubin.log`), and the run half of [mudler#2060](mudler#2060) is genuinely discharged. The failures were in the harness and in the records. **The KV guard was failing open on the real denominator, which is the one that would have hurt.** `assert_reported_kv_fits` grepped the server log for `KV self size = N` and, finding nothing, set the term to 0 and passed. `llama-server` at this pin prints no such line. Measured on this change's own production capture rather than reasoned about: `decode-proof/llama-server.log` is the COMPLETE unfiltered server output at 1,862 bytes, and it carries no `KV self size`, no `llama_kv_cache:` sizing line and no allocation summary of any kind, with sixteen minutes passing between `load_model:` and `threadpool init` and nothing printed in between; `/props` carries no KV bytes either, its only sizing fields being `n_ctx = 4096` and `total_slots = 1`. `KV_BYTES_PER_TOKEN` also defaults to 0, so the static check was weights-only by design. On the real box, therefore, NEITHER check carried a KV term, while the ladder configures `CTX_TOTAL=49152` over 32 slots against a 67.5 GiB model on a 119 GiB unified-memory device that reboots rather than swaps. Nothing caught it because every server stub in the suite emitted the line: the fixture and the measured denominator disagreed. The guard now prefers the engine's own number if a pin ever prints one, falls back to a supplied `KV_BYTES_PER_TOKEN` and names the fallback in its output, and REFUSES (`E_KV_UNREPORTED`, 21) when neither exists. The fixture server is now silent about KV exactly as the real one is, so every execute-path test feeds the guard a log without the line and passes only because the environment supplies a per-token cost; one test removes it and asserts the refusal with no leg written, and a mutation restores the zero-term fall-through and asserts the same scenario stops refusing. `/metrics` was read and rejected — it publishes a KV usage RATIO, not a size — and a post-launch RSS check was rejected because whether GB10's unified `cudaMalloc` appears in `smaps_rollup` cannot be settled without a lease. **Nothing binds, so the ladder must not run until a leased load supplies the number** ([mudler#2261](mudler#2261)). **The two arms were not commensurable, which is the one thing this harness exists to guarantee.** The previous commit claimed every timed request used "the flags `build_client_command` already issues". It was wrong twice. `--num-warmups` was hardcoded to `0` while `OnlineRun.num_warmups` defaults to the CONCURRENCY, so at c=32 our arm would have taken 32 warmups and llama.cpp none, and this repository's own `docs/benchmarks/vllm-online-serving.md` records warmup as exactly where two engines differ most. And ONE corpus file served all 18 legs while `prepare_corpus_views` builds provably disjoint per-(concurrency, repetition) partitions and refuses overlap, so with `--disable-shuffle` every repetition replayed the same prompts and warmed `llama-server`'s slot prefix cache in a way our arm never sees. Both are reconciled rather than documented: the harness passes the concurrency, and `--corpus-dir` now consumes the same `c<C>-r<R>.jsonl` partition set by the same names, with all 18 asserted before the lock. The claim is now a test: one case reads the flag SEQUENCE out of `build_client_command` and compares it to the invocation, two more read what the client was actually handed on all 18 legs, and a fourth pins `REPETITIONS`, the third grid axis, which had no assertion at all. **Three record repairs.** The oracle file said `gateable = yes` at the top and, forty lines lower and in the present tense, "Runs: NOT measured … the flag says `no`"; that paragraph is now dated, past-tense history. The cubin histogram reads `142 sm_121a`, not `142 sm_121` — and the root cause is in the tree: the build script ran `grep -o 'sm_[0-9]*'`, a pattern that CANNOT print the `a`. The error was in this project's own disfavour, because the histogram is the second, independent witness that the `121` asked for became the `121a` the GB10 wants. And G4 named `llama-cpp` at released `b10451`, which cannot name `qwen4exp` — this change's own oracle record documents that at rc=1 with a working `qwen3vl` control — so it named a tool that refuses the model. G4 now names `llama-cpp-qwen4exp`, links mudler#2060, and states the 1,024/128 ISL/OSL the published grid's values already used while the spec was silent. That spec is a keyed record under concurrent edit, so the edit was applied to the complete `origin/main` version and the result diffed back, and it has been re-applied that way again since — see the section below for the current counts. The stale "W6 owes a local sha256" is corrected in both places — this change computed all three and they agree with the Hub `lfs.oid` values three for three. **The fidelity caveat no longer rests on the reader.** The denominator is text-only (`modalities` all false) while this row is a multimodal port, and nothing enforced which configuration our arm must run. The rendered table now names it: same UD-IQ1_S artifact, text-only prompts, no vision or video tower resident, 1,024 in / 128 out, the cell's own corpus partition, warmups equal to the concurrency. **Reproducibility, decided rather than deferred.** The build and decode-proof scripts lived on the NAS, so the two jobs that flipped an oracle to `gateable = yes` were not reproducible from the repository: they are now `scripts/qwen4exp-llamacpp-build-cuda.sh` and `scripts/qwen4exp-llamacpp-decode-proof.sh`, pinned to the oracle record by a new test class. The toolchain is recorded but NOT pinned — `apt-get install -y cuda-toolkit-13-0` names a channel, and a rerun gets whatever apt serves — so the build now asserts `EXPECT_NVCC=13.0.88`, the version the evidence records, and exits 89 on a mismatch; drift is visible instead of silent, which is not the same as pinned. Only 2 of the 11 recorded mutations are committed as tests, because the sweep driver is not in the tree; that and the toolchain are owed under [mudler#2262](mudler#2262). The decode proof also printed the staged binaries' sha256 under a comment saying they "must equal BUILD-RECORD.txt" and compared nothing; it now reads the pin out of that record, asserts it, and hard-exits on any hash that differs. And the suite ran on no lane at all — absent from preflight's `SUITES` and from CI — so it has a CI line now. That closes the CI half only: `scripts/agent-preflight.sh` still does not list it, so a local preflight does not execute it and it has to be run by name. **No number here is a measurement.** One prompt, one repetition, no clock window, no contention control. The ladder has not run, cannot run until mudler#2261 supplies a KV term, and has nothing to compare against until `ModelRegistry::Forward` stops refusing `qwen4_exp` by name. No GPU lease, no GPU mutex and no clock pin was taken for this repair; every finding above is settled from the recorded evidence, the scripts and the tests. 47 tests pass, up from 32. **Brought current against a `main` that moved 39 commits, in four rounds.** A merge, not a rebase: this branch is published and now carries seven merge commits, and a rebase would need a force push. Only two files overlap, and both are records this protocol forbids resolving by automatic three-way merge. Everything else this branch touches — `tools/bench/online_gate.py`, which the harness reads for its anti-drift assertion, the oracle records, `ci.yml`, and every script and suite added here — is untouched by `main`, checked against the merge-base rather than assumed, so those assertions still hold against the same bytes they were written against. `.agents/issue-index.md` was resolved as a SET against the merge-base every round rather than left to the `merge=union` driver, and the two outcomes are worth separating. In the `d858fa9f9` round the driver was WRONG: it interleaved this branch's mudler#2261 and mudler#2262 AHEAD of `origin/main`'s tail row mudler#2329, so main's file was not a prefix of its result. That answer was discarded and the file rebuilt as main's complete bytes with this branch's two rows appended. In the final `c0fa299b1` round the driver was right, and it is kept on evidence rather than on trust: an independent rebuild of main's complete bytes plus the two rows is BYTE-IDENTICAL to what it produced. At the final head: merge-base 882 rows, `origin/main` 883, this branch 884, result 885, and 883 + (884 - 882) = 885 holds; all 882 merge-base rows present byte-for-byte with none missing; zero duplicate ids matched on the leading `^| [#NNNN]`; exactly one row each for mudler#2261 and mudler#2262 from this branch, one mudler#2249 row inherited from `main`, and one mudler#2241 row main appended last; and `origin/main`'s file a strict byte prefix of the result, by sha256 over its 1,378,584 bytes. mudler#2249 and mudler#2275 each appear exactly once and both arrived from `main`; this branch adds neither. One caveat for whoever merges this file next: `origin/main` is NOT itself a suffix-append of the merge-base, because mudler#2324 and mudler#2241 already sit interleaved ahead of main's own tail rows. The prefix property is a claim about the RESULT against main, never about main against the base. The spec grew from 3,853 to 4,192 lines on `main` in the meantime, so the three-way result was discarded and `main`'s complete version taken again, with this branch's six scoped edits re-applied onto it at zero fuzz, giving 4,259 lines. The final round does not touch it at all -- mudler#2334 is a different model's wave -- and the file is byte-identical across that merge at `55fdfa9ab5eb6dee83ac6b306f6e175a26e49638`, so this resolution stands as made. Each was located by its own unmodified text and each match was unique: all fourteen lines this change deletes were confirmed present exactly once in `main` before anything was applied. Against `origin/main` that is six hunks and fourteen deleted lines, byte-identical in content to the base-to-head diff, and every one of those lines is a sentence this change supersedes on purpose: the two `lfs.oid` "not locally computed" passages at 2 and 3 lines, the seven-line G4 clause naming the stock `llama-cpp` pin, and the two single lines the ladder grid and the corpus-partition requirement extend. The new `### The llama.cpp G4 denominator` subsection lands inside `## Owed` in both versions, which was checked rather than assumed, and it is the only heading the diff adds. **The sixth chance for mudler#2288 is a negative result, and looking for it caught a different defect instead.** The refusal has moved on again: mudler#2330 (W5c-2) makes `ForwardQwen4ExpForConditionalGeneration` say that ZERO of mudler#2249's five prerequisites remain and name the layer loop, `Qwen4ExpTextModel::Forward` (mudler#2031), as what is missing. This branch adds NO product code — its diff against `origin/main` under `src/` and `include/` is empty and `qwen4_exp_registry.cpp` is byte-identical to `main` — so it cannot have falsified a clause, and the statement above that there is nothing to compare until `ModelRegistry::Forward` stops refusing `qwen4_exp` by name was re-verified against the merged tree rather than carried over. **An earlier version of this body claimed the prose this branch adds "carries no `file:line` anchor at all". That was false, and one of the two anchors it does carry had gone stale inside this very pull request.** `docs/bench-evidence/qwen4exp-llamacpp-ladder-arm-20260829.md` cited `qwen4_exp_registry.cpp:174` for exactly the refusal discussed above. It was true when written, went stale at the `a920df81a` merge, which moved the refusal to line 229, and moved again to 266 here. Read back at the pushed head, line 174 is a comment about a dots3-note delegate hop — a line that still resolves and says nothing about a refusal, which is why no gate and no reader caught it. It is re-pinned to 266 by READING THE LINE BACK rather than by applying the offset, and the citation now also names the enclosing hook so the next shift leaves a recoverable anchor instead of a number pointing at unrelated prose. The other anchor, `tools/bench/online_gate.py:59-60,232`, was read back and is correct — `INPUT_LEN = 1024`, `OUTPUT_LEN = 128`, and the six-point `POINTS` grid — and that file is untouched by `main`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude-code:claude-opus-5 [Claude Code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
Five CI checks failed on the PR; four root causes are fixed here (the fifth -- commit trailers on old commits -- is repaired in the same push by a history rewrite). 1. MSVC C4456 (windows-msvc-cpu, windows-msvc-vulkan): glm5_next_weights declared 'const GgufValue* v' in three consecutive if/else-if branches. GCC does not warn across sibling scopes but MSVC does. Rename the second and third to 'gv'. 2. agent-record stale anchor: upstream advanced 10 commits since the PR base, shifting PrepareMarlinResident from line 8614 to 8624 in qwen3_5.cpp. The CI merges the PR head into latest upstream, so the citation at :8614 is stale. Updated to :8624. 3. pr-size unclassified path: t35r3-neartie.json was committed at the repo root, which the pr-size classifier does not recognise. Moved to tools/ where it is classified. 4. Latest upstream/main (10 commits) merged into the row branch to match what the CI sees. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM [omp]
ghazni101
force-pushed
the
row/GFX1100-TG200
branch
from
August 31, 2026 12:50
930f85a to
3c95548
Compare
localai-org-maint-bot
requested changes
Aug 31, 2026
localai-org-maint-bot
left a comment
localai-org-maint-bot
left a comment
Collaborator
There was a problem hiding this comment.
Reviewed exact head 3c95548ec4a7d0ea72f0f7eec35a21d829b72886. This accumulated branch is dirty and cannot merge:
src/vt/rocm/rocm_grouped_gemm.hip:coopis used after its definition was removed; the HIP translation unit does not compile.src/vt/rocm/rocm_paged_attn.hip: the launch uses undefinedkNwarpsGandkQgGidentifiers.- The fused quant token survives reuse of the same per-stream scratch by nonmatching consumers, allowing stale quant data.
src/vllm/model_executor/models/qwen4_exp_registry.cpp: this branch rolls the landed Qwen4-Exp forward path back to an unconditional refusal.src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp: retained f16 merged weights are sent through the quant-only loader.
#2168 and the other split PRs overlap this branch. Keep the split changes, remove superseded paths from this campaign branch, and rebase before another review.
localai-org-maint-bot
pushed a commit
that referenced
this pull request
Aug 31, 2026
## Row `BACKEND-ROCM` — one records-only unit for issue #2427. Historical upstream issue #2164 is deleted or unavailable. It remains only as attribution for the original integration request; issue #2427 owns the live landing. ## Before starting - Issue #2427 is open. Its first body line is exactly `Row: BACKEND-ROCM`, with `BACKEND-ROCM` delimited by backticks. - PR #1936 remains the immutable campaign source at `3a345b5ae5df7cf08f1383b6623b38db9a1335bd`. It mixes product changes with the historical records, so this split remains records-only. - `BACKEND-ROCM` remains `ACTIVE`. This pull request changes no product code, checker, test, public document, or live benchmark index. ## What changed Import the complete GFX1100-TG200 campaign evidence and exact gate prompt from PR #1936. The specification now names issue #2427 as the live landing owner and labels #2164 as historical and unavailable. The T5 record now distinguishes the fresh-build 841 assertions across 19 cases from the old 132,094 count read from a stale ROCm 7.14-era binary. Closes #2427 ## Evidence - [x] Final tracked head `6e25501d029a6efebcf44733a4ad5dea7eec02ab` contains `origin/main` `ffab931668f650cb79e9bf044ff1da34b7d11863` as an ancestor. The integration merges had no conflicts or manual resolutions. - [x] Source-identity verification proved 15 evidence files plus `tools/tg200-prompt.txt` byte-identical to `3a345b5ae5df7cf08f1383b6623b38db9a1335bd`. The three reconciled records are the spec, T20, and T5. - [x] Source commit `09da0553c880a9233dc80aba26ae8aab97aaa825` records the fresh-build 841-assertion, 19-case result and identifies the earlier 132,094 claim as stale-binary provenance. - [x] The semantic check was red before the edits for the issue/spec/T5/body invariants. Twenty-five independent negative mutations then failed for the intended guarantee and restored every fixture byte-for-byte. - [x] `scripts/check-agent-record.py`, `scripts/check-symbol-anchors.py`, `scripts/check-conflict-markers.py`, `scripts/check-benchmark-index.py`, `git diff --check`, the commit trailer check, and the commit style check exited 0 on that head. - [x] `.agents/completed/issue-index.md` stayed at blob `199a10828f22806a6f0fc3b89b01604bafb849a1`, matching the frozen PR head and current `origin/main`. Its historical #2164 row is unchanged. - [x] The landed #2401 repair removed the `ENG-PREFLIGHT-COMPILES` baseline failure. Exact head `6e25501d029a6efebcf44733a4ad5dea7eec02ab` passed the agent-record, symbol-anchor, conflict-marker, benchmark-index, gate-command, 59-case gate-command mutation, PR-size, commit-trailer, and commit-style gates before its normal push. - [x] Fresh independent review of exact head `6e25501d029a6efebcf44733a4ad5dea7eec02ab` returned findings `NONE` and verdict `PASS`. Its full preflight passed every changed-surface gate; only unchanged `test_check_windows_portability` hit `ENOSPC` while the filesystem was full, then passed 101/101 on the required serial rerun. Separate PR-size verification passed. Exact-head GitHub checks remain queued; no CI success is claimed. ## Speed claims - [x] This pull request makes no current speed claim. It preserves attributed historical measurements from an unmerged campaign and labels them as such. ## Honest gaps No product change from PR #1936 is included. The fresh current-position `rocprofv3` capture and the next dispatch-removal implementation remain campaign work, not part of this records landing. No GPU or oracle run applies to this records-only import. The exact-head full preflight is not reported as green because of the unchanged Windows-portability suite's transient `ENOSPC` result. Its 101-case serial rerun passed, and fresh immutable review returned `PASS`. No completed GitHub check result is claimed yet. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
Collaborator
|
Reviewed as part of an external-contributor sweep, alongside #2168, #2361 and #2363. #2168 is now merged, so part of this branch has landed by another route. The mechanical position:
A branch this size cannot be reviewed as one unit, and it cannot land as one either. The repository's rule is that size is a review decision rather than a counter: a change is split when the parts help the reviewer. Here they clearly would, and the campaign table in the description is already the natural seam. Each adopted lever in that table (T5a, T6a, T8, T9, T10, T11, T16, T18, T21, T24, T25, T27) is an independently reviewable change with its own measurement. What I would suggest, and I am happy to help land the pieces:
One thing worth stating plainly so it does not get lost: the description records ~103 tok/s against a 200 tok/s target. That gap is not a reason to reject the levers that are measured and bit-identical, and I would rather see those land individually than sit behind the campaign's headline number. For reference, |
Contributor
Author
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What changed
This PR lands the GFX1100-TG200 decode performance campaign and the
fp8 KV cache decode attention arm on the ROCm backend. The branch
accumulated both bodies of work on
row/GFX1100-TG200.The TG200 campaign
Raises Qwen3.5-4B Q4_K_M text-generation throughput on the RX 7900 XTX
(gfx1100, RDNA3, 24 GiB, ROCm 7.14.0) toward >= 200 tok/s, single
stream, batch 1, greedy decode. Spec:
.agents/specs/gfx1100-tg200.md.Tracking issue: ghazni101#5
Adopted levers and measured gains:
Current measured position: ~103 tok/s (T18 idle-host gate 100.46 tok/s
The remaining gap is launch/sync overhead (~4.2 ms/tok above the weight
read floor), not kernel internals. Next attack: HIP graph capture or
kernel fusion to close the overhead gap.
Closed negative (evidence in
docs/bench-evidence/): T5c MMVQnontemporal, T7 COALK wash, T12 gated-quant fusion, T13 async server
wash, T15 LDS bank conflicts, T17 v_dot2 memory-bound, T19 kGemvWarps
block-limited, T20 full-warp cooperative GEMV (kernel 2.4-3.1x on large
grids but engine wash at small grids).
fp8 KV cache decode attention
--kv-cache-dtype fp8set the KV cache toDType::kI8, which failedthe
kBF16guard on every optimized decode kernel inrocm_paged_attn.hip. The dispatch fell through toPagedAttnOnline-- the reference kernel with per-key
__syncthreadsreduction, no GQAamortization, no vectorized loads. Up to 7.5x decode regression at long
context on gfx1100. Spec:
.agents/specs/rocm-fp8-kv-decode-attn.md. Tracking issue:ghazni101#7
Fix: widen the
VT_ATTN_DECODE_GQA4dispatch guard to acceptkI8KVand add an fp8 dequant load path inside the GQA kernel.
PagedAttnDecodeGqaF32QonTKV(default__hip_bfloat16) withk_scale/v_scaleparamsLoadRowEplFp8<EPL>: vectorized uint8_t loads +F8E4M3ToF32Devdequant x scaleLoadRowEplKv<EPL, TKV>:if constexprdispatch between bf16and fp8
kBF16 || (kI8 && kv_cache_dtype != kAuto)The dequant arithmetic is identical to the existing
LoadKv(uint8_t *, ...)at line 176 -- the fast kernel now calls the same dequant thefallback already used.
Benchmark: Qwen3.5-4B Q4_K_M, RX 7900 XTX, ROCm 7.14.0, 128-token
greedy decode, 4 reps:
Correctness: 3/3 simple prompts produce token-exact identical output
vs the
PagedAttnOnlinefallback. Divergence at thinking-tag near-tiesis the reduction-order risk the bf16 GQA4 path already carries.
Relationship with #2080 and landing order
This PR cherry-picks the W6 ROCm fp8-e4m3 KV cache store and read from
#2080 (commit
87c6f2518516) so the decode attention arm can be testedend-to-end. The fp8 decode read side (this PR, commit
e43a39d45803)depends on the store side (#2080) -- both are needed for fp8 KV cache
to work.
Landing order:
(
ReshapeAndCacheFp8KernelRocm), closes KV-FP8 W6: the ROCm fp8-e4m3 KV cache store and read #2065.87c6f2518516(and its matrix/spec deltas) since the store work isthen in main. This PR keeps the TG200 campaign and the fp8 decode
read side.
All 13 files in #2080 are currently in this PR. After #2080 lands and
this PR rebases, the overlap is zero.
CI status
The current head
b058bb752a42has all four workflows (ci, containers,docs site, triton-aot-sync) in
action_requiredstate after theupstream-main integration merges. The prior CI approval was on
1033485e. Workflow re-approval is needed on the current head. Mergeshould be held until the full integrated head is green.
Pre-existing failures on main (not this PR's responsibility):
windows-msvc-cpu,windows-msvc-vulkan(#503/#584), and theagent-recordhugo missing-binary lane (#1722).How to verify
Build with HIP and run the test and gate suites:
cmake -B /tmp/build -DVLLM_CPP_HIP=ON make -j4 vllm test_rocm_fp8_kv_cache test_rocm_quant_dot test_rocm_backend ctest -R 'rocm|cross_device|quant'The acceptance gate: median of >= 5 reps, idle host, gpu-ctl lock held,
batch 1, one ~512-token real prompt, 256 generated tokens, greedy
(
--temperature 0 --seed 0), throughexamples/vllm-cli.Out of scope
fp8_e5m2 compute, per-attention-head scales, the Metal arm, fast-path
(tensor-core/rocWMMA) fp8 attention kernels, the bf16 decode-opt
kernels, the prefill fp8 path, non-gfx1100 architectures, and the
memory-halving e2e measurement on a ROCm gate model. The 200 tok/s
target is not yet met; the overhead gap is the next attack.
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:GLM-5-2 [OMP]