From 8bd5e343594ef380db9c431006f3e8880ad3f033 Mon Sep 17 00:00:00 2001
From: Kaden Schutt <151092359+Kaden-Schutt@users.noreply.github.com>
Date: Thu, 3 Sep 2026 00:56:17 +0000
Subject: [PATCH 1/5] fix(prefill): share MQ-V2 admit rule between llama and
 qwen35 gates
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

llama::is_batchable_la admitted MQ4G256V2/MQ6/5/3/2G256V2 for WMMA
prefill only on gfx1200/gfx1201 while qwen35::is_batchable_la admitted
them on gfx11+gfx12 behind HIPFIRE_MQV2_GFX11_WMMA, so plain Llama/Qwen3
dense qt=44 models prefetched per-token on gfx1100/1151 while Qwen3.5/3.8
took WMMA — despite both doc-comments claiming an exact match (audit
2026-09-02 Broken 1).

Move the dtype set + arch set + kill-switch helper into
llama::mqv2_wmma_batchable / llama::mqv2_gfx11_wmma_enabled_from_env in
hipfire-runtime and delegate from both callers, so the lockstep is
structural. MQ4CG256 (qt=45) stays gfx12-only in both by intent.

Tests: rename the two gfx12-only llama admit tests to gfx11+gfx12
expectations, repoint qwen35 env-escape test at the shared helper, and
add mqv2_admit_llama_qwen35_lockstep asserting both gates agree over
the MQ-V2 dtypes x {gfx1100, gfx1151, gfx1201, gfx1030, gfx1010}.
---
 .../hipfire-arch-qwen35/src/qwen35/prefill.rs |  87 +++++++------
 crates/hipfire-runtime/src/llama.rs           | 114 +++++++++++++-----
 2 files changed, 135 insertions(+), 66 deletions(-)

diff --git a/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs b/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs
index 418e23b44..01489eac9 100644
--- a/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs
+++ b/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs
@@ -1543,17 +1543,14 @@ pub(crate) fn is_batchable_la(dt: DType, arch: &str) -> bool {
     // (gfx1100/1101/1102/1150/1151 + gfx1200/1201) but gate the gfx11 half
     // behind HIPFIRE_MQV2_GFX11_WMMA != "0" — setting
     // HIPFIRE_MQV2_GFX11_WMMA=0 restores the per-token fallback ONLY on
-    // gfx11, leaving gfx12 untouched. Lockstep with the HasWmma predicate
-    // on GemmMq*G256V2* keys and with gemm_mq*g256v2's has_wmma() guard.
+    // gfx11, leaving gfx12 untouched. Delegates to the shared
+    // `hipfire_runtime::llama::mqv2_wmma_batchable` rule so this stays in
+    // lockstep with `llama::is_batchable_la` structurally. Lockstep with
+    // the HasWmma predicate on GemmMq*G256V2* keys and with
+    // gemm_mq*g256v2's has_wmma() guard.
     // MQ4CG256 (qt45) remains gfx12-only until its gfx11 sibling lands.
-    let mqv2_with_wmma = matches!(
+    let mqv2_with_wmma = llama::mqv2_wmma_batchable(
         dt,
-        DType::MQ4G256V2
-            | DType::MQ6G256V2
-            | DType::MQ5G256V2
-            | DType::MQ3G256V2
-            | DType::MQ2G256V2
-    ) && mqv2_gfx11_wmma_enabled_from_env(
         hipfire_config::developer_var("HIPFIRE_MQV2_GFX11_WMMA")
             .ok()
             .as_deref(),
@@ -1582,24 +1579,6 @@ pub(crate) fn is_batchable_la(dt: DType, arch: &str) -> bool {
         || bf16_with_gfx942
 }
 
-/// Helper for MQ2/3/4/5/6G256V2 (qt44,47-50) batched prefill admit: gfx12 always, gfx11
-/// gated by HIPFIRE_MQV2_GFX11_WMMA != "0". Public for testability, mirrors
-/// `mq6_batched_admit_enabled_from_env` / `q8_prefill_wmma_enabled_from_env`.
-/// `value` is the raw env var (None = unset → default ON); only Some("0")
-/// disables the gfx11 path. Gfx12 is unaffected by the env var.
-pub(crate) fn mqv2_gfx11_wmma_enabled_from_env(value: Option<&str>, arch: &str) -> bool {
-    let gfx11_enabled = value != Some("0");
-    if matches!(arch, "gfx1200" | "gfx1201") {
-        true
-    } else if matches!(
-        arch,
-        "gfx1100" | "gfx1101" | "gfx1102" | "gfx1150" | "gfx1151"
-    ) {
-        gfx11_enabled
-    } else {
-        false
-    }
-}
 /// Single source of truth for per-layer batchability and checked geometry.
 /// Called by `validate_ep_batch_compatibility`, `prefill_batch_pbs_eligible`,
 /// `fa_batched_ok` guard, and later EP state preflight. Validates every
@@ -8152,44 +8131,44 @@ mod tests {
     #[test]
     fn qwen35_is_batchable_la_mq4_v2_env_escape() {
         // HIPFIRE_MQV2_GFX11_WMMA=0 restores fallback ONLY on gfx11; gfx12
-        // remains admitted. Use the helper directly to avoid global env
+        // remains admitted. Use the shared helper directly to avoid global env
         // mutation flakiness in parallel tests — is_batchable_la delegates
-        // to this helper verbatim.
+        // to `llama::mqv2_wmma_batchable`, which calls this helper verbatim.
         for arch in ["gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151"] {
             assert!(
-                !mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch),
+                !llama::mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch),
                 "env=0 should disable {arch}"
             );
             assert!(
-                mqv2_gfx11_wmma_enabled_from_env(None, arch),
+                llama::mqv2_gfx11_wmma_enabled_from_env(None, arch),
                 "unset should enable {arch}"
             );
             assert!(
-                mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch),
+                llama::mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch),
                 "env=1 should enable {arch}"
             );
         }
         for arch in ["gfx1200", "gfx1201"] {
             assert!(
-                mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch),
+                llama::mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch),
                 "gfx12 unaffected by env=0 on {arch}"
             );
             assert!(
-                mqv2_gfx11_wmma_enabled_from_env(None, arch),
+                llama::mqv2_gfx11_wmma_enabled_from_env(None, arch),
                 "gfx12 enabled without env on {arch}"
             );
         }
         for arch in ["gfx1010", "gfx942", "gfx1030", "gfx1103", "gfx1152"] {
             assert!(
-                !mqv2_gfx11_wmma_enabled_from_env(None, arch),
+                !llama::mqv2_gfx11_wmma_enabled_from_env(None, arch),
                 "non-WMMA {arch} must never admit"
             );
             assert!(
-                !mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch),
+                !llama::mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch),
                 "non-WMMA {arch} with env=0"
             );
             assert!(
-                !mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch),
+                !llama::mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch),
                 "non-WMMA {arch} with env=1"
             );
         }
@@ -8241,6 +8220,40 @@ mod tests {
         assert_eq!(rdna_compute::MQ4V2_GROUP_BYTES, 136);
     }
 
+    #[test]
+    fn mqv2_admit_llama_qwen35_lockstep() {
+        // Audit 2026-09-02 Broken 1: `llama::is_batchable_la` admitted MQ-V2
+        // only on gfx12 while this module admitted gfx11+gfx12, despite both
+        // doc-comments claiming an exact match. Both now delegate to the
+        // shared `llama::mqv2_wmma_batchable` rule; this test iterates the
+        // MQ-V2 dtypes over gfx11, gfx12, and pre-WMMA arches and asserts the
+        // two gates agree. Both read `HIPFIRE_MQV2_GFX11_WMMA` from the
+        // environment identically, so equality holds in any env state
+        // without mutating globals.
+        let dts = [
+            DType::MQ4G256V2,
+            DType::MQ6G256V2,
+            DType::MQ5G256V2,
+            DType::MQ3G256V2,
+            DType::MQ2G256V2,
+            DType::MQ4CG256,
+        ];
+        for dt in dts {
+            for arch in ["gfx1100", "gfx1151", "gfx1201", "gfx1030", "gfx1010"] {
+                assert_eq!(
+                    llama::is_batchable_la(dt, arch),
+                    is_batchable_la(dt, arch),
+                    "lockstep drift for {dt:?} on {arch}"
+                );
+            }
+        }
+        // Absolute pins so the test also fails if the shared rule itself
+        // regresses, not just on caller drift.
+        assert!(is_batchable_la(DType::MQ4G256V2, "gfx1201"));
+        assert!(!is_batchable_la(DType::MQ4G256V2, "gfx1030"));
+        assert!(!is_batchable_la(DType::MQ4CG256, "gfx1100"));
+    }
+
     #[test]
     fn qwen35_v2_dense_keys_are_exact_no_hfq4_default() {
         // Contract: every admitted V2 dtype maps 1:1 to its exact V2 kernel
diff --git a/crates/hipfire-runtime/src/llama.rs b/crates/hipfire-runtime/src/llama.rs
index e9e909107..1251a95a8 100644
--- a/crates/hipfire-runtime/src/llama.rs
+++ b/crates/hipfire-runtime/src/llama.rs
@@ -1838,10 +1838,49 @@ pub fn prefill_forward(
 /// largest physical_cap any consumer sets up.
 pub const PREFILL_MAX_BATCH: usize = 256;
 
+/// Kill-switch for the MQ-V2 (qt44 + neutral qt47-50) gfx11 WMMA prefill path:
+/// gfx12 (`gfx1200`/`gfx1201`) is always admitted, gfx11
+/// (`gfx1100`/`gfx1101`/`gfx1102`/`gfx1150`/`gfx1151`) is admitted unless
+/// `HIPFIRE_MQV2_GFX11_WMMA=0`, anything else is rejected. Single definition
+/// shared by `llama::is_batchable_la` and `qwen35::is_batchable_la` so the two
+/// stay in lockstep structurally instead of by matching comments.
+/// `value` is the raw env var (None = unset → default ON); only `Some("0")`
+/// disables the gfx11 path. Gfx12 is unaffected by the env var.
+pub fn mqv2_gfx11_wmma_enabled_from_env(value: Option<&str>, arch: &str) -> bool {
+    let gfx11_enabled = value != Some("0");
+    if matches!(arch, "gfx1200" | "gfx1201") {
+        true
+    } else if matches!(
+        arch,
+        "gfx1100" | "gfx1101" | "gfx1102" | "gfx1150" | "gfx1151"
+    ) {
+        gfx11_enabled
+    } else {
+        false
+    }
+}
+
+/// Admit rule for the MQ-V2 family (`MQ4G256V2` + neutral `MQ6/5/3/2G256V2`)
+/// in batched WMMA prefill: dtype set × arch set × the
+/// `HIPFIRE_MQV2_GFX11_WMMA` kill-switch in one function. Both
+/// `llama::is_batchable_la` and `qwen35::is_batchable_la` delegate here;
+/// `MQ4CG256` (qt45) stays gfx12-only in both callers and is intentionally
+/// NOT part of this rule.
+pub fn mqv2_wmma_batchable(dt: DType, mqv2_gfx11_wmma: Option<&str>, arch: &str) -> bool {
+    matches!(
+        dt,
+        DType::MQ4G256V2
+            | DType::MQ6G256V2
+            | DType::MQ5G256V2
+            | DType::MQ3G256V2
+            | DType::MQ2G256V2
+    ) && mqv2_gfx11_wmma_enabled_from_env(mqv2_gfx11_wmma, arch)
+}
+
 /// Is this dtype/arch combination eligible for the batched WMMA prefill
-/// kernels? Matches `qwen35::is_batchable_la` exactly so plain Qwen3 and
-/// hybrid Qwen3.5 share one rule and stay in lockstep when new dtypes or
-/// arches gain WMMA support.
+/// kernels? Shares the MQ-V2 admit rule with `qwen35::is_batchable_la` via
+/// `mqv2_wmma_batchable`, so plain Qwen3 and hybrid Qwen3.5 stay in lockstep
+/// structurally when new dtypes or arches gain WMMA support.
 pub fn is_batchable_la(dt: DType, arch: &str) -> bool {
     let always_ok = matches!(
         dt,
@@ -1874,21 +1913,25 @@ pub fn is_batchable_la(dt: DType, arch: &str) -> bool {
             arch,
             "gfx1010" | "gfx1011" | "gfx1012" | "gfx1013" | "gfx1030" | "gfx1031" | "gfx1032"
         );
-    // MQ4G256V2 / MQ4CG256 batched prefill + batched lm_head GEMM exist only
-    // on gfx12 (gfx1200/gfx1201). Outside gfx12, fall back to per-token decode
-    // rather than dispatching a gfx12 WMMA kernel. Lockstep with
-    // qwen35::is_batchable_la (qt44/qt45).
-    // Extended to neutral V2 family qt47-50.
-    let mq4_v2_gfx12 = matches!(
+    // MQ-V2 family (qt44 + neutral qt47-50) batched prefill + batched lm_head
+    // GEMM: gfx11 + gfx12 via the shared rule, with the gfx11 half behind the
+    // `HIPFIRE_MQV2_GFX11_WMMA=0` kill-switch. Delegates to
+    // `mqv2_wmma_batchable` so this stays in lockstep with
+    // `qwen35::is_batchable_la` structurally, not by matching comments.
+    // Outside the admitted arches, fall back to per-token decode rather than
+    // dispatching a foreign-arch WMMA kernel.
+    let mq4_v2 = mqv2_wmma_batchable(
         dt,
-        DType::MQ4G256V2
-            | DType::MQ4CG256
-            | DType::MQ6G256V2
-            | DType::MQ5G256V2
-            | DType::MQ3G256V2
-            | DType::MQ2G256V2
-    ) && matches!(arch, "gfx1200" | "gfx1201");
-    wmma_only || mq3_gfx10_scalar || mq4_v2_gfx12
+        hipfire_config::developer_var("HIPFIRE_MQV2_GFX11_WMMA")
+            .ok()
+            .as_deref(),
+        arch,
+    );
+    // MQ4CG256 (qt45) remains gfx12-only until its gfx11 sibling lands —
+    // intentionally not part of the shared rule, in both callers.
+    let mq4cg256_gfx12 =
+        matches!(dt, DType::MQ4CG256) && matches!(arch, "gfx1200" | "gfx1201");
+    wmma_only || mq3_gfx10_scalar || mq4_v2 || mq4cg256_gfx12
 }
 
 /// Per-call scratch for `forward_prefill_batch`. Holds [N × ...] working
@@ -8247,24 +8290,32 @@ mod tests {
     }
 
     #[test]
-    fn is_batchable_la_mq4_v2_gfx12_only() {
-        // MQ4G256V2 / MQ4CG256 batched prefill is gfx12-only; other arches
-        // fall back to per-token decode.
-        for arch in ["gfx1200", "gfx1201"] {
+    fn is_batchable_la_mq4_v2_gfx11_and_gfx12() {
+        // MQ4G256V2 batched prefill admits gfx11 + gfx12 through the shared
+        // `mqv2_wmma_batchable` rule (gfx11 behind HIPFIRE_MQV2_GFX11_WMMA);
+        // MQ4CG256 stays gfx12-only. Lockstep with
+        // `qwen35::is_batchable_la` by construction.
+        for arch in [
+            "gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151", "gfx1200", "gfx1201",
+        ] {
             assert!(
                 is_batchable_la(DType::MQ4G256V2, arch),
                 "MQ4G256V2 should batch on {arch}"
             );
+        }
+        for arch in ["gfx1200", "gfx1201"] {
             assert!(
                 is_batchable_la(DType::MQ4CG256, arch),
                 "MQ4CG256 should batch on {arch}"
             );
         }
-        for arch in ["gfx1010", "gfx1100", "gfx942"] {
+        for arch in ["gfx1010", "gfx1030", "gfx942"] {
             assert!(
                 !is_batchable_la(DType::MQ4G256V2, arch),
                 "MQ4G256V2 must fall back on {arch}"
             );
+        }
+        for arch in ["gfx1010", "gfx1100", "gfx1151", "gfx942"] {
             assert!(
                 !is_batchable_la(DType::MQ4CG256, arch),
                 "MQ4CG256 must fall back on {arch}"
@@ -8273,14 +8324,19 @@ mod tests {
     }
 
     #[test]
-    fn is_batchable_la_v2_family_gfx12_only() {
-        for arch in ["gfx1200", "gfx1201"] {
-            assert!(is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 gfx12");
-            assert!(is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 gfx12");
-            assert!(is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 gfx12");
-            assert!(is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 gfx12");
+    fn is_batchable_la_v2_family_gfx11_and_gfx12() {
+        // Neutral V2 family (qt47-50) admits gfx11 + gfx12 through the shared
+        // `mqv2_wmma_batchable` rule, mirroring
+        // `qwen35_is_batchable_la_v2_family_gfx11_and_gfx12`.
+        for arch in [
+            "gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151", "gfx1200", "gfx1201",
+        ] {
+            assert!(is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 on {arch}");
+            assert!(is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 on {arch}");
+            assert!(is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 on {arch}");
+            assert!(is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 on {arch}");
         }
-        for arch in ["gfx1010", "gfx1100", "gfx942"] {
+        for arch in ["gfx1010", "gfx1030", "gfx942"] {
             assert!(!is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 fallback");
             assert!(!is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 fallback");
             assert!(!is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 fallback");

From db9c1b2497fb0966fc05e1cfa6e9f2231b387dcf Mon Sep 17 00:00:00 2001
From: Kaden Schutt <151092359+Kaden-Schutt@users.noreply.github.com>
Date: Thu, 3 Sep 2026 01:01:27 +0000
Subject: [PATCH 2/5] fix(mq4v2): discriminating disjoint-halves GEMM parity
 arm + doc corrections

mq4v2_gemm_parity's Gaussian weights give both halves near-identical
headers, so a wrong half-select hides in quantization noise despite the
header comment promising a systematic blow-up (audit 2026-09-02
Missing 1). Add arm 2 using the disjoint-halves construction from
mq4v2_residual_parity (half0 [-1,1], half1 [96,160]) over the same
batch-size sweep: v2 output must match the f32 reference within 5%
rel-RMS, and the swapped-headers negative control (as in
mq4v2_moe_parity) must DISAGREE. Keep the Gaussian v1-vs-v2 arm and fix
the header comment to state what each arm can and cannot detect. Add a
host-side test proving the fixture discriminates with no GPU.

Docs (audit Would-change 1-2): spec section 9 now records MoE as
production-wired for qt=44, the XBATCH single-row path as ported, and
the gfx11 kt+=2 / residual kt++ stepping; residual_mmq.hip loses its
stale Experimental tag; the gfx12 QKV kernel loses its
HYPOTHESIS/scaffold wording for the validated C-map statement.
---
 .../examples/mq4v2_gemm_parity.rs             | 181 ++++++++++++++++--
 docs/quant-formats/mq4-v2.md                  |  11 +-
 kernels/src/gemm_mq4g256v2_residual_mmq.hip   |   7 +-
 kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip |  38 ++--
 4 files changed, 193 insertions(+), 44 deletions(-)

diff --git a/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs b/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs
index 4a10c0461..3fbf3318e 100644
--- a/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs
+++ b/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs
@@ -1,4 +1,5 @@
-//! v1-vs-v2 cross-check for the **WMMA GEMM** path (qt=13 vs qt=44).
+//! v1-vs-v2 cross-check for the **WMMA GEMM** path (qt=13 vs qt=44),
+//! plus a discriminating disjoint-halves arm for the v2 residual path.
 //!
 //! `mq4v2_parity` verifies the decode GEMV against a host oracle. It does NOT
 //! cover the WMMA prefill GEMMs, which are what `--scoring-mode prefill` actually
@@ -6,7 +7,27 @@
 //! executing (8 v2 modules compiled), WT2 KLD came back 16.705139 against a
 //! 0.043776 baseline.
 //!
-//! ## Why cross-check instead of a host reference
+//! ## Arm 1: Gaussian v1-vs-v2 agreement — CANNOT catch a half-select bug
+//!
+//! Realistic post-FWHT weights (Gaussian, sigma ~0.011) give the two halves of
+//! every group near-identical `(scale, zero)` headers, so a wrong half-select
+//! predicate lands inside 4-bit quantization noise. This arm detects gross v2
+//! decode errors (agreement far above the quantization floor) but a wrong
+//! predicate passes it silently. See arm 2 for the discriminating fixture.
+//!
+//! ## Arm 2: disjoint halves + negative control — CATCHES a half-select bug
+//!
+//! Same construction as `mq4v2_residual_parity.rs`: half 0 in `[-1, 1]`, half 1
+//! in `[96, 160]`, packed through the same fp16 round-trip. A kernel that
+//! decodes half 1 with half 0's header reconstructs `~0` instead of `~128`,
+//! so the v2 output is asserted against an exact-dequant f32 reference within
+//! a tight tolerance (rel-RMS below 5%). The negative control — a reference
+//! computed with the halves' headers swapped — must DISAGREE by an order of
+//! magnitude more; if it ever agrees, the fixture has stopped separating the
+//! halves and the arm is vacuous (same control as
+//! `rdna-compute/examples/mq4v2_moe_parity.rs`).
+//!
+//! ## Why cross-check instead of a host reference (arm 1)
 //!
 //! Replicating a WMMA kernel on the host means reproducing fp16 activation
 //! conversion, 16x16 tiling, and accumulation order — a reference that is itself
@@ -16,10 +37,11 @@
 //! Both paths then share every stage except the 8 header bytes and their decode.
 //! v1 quantizes with one affine grid per 256 weights; v2 with one per 128. v2 is
 //! therefore slightly MORE accurate, so agreement should sit at the scale of
-//! 4-bit quantization noise. A systematic blow-up isolates the v2 header decode —
-//! in practice the half-select predicate, which the spec calls out as "the single
-//! highest-risk detail in the port" because a wrong one "compiles, runs, and
-//! silently applies the wrong scale to half of every tensor."
+//! 4-bit quantization noise. A systematic blow-up isolates a gross v2 header
+//! decode error — but NOT the half-select predicate, which the spec calls out as
+//! "the single highest-risk detail in the port" because a wrong one "compiles,
+//! runs, and silently applies the wrong scale to half of every tensor." Only
+//! arm 2's disjoint fixture can see that failure.
 //!
 //! ## Why sweep batch size
 //!
@@ -27,7 +49,7 @@
 //! body by batch size and flags. Scoring compiled `_bt8` and `_bt12`, so the BT
 //! bodies are live — and BT is b-transposed, which changes the nibble addressing
 //! the half-select must be derived from. Sweeping batch size tells us WHICH body
-//! is wrong rather than just that something is.
+//! is wrong rather than just that something is. Both arms sweep.
 //!
 //! Run: `cargo run --release -p hipfire-runtime --example mq4v2_gemm_parity`
 
@@ -51,6 +73,7 @@ fn prng(i: usize, salt: u32) -> f32 {
 /// Realistic post-FWHT weights: roughly Gaussian, sigma ~0.011 as measured on the
 /// Qwen3.8-27B parent. Deliberately NOT the disjoint-halves fixture -- here both
 /// containers must be individually reasonable so their outputs are comparable.
+/// (The discriminating fixture is `build_disjoint_halves` below, used by arm 2.)
 fn build_weights(m: usize, k: usize) -> Vec<f32> {
     let mut w = vec![0.0f32; m * k];
     for (i, v) in w.iter_mut().enumerate() {
@@ -62,6 +85,43 @@ fn build_weights(m: usize, k: usize) -> Vec<f32> {
     w
 }
 
+/// Discriminating fixture (same construction as `mq4v2_residual_parity.rs`):
+/// half 0 in `[-1, 1]`, half 1 in `[96, 160]`. The two halves occupy disjoint
+/// ranges, so a kernel that decodes half 1 with half 0's header reconstructs
+/// `~0` instead of `~128` and fails by >100% relative error instead of hiding
+/// inside quantization noise.
+fn build_disjoint_halves(m: usize, k: usize) -> Vec<f32> {
+    let mut w = vec![0.0f32; m * k];
+    for r in 0..m {
+        for c in 0..k {
+            let gi = c % GROUP;
+            let idx = r * k + c;
+            if gi < HALF {
+                // [-1, 1]
+                w[idx] = prng(idx, 0xA5A5_0001) * 2.0 - 1.0;
+            } else {
+                // [96, 160] — disjoint from half0 by two orders of magnitude
+                w[idx] = 96.0 + prng(idx, 0x5A5A_0002) * 64.0;
+            }
+        }
+    }
+    w
+}
+
+/// Swap the two 4-byte half-headers of every group in a packed v2 blob. The
+/// result decodes each half with the OTHER half's grid — the negative control:
+/// a reference built from this blob must DISAGREE with the correct reference.
+fn swap_v2_half_headers(blob: &[u8]) -> Vec<u8> {
+    let mut out = blob.to_vec();
+    for chunk in out.chunks_exact_mut(GROUP_BYTES) {
+        let mut tmp = [0u8; 8];
+        tmp.copy_from_slice(&chunk[0..8]);
+        chunk[0..4].copy_from_slice(&tmp[4..8]);
+        chunk[4..8].copy_from_slice(&tmp[0..4]);
+    }
+    out
+}
+
 /// qt=13 / HFQ4 container: `[0..4) f32 scale, [4..8) f32 zero` over all 256.
 fn pack_v1(w: &[f32], m: usize, k: usize) -> Vec<u8> {
     let gpr = k / GROUP;
@@ -250,6 +310,57 @@ fn main() {
         }
     }
 
+    // ── Arm 2: disjoint halves + negative control (residual path) ──────────
+    //
+    // Arm 1's Gaussian weights cannot discriminate a wrong half-select; this
+    // arm can. Same batch-size sweep, same v2 residual WMMA kernel, but the
+    // weights put half 0 in [-1, 1] and half 1 in [96, 160], so decoding half
+    // 1 with half 0's header is a ~100x scale error. The v2 output must match
+    // the exact-dequant f32 reference within a tight tolerance, AND the
+    // swapped-headers reference must DISAGREE — otherwise the fixture is
+    // vacuous and the arm proves nothing.
+    {
+        let wd = build_disjoint_halves(m, k);
+        let bd = pack_v2(&wd, m, k);
+        let bd_swapped = swap_v2_half_headers(&bd);
+        for &batch in &[1usize, 8, 12, 16, 32] {
+            let x: Vec<f32> = (0..batch * k)
+                .map(|i| prng(i, 0xC0FF_EE00) * 2.0 - 1.0)
+                .collect();
+            let want = ref_gemm(&bd, &x, m, k, batch, true);
+            let want_bug = ref_gemm(&bd_swapped, &x, m, k, batch, true);
+            let bug_rel = {
+                let bug_f32: Vec<f32> = want_bug.iter().map(|&v| v as f32).collect();
+                rel_rms(&bug_f32, &want)
+            };
+            // Host-side negative control: the swapped grid must be badly wrong
+            // before any GPU result is scored against it.
+            assert!(
+                bug_rel > 0.5,
+                "disjoint fixture not discriminating at batch {batch}: bug_rel {bug_rel:.3e} — halves overlap"
+            );
+            let d_a = gpu.upload_raw(&bd, &[bd.len()]).unwrap();
+            let d_x = gpu.upload_f32(&x, &[batch * k]).unwrap();
+            let d_y = gpu.zeros(&[batch * m], rdna_compute::DType::F32).unwrap();
+            gpu.gemm_hfq4g256_residual_wmma_gfx12_mq4v2(&d_a, &d_x, &d_y, m, k, batch)
+                .expect("v2 disjoint residual wmma launch");
+            gpu.hip.device_synchronize().unwrap();
+            let got = gpu.download_f32(&d_y).unwrap();
+            let e = rel_rms(&got, &want);
+            let verdict = if e < 0.05 && e < bug_rel * 0.1 {
+                "ok"
+            } else {
+                "FAIL"
+            };
+            eprintln!(
+                "disjoint batch {batch:>3}: v2 rel-rms {e:.4e}   bug {bug_rel:.3e}   {verdict}"
+            );
+            if verdict == "FAIL" {
+                failures.push((3000 + batch, e, bug_rel));
+            }
+        }
+    }
+
     // ── The fused multi-output GEMMs ────────────────────────────────────────
     //
     // These are the rest of the live v2 set. `gemm_qkvza` carries NINETEEN header
@@ -356,19 +467,63 @@ fn main() {
 
     if failures.is_empty() {
         eprintln!(
-            "\nmq4v2_gemm_parity: PASS — every live v2 WMMA GEMM matches its own exact dequant"
+            "\nmq4v2_gemm_parity: PASS — every live v2 WMMA GEMM matches its own exact dequant, and the disjoint-halves residual arm is half-select correct"
         );
     } else {
         eprintln!(
             "\nmq4v2_gemm_parity: FAIL — codes {:?}",
             failures.iter().map(|f| f.0).collect::<Vec<_>>()
         );
-        eprintln!("(1000 = gate_up, 2000 = qkvza, otherwise the residual batch size)");
+        eprintln!("(1000 = gate_up, 2000 = qkvza, 3000+batch = disjoint-halves residual, otherwise the Gaussian residual batch size)");
         eprintln!("The v1 row is the WMMA fp16 error floor; a v2 row far above it means that");
-        eprintln!("kernel mis-decodes its own header. Each body has its OWN nibble addressing,");
-        eprintln!(
-            "so its half-select predicate must be derived from that addressing, never copied."
-        );
+        eprintln!("kernel mis-decodes its own header. Only the disjoint arm (3000+batch) can");
+        eprintln!("see a half-select bug — the Gaussian arm hides one inside quantization noise.");
+        eprintln!("Each body has its OWN nibble addressing, so its half-select predicate must");
+        eprintln!("be derived from that addressing, never copied.");
         std::process::exit(1);
     }
 }
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn disjoint_fixture_discriminates_half_select() {
+        // Host-side proof that arm 2's negative control is load-bearing: the
+        // halves occupy disjoint ranges, header-swapping is an involution,
+        // and a swapped-headers reference DISAGREES with the correct one.
+        // Runs with no GPU.
+        let (m, k, batch) = (16usize, 256usize, 4usize);
+        let w = build_disjoint_halves(m, k);
+        for r in 0..m {
+            for c in 0..k {
+                let v = w[r * k + c];
+                if (c % GROUP) < HALF {
+                    assert!((-1.0..=1.0).contains(&v), "half0 out of range: {v}");
+                } else {
+                    assert!((96.0..=160.0).contains(&v), "half1 out of range: {v}");
+                }
+            }
+        }
+        let blob = pack_v2(&w, m, k);
+        let swapped = swap_v2_half_headers(&blob);
+        assert_ne!(swapped, blob, "swapping identical headers would be vacuous");
+        assert_eq!(
+            swap_v2_half_headers(&swapped),
+            blob,
+            "header swap must be an involution"
+        );
+        let x: Vec<f32> = (0..batch * k)
+            .map(|i| prng(i, 0xC0FF_EE00) * 2.0 - 1.0)
+            .collect();
+        let want = ref_gemm(&blob, &x, m, k, batch, true);
+        let want_bug = ref_gemm(&swapped, &x, m, k, batch, true);
+        let bug_f32: Vec<f32> = want_bug.iter().map(|&v| v as f32).collect();
+        let bug_rel = rel_rms(&bug_f32, &want);
+        assert!(
+            bug_rel > 0.5,
+            "swapped-headers reference must DISAGREE: bug_rel {bug_rel:.3e}"
+        );
+    }
+}
diff --git a/docs/quant-formats/mq4-v2.md b/docs/quant-formats/mq4-v2.md
index 7c5d25fbb..9374d3e0b 100644
--- a/docs/quant-formats/mq4-v2.md
+++ b/docs/quant-formats/mq4-v2.md
@@ -241,8 +241,11 @@ Measured across the 11 ported kernels, three distinct forms occur:
 | `kt < 8` | 8 | WMMA GEMM main + BT bodies | nibbles at `gp + 8 + kt*8 + k_grp*4` with `k_grp = tid>>4 ∈ {0,1}`; `kt*8 + 4 < 64` ⇔ `kt < 8` for both `k_grp` |
 | `quarter_in_group < 2` | 2 | WMMA `ldsstage` bodies | nibbles at `gp + 8 + quarter*32 + {0,8,16,24} + k_grp*4`; max offset at `quarter=1` is 60 < 64 |
 
-In the WMMA main path the loop steps `kt += 4` over tiles `kt..kt+3`, so all four tiles in one
-body share a half (kt=0,4 → half 0; kt=8,12 → half 1) and one select per body suffices.
+The step varies by variant — gfx11 main bodies step `kt += 2` (K2 unroll; both tiles in a
+pair share a half), the gfx11 residual body steps `kt++` (per-tile select), gfx12 bodies
+step `kt += 4` over tiles `kt..kt+3` (all four tiles in one body share a half: kt=0,4 →
+half 0; kt=8,12 → half 1) — but every body selects with `kt < 8`, so one select per
+body suffices in all three forms.
 
 A wrong predicate here **compiles, runs, and silently applies the wrong scale to half of every
 tensor**. It is the single highest-risk detail in the port. For any kernel not in the 11, redo
@@ -429,10 +432,10 @@ relatively less of the damage. "Codebooks are for the sub-4-bit tier" is defensi
 ### Not claimed / out of scope
 
 - **wave64** half-split (§ 4) — not verified; wave64 remains unsupported for these formats.
-- **MoE** paths (`gemv_hfq4g256_moe_*`, `gemm_*_moe_grouped_*`) — out of scope / fail-closed for V2 product tiers.
+- **MoE** paths (`gemv_mq4g256v2_moe_*`, `gemm_mq4g256v2_moe_grouped_*`) — production-wired for qt=44 (decode + prefill, gfx11 + gfx12; loader → `MoeResolution.routed_indexable_mq4v2` at `families/moe.rs:244` → `pipeline/mod.rs:1244-1262`).
 - **qt=45 on gfx11** — no gfx11 WMMA sibling; do not promote.
 - **gfx1030 default-R decision** (§ 6) — still open if/when gfx1030 ships these dtypes; not a dense-WMMA blocker.
-- Research-only surfaces: `muse_*`, dp4a / cpol / `ldscoop` / `ldsx` / `.v1`–`.v5` / `XBATCH` single-row path. (gfx11 base/BT WMMA for **qt=44** is production, not research-only.)
+- Research-only surfaces: `muse_*`, dp4a / cpol / `ldscoop` / `ldsx` / `.v1`–`.v5`. (The `XBATCH` single-row path IS ported for qt=44 — `gemv_mq4g256v2.hip:295-361` — and gfx11 base/BT WMMA for **qt=44** is production, not research-only.)
 
 ### Port surface that landed for dense HasWmma (qt=44) / dense gfx12 (qt=45)
 
diff --git a/kernels/src/gemm_mq4g256v2_residual_mmq.hip b/kernels/src/gemm_mq4g256v2_residual_mmq.hip
index 595eee5ff..43b3fcee4 100644
--- a/kernels/src/gemm_mq4g256v2_residual_mmq.hip
+++ b/kernels/src/gemm_mq4g256v2_residual_mmq.hip
@@ -7,7 +7,12 @@
 #include <stdint.h>
 #include <math.h>
 
-// Experimental HFQ4-G256 MMQ residual GEMM for RDNA3/RDNA3.5.
+// Production MQ4G256V2 (qt=44) MMQ residual GEMM for gfx1100/gfx1151.
+//
+// Selected as the fast path when `batch >= 128 && batch % 128 == 0`
+// (see `crates/rdna-compute/src/gemm.rs:17860-17909` and the qkv/qkvza/
+// gate_up/residual call sites); otherwise the WMMA residual path runs.
+// Applies the scale per 128-K half (`kt < 8` select), not a per-256 V1-ism.
 //
 // This mirrors the important parts of llama.cpp's AMD MMQ path:
 // - pre-quantize the activation matrix into block_q8_1_mmq DS4 layout
diff --git a/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip b/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip
index 5959ae87f..8ec6bfa8d 100644
--- a/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip
+++ b/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip
@@ -12,17 +12,10 @@
 // (which targets gfx11 / RDNA3). Compile this file with
 //   hipcc --offload-arch=gfx1200  (or gfx1201, ...).
 //
-// SCAFFOLD STATUS (2026-04-26):
-//   This file is NOT yet wired into dispatch.rs. It compiles for gfx12
-//   targets and is intended as the canonical pattern reference for the
-//   five remaining gfx12 WMMA ports (qkvza-hfq4, gate_up-hfq4, plus
-//   the three hfq6 variants). Runtime correctness must be validated on
-//   real RDNA4 hardware (9070 XT / R9700) via test_kernels before any
-//   dispatch.rs site is allowed to route here. The C-output mapping
-//   below is a HYPOTHESIS derived from CK trait math and has NOT been
-//   confirmed against silicon — see the channel-test note at the
-//   write-back loop. See `.skills/hipfire-arch-port/` for the port
-//   workflow.
+// STATUS: production gfx12 WMMA path, wired into dispatch. The C-output
+// mapping below was validated 2026-04-27 in this file's channel-test on
+// R9700 — the same map its `gemm_gate_up_mq4g256v2_wmma.gfx12.hip` sibling
+// documents.
 //
 // Differences from the gfx11 kernel:
 //   1. WMMA builtin: __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12
@@ -34,11 +27,12 @@
 //        lane (tid >> 4) = 0  -> carries K = [0..7]  of the 16-K tile
 //        lane (tid >> 4) = 1  -> carries K = [8..15] of the 16-K tile
 //      Each lane therefore loads HALF as much per WMMA tile as gfx11.
-//   4. C-output mapping (HYPOTHESIS — see below):
-//        gfx11 (validated, commit b7ac66a):
-//          acc[j] = C[2*j + (tid>>4)][tid & 15]
-//        gfx12 (UNVERIFIED, derived from CK kCM0/kCM1PerLane swap):
-//          acc[j] = C[8*(tid>>4) + j][tid & 15]
+//   4. C-output mapping (validated 2026-04-27 in the QKV scaffold's
+//      channel-test on R9700):
+//        gfx11: acc[j] = C[2*j + (tid>>4)][tid & 15]    (rows interleaved)
+//        gfx12: acc[j] = C[8*(tid>>4) + j][tid & 15]    (rows contiguous —
+//                                                        group 0: rows 0..7,
+//                                                        group 1: rows 8..15)
 //
 // Reference (ROCm 7.x):
 //   /opt/rocm/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_base_traits.hpp
@@ -165,10 +159,8 @@ extern "C" __global__ void gemm_qkv_mq4g256v2_wmma_gfx12(
     }
 
     // --- Output ---
-    //
-    // gfx12 wave32 WMMA C-mapping HYPOTHESIS — DO NOT TRUST WITHOUT
-    // CHANNEL-TEST ON REAL HARDWARE. Derived from the CK trait swap
-    // (warp_gemm_attribute_wmma_impl_base_traits.hpp):
+    // gfx12 wave32 WMMA C-mapping (validated 2026-04-27 in this file's
+    // channel-test on R9700; same map the gate_up gfx12 sibling documents):
     //
     //   gfx11: kCMLane=2, kCM0PerLane=8, kCM1PerLane=1
     //          -> acc[j] = C[2*j + (tid>>4)][tid & 15]  (rows interleaved)
@@ -176,12 +168,6 @@ extern "C" __global__ void gemm_qkv_mq4g256v2_wmma_gfx12(
     //          -> acc[j] = C[8*(tid>>4) + j][tid & 15]  (rows contiguous)
     //
     // i.e. lane group 0 holds output rows 0..7, lane group 1 rows 8..15.
-    // The gfx11 mapping was silently wrong for ~6 weeks before being fixed
-    // in commit b7ac66a — assume the same risk here. Validation recipe
-    // (from that fix): add `if (blockIdx.x == 0 && blockIdx.y == 0)
-    // printf("tid=%d j=%d row=%d col=%d acc=%f\\n", tid, j, out_row,
-    // out_col, acc[j]);` and compare against a CPU reference for a small
-    // golden case. Adjust the (j, k_grp) -> out_row formula until it matches.
     const int out_col = batch_start + m_lane;  // batch index
     if (out_col < N) {
         #pragma unroll

From 077aac54523a23b338abe8ca5836b77c81074399 Mon Sep 17 00:00:00 2001
From: Kaden Schutt <kaden@hipfire.dev>
Date: Thu, 3 Sep 2026 01:02:22 +0000
Subject: [PATCH 3/5] style: rustfmt changed files

---
 crates/hipfire-runtime/src/llama.rs | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/crates/hipfire-runtime/src/llama.rs b/crates/hipfire-runtime/src/llama.rs
index 1251a95a8..b16edb8e9 100644
--- a/crates/hipfire-runtime/src/llama.rs
+++ b/crates/hipfire-runtime/src/llama.rs
@@ -1929,8 +1929,7 @@ pub fn is_batchable_la(dt: DType, arch: &str) -> bool {
     );
     // MQ4CG256 (qt45) remains gfx12-only until its gfx11 sibling lands —
     // intentionally not part of the shared rule, in both callers.
-    let mq4cg256_gfx12 =
-        matches!(dt, DType::MQ4CG256) && matches!(arch, "gfx1200" | "gfx1201");
+    let mq4cg256_gfx12 = matches!(dt, DType::MQ4CG256) && matches!(arch, "gfx1200" | "gfx1201");
     wmma_only || mq3_gfx10_scalar || mq4_v2 || mq4cg256_gfx12
 }
 

From 8d4450b18610a7c44a28f7974c0be47f60c11f37 Mon Sep 17 00:00:00 2001
From: Kaden Schutt <kaden@hipfire.dev>
Date: Thu, 3 Sep 2026 01:56:24 +0000
Subject: [PATCH 4/5] chore: refresh crate maps (check-crate-maps --write)

---
 crates/hipfire-arch-qwen35/map.md | 4 ++--
 crates/hipfire-runtime/map.md     | 6 +++---
 2 files changed, 5 insertions(+), 5 deletions(-)

diff --git a/crates/hipfire-arch-qwen35/map.md b/crates/hipfire-arch-qwen35/map.md
index 94a8f3efd..f27912845 100644
--- a/crates/hipfire-arch-qwen35/map.md
+++ b/crates/hipfire-arch-qwen35/map.md
@@ -42,7 +42,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 | [`src/qwen35/ep_batch.rs`](src/qwen35/ep_batch.rs) | 4,800 | 20 | 7 |
 | [`src/qwen35/forward.rs`](src/qwen35/forward.rs) | 6,255 | 31 | 12 |
 | [`src/qwen35/load.rs`](src/qwen35/load.rs) | 4,906 | 10 | 0 |
-| [`src/qwen35/prefill.rs`](src/qwen35/prefill.rs) | 9,312 | 11 | 48 |
+| [`src/qwen35/prefill.rs`](src/qwen35/prefill.rs) | 9,325 | 11 | 49 |
 | [`src/qwen35/weights.rs`](src/qwen35/weights.rs) | 1,971 | 43 | 10 |
 | [`src/qwen35.rs`](src/qwen35.rs) | 63 | 7 | 0 |
 | [`src/scheduler.rs`](src/scheduler.rs) | 142 | 3 | 4 |
@@ -97,6 +97,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 
 ### Totals
 
-- 29 modules · 57,281 lines · 436 public items · 189 tests · 4 examples
+- 29 modules · 57,294 lines · 436 public items · 190 tests · 4 examples
 
 <!-- crate-map:generated:end -->
diff --git a/crates/hipfire-runtime/map.md b/crates/hipfire-runtime/map.md
index 1995fd991..1c8b741fe 100644
--- a/crates/hipfire-runtime/map.md
+++ b/crates/hipfire-runtime/map.md
@@ -53,7 +53,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 | [`src/kv_backend.rs`](src/kv_backend.rs) | 129 | 1 | 7 |
 | [`src/kv_mode.rs`](src/kv_mode.rs) | 298 | 10 | 7 |
 | [`src/lib.rs`](src/lib.rs) | 80 | 55 | 0 |
-| [`src/llama.rs`](src/llama.rs) | 8,738 | 83 | 42 |
+| [`src/llama.rs`](src/llama.rs) | 8,793 | 85 | 42 |
 | [`src/llama_spec.rs`](src/llama_spec.rs) | 617 | 6 | 1 |
 | [`src/loader_api.rs`](src/loader_api.rs) | 256 | 10 | 4 |
 | [`src/loop_guard.rs`](src/loop_guard.rs) | 194 | 8 | 4 |
@@ -114,7 +114,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 - [`src/kv_backend.rs`](src/kv_backend.rs): `saddle_core`
 - [`src/kv_mode.rs`](src/kv_mode.rs): `saddle_core`, `KvModePolicy`, `ResolveResult`, `QWEN35_HFQ_POLICY`, `QWEN35_PARO_POLICY`, `DIR_SAFETENSORS_POLICY`, `LLAMA_HFQ_POLICY`, `HFQ_Q8_ONLY_POLICY`, `QWEN35_PP_POLICY`, `resolve`
 - [`src/lib.rs`](src/lib.rs): `admission`, `arch`, `arch_mapping`, `arch_model`, `arch_spec`, `augmentor`, `bf16_loader`, `cache_plan`, `cask`, `config`, `cpu_router`, `ddtree`, +43 more
-- [`src/llama.rs`](src/llama.rs): `ModelArch`, `LlamaConfig`, `from_gguf`, `dequantize_q4_0`, `dequantize_q8_0`, `f16_to_f32`, `f32_to_f16`, `dequantize_q4_k`, `convert_q4k_to_q4f16_g64`, `convert_q4k_to_q4f16_g32`, `dequantize_q6_k`, `ParoRotation`, +71 more
+- [`src/llama.rs`](src/llama.rs): `ModelArch`, `LlamaConfig`, `from_gguf`, `dequantize_q4_0`, `dequantize_q8_0`, `f16_to_f32`, `f32_to_f16`, `dequantize_q4_k`, `convert_q4k_to_q4f16_g64`, `convert_q4k_to_q4f16_g32`, `dequantize_q6_k`, `ParoRotation`, +73 more
 - [`src/llama_spec.rs`](src/llama_spec.rs): `verify_block_argmax`, `verify_block_logits`, `verify_block_argmax_capture_gpu`, `verify_block_sampled_capture_gpu`, `verify_tree_logits`, `lm_head_logits_n_rows`
 - [`src/loader_api.rs`](src/loader_api.rs): `ModelSource`, `from_path`, `arch_id`, `is_dir`, `describe`, `LoadCtx`, `SpecLoadCfg`, `CaskConfig`, `physical_cap`, `physical_cap_with_override`
 - [`src/loop_guard.rs`](src/loop_guard.rs): `StopReason`, `LoopGuard`, `from_config`, `new`, `off`, `enabled`, `check`, `window_len`
@@ -156,6 +156,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 
 ### Totals
 
-- 58 modules · 53,079 lines · 893 public items · 617 tests · 132 examples
+- 58 modules · 53,134 lines · 895 public items · 617 tests · 132 examples
 
 <!-- crate-map:generated:end -->

From 4e360fb2a8308e7522c3ea4e28b251b5b35da253 Mon Sep 17 00:00:00 2001
From: Kaden Schutt <kaden@hipfire.dev>
Date: Fri, 4 Sep 2026 17:24:32 +0000
Subject: [PATCH 5/5] fix(runtime): llama never admits MQ-V2 to batched
 prefill; forward_prefill_chunk has no V2 arms
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

hw-gate Fable seat on #690 (run 33895641944), source trace verified: the shared MQ-V2 admit rule made llama::is_batchable_la admit plain Llama/Qwen3-dense qt44/47-50 artifacts to WMMA prefill on gfx11 and gfx12, but llama.rs::forward_prefill_chunk's per-layer matchers (qkv_is_mq ~:2570, wo_is_mq ~:3025, ffn_is_mq ~:3117, w_down_is_mq ~:3248) list only MQ4G256|MQ6G256|MQ3G256|MFP4G32 — an admitted V2 model skips the FWHT rotate and runs the V1 hfq4g256 launchers on V2 blobs: silently incoherent prefill. master's pre-existing mq4_v2_gfx12 arm had the same hole on gfx12; no gfx12 Llama-V2 artifact has ever tripped it.

llama::is_batchable_la now refuses every *G256V2 dtype and MQ4CG256 on every arch, with the reason at the site. qwen35::is_batchable_la keeps the shared mqv2_wmma_batchable rule (its chunk path has the V2 arms; gfx11 kill-switch intact). llama_spec::batched_verify_eligible routes all seven weights through is_batchable_la, so it is covered without an edit. Lockstep test now asserts the true contract: agreement on every non-V2 dtype across 5 arches; for V2, qwen35 admits on gfx11/gfx12 and llama refuses everywhere. Spec §9 row and crate maps corrected.

hipfire-runtime is_batchable_la: 7 passed; qwen35 is_batchable + lockstep: 9 passed.
---
 crates/hipfire-arch-qwen35/map.md             |   4 +-
 .../hipfire-arch-qwen35/src/qwen35/prefill.rs |  90 ++++++++++++---
 crates/hipfire-runtime/map.md                 |   4 +-
 crates/hipfire-runtime/src/llama.rs           | 109 ++++++++----------
 docs/quant-formats/mq4-v2.md                  |   2 +-
 5 files changed, 129 insertions(+), 80 deletions(-)

diff --git a/crates/hipfire-arch-qwen35/map.md b/crates/hipfire-arch-qwen35/map.md
index f27912845..ea268919d 100644
--- a/crates/hipfire-arch-qwen35/map.md
+++ b/crates/hipfire-arch-qwen35/map.md
@@ -42,7 +42,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 | [`src/qwen35/ep_batch.rs`](src/qwen35/ep_batch.rs) | 4,800 | 20 | 7 |
 | [`src/qwen35/forward.rs`](src/qwen35/forward.rs) | 6,255 | 31 | 12 |
 | [`src/qwen35/load.rs`](src/qwen35/load.rs) | 4,906 | 10 | 0 |
-| [`src/qwen35/prefill.rs`](src/qwen35/prefill.rs) | 9,325 | 11 | 49 |
+| [`src/qwen35/prefill.rs`](src/qwen35/prefill.rs) | 9,381 | 11 | 49 |
 | [`src/qwen35/weights.rs`](src/qwen35/weights.rs) | 1,971 | 43 | 10 |
 | [`src/qwen35.rs`](src/qwen35.rs) | 63 | 7 | 0 |
 | [`src/scheduler.rs`](src/scheduler.rs) | 142 | 3 | 4 |
@@ -97,6 +97,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 
 ### Totals
 
-- 29 modules · 57,294 lines · 436 public items · 190 tests · 4 examples
+- 29 modules · 57,350 lines · 436 public items · 190 tests · 4 examples
 
 <!-- crate-map:generated:end -->
diff --git a/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs b/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs
index 01489eac9..9c65fb3b1 100644
--- a/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs
+++ b/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs
@@ -1544,8 +1544,10 @@ pub(crate) fn is_batchable_la(dt: DType, arch: &str) -> bool {
     // behind HIPFIRE_MQV2_GFX11_WMMA != "0" — setting
     // HIPFIRE_MQV2_GFX11_WMMA=0 restores the per-token fallback ONLY on
     // gfx11, leaving gfx12 untouched. Delegates to the shared
-    // `hipfire_runtime::llama::mqv2_wmma_batchable` rule so this stays in
-    // lockstep with `llama::is_batchable_la` structurally. Lockstep with
+    // `hipfire_runtime::llama::mqv2_wmma_batchable` rule (shared home for
+    // the dtype/arch/kill-switch set). NOTE: `llama::is_batchable_la` does
+    // NOT delegate to it — the llama chunk path has no V2 arms, so llama
+    // refuses V2 everywhere; only this qwen35 caller admits V2. Lockstep with
     // the HasWmma predicate on GemmMq*G256V2* keys and with
     // gemm_mq*g256v2's has_wmma() guard.
     // MQ4CG256 (qt45) remains gfx12-only until its gfx11 sibling lands.
@@ -8222,15 +8224,39 @@ mod tests {
 
     #[test]
     fn mqv2_admit_llama_qwen35_lockstep() {
-        // Audit 2026-09-02 Broken 1: `llama::is_batchable_la` admitted MQ-V2
-        // only on gfx12 while this module admitted gfx11+gfx12, despite both
-        // doc-comments claiming an exact match. Both now delegate to the
-        // shared `llama::mqv2_wmma_batchable` rule; this test iterates the
-        // MQ-V2 dtypes over gfx11, gfx12, and pre-WMMA arches and asserts the
-        // two gates agree. Both read `HIPFIRE_MQV2_GFX11_WMMA` from the
-        // environment identically, so equality holds in any env state
-        // without mutating globals.
-        let dts = [
+        // True contract (PR #690 hw-gate regression): llama and qwen35 agree
+        // on every NON-V2 dtype, but for the V2 family they deliberately
+        // diverge — qwen35's `forward_prefill_chunk` has V2 dispatch arms
+        // (206 hits) so it admits V2 via the shared
+        // `llama::mqv2_wmma_batchable` rule, while llama's chunk path has no
+        // V2 arms (`qkv_is_mq`/`wo_is_mq`/`ffn_is_mq`/`w_down_is_mq` list
+        // only V1 dtypes) so `llama::is_batchable_la` refuses V2 everywhere
+        // and stays on per-token decode. Admitting V2 to the llama path
+        // would skip the FWHT rotate and run V1 `hfq4g256` launchers on V2
+        // blobs — silently incoherent prefill.
+        // Non-V2 agreement across the 5-arch sample.
+        let non_v2 = [
+            DType::MQ4G256,
+            DType::HFQ4G256,
+            DType::MQ6G256,
+            DType::MQ3G256,
+            DType::MFP4G32,
+            DType::Q8_0,
+        ];
+        for dt in non_v2 {
+            for arch in ["gfx1100", "gfx1151", "gfx1201", "gfx1030", "gfx1010"] {
+                assert_eq!(
+                    llama::is_batchable_la(dt, arch),
+                    is_batchable_la(dt, arch),
+                    "lockstep drift for {dt:?} on {arch}"
+                );
+            }
+        }
+        // V2 divergence: qwen35 admits on gfx11/gfx12 (kill-switch at its
+        // default ON here — both gates read `HIPFIRE_MQV2_GFX11_WMMA`
+        // identically, so with the var unset gfx11 admits), refuses
+        // pre-WMMA; llama refuses on all 5 arches.
+        let v2 = [
             DType::MQ4G256V2,
             DType::MQ6G256V2,
             DType::MQ5G256V2,
@@ -8238,12 +8264,41 @@ mod tests {
             DType::MQ2G256V2,
             DType::MQ4CG256,
         ];
-        for dt in dts {
-            for arch in ["gfx1100", "gfx1151", "gfx1201", "gfx1030", "gfx1010"] {
-                assert_eq!(
-                    llama::is_batchable_la(dt, arch),
-                    is_batchable_la(dt, arch),
-                    "lockstep drift for {dt:?} on {arch}"
+        for dt in v2 {
+            for arch in ["gfx1100", "gfx1151"] {
+                // MQ4CG256 is gfx12-only by intent in BOTH callers.
+                if dt == DType::MQ4CG256 {
+                    assert!(
+                        !is_batchable_la(dt, arch),
+                        "qwen35 must refuse {dt:?} on {arch}"
+                    );
+                } else {
+                    assert!(
+                        is_batchable_la(dt, arch),
+                        "qwen35 should admit {dt:?} on {arch}"
+                    );
+                }
+                assert!(
+                    !llama::is_batchable_la(dt, arch),
+                    "llama must refuse {dt:?} on {arch}"
+                );
+            }
+            assert!(
+                is_batchable_la(dt, "gfx1201"),
+                "qwen35 should admit {dt:?} on gfx1201"
+            );
+            assert!(
+                !llama::is_batchable_la(dt, "gfx1201"),
+                "llama must refuse {dt:?} on gfx1201"
+            );
+            for arch in ["gfx1030", "gfx1010"] {
+                assert!(
+                    !is_batchable_la(dt, arch),
+                    "qwen35 must refuse {dt:?} on {arch}"
+                );
+                assert!(
+                    !llama::is_batchable_la(dt, arch),
+                    "llama must refuse {dt:?} on {arch}"
                 );
             }
         }
@@ -8252,6 +8307,7 @@ mod tests {
         assert!(is_batchable_la(DType::MQ4G256V2, "gfx1201"));
         assert!(!is_batchable_la(DType::MQ4G256V2, "gfx1030"));
         assert!(!is_batchable_la(DType::MQ4CG256, "gfx1100"));
+        assert!(!llama::is_batchable_la(DType::MQ4G256V2, "gfx1201"));
     }
 
     #[test]
diff --git a/crates/hipfire-runtime/map.md b/crates/hipfire-runtime/map.md
index 1c8b741fe..e7ad6225b 100644
--- a/crates/hipfire-runtime/map.md
+++ b/crates/hipfire-runtime/map.md
@@ -53,7 +53,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 | [`src/kv_backend.rs`](src/kv_backend.rs) | 129 | 1 | 7 |
 | [`src/kv_mode.rs`](src/kv_mode.rs) | 298 | 10 | 7 |
 | [`src/lib.rs`](src/lib.rs) | 80 | 55 | 0 |
-| [`src/llama.rs`](src/llama.rs) | 8,793 | 85 | 42 |
+| [`src/llama.rs`](src/llama.rs) | 8,786 | 85 | 42 |
 | [`src/llama_spec.rs`](src/llama_spec.rs) | 617 | 6 | 1 |
 | [`src/loader_api.rs`](src/loader_api.rs) | 256 | 10 | 4 |
 | [`src/loop_guard.rs`](src/loop_guard.rs) | 194 | 8 | 4 |
@@ -156,6 +156,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside
 
 ### Totals
 
-- 58 modules · 53,134 lines · 895 public items · 617 tests · 132 examples
+- 58 modules · 53,127 lines · 895 public items · 617 tests · 132 examples
 
 <!-- crate-map:generated:end -->
diff --git a/crates/hipfire-runtime/src/llama.rs b/crates/hipfire-runtime/src/llama.rs
index b16edb8e9..ece87326f 100644
--- a/crates/hipfire-runtime/src/llama.rs
+++ b/crates/hipfire-runtime/src/llama.rs
@@ -1841,9 +1841,9 @@ pub const PREFILL_MAX_BATCH: usize = 256;
 /// Kill-switch for the MQ-V2 (qt44 + neutral qt47-50) gfx11 WMMA prefill path:
 /// gfx12 (`gfx1200`/`gfx1201`) is always admitted, gfx11
 /// (`gfx1100`/`gfx1101`/`gfx1102`/`gfx1150`/`gfx1151`) is admitted unless
-/// `HIPFIRE_MQV2_GFX11_WMMA=0`, anything else is rejected. Single definition
-/// shared by `llama::is_batchable_la` and `qwen35::is_batchable_la` so the two
-/// stay in lockstep structurally instead of by matching comments.
+/// `HIPFIRE_MQV2_GFX11_WMMA=0`, anything else is rejected. Defined here as
+/// the shared home for the qwen35 caller (`qwen35::is_batchable_la`);
+/// `llama::is_batchable_la` does NOT delegate to it (see below).
 /// `value` is the raw env var (None = unset → default ON); only `Some("0")`
 /// disables the gfx11 path. Gfx12 is unaffected by the env var.
 pub fn mqv2_gfx11_wmma_enabled_from_env(value: Option<&str>, arch: &str) -> bool {
@@ -1862,9 +1862,11 @@ pub fn mqv2_gfx11_wmma_enabled_from_env(value: Option<&str>, arch: &str) -> bool
 
 /// Admit rule for the MQ-V2 family (`MQ4G256V2` + neutral `MQ6/5/3/2G256V2`)
 /// in batched WMMA prefill: dtype set × arch set × the
-/// `HIPFIRE_MQV2_GFX11_WMMA` kill-switch in one function. Both
-/// `llama::is_batchable_la` and `qwen35::is_batchable_la` delegate here;
-/// `MQ4CG256` (qt45) stays gfx12-only in both callers and is intentionally
+/// `HIPFIRE_MQV2_GFX11_WMMA` kill-switch in one function. Only
+/// `qwen35::is_batchable_la` delegates here — qwen35's
+/// `forward_prefill_chunk` has V2 dispatch arms, while the llama chunk path
+/// does not (see `llama::is_batchable_la`).
+/// `MQ4CG256` (qt45) stays gfx12-only in its caller and is intentionally
 /// NOT part of this rule.
 pub fn mqv2_wmma_batchable(dt: DType, mqv2_gfx11_wmma: Option<&str>, arch: &str) -> bool {
     matches!(
@@ -1878,9 +1880,8 @@ pub fn mqv2_wmma_batchable(dt: DType, mqv2_gfx11_wmma: Option<&str>, arch: &str)
 }
 
 /// Is this dtype/arch combination eligible for the batched WMMA prefill
-/// kernels? Shares the MQ-V2 admit rule with `qwen35::is_batchable_la` via
-/// `mqv2_wmma_batchable`, so plain Qwen3 and hybrid Qwen3.5 stay in lockstep
-/// structurally when new dtypes or arches gain WMMA support.
+/// kernels? NOTE: unlike `qwen35::is_batchable_la`, this does NOT admit the
+/// MQ-V2 family — see the `never_v2` refusal below.
 pub fn is_batchable_la(dt: DType, arch: &str) -> bool {
     let always_ok = matches!(
         dt,
@@ -1913,24 +1914,30 @@ pub fn is_batchable_la(dt: DType, arch: &str) -> bool {
             arch,
             "gfx1010" | "gfx1011" | "gfx1012" | "gfx1013" | "gfx1030" | "gfx1031" | "gfx1032"
         );
-    // MQ-V2 family (qt44 + neutral qt47-50) batched prefill + batched lm_head
-    // GEMM: gfx11 + gfx12 via the shared rule, with the gfx11 half behind the
-    // `HIPFIRE_MQV2_GFX11_WMMA=0` kill-switch. Delegates to
-    // `mqv2_wmma_batchable` so this stays in lockstep with
-    // `qwen35::is_batchable_la` structurally, not by matching comments.
-    // Outside the admitted arches, fall back to per-token decode rather than
-    // dispatching a foreign-arch WMMA kernel.
-    let mq4_v2 = mqv2_wmma_batchable(
+    // MQ-V2 family (`MQ4G256V2` + neutral `MQ6/5/3/2G256V2`, qt44/qt47-50)
+    // plus `MQ4CG256` (qt45): REFUSED on every arch. `forward_prefill_chunk`
+    // has no V2 arms — its per-layer dtype matchers (`qkv_is_mq` ~:2570,
+    // `wo_is_mq` ~:3025, `ffn_is_mq` ~:3117, `w_down_is_mq` ~:3248) list only
+    // `MQ4G256|MQ6G256|MQ3G256|MFP4G32`, so an admitted V2 model would skip
+    // the FWHT rotate and run the V1 `hfq4g256` launchers
+    // (`gemm_qkv_hfq4g256`, `gemm_hfq4g256_residual`, `gemm_gate_up_hfq4g256`)
+    // on V2 blobs — silently incoherent prefill. Per-token decode is the
+    // only correct llama path for V2 until those arms exist. qwen35's chunk
+    // path DOES have the V2 arms, so `qwen35::is_batchable_la` keeps
+    // admitting V2 via the shared `mqv2_wmma_batchable` rule above.
+    let never_v2 = matches!(
         dt,
-        hipfire_config::developer_var("HIPFIRE_MQV2_GFX11_WMMA")
-            .ok()
-            .as_deref(),
-        arch,
+        DType::MQ4G256V2
+            | DType::MQ6G256V2
+            | DType::MQ5G256V2
+            | DType::MQ3G256V2
+            | DType::MQ2G256V2
+            | DType::MQ4CG256
     );
-    // MQ4CG256 (qt45) remains gfx12-only until its gfx11 sibling lands —
-    // intentionally not part of the shared rule, in both callers.
-    let mq4cg256_gfx12 = matches!(dt, DType::MQ4CG256) && matches!(arch, "gfx1200" | "gfx1201");
-    wmma_only || mq3_gfx10_scalar || mq4_v2 || mq4cg256_gfx12
+    if never_v2 {
+        return false;
+    }
+    wmma_only || mq3_gfx10_scalar
 }
 
 /// Per-call scratch for `forward_prefill_batch`. Holds [N × ...] working
@@ -8289,32 +8296,22 @@ mod tests {
     }
 
     #[test]
-    fn is_batchable_la_mq4_v2_gfx11_and_gfx12() {
-        // MQ4G256V2 batched prefill admits gfx11 + gfx12 through the shared
-        // `mqv2_wmma_batchable` rule (gfx11 behind HIPFIRE_MQV2_GFX11_WMMA);
-        // MQ4CG256 stays gfx12-only. Lockstep with
-        // `qwen35::is_batchable_la` by construction.
+    fn is_batchable_la_mq4_v2_refused_everywhere() {
+        // `forward_prefill_chunk` has no V2 arms (its `qkv_is_mq` ~:2570,
+        // `wo_is_mq` ~:3025, `ffn_is_mq` ~:3117, `w_down_is_mq` ~:3248
+        // matchers list only V1 dtypes), so llama must refuse MQ4G256V2 and
+        // MQ4CG256 on EVERY arch — including gfx11/gfx12 — and stay on
+        // per-token decode. qwen35's chunk path has the arms and keeps the
+        // shared `mqv2_wmma_batchable` rule; see
+        // `qwen35_is_batchable_la_mq4_v2_gfx11_and_gfx12` for the admit side.
         for arch in [
             "gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151", "gfx1200", "gfx1201",
+            "gfx1010", "gfx1030", "gfx942",
         ] {
-            assert!(
-                is_batchable_la(DType::MQ4G256V2, arch),
-                "MQ4G256V2 should batch on {arch}"
-            );
-        }
-        for arch in ["gfx1200", "gfx1201"] {
-            assert!(
-                is_batchable_la(DType::MQ4CG256, arch),
-                "MQ4CG256 should batch on {arch}"
-            );
-        }
-        for arch in ["gfx1010", "gfx1030", "gfx942"] {
             assert!(
                 !is_batchable_la(DType::MQ4G256V2, arch),
                 "MQ4G256V2 must fall back on {arch}"
             );
-        }
-        for arch in ["gfx1010", "gfx1100", "gfx1151", "gfx942"] {
             assert!(
                 !is_batchable_la(DType::MQ4CG256, arch),
                 "MQ4CG256 must fall back on {arch}"
@@ -8323,23 +8320,19 @@ mod tests {
     }
 
     #[test]
-    fn is_batchable_la_v2_family_gfx11_and_gfx12() {
-        // Neutral V2 family (qt47-50) admits gfx11 + gfx12 through the shared
-        // `mqv2_wmma_batchable` rule, mirroring
-        // `qwen35_is_batchable_la_v2_family_gfx11_and_gfx12`.
+    fn is_batchable_la_v2_family_refused_everywhere() {
+        // Neutral V2 family (qt47-50): same refusal as MQ4G256V2 — no V2 arms
+        // in the llama chunk path, so refuse on every arch (including
+        // gfx11/gfx12). Mirrors `qwen35_is_batchable_la_v2_family_gfx11_and_gfx12`
+        // on the admit side.
         for arch in [
             "gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151", "gfx1200", "gfx1201",
+            "gfx1010", "gfx1030", "gfx942",
         ] {
-            assert!(is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 on {arch}");
-            assert!(is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 on {arch}");
-            assert!(is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 on {arch}");
-            assert!(is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 on {arch}");
-        }
-        for arch in ["gfx1010", "gfx1030", "gfx942"] {
-            assert!(!is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 fallback");
-            assert!(!is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 fallback");
-            assert!(!is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 fallback");
-            assert!(!is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 fallback");
+            assert!(!is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 fallback on {arch}");
+            assert!(!is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 fallback on {arch}");
+            assert!(!is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 fallback on {arch}");
+            assert!(!is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 fallback on {arch}");
         }
         assert_ne!(DType::MQ6G256, DType::MQ6G256V2);
         assert_ne!(DType::MQ3G256, DType::MQ3G256V2);
diff --git a/docs/quant-formats/mq4-v2.md b/docs/quant-formats/mq4-v2.md
index 9374d3e0b..3943cc21e 100644
--- a/docs/quant-formats/mq4-v2.md
+++ b/docs/quant-formats/mq4-v2.md
@@ -421,7 +421,7 @@ relatively less of the damage. "Codebooks are for the sub-4-bit tier" is defensi
 | **qt=45** batched prefill WMMA GEMM + batched lm_head GEMM | **gfx12-only** (`HasWmmaGfx12`); **not** admitted on gfx11 |
 | Unsupported batched prefill | **per-token decode fallback** (does not dispatch a foreign-arch WMMA kernel) |
 | FusedQkv / FusedGateUp decode registrations | remain **cross-arch** (not narrowed to gfx12) |
-| Prefill LA admission (qt=44) | gfx1100/1101/1102/1150/1151 + gfx1200/1201; gfx11 opt-out `HIPFIRE_MQV2_GFX11_WMMA=0` |
+| Prefill LA admission (qt=44) | **qwen35 only**: gfx1100/1101/1102/1150/1151 + gfx1200/1201; gfx11 opt-out `HIPFIRE_MQV2_GFX11_WMMA=0`. Plain-Llama dense (`llama::is_batchable_la`) **refuses V2 on every arch** (per-token decode) until `forward_prefill_chunk` grows V2 arms |
 | Exact parity examples | `mq4v2_parity`, `mq4v2_gemm_parity`, `mq4v2_fused_parity`, `mq4v2_residual_parity`, `mq4c_parity`; BT screens `test_mq4v2_*_bt_gfx{1100,1151,1201}.rs` |
 | Qwen3.8 fixture-bound KLD | qt=44 `ctl` WT2 0.039033 / v6 0.544517; `ctl2` WT2 0.032495; `attn` WT2 0.025437 (§ 5) |
 | gfx1010 | scalar fused decode TUs compile; batched prefill falls back as above |