diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 458b2bd41..b5d17c2de 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -151,8 +151,9 @@ host/sched. Detail: state `KERNEL-FA2-GQA-SWAP-FLIP`. | `KERNEL-MOE-UNQUANTIZED` | Unquantized grouped/batched MoE GEMM | core MoE sources `CMakeLists.txt:1135-1157`; upstream `tests/kernels/moe/test_unquantized_backend_selection.py` | activation/combine subset [cuda_moe.cu:349](../src/vt/cuda/cuda_moe.cu#L349); grouped execution remains NVFP4-specialized | [MoE tests](../tests/vt/test_ops_moe.cpp#L193), [grouped tests](../tests/vt/test_ops_moe_grouped.cpp#L160) | [inventory](specs/kernel-family-inventory.md) | `PARTIAL` | - | | `KERNEL-MOE-QUANTIZED` | FP8/INT8/NVFP4/MXFP4 grouped MoE | CUTLASS/FP4 builds `CMakeLists.txt:865-1002`; NVFP4 oracle `fused_moe/oracle/nvfp4.py:38-276` | NVFP4 fallback [cuda_matmul_nvfp4.cu:761](../src/vt/cuda/cuda_matmul_nvfp4.cu#L761), Marlin [cuda_moe_marlin.cu:156](../src/vt/cuda/cuda_moe_marlin.cu#L156) | [NVFP4 grouped tests](../tests/vt/test_ops_moe_grouped.cpp#L160); 35B gate | [inventory](specs/kernel-family-inventory.md) | `PARTIAL` | - | | `KERNEL-MOE-SPECIAL` | Cooperative top-k, DSV3 router, Hadacore, DeepEP/AllToAll integrations | target/build sources `CMakeLists.txt:408-424,1063-1072,1282-1299`; FlashInfer comm probes `vllm/utils/flashinfer.py:171-203` | - | - | [inventory](specs/kernel-family-inventory.md) | `INVENTORIED` | - | -| `KERNEL-GDN-PACKED-DECODE` | Pure non-spec packed GDN recurrent decode for FP16/BF16/F32 | default and dispatch `vllm/envs.py:117,1123-1125`, `qwen_gdn_linear_attn.py:1286-1298,1644-1695`; packed body/launch `fla/ops/fused_recurrent.py:255-478`; upstream test `tests/kernels/test_fused_recurrent_packed_decode.py:13-98` | public [API](../include/vt/ops.h#L518), [validation](../src/vt/ops.cpp#L1466), portable [CPU recurrence](../src/vt/cpu/cpu_ops.cpp), registered [CUDA kernel](../src/vt/cuda/cuda_gdn.cu#L156), **register-resident perf kernel** [GdnPackedDecodeRegTileKernel](../src/vt/cuda/cuda_gdn.cu#L1183) + [rollback flag](../src/vt/cuda/gdn_packed_reg_tile.h) + [CPU flag test](../tests/vt/test_gdn_packed_reg_tile.cpp), exact [model dispatch](../src/vllm/model_executor/models/qwen3_5.cpp#L1555), [request-identity slot pool](../src/vllm/v1/worker/gpu/runner.cpp), [trace finalizer](../tools/bench/finalize_gdn_packed_trace.py#L218), production [component driver](../scripts/dgx-gdn-packed-component.sh), and [component finalizer](../tools/bench/gdn_packed_component.py) | Clean `f18ca23`/`9ad8fb7` close G0/G1; `f344dec` closes W1D2/G2; `7ff713e` + `24cea4f` close W1D3 structure; the identity-keyed slot fix at `c172336` proved the c16 repair (model gates 235/235, `--diagnostic-c16` 3/3). **W1D3 CLOSES on EQUIVALENCE** over eight sealed component roots + the 8-pair locked c16 A/B (**−0.205% ± 0.30, <1σ**) + the 24-window trace attribution (packed GPU-cheaper, kernel compute −1.30..−1.58%/step, no attributable packed-side cost): no STABLE regression on any axis; every failing axis is a sign-flipping band-edge statistic of a true-zero effect. The eighth (first 22-leg: cold-discard pair + 5 reps) seal `complete-failed` at **38/40 + 8/8 memory**, stability clean, `validation_error=None`, paired-consistency PASS at BOTH c2/c16 (the two fails are c2 `median_tpot_ms` 0.9899 and c2 pooled `p99_ttft_ms` 0.8464). Packed stays the **default** (exact-upstream semantics, 48-launch reduction); `VT_GDN_PACKED_DECODE=0` is the rollback; **no `complete-pass` marker exists and NO speed credit is claimed**. [Component tests](../tests/tools/test_gdn_packed_component.py#L1) **79/79**, all tools **162/162**; [runner test](../tests/vllm/v1/worker/test_runner.cpp) 8/8. Correctness blast radius in [BENCHMARKS](../docs/BENCHMARKS.md); closure [ledger](parity-ledger.md#L469). qkvz (`KERNEL-GEMM-BF16` W2) is UNBLOCKED and the exact grid is authorized. **2026-07-16 PERF LEVER (test-first, CPU-gated, DGX-pending):** the named +2.06 ms/step recurrence-tiling lever is ported into `GdnPackedDecodeRegTileKernel` — one warp per `[BV=32,BK]` tile, state block REGISTER-RESIDENT (no shared-state round-trip, no cross-warp `__shfl` reduction, no `__syncthreads`; mirrors vLLM FLA `num_warps=1`/`num_stages=3` `fused_recurrent.py:256-336`), default-on behind `VT_GDN_PACKED_REG_TILE` (=0 → legacy bit-for-bit, same binary), selected for `bv==32 && dk in {32,128}`. Boundary-fixture bit-exactness preserved (sequential per-row Dk reduction). **PROOF FAILED (dgx 54f0541): oracle boundary FAIL + c16 700.5 vs 793.6 tok/s → default flipped OFF (`309c218`).** **2026-07-16 RESOLVED via the sanctioned vendored Triton cubin (`CLAIM-GDN-DECODE-TRITON`):** Phase-1 cuobjdump MEASURED the codegen cause — vLLM FLA decode cubin REG:205/0-spill vs the naive hand port REG:255+STACK:48 (spills); DECISION = vendored cubin (`gdn_decode_h48`, 27B-only, `VT_GDN_PACKED_DECODE_TRITON` default OFF, hand kernel stays default). DGX gates: AOT op test 28/28, full GDN 49/49, oracle boundary 12/12, **27B model gate 235/235 token-exact with the Triton path ON**, memcheck 0/0. c16 A/B: triton [817.51, 821.06, 822.55] vs legacy [813.77, 815.62, 815.30] tok/s — paired mean **+5.48 tok/s (+0.67%)**, monotone (+3.74/+5.44/+7.25), 3/3 pairs positive; mean TPOT triton [161.04, 160.49, 160.35] vs legacy [162.09, 161.65, 161.93] = **-1.26 ms (-0.78%)** (median TPOT -1.13 ms); w0 cold-discard (triton 821.48/160.44) excluded. Anchors: [shim](../triton_kernels/fused_recurrent_packed_decode.py), [vendored cubin](../src/vt/cuda/triton_aot_vendored/sm_121a/gdn_decode_h48.h), [launcher `TryTritonPackedDecode`](../src/vt/cuda/cuda_gdn.cu), [default-ON flag header](../src/vt/cuda/gdn_packed_decode_triton.h) + [CPU flag test](../tests/vt/test_gdn_packed_decode_triton.cpp), [AOT test](../tests/vt/test_ops_gdn.cpp). ACCEPTANCE MET (oracle PASS + consistent c16 TPOT improvement + no throughput regression). **2026-07-16 DEFAULT FLIP ON (`CLAIM-GDN-DECODE-TRITON-FLIP`):** per MIRROR policy (the vendored kernel IS vLLM's exact token-identical FLA kernel, run by default upstream) `VT_GDN_PACKED_DECODE_TRITON` flipped OFF→ON, joining the sibling GDN Triton kernels; `=0` is the same-binary rollback. Test-first default-ON pure-header predicate + CPU flag test (RED→GREEN 10/10). **35B: no specialization** — excluded at the model level by the dense-only `ShouldUsePackedGdnDecode` (`qwen3_5.cpp:49`) and the launcher guard rejects its `Hv=32` shape (`cmake` H=32=35B) anyway (clean fallback; a 35B cubin would be dead code). Flip gates ALL EIGHT PASS exit 0 (root `~/work/vllm.cpp-gdn-decode-triton-flip` `gates.verdict`, `-DVLLM_CPP_TRITON=ON` + CUTLASS-4.5.0/nvcc-13.0, CUTLASS/FA2 configure-log lines verified, one flock): 27B DEFAULT (Triton) **235/235** + `=0` rollback **235/235**; 35B DEFAULT **315/315** + `=0` rollback **315/315** (inert); AOT op test **28/28** (default fires cubin, `=0` fires legacy); full GDN **49/49 (2,343/2,343)**; oracle boundary **12/12**; memcheck **28/28, 0 errors**. No new A/B (9dd7d3f's +5.48 tok/s / −1.26 ms TPOT stands). The next binding grid runs the Triton decode path by default; no separate flip speed credit is claimed. | [packed decode](specs/gdn-packed-decode.md) | `DONE` | `e47b4d6` | -| `KERNEL-GDN-AOT-BF16` | GDN prefill/decode and BF16-output Triton/CuTe fast-path coverage | backend resolution `qwen_gdn_linear_attn.py:150-411`; FLA kernels `fla/ops/chunk_delta_h.py:43-371`; FlashInfer GDN entry `flashinfer-ref/gdn_prefill.py:37-87` | [canonical AOT contract](../cmake/TritonAOTKernels.cmake#L47), [BF16 manifest bases](../src/vt/cuda/triton_aot_vendored/sm_121a/MANIFEST#L21), [typed dispatch](../src/vt/cuda/cuda_gdn.cu#L2910), [27B-only default/f32 fallback](../src/vllm/model_executor/models/qwen3_5.cpp#L1555) | Existing AOT/safety/native gates are green. The BF16/f32 component is **1.007989×**, 16/20 timing and 2/4 memory. Immutable `3f256ab` is the exact v0.25 denominator: c16 total throughput passes at **1.027889×**, but mean TPOT/ITL remains **0.987450× normalized**; c1-c8 total and decode-shaped axes fail. Diagnostic cross-profiler totals are local `GdnDecodeFusedKernel` 19.101 s/73,578 calls and vLLM fused recurrent 28.659 s/70,848 calls; they are not a valid speed ratio. The independently gateable pure-decode mismatch is now owned by `KERNEL-GDN-PACKED-DECODE`; this row retains prefill AOT, BF16 default and scratch/teardown scope. Every 35B path stays f32; strict component/pool/teardown remain open | [inventory](specs/kernel-family-inventory.md) | `READY` | - | +| `KERNEL-GDN-PACKED-DECODE` | Pure non-spec packed GDN recurrent decode for FP16/BF16/F32 | default and dispatch `vllm/envs.py:117,1123-1125`, `qwen_gdn_linear_attn.py:1286-1298,1644-1695`; packed body/launch `fla/ops/fused_recurrent.py:255-478`; upstream test `tests/kernels/test_fused_recurrent_packed_decode.py:13-98` | public [API](../include/vt/ops.h#L518), [validation](../src/vt/ops.cpp#L1466), portable [CPU recurrence](../src/vt/cpu/cpu_ops.cpp), registered [CUDA kernel](../src/vt/cuda/cuda_gdn.cu#L156), **register-resident perf kernel** [GdnPackedDecodeRegTileKernel](../src/vt/cuda/cuda_gdn.cu#L1183) + [rollback flag](../src/vt/cuda/gdn_packed_reg_tile.h) + [CPU flag test](../tests/vt/test_gdn_packed_reg_tile.cpp), exact [model dispatch (`ShouldUsePackedGdnDecode`, `GdnBlockPaged`)](../src/vllm/model_executor/models/qwen3_5.cpp#L4505), [request-identity slot pool](../src/vllm/v1/worker/gpu/runner.cpp), [trace finalizer](../tools/bench/finalize_gdn_packed_trace.py#L218), production [component driver](../scripts/dgx-gdn-packed-component.sh), and [component finalizer](../tools/bench/gdn_packed_component.py) | Clean `f18ca23`/`9ad8fb7` close G0/G1; `f344dec` closes W1D2/G2; `7ff713e` + `24cea4f` close W1D3 structure; the identity-keyed slot fix at `c172336` proved the c16 repair (model gates 235/235, `--diagnostic-c16` 3/3). **W1D3 CLOSES on EQUIVALENCE** over eight sealed component roots + the 8-pair locked c16 A/B (**−0.205% ± 0.30, <1σ**) + the 24-window trace attribution (packed GPU-cheaper, kernel compute −1.30..−1.58%/step, no attributable packed-side cost): no STABLE regression on any axis; every failing axis is a sign-flipping band-edge statistic of a true-zero effect. The eighth (first 22-leg: cold-discard pair + 5 reps) seal `complete-failed` at **38/40 + 8/8 memory**, stability clean, `validation_error=None`, paired-consistency PASS at BOTH c2/c16 (the two fails are c2 `median_tpot_ms` 0.9899 and c2 pooled `p99_ttft_ms` 0.8464). Packed stays the **default** (exact-upstream semantics, 48-launch reduction); `VT_GDN_PACKED_DECODE=0` is the rollback; **no `complete-pass` marker exists and NO speed credit is claimed**. [Component tests](../tests/tools/test_gdn_packed_component.py#L1) **79/79**, all tools **162/162**; [runner test](../tests/vllm/v1/worker/test_runner.cpp) 8/8. Correctness blast radius in [BENCHMARKS](../docs/BENCHMARKS.md); closure [ledger](parity-ledger.md#L469). qkvz (`KERNEL-GEMM-BF16` W2) is UNBLOCKED and the exact grid is authorized. **2026-07-16 PERF LEVER (test-first, CPU-gated, DGX-pending):** the named +2.06 ms/step recurrence-tiling lever is ported into `GdnPackedDecodeRegTileKernel` — one warp per `[BV=32,BK]` tile, state block REGISTER-RESIDENT (no shared-state round-trip, no cross-warp `__shfl` reduction, no `__syncthreads`; mirrors vLLM FLA `num_warps=1`/`num_stages=3` `fused_recurrent.py:256-336`), default-on behind `VT_GDN_PACKED_REG_TILE` (=0 → legacy bit-for-bit, same binary), selected for `bv==32 && dk in {32,128}`. Boundary-fixture bit-exactness preserved (sequential per-row Dk reduction). **PROOF FAILED (dgx 54f0541): oracle boundary FAIL + c16 700.5 vs 793.6 tok/s → default flipped OFF (`309c218`).** **2026-07-16 RESOLVED via the sanctioned vendored Triton cubin (`CLAIM-GDN-DECODE-TRITON`):** Phase-1 cuobjdump MEASURED the codegen cause — vLLM FLA decode cubin REG:205/0-spill vs the naive hand port REG:255+STACK:48 (spills); DECISION = vendored cubin (`gdn_decode_h48`, 27B-only, `VT_GDN_PACKED_DECODE_TRITON` default OFF, hand kernel stays default). DGX gates: AOT op test 28/28, full GDN 49/49, oracle boundary 12/12, **27B model gate 235/235 token-exact with the Triton path ON**, memcheck 0/0. c16 A/B: triton [817.51, 821.06, 822.55] vs legacy [813.77, 815.62, 815.30] tok/s — paired mean **+5.48 tok/s (+0.67%)**, monotone (+3.74/+5.44/+7.25), 3/3 pairs positive; mean TPOT triton [161.04, 160.49, 160.35] vs legacy [162.09, 161.65, 161.93] = **-1.26 ms (-0.78%)** (median TPOT -1.13 ms); w0 cold-discard (triton 821.48/160.44) excluded. Anchors: [shim](../triton_kernels/fused_recurrent_packed_decode.py), [vendored cubin](../src/vt/cuda/triton_aot_vendored/sm_121a/gdn_decode_h48.h), [launcher `TryTritonPackedDecode`](../src/vt/cuda/cuda_gdn.cu), [default-ON flag header](../src/vt/cuda/gdn_packed_decode_triton.h) + [CPU flag test](../tests/vt/test_gdn_packed_decode_triton.cpp), [AOT test](../tests/vt/test_ops_gdn.cpp). ACCEPTANCE MET (oracle PASS + consistent c16 TPOT improvement + no throughput regression). **2026-07-16 DEFAULT FLIP ON (`CLAIM-GDN-DECODE-TRITON-FLIP`):** per MIRROR policy (the vendored kernel IS vLLM's exact token-identical FLA kernel, run by default upstream) `VT_GDN_PACKED_DECODE_TRITON` flipped OFF→ON, joining the sibling GDN Triton kernels; `=0` is the same-binary rollback. Test-first default-ON pure-header predicate + CPU flag test (RED→GREEN 10/10). **35B: no specialization** — excluded at the model level by the dense-only `ShouldUsePackedGdnDecode` (`qwen3_5.cpp:49`) and the launcher guard rejects its `Hv=32` shape (`cmake` H=32=35B) anyway (clean fallback; a 35B cubin would be dead code). Flip gates ALL EIGHT PASS exit 0 (root `~/work/vllm.cpp-gdn-decode-triton-flip` `gates.verdict`, `-DVLLM_CPP_TRITON=ON` + CUTLASS-4.5.0/nvcc-13.0, CUTLASS/FA2 configure-log lines verified, one flock): 27B DEFAULT (Triton) **235/235** + `=0` rollback **235/235**; 35B DEFAULT **315/315** + `=0` rollback **315/315** (inert); AOT op test **28/28** (default fires cubin, `=0` fires legacy); full GDN **49/49 (2,343/2,343)**; oracle boundary **12/12**; memcheck **28/28, 0 errors**. No new A/B (9dd7d3f's +5.48 tok/s / −1.26 ms TPOT stands). The next binding grid runs the Triton decode path by default; no separate flip speed credit is claimed. | [packed decode](specs/gdn-packed-decode.md) | `DONE` | `e47b4d6` | +| `KERNEL-GDN-AOT-BF16` | GDN prefill/decode and BF16-output Triton/CuTe fast-path coverage | backend resolution `qwen_gdn_linear_attn.py:150-411`; FLA kernels `fla/ops/chunk_delta_h.py:43-371`; FlashInfer GDN entry `flashinfer-ref/gdn_prefill.py:37-87` | [canonical AOT contract](../cmake/TritonAOTKernels.cmake#L47), [BF16 manifest bases](../src/vt/cuda/triton_aot_vendored/sm_121a/MANIFEST#L21), [typed dispatch](../src/vt/cuda/cuda_gdn.cu#L2910), [`GdnOutDType()` — BF16 default on EVERY arm, `VT_GDN_OUT_BF16=0` the f32 rollback](../src/vllm/model_executor/models/qwen3_5.cpp#L172) | Existing AOT/safety/native gates are green. The BF16/f32 component is **1.007989×**, 16/20 timing and 2/4 memory. Immutable `3f256ab` is the exact v0.25 denominator: c16 total throughput passes at **1.027889×**, but mean TPOT/ITL remains **0.987450× normalized**; c1-c8 total and decode-shaped axes fail. Diagnostic cross-profiler totals are local `GdnDecodeFusedKernel` 19.101 s/73,578 calls and vLLM fused recurrent 28.659 s/70,848 calls; they are not a valid speed ratio. The independently gateable pure-decode mismatch is now owned by `KERNEL-GDN-PACKED-DECODE`; this row retains prefill AOT, BF16 default and scratch/teardown scope. The 35B GDN recurrence output and z gate are no longer f32 — `GDN-MOE-BF16-OUT` (#1168) removed the model-shape branch from `GdnOutDType`, so `VT_GDN_OUT_BF16` is the whole decision on both arms; this row's own prefill AOT and scratch scope is unchanged by that. Strict component/pool/teardown remain open | [inventory](specs/kernel-family-inventory.md) | `READY` | - | +| `GDN-MOE-BF16-OUT` | The GDN recurrence output `dcore` `[T,Hv,Dv]`, the `z` output gate and the gated-RMSNorm weight that has to match them follow the ONE model dtype on every arm, dense and MoE alike, instead of resolving bf16 from a dense checkpoint and f32 from a MoE one | vLLM branches on no model shape here: `core_attn_out = torch.zeros(..., dtype=hidden_states.dtype)` `qwen_gdn_linear_attn.py:870-873`, `z` a split of the bf16 `in_proj_qkvz` `:843`, `:859-860`, `RMSNormGated` with no dtype override `:459-465`, one shared `Qwen3_5ForCausalLMBase` `models/qwen3_5.py:280-297` @ `5559679`; `VLLM_ENABLE_FLA_PACKED_RECURRENT_DECODE` defaults True with no shape term `envs.py:124`. SGLang agrees (`qwen3_5.py:522-536`, `gdn_triton.py:43` @ `f63458b5be`) | resolver + parser [qwen3_5.cpp](../src/vllm/model_executor/models/qwen3_5.cpp), contract [qwen3_5_internal.h](../src/vllm/model_executor/models/qwen3_5_internal.h); `detail::ShouldUsePackedGdnDecode` lost its `e.dense_model` term in the same change | CPU tier only. `test_qwen35_paged_forward` enters through `ModelRegistry::Forward` on a MoE config and reads the dtypes off `dcore` and the projected gate, and runs a second time under `VT_GDN_OUT_BF16=0` so the documented rollback is gated too; `test_qwen27_paged_forward` pins the parser truth table and the shape-free eligibility. **Nothing is measured.** The 35B correctness gate, `test_qwen36_paged_engine` 315/315, `test_qwen27_paged_engine` 235/235 inertness, the `VT_GDN_OUT_BF16` 0-versus-1 same-binary A/B and the `nsys` memory-format confirmation are all `PENDING` on a GPU host and are listed in the spec's `## Now` | [gdn-moe-bf16-out.md](specs/gdn-moe-bf16-out.md) | `GATING` | [#1168](https://github.com/mudler/vllm.cpp/issues/1168) | | `KERNEL-GDN-SCRATCH` | Stream-ordered GDN AOT scratch allocation, reuse, and capture safety | FLA/FlashInfer runtime path `qwen_gdn_linear_attn.py:1073-1161`; CuTe Blackwell scheduler `flashinfer-ref/gdn_kernels/blackwell/gated_delta_net_chunked.py:177-182,379-846` | [stream-owned pool](../src/vt/cuda/cuda_gdn.cu#L2662), [queue cleanup](../src/vt/cuda/cuda_gdn.cu#L3641), [test poison hook](../src/vt/cuda/cuda_gdn.cu#L3679) | [same-stream poison/reuse/growth/pool-off assertions](../tests/vt/test_ops_gdn.cpp#L1494); all 11 buffers are filled with `0xff` before reuse and dispatch counters reject hand fallback; integrated current-main GB10 graph gate, sanitizer, trace and pool A/B pending | [inventory](specs/kernel-family-inventory.md) | `READY` | - | | `KERNEL-GDN-REPLAYSSM` | ReplaySSM buffered output-only decode for the GDN path: a per-slot ring of the last `L` steps' rank-1 factors `(d, k, g)`, with the full `[HV,V,K]` state written back only every `L` steps | **Algorithm (vLLM, MIRROR SOURCE) @ `555967922`:** kernel `layers/mamba/ops/selective_state_update_replayssm_output_only.py` (precompute `:22-129`, main `:131-472`, non-flush route `:275-279`, flush route `:358-362`, entry `:474-523`), launch heuristic `ops/replayssm_config.py`, ring dtypes/shapes `layers/mamba/mamba_utils.py:84-93,202-221`, layer wiring `mamba_mixer2.py:1105-1140`, cursor `v1/attention/backends/mamba_attn.py:77-82,575-620`, ring origin `v1/attention/backend.py:483-486`, `replayssm_buffer_len` default 16 `config/cache.py:148`, `use_replayssm` default `False` `:152`, refusals `config/vllm.py:2318-2337` (spec decode `:2329`), introduced `866fea2b` (vllm#48018). **Mamba2 ONLY — four walls to GDN:** `config/vllm.py:2318-2322` refuses a model without `supports_replayssm`; only `NemotronHForCausalLM` sets it (`models/nemotron_h.py:711,755,792`); `GDNAttentionMetadataBuilder` (`v1/attention/backends/gdn_attn.py:82`) does not subclass the builder that derives the cursor; the kernel asserts scalar-per-head `A` (`:540-542`) + Mamba2 `(B,C)` groups (`:528-530`). `gated_delta_net_state_dtype` (`mamba_utils.py:120-128`) appends no ring. Still true 877 commits past the pin. **GDN application (SGLang, SECONDARY ORACLE) @ `f63458b5`:** `layers/attention/fla/fused_recurrent_linear_replayssm.py` (`:50` credits vLLM, `:45-47` the `L=1` algebraic identity), `--enable-linear-replayssm` default `False` + `--linear-replayssm-cache-len` default 16 `server_args.py:1972-1987`, rings `mem_cache/memory_pool.py:465-487`, cursor `:641-645`, dispatch `linear/gdn_backend.py:317-327` + `linear/kernels/gdn_triton.py:84-119`, introduced `a10a24e9` (sglang#28451) | NOT IMPLEMENTED. `grep -rniE 'replayssm' src include tests` is empty. The path this would join: legacy packed decode [cuda_gdn.cu:2312-2427](../src/vt/cuda/cuda_gdn.cu) (state read `:2393`, write `:2425`), reg-tiled `:2505-2630` (default OFF), fused `:2793`, scan `:2209`, dispatch `:2668-2760`, vendored Triton AOT bridge `TryTritonPackedDecode` `:5189-5245`; state spec [qwen3_5_common.cpp:37-106](../src/vllm/model_executor/models/qwen3_5_common.cpp) (shape/dtype vectors `:84-87`), SSM dtype resolution [qwen3_5.cpp:460-470](../src/vllm/model_executor/models/qwen3_5.cpp), speculative arm `vt::GdnSpecDecode` [ops.h:2380-2384](../include/vt/ops.h) | **NO MEASUREMENT EXISTS AND NONE IS CLAIMED.** Source-verified anchors only, read at both pins. §5 of the spec DERIVES from the 27B shape (`HV=32, V=128, K=128, H_k=16`, fixture `tests/vllm/test_hf_config.cpp:80-82`): honest state-traffic ratio **0.61-0.63x** at `L=16`, NOT SGLang's published 0.53x (`fla/bench_gdn_replayssm_decode.py:58-63`), which models `(1+1/L)/2` and omits both the flush re-read and every ring access; ring costs **+18.9% mamba page** (395,264 B against a 2,097,152 B state) so ~19% fewer GDN slots, against vLLM's ~7% on Nemotron shapes, because GDN's state is `V*K` while the ring is `L*(V+K)`; `L* = sqrt(4*HV*V*K/(HV*V + a*H_k*K))` = **exactly 16** at this shape, matching both upstream defaults, with `L=8` and `L=32` both worse at 0.658x. Neither upstream is bit-exact and neither claims to be (vLLM fp32 `rtol=1e-4, atol=1e-3` / bf16 `6e-2, 2e-1`, `tests/kernels/mamba/test_replayssm_standard_decode_mamba2.py:47-57`; SGLang fp32 `atol=1e-4, rtol=1e-3`, `atol=2e-6, rtol=1e-5` at `L=1` alone, `test_linear_replayssm_decode.py:83-88,226-231` — atol/rtol roles are SWAPPED between the two). **TOP RISK, already measured in this tree:** the 27B decode default is the vendored FLA cubin, which hard-guards `state.dtype == kF32` (`cuda_gdn.cu:5212`), and the identical register-resident `[BV=32,BK=128]` fp32 tile is REG:205/0-spill under Triton but REG:255+STACK:48 (spills) as hand-CUDA, the hand kernel built on it losing c16 700 vs 794 tok/s (`gdn_packed_decode_triton.h:9-14`, `cuda_gdn.cu:2710`) — a ReplaySSM kernel holds strictly more live state. Motivation is the OPEN Qwen3.8-27B bf16 decode gap (c4 total **0.918x**, output **0.963x**, [BENCHMARKS](../docs/BENCHMARKS.md) `:192-205`). Payoff is UNMEASURED; the spec §6 names the A/B (same binary, default-OFF lever, the SHIPPED cubin path as denominator, interleaved arms, clocks pinned under the GPU lock, `failed == 0` asserted, capacity as an axis) | [gdn-replayssm](specs/gdn-replayssm.md) | `READY` | - | | `KERNEL-SSM-MAMBA` | General Mamba selective scan, causal convolution, SSD, and linear attention kernels | selective scan source `CMakeLists.txt:401`, `csrc/libtorch_stable/mamba/selective_scan_fwd.cu:394-820`; upstream mamba tests; decode conv-update FLA Triton `causal_conv1d.py:15-192` @ `e24d1b24` | only Qwen GDN-specific convolution/state path, including indexed cache gather/scatter [cuda_gdn.cu:156](../src/vt/cuda/cuda_gdn.cu#L156) and GDN dispatch [cuda_gdn.cu:3695](../src/vt/cuda/cuda_gdn.cu#L3695); **decode conv-update decode-fast** [CausalConv1dUpdateFastKernel + TryLaunchConvUpdateFast](../src/vt/cuda/cuda_gdn.cu) + [flag header](../src/vt/cuda/conv_update_fast.h); **prefill conv-fwd register-window** [CausalConv1dFwdRegKernel + LaunchConvFwdReg](../src/vt/cuda/cuda_gdn.cu) + **fused post-conv per-V-head split** [GdnPostConvSplitKernel](../src/vt/cuda/cuda_gdn.cu) + **fused post-conv fast (128-thread + 128-bit V copy)** [GdnPostConvFastKernel + GdnVecCopy8](../src/vt/cuda/cuda_gdn.cu) + [prefill flag header](../src/vt/cuda/gdn_prefill_conv.h) + [CPU flag test](../tests/vt/test_gdn_prefill_conv.cpp) | Indexed BF16/F32 cache-boundary + i8-mask CUDA tests [test_ops_gdn.cpp:994](../tests/vt/test_ops_gdn.cpp#L994), focused 7/7 memcheck and W1 structural trace pass; these GDN tests do not prove generic Mamba support. **2026-07-18 (`CLAIM-CONV-UPDATE-FAST-1`): decode conv-update kernel-efficiency sub-lever landed, DEFAULT ON** (the c16-trace scan lever #5; row stays `INVENTORIED` — broad Mamba support unchanged). `CausalConv1dUpdateFastKernel` behind `VT_CONV_UPDATE_FAST` (default ON / `=0` rollback) is BIT-IDENTICAL (0-ulp) to the shipped `CausalConv1dUpdateKernel`: same float op order, only a 2D grid (removes two int64 div/mod per thread) + a WIDTH-templated register-cached state row reused for the conv accumulation AND the roll (mirrors FLA's per-`KERNEL_WIDTH` `col0..col3` register specialization). DGX-PROVEN (flock, clean `-Werror`, CUTLASS+FA2 hard-verified): `test_ops_gdn` conv-update decode-fast `fast==shipped` **BYTE-EXACT (0-ulp)** on `out`+`conv_state` [330/330](../tests/vt/test_ops_gdn.cpp) (k∈{3,4,5}, bf16+f32 state, ±bias, silu/identity, compact+scattered cache incl. NULL-block); full GDN 51/51; [CPU flag test](../tests/vt/test_conv_update_fast.cpp) 10/10. Isolated nsys pure-kernel 27B c16 shape (batch=16 conv_dim=10240 k=4 bf16): shipped 7,072 ns vs fast 3,680 ns median = **1.92×** (clears ≥1.3× bar). Full default set (all fast kernels ON) 27B 235/235 + 35B 315/315; `=0` rollback arms 235/235 + 315/315. `benchmark_binding=false`; binding grid re-measures. **2026-07-18 (`CLAIM-GDN-PREFILL-CONV-1`): PREFILL conv-fwd + fused post-conv kernel-efficiency sub-lever landed** (row stays `INVENTORIED`). `CausalConv1dFwdRegKernel` (`VT_CONV_REG` DEFAULT ON / `=0`→tiled) is the register-resident sliding-window mirror of vLLM FLA `_causal_conv1d_fwd_kernel` (`causal_conv1d.py:397-452`): per-channel weights preloaded to registers, `(k-1)`-tap register sliding window (each x loaded once, coalesced), token-axis chunked over grid.z for `n<=4` low-batch prefill. `GdnPostConvSplitKernel` (`VT_GDN_POSTCONV_SPLIT` OPT-IN) mirrors `_fused_post_conv_kernel` grid `(cdiv(L,BLOCK_T), H+HV)` (per-V-head). BIT-IDENTICAL (0-ulp) to the shipped tiled/megablock kernels — same tap-order f32 sum + q/k L2-norm byte-for-byte. DGX-PROVEN: byte-exact reg==tiled + split==megablock 268 GPU assertions + full GDN 3081/3081; memcheck 0; 27B **235/235** + 35B **315/315** (final defaults reg ON/split OFF). nsys 35B (`~/work/prefill-attr-conv-35b`): conv c1 −4.7% / c6 −7.3% (bandwidth-bound) ⇒ DEFAULT ON; post-conv split c1 −3.8% / c6 +4.7% (near-neutral — GdnPostConv is q/k-L2-norm-dominated) ⇒ OPT-IN. TTFT c1 −0.39% within run-noise (conv ~2.5% of GPU). Residual vLLM conv gap = bf16 traffic (VT_GDN_IN_BF16, task #40 sibling), not structure. `benchmark_binding=false`. **2026-07-19 (`CLAIM-GDN-POSTCONV-FAST-1`): fused post-conv FAST sub-lever landed, DEFAULT ON** (row stays `INVENTORIED`). A fresh production-path nsys (`--cuda-graph-trace=node`, VLLM_CPP_TRITON=ON, input-1024) confirmed the GDN chunk **compute** runs the vendored FLA Triton AOT cubins by default (`chunk_gated_delta_rule_fwd_kernel_h_blockdim64` / `chunk_fwd_kernel_o` / `recompute_w_u_fwd_kernel` / `chunk_scaled_dot_kkt_fwd_kernel` = at FLA parity), so the #1 remaining NON-AOT GDN kernel on BOTH models is the fused post-conv prep. The split (above) measured neutral/slower, so `GdnPostConvFastKernel` (`VT_GDN_POSTCONV_FAST` DEFAULT ON / `=0`→megablock) keeps the megablock grid `(T,Hk+1)` but makes two BYTE-IDENTICAL changes for Dk==Dv==128: 128 threads/block (the 128-wide L2-norm tree is the 256-wide tree minus a leading `+0` step; better reduction occupancy) + a 128-bit-staged V copy (`GdnVecCopy8`: raw `int4` when dtypes match, else the SAME `__bfloat162float`/`__float2bfloat16` converts). No arithmetic reordered; grounded in FLA `_fused_post_conv_kernel` (`fused_gdn_prefill_post_conv.py:57-149`). Isolated nsys per-call: 27B **175.6→133.0 ms (−24.3%)**, 35B **93.5→70.3 ms (−24.8%)**; in-situ TTFT (input-1024, 3 reps): 27B c1 −1.14% / c2 −1.31%, 35B c1 −0.72% / c2 −0.99% (all reps positive). BIT-EXACT: 27B 235/235 + 35B 315/315 on default(fast)+`=0`(megablock) arms; CPU flag test 28/28; clean `-Werror`. Evidence `dgx:~/work/vllm.cpp-gdn-chunk/prof`. `benchmark_binding=false`. **2026-08-12 (#496): the SSD half of this row's title is named and specced.** Everything above is the GDN arm; the generic Mamba2 selective scan has no local symbol at all — `include/vt/ops.h` has `kCausalConv1dFwd/Update/SpecUpdate`, `kRmsNormGated`, `kGdnPrefill/Decode`, `kKdaGatedDeltaRule` and **no** chunked state-space scan and no selective state update. Gated delta rule is not Mamba2: GDN carries the `(I − β kkᵀ)` removal term and a per-head scalar decay, Mamba2's SSD is a diagonally-decayed gated linear recurrence driven by `A_log`/`dt`/`D` with `B`/`C` shared across `n_groups`. Owed: `mamba_chunk_scan_combined_varlen` (`mamba/ops/ssd_combined.py:157-235` @ `555967922`), `selective_state_update` (`mamba/ops/mamba_ssm.py:497`), and the silu-gated GROUP RMS norm (`mamba_mixer2.py:100-149`, a sibling of our sigmoid `kRmsNormGated`, not a parameter of it). Blocks 10 `INVENTORIED` model rows; immediate driver is `NemotronHForCausalLM` (Nemotron-3.5-Lightning-30B-A3B-NVFP4: 23 mamba / 6 attention / 23 MoE, `n_groups=8`, `ssm_state_size=128`, `chunk_size=128`). **2026-08-12 (#496): W1 LANDED** at `47960a009` — CPU host references `vt::Mamba2ChunkScan` (varlen chunked SSD prefill), `vt::Mamba2StateUpdate` (single-token selective update) and `vt::RmsNormGatedGroup` (silu-gated GROUP RMS norm), gated against a sequential per-token recurrence transcribed independently in `double` from upstream's OWN non-chunked CPU kernel (`csrc/cpu/mamba_kernels.hpp`), plus chunk-size invariance over {8,16,32,64,128}. Two fresh reviews: round 1 FAIL (the gated-norm `weight` was read through `Ptr()` while the validator accepted bf16/f16 — a 2x heap over-read on the dtype upstream actually passes, since `Mixer2RMSNormGated.weight` is `torch.ones(...)` at model dtype; plus four unpinned guarantees and a refuted "equivalent mutant"), round 2 PASS (13 mutations, F1 swept repo-wide for a second instance — none). Row stays `INVENTORIED`: a host reference is not generic Mamba support. **Owed:** the CUDA arm (W2), the `MambaSpec` producer for Mamba2 shapes (W3), and one decode refusal `SUBCASE` (spec §8.2). No performance claim is made | [Mamba2 SSD spec](specs/mamba2-ssd.md); [prefill conv spec](specs/gdn-prefill-conv-reg-2026-07-18.md); [conv-update decode-fast spec](specs/conv-update-decode-fast-2026-07-18.md); [inventory](specs/kernel-family-inventory.md) | `INVENTORIED` | - | diff --git a/.agents/specs/gdn-moe-bf16-out.md b/.agents/specs/gdn-moe-bf16-out.md index f1c4ca57e..044f54a08 100644 --- a/.agents/specs/gdn-moe-bf16-out.md +++ b/.agents/specs/gdn-moe-bf16-out.md @@ -14,21 +14,81 @@ Secondary oracle: `sglang` @ `f63458b5be`, see [`.agents/oracles/sglang.md`](../ ## Now -Spec only. No product code is in this change, and nothing is measured. The next -step belongs to a fresh implementer: write the red tests in `## Tests`, make the -two coupled edits in `## Design`, and hand the result to a fresh reviewer. The -speed and correctness gates in `## Gates` need the GPU host and belong to the -operator. - -The row is `SPEC` until a fresh implementer claims the implementation. It cannot -reach `DONE` on the 2.4T arm at all; see `## Scope`. - -Read `## Evidence` before starting. Four unmerged local branches touch activation -dtypes on this model and none of them covers this row, but two of them collide -with it: `row/REFACTOR-DTYPE-CONSISTENCY` relocates `GdnOutDType` verbatim into -`activation_dtype.h`, and `row/PERF-GDN-BF16-CHAIN` uses the MoE f32 default as -an eligibility bound that Edit 1 removes. Check which of them has landed before -you edit. +Both edits in `## Design` are implemented, on `row/GDN-MOE-BF16-OUT-IMPL` off +`fd64c76ee`. `GdnOutDType()` takes no model-shape argument and resolves bf16 from +`VT_GDN_OUT_BF16` alone, the three call sites drop `cfg.num_experts == 0`, and +`detail::ShouldUsePackedGdnDecode` no longer carries `e.dense_model`. The env +decision is extracted into the pure `detail::GdnOutBf16FlagIsOn(const char*)`, +the second shape `## Tests` offers, because the resolver caches its `getenv` and +a process can only observe one value of it. + +**Nothing here is measured, and no GPU gate has run.** The row cannot reach +`DONE` on this evidence. What ran and what is owed: + +| Gate | Result | +|---|---| +| `test_qwen35_paged_forward` (CPU, the MoE reachability vehicle) | 5/5 cases, 13/13 assertions | +| `test_qwen27_paged_forward` (CPU, the two predicates) | 31/31 cases, 770/770 assertions | +| `test_ops_gdn`, `test_qwen27_dense_forward`, `test_qwen3_5_gdn_spec_routing`, `test_model_registry`, `test_runner`, the three expert-stream suites | all green on the CPU tier | +| `scripts/agent-preflight.sh` | `PASS`, exit 0 | +| full `ctest` | **not run.** The implementer's host had 6 GB free, and a full test build does not fit. `PENDING`, named resource | +| `test_qwen27_paged_engine` 235/235 | **not run.** Skipped with `assertions: 0` — no 27B checkpoint and no GPU on the implementer's host. `PENDING` | +| `test_qwen36_paged_engine` 315/315 | **not run.** `PENDING`, same reason. This is the row's ONE measurable checkpoint | +| correctness on `nvidia/Qwen3.6-35B-A3B-NVFP4` | `PENDING`, GPU host | +| `VT_GDN_OUT_BF16=0|1` same-binary speed A/B, per leaf | `PENDING`, GPU host | +| `nsys --cuda-graph-trace=node` memory-format confirmation | `PENDING`, GPU host | +| bf16 into `vt::RmsNormGatedQuantFp8` (the 35B's fp8 `out_proj` gated norm) | **not run.** Unreachable on the CPU tier: no fp8 `out_proj` on the synthetic model and no fp8 platform. `PENDING`, GPU host | +| GGUF MoE | `PENDING`. See `## Stop conditions`: if it cannot be gated, do NOT reintroduce a shape term | +| 2.4T | `PENDING`, named resource, by construction | + +The GPU gates belong to the operator, and the correctness gate comes first. A +fresh reviewer sees the immutable head next; the mutation `## Design` names is to +restore the `dense_model` form of `GdnOutDType` in a scratch copy, which must +turn `test_qwen35_paged_forward`'s bf16 case red. + +**Fresh review 1 -> repair.** The reviewer confirmed both edits correct and bound +by tests, and found the LEVER unpinned in one direction and a false-red generator +in the other. `GdnOutDType()` mutated to `return DType::kBF16;` — severed from its +parser and from `VT_GDN_OUT_BF16` entirely — left `test_qwen35_paged_forward` 5/5 +13/13 and `test_qwen27_paged_forward` 31/31 770/770 GREEN, while +`VT_GDN_OUT_BF16=0` on the unmutated binary FAILED 4/5 cases 11/13 assertions +because the new case asserted bf16 unconditionally and the resolver's cached +`getenv` cannot be neutralised in-process. The rollback that the A/B above is +built on was therefore both ungated and a red an operator would read as a +regression. Repaired by giving the resolver a `detail::` declaration, asserting +against the environment as the test file reads it directly, and registering the +same binary a second time under `VT_GDN_OUT_BF16=0` — the shape +`tests/CMakeLists.txt` already uses for a read-once lever. Both arms are now +6/6 cases 15/15 assertions, and the severing mutation is RED in the `=0` arm +(3 assertions across 2 cases). + +**The row now has a lifecycle surface.** `GDN-MOE-BF16-OUT` existed only in +`.agents/issue-index.md` and in this file, so there was nothing to move when its +state changed. It is recorded in [`.agents/kernel-matrix.md`](../kernel-matrix.md) +beside `KERNEL-GDN-AOT-BF16`, at `GATING`: implemented, CPU tier green, every GPU +gate `PENDING`. That gap came from the spec pull request rather than from the +implementation. `KERNEL-GDN-AOT-BF16`'s own "Every 35B path stays f32" is +corrected in the same edit, because Edit 1 is what makes it false — the same +obligation as `## A record consequence outside this row`. `docs/STATUS.md` and +`docs/BENCHMARKS.md` are written with it, because a new lifecycle state is a +claim about the project and `scripts/check-doc-checkpoint.py` says so at the +gate. Both entries say the same thing this section does: **no number, none +claimed, every measurable axis owed to a GPU host.** An entry that recorded a +capability without recording that nothing is measured would be the more +misleading half. + +**The fp8-fused gated norm is OWED, not tested.** See `## Risks`: Test 3 cannot +reach `vt::RmsNormGatedQuantFp8`, so bf16 has never entered that op on a real +checkpoint. It belongs to the 35B GPU gate together with the rows above. + +**Collision check, run against `origin/main` at `fd64c76ee` before the edits:** +neither collides yet. `row/REFACTOR-DTYPE-CONSISTENCY` has not landed — +`include/vllm/model_executor/models/activation_dtype.h` does not exist on `main`, +and the branch is still local. `row/PERF-GDN-BF16-CHAIN` has not landed either; +the only `main` commit whose message names it is `918d4546e` +(`PERF-FP8-ALPHA-FOLD`), which merely says it composes with that branch. Both +still have to be reconciled onto this row's result when they land, which is what +`## Owed` records. ## Scope @@ -273,6 +333,13 @@ RED first, in this order. Capture each red result before the change. Assert: the default resolves bf16 with no model-shape input; `VT_GDN_OUT_BF16=0` resolves f32; `=1` resolves bf16. RED today for the MoE case. Mutation proof: restoring the `dense_model` branch must turn it red. + **Fresh review 1 added the missing half.** Pinning the extracted parser does + not pin that the RESOLVER reads it, and the two are separable in silence: see + `## Now`. The resolver therefore gets a `detail::` declaration and its own + case, and because it caches its `getenv` the two env values are two ctest + registrations of the same binary rather than two assertions in one process — + the `_glue_fuse_off` / `test_dense_gateup_fused_marlin_off_*` shape. Neither + case computes its expectation through the parser under test. 2. **CPU tier, the eligibility predicate.** In `tests/vllm/models/test_qwen27_paged_forward.cpp`, a `GdnPackedDecodeEligibility` with every remaining term true selects packed decode without any model-shape @@ -340,9 +407,16 @@ than as an omission. the merged arm slower while the split arm gets faster. It is a measurement, not a blocker, and it is why the A/B reports the two leaves separately. - **First bf16 exposure of the fp8-fused gated norm.** The 35B's `out_proj` is - W8A8 fp8, so it takes `vt::RmsNormGatedQuantFp8` (`:4768`); the 27B's is fp4 and - never did (`:4747-4748`). The op accepts bf16 and the CUDA kernel dispatches it, - but no gate has run that combination on a real checkpoint. Test 3 must reach it. + W8A8 fp8, so it takes `vt::RmsNormGatedQuantFp8` (`:4873`); the 27B's is fp4 and + never did (`:4896-4897`). The op accepts bf16 and the CUDA kernel dispatches it, + but no gate has run that combination on a real checkpoint. Test 3 was assigned + to reach it and **cannot**: the call is guarded by FOUR terms, + `!w.out_proj_fp8.Empty() && GdnOutFp8FuseEnabled() && GlueFuseEnabled() && + supports_fp8()` (`:4854-4855`), the synthetic CPU model carries no + fp8 `out_proj`, and the CPU platform does not report fp8. So Test 3 exercises + the plain `vt::RmsNormGated` arm and this risk is UNMET on the CPU tier by + construction, not by omission. It is owed to the 35B GPU gate and recorded + under `## Now` and `## Owed`; do not read the sentence above as satisfied. - **GGUF MoE is unmeasured.** The comment being replaced named GGUF explicitly as an arm to leave alone. If the GGUF MoE arm cannot be gated in this row, see `## Stop conditions`. @@ -451,6 +525,15 @@ mistake the silence for absence, and contradicts none of it. AOT arms are pinned to 48 or 32 linear V-heads, so `Qwen3.8-2.4T-A95B` (128) runs the hand CUDA kernels the vendoring exists to replace. Not measurable on this hardware. +- **The first bf16 pass through `vt::RmsNormGatedQuantFp8`** (guard + `qwen3_5.cpp:4854-4855`, call `:4873`; reached only when + `!w.out_proj_fp8.Empty() && GdnOutFp8FuseEnabled() && GlueFuseEnabled() && + supports_fp8()`). `## Risks` names + it and `## Tests` assigned it to Test 3, which cannot reach it on the CPU tier + for the two structural reasons recorded there. The op accepts bf16 and the CUDA + kernel dispatches it, so this is an unexercised combination rather than a known + defect — but nothing here has run it, and the 35B correctness gate is where it + first will. Owed to that gate; do not treat the Risks entry as discharged. - `row/PERF-GDN-BF16-CHAIN` (local, unmerged) — its lever's third term is `outdt == BF16` and it uses the MoE f32 default as an eligibility bound. Edit 1 removes that bound. The branch's eligibility rule and Risks section need diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 79ee28b93..065d8ee98 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -469,6 +469,7 @@ built on it rather than keeping the flattering one. | 35B low-batch MoE decode | CLOSED at low batch (c1 0.975x, c4 wins); c16 0.93x. `VT_ASYNC_DEVICE_MIRROR` **default ON for correctness**. `VT_ASYNC_EXECUTOR` Option A (H2D out of capture) A/B'd speed-NEUTRAL | c16 lever is prefill glue (task #61), not the decode drain. `test_qwen36_async_serving` GREEN | | CPU keep-quant MoE decode | **No number owed**: correctness-only P0. The grouped keep-quant GEMM read activations as f32 whatever their dtype, so CPU MoE decode emitted token-0 garbage from `b4f5610a` (2026-07-31) | Speed unmeasured and unclaimed; `test_ops_quant_dot` GREEN (150224 assertions) | | Accepted-and-inert serve args (`SERVE-RECIPE-ARGS`, #606) | **No number owed**: argument parsing only, so nothing to time and no oracle leg. Correctness gate 4 cases / 58 asserts GREEN, RED-first, mutation-proven | None. A speed axis would be fabricated; closes on review plus the operator gate rerun | +| GDN core/z at the model dtype, MoE arms (`GDN-MOE-BF16-OUT`, [#1168](https://github.com/mudler/vllm.cpp/issues/1168)) | **No number, none claimed.** `GATING`, CPU tier only. A token gate cannot see this axis: f32 is the MORE precise deviation ([spec](../.agents/specs/gdn-moe-bf16-out.md)) | 35B correctness first, then 315/315, 235/235 dense inertness, the `VT_GDN_OUT_BF16` 0-versus-1 A/B per leaf, and `nsys` for the memory format. All `PENDING` a GPU host | | DeepSeek-V2-Lite MLA | Attributed miss, `ACTIVE` | Throughput at every concurrency | | Qwen3.5 text-only arms (#490) | **No number; run gates OWED**, both `PARTIAL`. The loader half is CLOSED (#740, #864 `DONE`), so what blocks these is hardware, not a refusal | No fitting ckpt for either causal-LM arm: no denominator. `Qwen3.8-2.4T-A95B` is ~4.8 TB vs 128 GB; its load plan resolves, which is not a token | | Darwin Qwen3.5 build repair (#1054, 2026-08-16) | **NOT APPLICABLE.** Removing a redundant namespace-scope lambda capture that Apple Clang rejects under `-Werror` changes no generated refusal text, no model math and no runtime path | None. The binding gate is the Apple Clang build | diff --git a/docs/STATUS.md b/docs/STATUS.md index a8bc9ec89..050469b2d 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -1573,6 +1573,25 @@ regression. ## Performance detail +**GDN recurrence output and z gate at the model dtype on every arm +(`GDN-MOE-BF16-OUT`, `GATING`, +[#1168](https://github.com/mudler/vllm.cpp/issues/1168)):** `GdnOutDType` +resolved bf16 from a dense checkpoint and f32 from a MoE one, so every MoE +checkpoint held the recurrence output `dcore`, the `z` gate and the gated-RMSNorm +weight at double width. vLLM branches on no model shape here. The shape term is +gone from the resolver and from packed-decode eligibility, and +`VT_GDN_OUT_BF16=0` is the f32 rollback for both arms now, not the dense one +alone ([spec](../.agents/specs/gdn-moe-bf16-out.md)). + +**Nothing is measured and no GPU gate has run**, so no axis is claimed in either +direction. The CPU tier enters through `ModelRegistry::Forward` on a MoE config +and reads the dtypes off the tensors, in the default and the `=0` arm alike. The +35B correctness gate, the `315/315` and `235/235` engine counts, the same-binary +A/B and the `nsys` memory-format confirmation are owed to a GPU host. Dropping +the shape term reaches packed GDN decode on no MoE checkpoint either: +`in_proj_ba` has one writer, the dense loader +([#1169](https://github.com/mudler/vllm.cpp/issues/1169), owed). + **Local Qwen3.5-4B plain BF16 direct loader; throughput ahead, acceptance `PENDING`, latency and VRAM open:** production `AsyncLLM` uses default-ON exact `(sequence, 8-token chunk)` causal-conv dispatch. Graph-node `nsys` measures diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 09155c350..0fc55e79c 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -74,9 +74,30 @@ void ResetQwen3_5MixedSpecInvocations() { g_mixed_spec_invocations.store(0, std::memory_order_relaxed); } +// GDN-MOE-BF16-OUT (#1168) Edit 2 dropped the `e.dense_model` term. It entered at +// f344decf4 ("dispatch exact packed decode") as one of that change's "real-model +// safety gates", was never revisited, and neither reference has an equivalent: +// VLLM_ENABLE_FLA_PACKED_RECURRENT_DECODE defaults True with no shape term +// (vllm/envs.py:124 @ 5559679) and SGLang keys supports_packed_decode on the +// platform alone (gdn_triton.py:43 @ f63458b5be). It became REDUNDANT once +// GdnOutDType stopped branching on model shape: `core_out` is `outdt`, and +// GdnPackedDecodeDTypesCompatible below already pins it to BF16, so an f32 +// recurrence output is what deselects packed decode on either arm. Removing it +// BEFORE the dtype change would have removed a term that the dtype rule did not +// yet subsume, which is why the two edits are one change and in this order. +// +// Do not read that as "and now the removal is observable in production", because +// it is not, in either order (fresh-review finding). `has_packed_ba` needs +// `in_proj_ba`, written at exactly one site in the tree — the dense loader, +// qwen3_5_dense_weights.cpp:431 — so on a MoE checkpoint the eligibility is +// false before the shape term is ever read. Removing it therefore reaches packed +// decode on NO checkpoint; it removes a contradiction with both references and a +// second answer to a question the dtype rule already answers. Reaching packed +// decode on a MoE arm needs the merged `in_proj_ba` owner in the MoE loader, +// which is #1169, and it is owed. bool detail::ShouldUsePackedGdnDecode( const GdnPackedDecodeEligibility& e) { - return e.runtime_enabled && e.cuda && e.dense_model && e.has_packed_ba && + return e.runtime_enabled && e.cuda && e.has_packed_ba && e.merged_ba_enabled && e.dtype_compatible && e.has_state_indices && e.num_prefills == 0 && e.num_prefill_tokens == 0 && e.num_spec_decodes == 0 && e.num_spec_decode_tokens == 0 && @@ -101,7 +122,12 @@ vt::DType detail::GdnProjectedMixedQkvDType(const GdnMixedQkvDTypeInputs& in) { // in the ONE place both the producer and the predictor read. See the header for // why each term is required; the short version is that the toggle is the opt-in, // `indt` keeps VT_GDN_IN_BF16's rollback honest on this arm too, and `outdt` -// confines the narrowing to the dense 27B. +// keeps the chain dtype-uniform. `outdt` used to be described as what "confines +// the narrowing to the dense 27B"; GDN-MOE-BF16-OUT (#1168) removed the +// model-shape argument from `GdnOutDType`, so `outdt` is BF16 on BOTH arms at +// the default and confines nothing. The DEFAULT-OFF `VT_GDN_FP8_IN_BF16` toggle +// (`GdnFp8InBf16Enabled`, which requires a leading '1') is now the only term +// keeping this inert on the 35B. vt::DType detail::GdnFp8MergedMixedQkvDType(bool fp8_in_bf16_enabled, vt::DType in_dtype, vt::DType out_dtype) { @@ -127,6 +153,28 @@ bool detail::PackedGdnDecodeFp8TowerFlagIsOn(const char* env_value) { return env_value != nullptr && env_value[0] == '1'; } +// GDN-MOE-BF16-OUT (#1168) — VT_GDN_OUT_BF16, default ON, parsed here so the CPU +// tier can pin the truth table that GdnOutDType() caches. There is no model-shape +// term: the environment is the whole decision, on the dense and the MoE arms +// alike, exactly as upstream resolves one model dtype for every layer. +bool detail::GdnOutBf16FlagIsOn(const char* env_value) { + return env_value == nullptr || env_value[0] != '0'; +} + +// ...and the RESOLVER that consumes it, here rather than in the anonymous +// namespace below so the CPU tier can call the thing the model calls. Pinning +// the parser alone does not pin that anything reads it: a `GdnOutDType()` +// hardwired to BF16 keeps every default-environment gate green, and the +// documented `VT_GDN_OUT_BF16=0` rollback then silently stops rolling back. See +// the header for why that matters more here than coverage — the variable is the +// denominator of this row's same-binary A/B. The full derivation of WHAT this +// dtype is stays at the `using` declaration below, next to the call sites. +vt::DType detail::GdnOutDType() { + static const bool bf16 = + detail::GdnOutBf16FlagIsOn(std::getenv("VT_GDN_OUT_BF16")); + return bf16 ? vt::DType::kBF16 : vt::DType::kF32; +} + bool detail::ShouldUseMergedGdnQkvz(const GdnMergedQkvzEligibility& e) { return e.runtime_enabled && e.cuda && e.has_packed_qkvz && e.uniform_dtype; } @@ -182,6 +230,40 @@ void detail::DisableGdnFp8InProjDebugStats() { g_gdn_fp8_inproj_debug_enabled.store(false, std::memory_order_release); } +namespace { +// GDN-MOE-BF16-OUT (#1168). What the last NON-MIXED-SPEC paged GDN layer actually +// allocated and projected, recorded off the tensors themselves rather than off +// the predicate. `GdnBlockPagedMixedSpec` is a paged GDN layer too and records +// nothing, the stores are unconditional (the fp8 sibling above is default-off), +// and they happen at graph CAPTURE and not at replay. The header states all +// three; none of them is what the shape of this code suggests. +std::atomic g_gdn_out_dtypes_observed{false}; +std::atomic g_gdn_out_core_dtype{static_cast(vt::DType::kF32)}; +std::atomic g_gdn_out_z_dtype{static_cast(vt::DType::kF32)}; + +void RecordGdnOutActivationDTypes(vt::DType core_out, vt::DType z_gate) { + g_gdn_out_core_dtype.store(static_cast(core_out), std::memory_order_relaxed); + g_gdn_out_z_dtype.store(static_cast(z_gate), std::memory_order_relaxed); + g_gdn_out_dtypes_observed.store(true, std::memory_order_release); +} +} // namespace + +void detail::ResetGdnOutActivationDTypes() { + g_gdn_out_dtypes_observed.store(false, std::memory_order_release); + g_gdn_out_core_dtype.store(static_cast(vt::DType::kF32), std::memory_order_relaxed); + g_gdn_out_z_dtype.store(static_cast(vt::DType::kF32), std::memory_order_relaxed); +} + +detail::GdnOutActivationDTypes detail::LastGdnOutActivationDTypes() { + GdnOutActivationDTypes out; + out.observed = g_gdn_out_dtypes_observed.load(std::memory_order_acquire); + out.core_out = + static_cast(g_gdn_out_core_dtype.load(std::memory_order_relaxed)); + out.z_gate = + static_cast(g_gdn_out_z_dtype.load(std::memory_order_relaxed)); + return out; +} + bool detail::PackedGdnDecodeEnvSelected(const GdnPackedDecodeEnvConfig& env) { // Mirror PackedGdnDecodeRuntimeEnabled: enabled unless first char is '0'. const bool runtime_enabled = @@ -191,8 +273,8 @@ bool detail::PackedGdnDecodeEnvSelected(const GdnPackedDecodeEnvConfig& env) { !(env.merged_proj != nullptr && env.merged_proj[0] == '0') && (env.merged_ba == nullptr || env.merged_ba[0] != '0'); // Mirror the dtype_compatible expression on the real 27B dense gate: - // GdnInDType (default BF16), GdnOutDType dense default (BF16; override - // '0' -> F32), MergedGdnBaOutputDType(packed) (default BF16 under packed; + // GdnInDType (default BF16), GdnOutDType (default BF16 on every arm since + // #1168; override '0' -> F32), MergedGdnBaOutputDType(packed) (default BF16 under packed; // override '0' -> F32). The SSM cache dtype term is always a float dtype. const bool in_bf16 = env.in_bf16 == nullptr || env.in_bf16[0] != '0'; const bool out_bf16 = env.out_bf16 == nullptr || env.out_bf16[0] != '0'; @@ -3169,8 +3251,8 @@ bool GdnFp8InBf16Enabled() { return on; } -// GDN recurrence-OUTPUT + z-gate in bf16 (27B default ON; 35B keeps its former -// f32 default; VT_GDN_OUT_BF16=0/1 overrides both for diagnostics). +// GDN recurrence-OUTPUT + z-gate in bf16, on EVERY arm (default ON; +// VT_GDN_OUT_BF16=0 is the same-binary f32 rollback). // vLLM keeps core_attn_out and the z gate bf16 (the gated-RMSNorm consumes them): // FLA chunk_o.py stores o bf16, and Qwen3NextGatedRMSNorm reads bf16 core/gate, // upcasting to f32 only for the variance reduction (layernorm_guard.py). Our @@ -3185,18 +3267,26 @@ bool GdnFp8InBf16Enabled() { // this lever is the f32 `dcore` recurrence output that attempt left untouched. // This is correctness-significant for the 27B: with the repaired full NVFP4 // tactic stack, f32 core/z takes the alternate whitespace near-tie branch while -// bf16 reproduces native vLLM 16/16. Keep every unmeasured 35B arm, including -// GGUF, on its prior f32 default; the explicit env override remains available -// for its later independently gated campaign. -DType GdnOutDType(bool dense_model) { - static const int override = [] { - const char* e = std::getenv("VT_GDN_OUT_BF16"); - if (e == nullptr) return -1; - return e[0] == '0' ? 0 : 1; - }(); - const bool bf16 = override >= 0 ? override != 0 : dense_model; - return bf16 ? DType::kBF16 : DType::kF32; -} +// bf16 reproduces native vLLM 16/16. +// +// GDN-MOE-BF16-OUT (#1168) removed the `bool dense_model` parameter this used to +// default to. It resolved bf16 for a dense checkpoint and f32 for a MoE one, and +// all three call sites passed `cfg.num_experts == 0`, so every MoE checkpoint +// carried `dcore`, `z` and the gated-norm weight at double width. Upstream does +// not branch on model shape anywhere on this path — `Qwen3_5ForCausalLMBase` +// (vllm/model_executor/models/qwen3_5.py:280-297 @ 5559679) is the shared base of +// the dense and MoE causal-LM arms — and the deferral quoted above ("keep every +// unmeasured 35B arm on its prior f32 default") named its own successor campaign, +// which this is. The parameter is gone rather than defaulted because a signature +// that accepts a model shape makes the default unreadable at the definition and +// lets a new call site reintroduce the split silently. VT_GDN_OUT_BF16=0 is now +// the f32 rollback for BOTH arms rather than for the dense one alone. +// +// Fresh-review repair: the definition moved up beside `GdnOutBf16FlagIsOn` and +// into `detail::`, so that a gate can observe the RESOLVER and not only its +// parser. Nothing about the resolution changed. The call sites below are +// unqualified and keep reading it through this declaration. +using detail::GdnOutDType; // bf16 residual stream (default ON). vLLM runs the 35B in bf16 // (model_config.dtype=bfloat16): qwen3_next.py keeps `residual` as the bf16 hidden @@ -3530,8 +3620,11 @@ DBuf MergedFp8QkvzD(Dev d, const Tensor& x, const Tensor* h_fp8, // resident owner is shared by both arms: the fallback slices its output ROWS // (dim-0 raw-NK slices stay contiguous) and issues the two legacy GEMMs at // their independent dtypes, never retaining duplicate split weights. The -// merged arm requires one uniform output dtype (GdnInDType == GdnOutDType; the -// 27B default is BF16/BF16, matching vLLM's model-dtype projection). +// merged arm requires one uniform output dtype (GdnInDType == GdnOutDType, and +// since GDN-MOE-BF16-OUT (#1168) that is BF16/BF16 on every arm, matching vLLM's +// model-dtype projection). The uniformity term therefore no longer excludes a +// MoE checkpoint; `has_packed_qkvz` still does, because no MoE or GGUF loader +// builds the merged `in_proj_qkvz` owner either — the same gap as #1169's. // VT_GDN_MERGED_QKVZ=0 (or master VT_GDN_MERGED_PROJ=0) restores the split // GEMMs from the same binary and the same resident owner. bool MergedGdnQkvzEnabled(Dev d) { @@ -3596,11 +3689,17 @@ GdnQkvzOutput ProjectGdnQkvz(Dev d, const GdnLayerWeights& w, const Tensor& h, // narrowing is not value-neutral), // 2. indt == BF16, i.e. VT_GDN_IN_BF16 is on (its default) — this IS the // lever being unblocked, so honouring its rollback is mandatory, - // 3. outdt == BF16. This is what confines the change to the 27B: the 35B is - // MoE, so GdnOutDType(dense_model=false) is f32 by default and this stays - // inert there. It also keeps the whole chain dtype-uniform, which the - // downstream contracts need — vt::RmsNormGatedQuantFp8 requires + // 3. outdt == BF16, which keeps the whole chain dtype-uniform — the + // downstream contracts need that: vt::RmsNormGatedQuantFp8 requires // gate.dtype == x.dtype (ops.cpp), and `z` below is exactly that gate. + // This term USED to read "confines the change to the 27B, because the + // 35B is MoE and GdnOutDType(dense_model=false) is f32 there". #521 + // asked for that correction and GDN-MOE-BF16-OUT (#1168) is what makes + // it wrong: outdt is BF16 on every arm now. What still keeps this leaf + // inert on the 35B is condition 1, the DEFAULT-OFF VT_GDN_FP8_IN_BF16 + // toggle — a toggle term, not a model-shape one. Nothing moves by + // default in either merge order, but whoever turns that toggle on owns + // measuring the 35B too (#417). // PERF-GDN-PACKED-BRIDGE (#365): PERF-FP8-ALPHA-FOLD's three-term decision, // moved into the shared `GdnFp8MergedInProjDType` so the PRODUCER (this line, // which reaches MergedFp8QkvzD and allocates the buffer) and the PREDICTOR @@ -3747,7 +3846,7 @@ DBuf GdnBlock(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, // and on the merged fp8 branch too under VT_GDN_FP8_IN_BF16 (default OFF, see // GdnFp8InBf16Enabled). See GdnInDType(). const DType indt = GdnInDType(); - const DType outdt = GdnOutDType(cfg.num_experts == 0); + const DType outdt = GdnOutDType(); GdnQkvzOutput qkvz = ProjectGdnQkvz(d, w, h, conv_dim, value_dim, indt, outdt, h_fp8); Tensor mixed = qkvz.mixed; // [T,conv_dim], contiguous or row-strided view @@ -4153,7 +4252,7 @@ DBuf GdnBlockPagedMixedSpec(Dev d, const GdnLayerWeights& w, const HfConfig& cfg const int64_t conv_dim = 2 * key_dim + value_dim; const float eps = static_cast(cfg.rms_norm_eps); const DType convdt = mixed.dtype; - const DType outdt = GdnOutDType(cfg.num_experts == 0); + const DType outdt = GdnOutDType(); const DType actdt = GdnActDType(); const float scale = 1.0F / std::sqrt(SizeF(Dk)); g_mixed_spec_invocations.fetch_add(1, std::memory_order_relaxed); @@ -4384,7 +4483,7 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, const bool mixed_spec = spec && np > 0; const DType indt = GdnInDType(); - const DType outdt = GdnOutDType(cfg.num_experts == 0); + const DType outdt = GdnOutDType(); // PERF-27B-GDN-PACKED-REACHABLE (#365). `dtype_compatible` is decided by the // ACTIVATION dtypes vt::GdnPackedDecode requires, not by how the GDN weights // are stored. `mixed_qkv` has to be PREDICTED because this decision runs @@ -4407,7 +4506,6 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, detail::GdnPackedDecodeEligibility{ PackedGdnDecodeRuntimeEnabled(), vllm::platforms::GetPlatform(d.q.device.type).needs_weight_staging(), - cfg.num_experts == 0, !w.in_proj_ba.Empty(), MergedGdnBaEnabled(d), detail::GdnPackedDecodeDTypesCompatible( @@ -4586,6 +4684,12 @@ DBuf GdnBlockPaged(Dev d, const GdnLayerWeights& w, const HfConfig& cfg, // intermediates are allocated. This mirrors vLLM v0.25.0 // _forward_core_decode_non_spec:1644-1695. DBuf dcore(d, outdt, {T, Hv, Dv}); + // GDN-MOE-BF16-OUT (#1168): read off the tensors, not off GdnOutDType, so a + // gate entering through ModelRegistry::Forward observes what this layer RAN. + // This is the ONLY recording site. The mixed spec+non-spec batch returns into + // GdnBlockPagedMixedSpec above and never reaches it, so a mixed step leaves + // the record untouched rather than stale-free — see the header's limits. + RecordGdnOutActivationDTypes(dcore.t().dtype, z.dtype); const float scale = 1.0F / std::sqrt(SizeF(Dk)); if (packed_decode) { Tensor gidx = SubView(sdi.gdn_state_idx.t(), 0, nd); @@ -10353,8 +10457,11 @@ ForwardLogits Qwen3_5DenseDecodeGraph::Step( // Drain the capture if the forward throws — see the 35B driver above for why // a skipped EndCaptureGraph poisons the stream permanently (#339, F-B). This // is the 27B DENSE driver, and it is the one that matters most here: the - // bf16-D fp8 lever this row guards (VT_GDN_FP8_IN_BF16) is confined to the - // 27B by construction, since the 35B is MoE and its GdnOutDType is f32. + // bf16-D fp8 lever this row guards (VT_GDN_FP8_IN_BF16) is off by DEFAULT + // (GdnFp8InBf16Enabled), which is the only thing keeping it inert on the + // 35B. It used to read "confined to the 27B by construction, since the 35B + // is MoE and its GdnOutDType is f32"; GDN-MOE-BF16-OUT (#1168) made outdt + // BF16 on both arms, so the bound is a toggle now, not a model shape. DBuf lg = [&] { try { return DenseForwardLayers(d, s.hidden->t(), s.positions, s.attn_meta, diff --git a/src/vllm/model_executor/models/qwen3_5_internal.h b/src/vllm/model_executor/models/qwen3_5_internal.h index 86ee7b40e..1e0bde29d 100644 --- a/src/vllm/model_executor/models/qwen3_5_internal.h +++ b/src/vllm/model_executor/models/qwen3_5_internal.h @@ -34,8 +34,8 @@ struct GDNAttentionMetadata; // the real forward's embed target) through EITHER the merged single-GEMM arm // (`merged=true`) or the exact two legacy fp8 GEMMs, and return the // concatenation [mixed_qkv | z] as [T*(conv_dim+value_dim)] f32. `z` is -// produced at BF16 when `z_bf16` (the 27B dense default `GdnOutDType`) and -// upcast on the way out, so the two arms are directly byte-comparable. The +// produced at BF16 when `z_bf16` (the `GdnOutDType` default, on every arm since +// #1168) and upcast on the way out, so the two arms are directly byte-comparable. The // merged arm is byte-identical to the split arm by construction; this proves it // on the real fp8 GEMM. CUDA-only (the fp8 W8A8 path is). Defined in // qwen3_5.cpp. @@ -62,7 +62,6 @@ namespace vllm::detail { struct GdnPackedDecodeEligibility { bool runtime_enabled = false; bool cuda = false; - bool dense_model = false; bool has_packed_ba = false; bool merged_ba_enabled = false; bool dtype_compatible = false; @@ -84,8 +83,8 @@ bool ShouldUsePackedGdnDecode(const GdnPackedDecodeEligibility& eligibility); // wires into GdnPackedDecodeEligibility (qwen3_5.cpp GdnBlockPaged): // PackedGdnDecodeRuntimeEnabled (VT_GDN_PACKED_DECODE), MergedGdnBaEnabled // (master VT_GDN_MERGED_PROJ + leaf VT_GDN_MERGED_BA), and dtype_compatible — -// GdnInDType (VT_GDN_IN_BF16) == BF16, GdnOutDType dense default -// (VT_GDN_OUT_BF16) == BF16, MergedGdnBaOutputDType(packed) +// GdnInDType (VT_GDN_IN_BF16) == BF16, GdnOutDType (VT_GDN_OUT_BF16) == BF16 — +// its default on every arm since #1168 — MergedGdnBaOutputDType(packed) // (VT_GDN_BA_OUT_BF16) == BF16. Splitting BA (master OR leaf) or reverting any // coupled dtype therefore deselects packed and runs the decomposed recurrence: // the gate's dispatch-count contract must expect ZERO packed launches on those @@ -140,8 +139,10 @@ vt::DType GdnProjectedMixedQkvDType(const GdnMixedQkvDTypeInputs& in); // Terms are PERF-FP8-ALPHA-FOLD's own, verbatim: the opt-in toggle // (`VT_GDN_FP8_IN_BF16`, default OFF), `indt == BF16` so `VT_GDN_IN_BF16`'s // documented rollback still restores f32 on this arm too, and `outdt == BF16`, -// which confines the narrowing to the dense 27B (the 35B is MoE, so -// `GdnOutDType` is f32 there and this stays inert). +// which keeps the chain dtype-uniform. That third term used to be described as +// what "confines the narrowing to the dense 27B, since the 35B is MoE and +// `GdnOutDType` is f32 there"; #1168 made `outdt` BF16 on both arms, so the +// DEFAULT-OFF toggle is now the only thing keeping this inert on the 35B. vt::DType GdnFp8MergedMixedQkvDType(bool fp8_in_bf16_enabled, vt::DType in_dtype, vt::DType out_dtype); @@ -164,7 +165,7 @@ vt::DType GdnFp8MergedMixedQkvDType(bool fp8_in_bf16_enabled, vt::DType in_dtype struct GdnPackedDecodeDTypes { vt::DType mixed_qkv = vt::DType::kF32; // post-conv activation (== mixed.dtype) vt::DType ba_out = vt::DType::kF32; // MergedGdnBaOutputDType(packed) - vt::DType core_out = vt::DType::kF32; // dcore, GdnOutDType(dense) + vt::DType core_out = vt::DType::kF32; // dcore, GdnOutDType() vt::DType ssm_state = vt::DType::kF32; // the recurrent cache, independent }; @@ -185,12 +186,47 @@ bool GdnPackedDecodeDTypesCompatible(const GdnPackedDecodeDTypes& dtypes); // unaffected. bool PackedGdnDecodeFp8TowerFlagIsOn(const char* env_value); +// GDN-MOE-BF16-OUT (#1168). VT_GDN_OUT_BF16, DEFAULT ON — the whole decision +// behind `GdnOutDType()`, which resolves the dtype of the GDN recurrence output +// `dcore` [T,Hv,Dv], the `z` output gate, and the gated-RMSNorm weight that has +// to match them. Parsed here rather than at the getenv so the CPU tier can pin +// the truth table, the same way PackedGdnDecodeFp8TowerFlagIsOn is; the resolver +// caches its getenv in a function-local static, so a process observes exactly +// one value of it and the table cannot be asked of the resolver itself. +// +// Note the POLARITY differs from the flag above: this one is default ON and only +// a leading '0' turns it off, matching VT_GDN_IN_BF16 and the `out_bf16` term +// PackedGdnDecodeEnvSelected already mirrors. There is no model-shape argument. +// `GdnOutDType` used to take a `bool dense_model` and default to it, which made +// every MoE checkpoint f32 while vLLM keeps both tensors at the one model dtype +// it resolves for every layer (qwen_gdn_linear_attn.py:870-873, :843, :459-465 @ +// 5559679). +bool GdnOutBf16FlagIsOn(const char* env_value); + +// GDN-MOE-BF16-OUT (#1168), fresh-review repair. The RESOLVER the model actually +// calls, declared here so a gate can observe what it answers and not only what +// its parser answers. The two are separable, and severing them is silent: a +// `GdnOutDType()` hardwired to BF16 keeps every default-environment gate green +// while `VT_GDN_OUT_BF16=0` stops rolling anything back — and that variable is +// the denominator of this row's whole same-binary A/B, so a disconnected lever +// would invalidate the measurement that decides the row rather than merely lose +// coverage. +// +// It caches its getenv in a function-local static, so ONE process observes ONE +// value. The gate is therefore the shape tests/CMakeLists.txt already uses for a +// read-once lever (`_glue_fuse_off`, `test_dense_gateup_fused_marlin_off_*`): the +// SAME binary is registered a second time with `VT_GDN_OUT_BF16=0`, and the case +// asserts against the environment as it reads it directly. +vt::DType GdnOutDType(); + // W2 merged-qkvz dispatch. vLLM always issues one in_proj_qkvz GEMM // (qwen_gdn_linear_attn.py:923-936 @ 702f4814); locally the single GEMM is // selected only on CUDA with the packed 27B owner resident, the runtime // toggles on (VT_GDN_MERGED_PROJ master, VT_GDN_MERGED_QKVZ leaf) and one // uniform output dtype (mixed_qkv and z leave one GEMM, so GdnInDType must -// equal GdnOutDType — the 27B default is BF16/BF16). Every other combination +// equal GdnOutDType — BF16/BF16 by default on every arm since #1168, so this +// term no longer excludes a MoE checkpoint; `has_packed_qkvz` still does). +// Every other combination // issues the exact two split GEMMs sliced from the same resident owner. struct GdnMergedQkvzEligibility { bool runtime_enabled = false; @@ -270,6 +306,61 @@ void ResetGdnFp8InProjDebugStats(); GdnFp8InProjDebugStats GetGdnFp8InProjDebugStats(); void DisableGdnFp8InProjDebugStats(); +// GDN-MOE-BF16-OUT (#1168) — the dtypes the GDN block's two OUTPUT-side +// activations ACTUALLY carried on the last NON-MIXED-SPEC paged GDN layer this +// process ran. +// +// `GdnOutDType()` is a predicate, and a test that calls it proves only that the +// predicate answers. This records what `GdnBlockPaged` allocated and projected: +// `core_out` is `dcore`'s own `[T,Hv,Dv]` tensor dtype and `z_gate` is the dtype +// `ProjectGdnQkvz` actually emitted for the gate, both read off the tensors +// after they exist. So a gate can enter through `ModelRegistry::Forward` on a +// real MoE config and ask what the model ran, which is the question +// .agents/reachability.md says a unit test over the predicate cannot answer. +// +// `observed` is false until such a layer has run since the last reset. It is the +// term that separates "the forward produced f32" from "the forward never reached +// that path", which are the same silence to a dtype comparison. +// +// FOUR LIMITS, stated because the sentence above is narrower than it first +// reads and a reader who takes `observed` for a complete answer gets a wrong one +// (fresh-review finding). +// +// 1. It does NOT cover `GdnBlockPagedMixedSpec`. That is also a paged GDN +// layer, at the same `outdt`, and it records nothing: `GdnBlockPaged` +// returns into it before reaching the one recording site. On a step whose +// batch mixes speculative and non-speculative requests `observed` therefore +// stays false, or worse stays TRUE from an earlier step — which is the +// ambiguity `observed` exists to remove, reappearing one level up. Read it +// only from a gate that controls its own batch shape, as +// test_qwen35_paged_forward does. The recording is not extended to the mixed +// path because nothing gates that path's dtypes today, and a second call +// site no test enters is the shape .agents/reachability.md warns about. +// 2. It is UNCONDITIONAL, unlike `GdnFp8InProjDebugStats` above, whose shape it +// otherwise borrows. That sibling is default-off behind an enable flag and +// counts only after `ResetGdnFp8InProjDebugStats`. This one stores three +// atomics on every non-mixed-spec paged GDN layer of every forward of every +// model, whether or not anybody reads it. Relaxed stores off the critical +// path, so nothing observable changes — but do not read the resemblance as +// "off by default", because it is not. +// 3. It records at CUDA-graph CAPTURE, not at replay. Replay does no host +// dispatch, so under graphed decode the values are whatever the capturing +// step wrote. Same caveat as the fp8 sibling's "read on an eager step". +// 4. It does not cover the NON-PAGED `GdnBlock` either (qwen3_5.cpp), which +// allocates at the same `outdt` and records nothing. The word "paged" above +// already excludes it, but a list a reader takes for complete has to say so: +// a forward that runs the non-paged block leaves `observed` exactly as +// limit 1 does, and this list is complete only with both uncovered sites on +// it. +struct GdnOutActivationDTypes { + bool observed = false; + vt::DType core_out = vt::DType::kF32; // dcore [T,Hv,Dv] + vt::DType z_gate = vt::DType::kF32; // the output gate handed to the norm +}; + +void ResetGdnOutActivationDTypes(); +GdnOutActivationDTypes LastGdnOutActivationDTypes(); + // Validate the exact prefix that will be uploaded. Negative rows are inert // padding; every live slot must be unique and in range. This runs on host // metadata before the device buffer is constructed, keeping CUDA capture free diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index 5dcfb22e8..def698918 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -1876,8 +1876,9 @@ __global__ void RmsNormGatedRowKernel(Tout* out, const Tin* x, const Tin* gate, // --------------------------------------------------------------------------- // Fast gated-RMSNorm variant (VT_RMSNORM_GATED_FAST). BIT-IDENTICAL (0-ulp) to the // shipped RmsNormGatedRowKernel above for the d==128 gated-norm path — covering BOTH -// the 27B dense bf16 core/z/weight AND the 35B MoE f32 core/z/weight (GdnOutDType is -// f32 for the MoE), mirroring the RMSNorm decode-fast bit-safety technique (cuda_ops.cu +// the bf16 core/z/weight both arms carry by default AND the f32 core/z/weight the +// VT_GDN_OUT_BF16=0 rollback restores (before #1168 that f32 arm was the MoE 35B's +// DEFAULT), mirroring the RMSNorm decode-fast bit-safety technique (cuda_ops.cu // RmsNormRowFastKernel, 348d12d). Templated over (Tin,Tout) via the shared Load/Store // helpers so a single kernel is bit-exact for every dispatched dtype pair. // @@ -1916,10 +1917,11 @@ __global__ void RmsNormGatedRowKernel(Tout* out, const Tin* x, const Tin* gate, // // Scope: d==128 (both gate models' Dv), core/gate/weight one shared dtype in {f32,bf16}, // out in {f32,bf16}. Templated over (Tin,Tout) so ONE bit-identical kernel covers BOTH -// the 27B dense bf16 core/z/weight path AND the 35B MoE f32 core/z/weight path: -// GdnOutDType returns bf16 for the dense 27B but f32 for the MoE 35B, so the 35B's -// gated-norm inputs are f32 and the former bf16-only guard sent it to the slow shipped -// kernel. The kernel below uses the SAME Load/Store helpers as the shipped +// the default bf16 core/z/weight path AND the f32 one: GdnOutDType returns bf16 on +// every arm since #1168 and f32 under the VT_GDN_OUT_BF16=0 rollback, and before #1168 +// f32 was the MoE 35B's DEFAULT — either way f32 gated-norm inputs were sent to the +// slow shipped kernel by the former bf16-only guard. +// The kernel below uses the SAME Load/Store helpers as the shipped // RmsNormGatedRowKernel, so for EVERY (Tin,Tout) the float op sequence is // byte-for-byte identical to shipped (Load(float) is the identity; Load(bf16) is the // same __bfloat162float; Store the same rounding) — the bit-identity argument above @@ -1970,8 +1972,18 @@ inline void LaunchGatedFast(cudaStream_t s, Tensor& out, const Tensor& x, const // Runtime predicate + launch for the gated fast path. Returns true iff it ran. // Guard: d==128 (the GDN Dv used by both gate models); core/gate/weight share one dtype // (the caller VT_CHECKs gate.dtype==x.dtype==w.dtype) in {f32,bf16}; out in {f32,bf16}. -// This covers the 27B (bf16 core/z, bf16 out) AND the 35B (f32 core/z; bf16 out under -// GlueFuse, else f32 out). Out-of-scope shapes keep RmsNormGatedRowKernel. The launch +// This covers BOTH the bf16 core/z/weight both gate models carry by default since +// #1168 (GdnOutDType stopped branching on model shape, so the 35B is no longer the +// f32 arm) AND the f32 core/z/weight the VT_GDN_OUT_BF16=0 rollback restores on +// either of them; `out` is bf16 exactly when GlueFuseEnabled() and f32 +// UNCONDITIONALLY otherwise -- it does NOT follow the input. All three +// non-GlueFuse gated-norm call sites allocate `DBuf dgated(d, DType::kF32, ...)` +// whatever the input dtype (qwen3_5.cpp:3965, :4435, :4886), so this row's own +// VT_GDN_OUT_BF16=1 default plus the documented VT_GLUE_FUSE=0 rollback makes +// bf16-in/f32-out a REACHABLE pair. That is why all four (Tin,Tout) +// instantiations below are live: the dispatch reads out.dtype independently of +// x.dtype, and narrowing it to the diagonal would drop a shipped combination. +// Out-of-scope shapes keep RmsNormGatedRowKernel. The launch // uses kGatedFastBlock (=128) threads and reproduces shipped's reduction ORDER, so the // output is bit-identical for every dispatched (Tin,Tout). inline bool TryLaunchRmsNormGatedFast(cudaStream_t s, Tensor& out, const Tensor& x, diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index bad02c87d..b865d484f 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -404,6 +404,22 @@ vllm_cpp_add_test(test_linear_method vllm/model_executor/layers/test_linear_meth target_include_directories(test_linear_method PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_decode_graph_sizes vllm/models/test_decode_graph_sizes.cpp) vllm_cpp_add_test(test_qwen35_paged_forward vllm/models/test_qwen35_paged_forward.cpp) +# GDN-MOE-BF16-OUT (#1168): the MoE forward's GDN output-activation dtype probe +# lives in the not-installed qwen3_5_internal.h, beside the other detail:: seams. +target_include_directories(test_qwen35_paged_forward PRIVATE ${CMAKE_SOURCE_DIR}/src) +# GDN-MOE-BF16-OUT (#1168), fresh-review repair. `GdnOutDType()` caches its +# getenv in a function-local static, so one process observes exactly ONE value of +# VT_GDN_OUT_BF16 and a single run cannot tell a resolver that READS the lever +# from one that was hardwired to its default answer — the mutation +# `return DType::kBF16;` left this binary 5/5 13/13. Re-run the SAME binary with +# the lever forced to 0, the shape `_glue_fuse_off` and +# `test_dense_gateup_fused_marlin_off_*` above already use. This is also the +# documented f32 rollback and the denominator of the row's same-binary A/B, so +# without this entry the arm an operator is told to run is ungated. +add_test(NAME test_qwen35_paged_forward_gdn_out_f32 + COMMAND test_qwen35_paged_forward) +set_tests_properties(test_qwen35_paged_forward_gdn_out_f32 PROPERTIES + ENVIRONMENT "VT_GDN_OUT_BF16=0") # MoE vision tower (#891): the tower is loaded (and its ABSENCE refused by name) # and the forked VL forward actually CONSUMES the merger rows. vllm_cpp_add_test(test_qwen3_5_moe_vision vllm/models/test_qwen3_5_moe_vision.cpp) diff --git a/tests/vllm/models/test_qwen27_paged_forward.cpp b/tests/vllm/models/test_qwen27_paged_forward.cpp index c6317456b..6ea1854a9 100644 --- a/tests/vllm/models/test_qwen27_paged_forward.cpp +++ b/tests/vllm/models/test_qwen27_paged_forward.cpp @@ -367,14 +367,19 @@ GDNAttentionMetadata ChunkGdnMeta(int64_t qlen, int32_t sidx, bool has_initial) // Mirrors vLLM v0.25.0 // qwen_gdn_linear_attn.py::_forward_core:1286-1298. The packed branch is // selected only for enabled, pure non-spec decode; every prefill/mixed/spec -// shape remains on the standard recurrence. Local W1D2 additionally scopes the -// first model consumer to the real dense packed-BA path and requires compatible -// BF16 storage plus persistent device indices. +// shape remains on the standard recurrence. Locally it additionally requires the +// packed-BA owner, compatible BF16 activation dtypes, and persistent device +// indices. +// +// W1D2's `dense_model` term is gone (GDN-MOE-BF16-OUT, #1168). The `rejects` +// case that pinned it is replaced below by the argument for its removal rather +// than deleted: the term is redundant behind the dtype rule, and what still +// keeps a MoE checkpoint off this path is `has_packed_ba` (#1169), which the +// case below now names. TEST_CASE("qwen27 packed GDN selection is pure non-spec decode only") { vllm::detail::GdnPackedDecodeEligibility e; e.runtime_enabled = true; e.cuda = true; - e.dense_model = true; e.has_packed_ba = true; e.merged_ba_enabled = true; e.dtype_compatible = true; @@ -399,11 +404,16 @@ TEST_CASE("qwen27 packed GDN selection is pure non-spec decode only") { rejects(x); } { - auto x = e; - x.dense_model = false; - rejects(x); // 35B/MoE is deliberately inert in W1D2. - } - { + // W1D2 asserted here that `dense_model = false` deselects, which was the + // record of f344decf4's day-one staging gate. GDN-MOE-BF16-OUT (#1168) + // removed that term, so the assertion is replaced by the two terms that + // ACTUALLY keep a MoE checkpoint off this path, neither of which is a model + // shape. The merged BA owner below is built only by the dense loader + // (qwen3_5_dense_weights.cpp:436) — that is #1169, and it is owed. The + // second is the dtype rule, already pinned by the `dtype_compatible` case + // further down: `core_out` is `outdt` and GdnPackedDecodeDTypesCompatible + // pins it to BF16, so an f32 recurrence output deselects on EITHER arm, + // which is what made the shape term redundant rather than merely unmeasured. auto x = e; x.has_packed_ba = false; rejects(x); @@ -520,9 +530,17 @@ TEST_CASE("qwen27 packed GDN predicts the mixed_qkv dtype ProjectGdnQkvz emits") // // The three terms are PERF-FP8-ALPHA-FOLD's own, kept verbatim: the opt-in // toggle, `indt == BF16` (honouring VT_GDN_IN_BF16's documented rollback, since -// that lever is the one being unblocked), and `outdt == BF16` (which is what -// confines this to the dense 27B — the 35B is MoE, so GdnOutDType is f32 there -// and the whole arm stays inert). +// that lever is the one being unblocked), and `outdt == BF16`, which keeps the +// chain dtype-uniform. +// +// GDN-MOE-BF16-OUT (#1168), and #521 asked for exactly this correction. That +// third term USED to be described here as what "confines this to the dense 27B — +// the 35B is MoE, so GdnOutDType is f32 there and the whole arm stays inert". +// `GdnOutDType()` no longer branches on model shape, so `outdt == BF16` is now +// unconditional at the default and the term excludes NO checkpoint. What keeps +// this arm inert on the 35B is the DEFAULT-OFF `VT_GDN_FP8_IN_BF16` toggle, and +// that is the only remaining bound — which is why the toggle case below is the +// one a reader must not mistake for a shape guard. TEST_CASE("qwen27 fp8 merged in_proj dtype is one decision, all three terms") { using vllm::detail::GdnFp8MergedMixedQkvDType; using vt::DType; @@ -539,7 +557,7 @@ TEST_CASE("qwen27 fp8 merged in_proj dtype is one decision, all three terms") { CHECK(GdnFp8MergedMixedQkvDType(true, DType::kF32, DType::kBF16) == DType::kF32); // VT_GDN_IN_BF16=0 rollback. CHECK(GdnFp8MergedMixedQkvDType(true, DType::kBF16, DType::kF32) == - DType::kF32); // VT_GDN_OUT_BF16=0 / the 35B MoE default. + DType::kF32); // the VT_GDN_OUT_BF16=0 rollback, on EITHER arm (#1168). CHECK(GdnFp8MergedMixedQkvDType(true, DType::kF32, DType::kF32) == DType::kF32); CHECK(GdnFp8MergedMixedQkvDType(false, DType::kF32, DType::kF32) == @@ -652,6 +670,67 @@ TEST_CASE("qwen27 packed GDN fp8-tower toggle defaults OFF") { CHECK(PackedGdnDecodeFp8TowerFlagIsOn("1x")); } +// GDN-MOE-BF16-OUT (#1168), Edit 1. The GDN recurrence-output dtype is resolved +// from the ENVIRONMENT alone, and the resolver has no model-shape input to give +// it a second answer. `GdnOutDType()` used to take a `bool dense_model` and +// default to it, so a reader had to visit every call site to learn what the +// default was and a new call site could reintroduce the dense/MoE split +// silently. vLLM has no such parameter because it has no such decision: it +// resolves ONE model dtype and every layer inherits it +// (qwen_gdn_linear_attn.py:870-873, :843, :459-465 @ 5559679). +// +// The truth table is asked of the pure helper the production resolver reads, +// mirroring PackedGdnDecodeFp8TowerFlagIsOn: `GdnOutDType()` caches its getenv +// in a function-local static, so one process can only ever observe one value of +// it. `nullptr` IS the production default, and it answers bf16 with nothing in +// the signature that could say dense or MoE. Default-ON — anything that is not +// a leading '0' is bf16 — the polarity VT_GDN_IN_BF16 uses and the one +// PackedGdnDecodeEnvSelected already mirrors for VT_GDN_OUT_BF16. +// +// What this case CANNOT say is what the model runs; that is +// `test_qwen35_paged_forward`'s "the GDN recurrence output and z gate are bf16", +// which enters through ModelRegistry::Forward on a MoE config. +TEST_CASE("qwen27 GDN out dtype is bf16 by default and keys on no model shape") { + using vllm::detail::GdnOutBf16FlagIsOn; + + CHECK(GdnOutBf16FlagIsOn(nullptr)); // unset -> bf16, on EITHER arm. + CHECK_FALSE(GdnOutBf16FlagIsOn("0")); + CHECK_FALSE(GdnOutBf16FlagIsOn("0x")); // leading char decides, as elsewhere. + CHECK(GdnOutBf16FlagIsOn("1")); + CHECK(GdnOutBf16FlagIsOn("")); // not a leading '0'. +} + +// GDN-MOE-BF16-OUT (#1168), Edit 2. The eligibility carries NO model-shape term: +// an eligibility whose every remaining term is true selects packed decode, and +// nothing in it can say whether the checkpoint is dense or MoE. +// +// The `dense_model` term this replaces entered at f344decf4 as one of that +// change's "real-model safety gates" and was never revisited. Neither reference +// has an equivalent: VLLM_ENABLE_FLA_PACKED_RECURRENT_DECODE defaults True with +// no shape term (vllm/envs.py:124 @ 5559679), and SGLang keys +// `supports_packed_decode` on the platform alone (gdn_triton.py:43 @ f63458b5be). +// It became REDUNDANT rather than merely unsupported once GdnOutDType stopped +// branching on model shape: `core_out` is `outdt` and +// GdnPackedDecodeDTypesCompatible already pins it to BF16, so the dtype rule is +// what excludes an f32 recurrence output on either arm. +// +// This case deliberately never NAMES a model-shape field, so it says the same +// thing before and after the field exists. +TEST_CASE("qwen27 packed GDN selection carries no model-shape term") { + vllm::detail::GdnPackedDecodeEligibility e; + e.runtime_enabled = true; + e.cuda = true; + e.has_packed_ba = true; + e.merged_ba_enabled = true; + e.dtype_compatible = true; + e.has_state_indices = true; + e.num_actual_tokens = 4; + e.num_decodes = 4; + e.num_decode_tokens = 4; + + CHECK(vllm::detail::ShouldUsePackedGdnDecode(e)); +} + // THE ROW. The eligibility must key on the dtypes the packed op needs and NOT // on "the GDN weights happen to be fp8". Composing the three helpers exactly as // GdnBlockPaged does, an fp8 tower and a bf16 tower that produce the SAME @@ -678,7 +757,6 @@ TEST_CASE("qwen27 packed GDN selection keys on dtypes, not on fp8 weights") { GdnPackedDecodeEligibility e; e.runtime_enabled = true; e.cuda = true; - e.dense_model = true; e.has_packed_ba = true; e.merged_ba_enabled = true; e.dtype_compatible = diff --git a/tests/vllm/models/test_qwen35_paged_forward.cpp b/tests/vllm/models/test_qwen35_paged_forward.cpp index d7d5dce4b..28ca0e197 100644 --- a/tests/vllm/models/test_qwen35_paged_forward.cpp +++ b/tests/vllm/models/test_qwen35_paged_forward.cpp @@ -19,11 +19,16 @@ #include #include +#include #include #include +#include + +#include "vllm/model_executor/models/model_registry.h" #include "vllm/model_executor/models/qwen3_5.h" #include "vllm/model_executor/models/qwen3_5_weights.h" +#include "vllm/model_executor/models/qwen3_5_internal.h" #include "vllm/transformers_utils/hf_config.h" #include "vllm/v1/attention/backend.h" #include "vllm/v1/attention/backends/gdn_attn.h" @@ -37,12 +42,26 @@ using vllm::OwnedTensor; using vllm::PagedKvCache; using vllm::Qwen3_5Model; using vllm::Qwen3_5MoeWeights; +using vllm::ModelForwardInput; +using vllm::ModelRegistry; using vllm::v1::CommonAttentionMetadata; using vllm::v1::GDNAttentionMetadata; using vt::DType; namespace { +// GDN-MOE-BF16-OUT (#1168), fresh-review repair. Read VT_GDN_OUT_BF16 DIRECTLY +// from the process environment, deliberately NOT through +// `detail::GdnOutBf16FlagIsOn`: these cases exist to catch a resolver that was +// hardwired or severed from its parser, and re-using either of them to compute +// the expectation would make exactly that mutation self-consistent. Same reason +// `LeverOn` exists in test_dense_gateup_fused_marlin.cpp. Default ON; only a +// leading '0' turns it off. +bool GdnOutBf16LeverOn() { + const char* e = std::getenv("VT_GDN_OUT_BF16"); + return !(e != nullptr && e[0] == '0'); +} + // splitmix64-based small deterministic weight values in [-0.08, 0.08). uint64_t Mix(uint64_t x) { x += 0x9E3779B97F4A7C15ULL; @@ -462,3 +481,103 @@ TEST_CASE("qwen35 paged: GDN state zeroing protects a fresh req in a mixed batch << " (garbage-seeded mamba block, zeroing must scrub)"); CHECK(d < 1e-2); } + +// GDN-MOE-BF16-OUT (#1168). The GDN recurrence output `dcore` and the `z` gate +// are the two largest OUTPUT-side activations of a GDN layer, and on a MoE +// checkpoint both were f32 while vLLM keeps them at the bf16 model dtype: +// `core_attn_out = torch.zeros(..., dtype=hidden_states.dtype)` +// (vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py:870-873 @ +// 5559679) and `z` is a split of the bf16 in_proj_qkvz output (:843, :859-860). +// Nothing upstream branches on dense vs MoE; vLLM resolves one model dtype and +// every layer inherits it. +// +// This case enters through ModelRegistry::Forward on a MoE config — the +// production entry point, over the registered MoE factory's own forward — and +// asks what the GDN block RAN, read off `dcore`'s tensor and the projected gate +// rather than off the GdnOutDType predicate. A test that called the predicate +// would prove the predicate answers; AGENTS.md's "Nothing lands dead" wants the +// capability, and the reviewer's mutation is to restore the `dense_model` form +// of GdnOutDType, which must turn this red. +// +// No token gate can see this axis: f32 is the MORE precise deviation, so the +// goldens pass either way while the path moves twice the bytes. +TEST_CASE("qwen35 paged MoE: the GDN recurrence output and z gate follow VT_GDN_OUT_BF16") { + const HfConfig c = MakeConfig(); + REQUIRE(c.num_experts > 0); // the arm that used to resolve f32. + const Qwen3_5MoeWeights w = MakeWeights(c); + vt::Queue q = Q(); + const int64_t T = 6; + const std::vector ids = {5, 9, 2, 31, 17, 3}; + const std::vector pos = {0, 1, 2, 3, 4, 5}; + const std::vector logits_indices; + + CachePool pool(c, /*num_blocks=*/8, /*block_size=*/8); + const CommonAttentionMetadata am = PrefillAttnMeta(T, {0, 1}, 8, 0); + const GDNAttentionMetadata gm = PrefillGdnMeta(T, 0); + + vllm::detail::ResetGdnOutActivationDTypes(); + std::unique_ptr model = + vllm::BorrowQwen3_5MoeLoadedModel(w); + ModelForwardInput in{ids, pos, am, gm, pool.attn_kv, pool.gdn_state, + c, q, logits_indices}; + in.num_reqs = 1; + const vllm::ForwardLogits logits = ModelRegistry::Forward(*model, in); + CHECK((logits.host.size() == static_cast(T * c.vocab_size) || + logits.on_device())); + + // `observed` separates "the forward produced f32" from "the forward never + // reached a paged GDN layer", which look identical to a dtype comparison. + const vllm::detail::GdnOutActivationDTypes dt = + vllm::detail::LastGdnOutActivationDTypes(); + REQUIRE(dt.observed); + + // Fresh-review repair, both halves. The case used to assert BF16 + // UNCONDITIONALLY, which made it two things it should not be: + // + // - a FALSE RED under the documented rollback. `VT_GDN_OUT_BF16=0` is what + // the row's own `## Gates` A/B sets, and an operator who exported it and + // re-ran this suite got 4/5 cases, 11/13 assertions and a FAILURE that was + // not a regression. The resolver caches its getenv, so the case could not + // neutralise the variable in-process either. + // - blind to a SEVERED resolver. `GdnOutDType()` hardwired to BF16 — cut from + // its parser and from the variable entirely — left this suite 5/5 13/13 and + // test_qwen27_paged_forward 31/31 770/770. + // + // Asserting against the environment as this file reads it directly fixes both: + // the rollback arm now INVERTS rather than fails, and it is the arm in which a + // hardwired BF16 reads BF16 where F32 was ordered. tests/CMakeLists.txt + // registers this same binary a second time with VT_GDN_OUT_BF16=0 so both arms + // actually run, the shape `_glue_fuse_off` and + // `test_dense_gateup_fused_marlin_off_*` already use for a read-once lever. + const bool lever = GdnOutBf16LeverOn(); + CAPTURE(lever); + const DType expect = lever ? DType::kBF16 : DType::kF32; + CHECK(dt.core_out == expect); + CHECK(dt.z_gate == expect); +} + +// GDN-MOE-BF16-OUT (#1168), fresh-review repair. The RESOLVER, not the parser. +// +// The CPU tier already pins `detail::GdnOutBf16FlagIsOn`'s truth table +// (test_qwen27_paged_forward), and the case above pins what the model RAN. Both +// stayed green when `GdnOutDType()` was replaced with `return DType::kBF16;`, +// because on the default environment BF16 is also the right answer. Nothing +// asserted that the production resolver consumes the parser at all, and this +// row's `## Gates` rest on a `VT_GDN_OUT_BF16=0|1` same-binary A/B: a lever +// wired to nothing does not lose coverage, it invalidates the measurement. +// +// One process observes one value of the cached getenv, so the two env values are +// two ctest registrations of this binary rather than two assertions here. +TEST_CASE("qwen35: GdnOutDType resolves from VT_GDN_OUT_BF16, not from a constant") { + const bool lever = GdnOutBf16LeverOn(); + CAPTURE(lever); + CHECK(vllm::detail::GdnOutDType() == + (lever ? DType::kBF16 : DType::kF32)); + + // And the parser it is supposed to read agrees with the environment on this + // process's own value. Two separate statements: the one above says the + // resolver answers what the environment says, this one says the parser does + // too, so a divergence names WHICH of the two moved. + CHECK(vllm::detail::GdnOutBf16FlagIsOn(std::getenv("VT_GDN_OUT_BF16")) == + lever); +} diff --git a/tests/vt/test_ops_gdn.cpp b/tests/vt/test_ops_gdn.cpp index 7f47d177c..26b28b2c9 100644 --- a/tests/vt/test_ops_gdn.cpp +++ b/tests/vt/test_ops_gdn.cpp @@ -1754,8 +1754,9 @@ void RunRmsNormGatedFastContiguous(int64_t rows, bool sigmoid_gate, uint32_t see DType in_dt = DType::kBF16, DType out_dt = DType::kBF16) { constexpr int64_t d = 128; // Dv, the only production gated-norm shape // Adversarial: wide x/z ranges exercise variance magnitude and silu/sigmoid - // saturation; w spans [-1,1] like the real norm weight. in_dt=bf16 is the 27B - // dense core/z/weight; in_dt=f32 is the 35B MoE core/z/weight (GdnOutDType f32). + // saturation; w spans [-1,1] like the real norm weight. in_dt=bf16 is the + // core/z/weight both arms carry by default; in_dt=f32 is the VT_GDN_OUT_BF16=0 + // rollback arm, which was the 35B MoE's DEFAULT before #1168. const auto xf = RandomF32(static_cast(rows * d), seed, -6.0f, 6.0f); const auto zf = RandomF32(static_cast(rows * d), seed + 1, -6.0f, 6.0f); const auto wf = RandomF32(static_cast(d), seed + 2, -1.0f, 1.0f); @@ -3413,7 +3414,7 @@ TEST_CASE("CUDA rmsnorm_gated decode-fast (VT_RMSNORM_GATED_FAST) matches rollba // Real GDN decode shapes: [T*Hv, Dv=128] (27B Hv=48, 35B Hv=32) at c1-c16, both // silu (sigmoid_gate=false, the production gate) and sigmoid. fast==shipped 0-ulp. // dtype combos: bf16->bf16 (27B dense core/z), f32->bf16 (35B MoE core/z under - // GlueFuse — GdnOutDType f32, gated-norm stores bf16), f32->f32 (35B GlueFuse off). + // GlueFuse — the VT_GDN_OUT_BF16=0 f32 arm, gated-norm stores bf16), f32->f32. struct DtCombo { DType in, out; };