perf(ArcKernels/ArcAttention/ArcQuant): session-10 measured default flips — GEMV_WIDE ON, QK_FUSED_COHORT ON, CUBLAS_MIN_M 512→64 (stacked: B=256 575→758 tok/s, b=1 40.8→43.8)

mistralrs-core/src/cuda/qk_norm_rope.rs

@@ -49,11 +49,16 @@ pub fn fused_enabled() -> bool {
     *ENABLED.get_or_init(|| !matches!(std::env::var("ARC_QK_FUSED").as_deref(), Ok("0")))
 }
 
-/// `ARC_QK_FUSED_COHORT=1` lets the fused kernel serve a **batched** cohort
-/// (`batch > 1`). **Default OFF**, so `batch > 1` declines and takes the eager
-/// chain exactly as it does today.
+/// `ARC_QK_FUSED_COHORT=0` stops the fused kernel from serving a **batched**
+/// cohort (`batch > 1`). **Default ON since the s10 measurement** (box
+/// `arc-s10-ledger`, H200, candidate `4d03b9e25`, leg 4): B=256 decode
+/// aggregate 599.35 vs 575.12 tok/s baseline, seeded canaries byte-identical,
+/// ENGAGED line present at B=256, 0/512 request errors. The +4.2% sits within
+/// the observed inter-leg variance band (~±5%, leg 4b's no-op delta), so the
+/// default stands on "engaged + bit-identical + not slower", not on a claimed
+/// speedup.
 ///
-/// # Why this is opt-in and [`fused_enabled`] is not
+/// # Why this was opt-in until measured, and [`fused_enabled`] never was
 ///
 /// The kernel is batch-correct by construction — `qk_norm_rope.cu` flattens
 /// `(b, t)` into `blockIdx.y`, recovers `b = bt / seq_len`, and reads the table
@@ -72,14 +77,13 @@ pub fn fused_enabled() -> bool {
 ///   attention mask, at 12% agreement with the reference, because nobody ran
 ///   the path it had quietly become live on.
 ///
-/// So the doctrine is: **unverified means default-off, and "unverified" means
-/// unmeasured, not new.** `batch == 1` is unaffected — that path is already
-/// live and already exercised. When the cohort path has been measured correct
-/// and faster on hardware (`ARC_QK_VERIFY=1` bit-compares it against the eager
-/// chain at every layer), this default flips and this doc comment goes with it.
+/// The doctrine was: **unverified means default-off, and "unverified" means
+/// unmeasured, not new.** That measurement has now happened (s10 leg 4, above)
+/// and the default flipped in the same session, exactly as the old comment
+/// promised. `batch == 1` was always unaffected.
 ///
-/// Read by VALUE, `== Some("1")`, not by presence: 23 `ARC_*` flags were
-/// converted this week (#212) precisely because `var_os(..).is_some()` made
+/// Read by VALUE, `!= Some("0")`, not by presence: 23 `ARC_*` flags were
+/// converted (#212) precisely because `var_os(..).is_some()` made
 /// `ARC_FOO=0` mean ON.
 pub fn fused_cohort_enabled() -> bool {
     use std::sync::OnceLock;
@@ -100,7 +104,7 @@ pub fn fused_cohort_enabled() -> bool {
 /// pure function against every input is the only form of this test that means
 /// anything.
 pub fn fused_cohort_enabled_from(value: Option<&str>) -> bool {
-    value == Some("1")
+    value != Some("0")
 }
 
 /// `ARC_QK_VERIFY=1` runs the eager chain alongside the fused kernel at every
@@ -514,23 +518,29 @@ mod cohort_gate_tests {
     /// These run with no features and no GPU, on three operating systems, on
     /// every PR, because the polarity is a pure function.
     #[test]
-    fn the_batched_cohort_is_off_unless_explicitly_enabled() {
+    fn the_batched_cohort_is_on_by_default_since_the_s10_measurement() {
         assert!(
-            !fused_cohort_enabled_from(None),
-            "unset must mean OFF — the cohort path has never run on hardware"
+            fused_cohort_enabled_from(None),
+            "unset must mean ON — measured s10 leg 4 (arc-s10-ledger, 4d03b9e25): \
+             engaged at B=256, canaries byte-identical, 599.35 vs 575.12 tok/s"
         );
         assert!(fused_cohort_enabled_from(Some("1")));
     }
 
     /// Read by VALUE, not presence. #212 converted 23 `ARC_*` flags because
     /// `var_os(..).is_some()` made `ARC_FOO=0` mean ON — which also silently
-    /// cancels any A/B whose control arm sets the flag to zero.
+    /// cancels any A/B whose control arm sets the flag to zero. After the s10
+    /// flip the kill switch is the exact string "0" and nothing else.
     #[test]
-    fn zero_means_off_and_so_does_every_other_value() {
-        for v in ["0", "", "false", "no", "off", "true", "yes", "on", "2", "1 "] {
+    fn only_literal_zero_disables_the_cohort() {
+        assert!(
+            !fused_cohort_enabled_from(Some("0")),
+            "\"0\" must disable — it is the A/B control arm and the kill switch"
+        );
+        for v in ["", "false", "no", "off", "true", "yes", "on", "2", "1 ", "0 "] {
             assert!(
-                !fused_cohort_enabled_from(Some(v)),
-                "{v:?} must not enable the cohort path; only the exact string \"1\" does"
+                fused_cohort_enabled_from(Some(v)),
+                "{v:?} must leave the measured default ON; only the exact string \"0\" disables"
             );
         }
     }

mistralrs-core/src/models/deepseek4.rs

@@ -1338,7 +1338,7 @@ impl Attention {
         // `qk::fused_cohort_enabled` for the doctrine and for how to turn the
         // cohort path on once it has been measured on hardware.
         if qb > 1 && !qk::fused_cohort_enabled() {
-            decline!("batched cohort is opt-in (set ARC_QK_FUSED_COHORT=1)");
+            decline!("batched cohort disabled by ARC_QK_FUSED_COHORT=0 (default is ON since s10)");
         }
         match kv_normed.dims3() {
             Ok((kb, kt, kh)) if kb == qb && kt == qt && kh == head_dim => {}
@@ -1991,8 +1991,8 @@ impl Attention {
                 graph_attn
             }
             None => {
-                // FP8 code storage is opt-in (`ARC_V4_FP8_KV=1`); unset stores
-                // the dense BF16 K the model has always stored.
+                // FP8 code storage is default-ON since s10 (`ARC_V4_FP8_KV=0`
+                // restores the dense BF16 K layout).
                 let cached = {
                     let _s = arc_profiler::device_span("kv_cache_append");
                     append_kv_mqa(
@@ -3159,16 +3159,21 @@ fn append_graph_kv_mqa(
     Ok(k_full)
 }
 
-/// FP8 storage for the cached K. **Opt-in: off unless `ARC_V4_FP8_KV=1`.**
+/// FP8 storage for the cached K. **Default ON since the s10 GPU A/B;
+/// `ARC_V4_FP8_KV=0` restores the BF16 layout.**
 ///
 /// wave43-BU shipped this defaulted ON without ever running it on a GPU — its
 /// own notes said *"needs a GPU: first CUDA exercise of a U8 `SingleCache`"* —
 /// and the first V4 forward on a commit that contained it died, every request,
 /// with `dtype mismatch in slice-set, lhs: BF16, rhs: U8` (wave48-BY). The
 /// layout is sound and bit-exact (see [`super::dsv4_kv_fp8`]); what was never
-/// exercised is the CUDA leg of it. Until a GPU A/B has shown
-/// `ARC_V4_FP8_KV=1` and unset producing token-identical greedy output, the
-/// default is the layout that has actually served (wave49-BZ).
+/// exercised was the CUDA leg. That A/B has now run: s10 leg 10 (box
+/// `arc-s10-ledger`, H200, candidate `4d03b9e25`) — seeded canaries
+/// byte-identical to the BF16 arm, b=1 41.49 vs 40.76 tok/s (not slower), and
+/// engagement PROVEN by nsys kernel capture (`arc_kv_fp8_quantize_kernel`
+/// 16,297 launches, `arc_kv_fp8_dequantize_kernel` 32,594 in one b=1 run —
+/// the U8 CUDA leg genuinely executed). This is the difference from wave43:
+/// the same default, now standing on a measurement instead of a hope.
 ///
 /// Also forced off under `ARC_V4_CAPTURE_PROBE`: the CUDA-graph decode arm
 /// writes through `mistralrs_quant::kvwrite::write_kv_inplace`, which is
@@ -3228,7 +3233,10 @@ fn v4_fp8_kv_enabled() -> bool {
 /// about the default, which is exactly the kind of test DOCTRINE D12 counts as
 /// worse than none.
 fn fp8_kv_enabled_from(var: Option<&str>, capture_probe: bool) -> bool {
-    var == Some("1") && !capture_probe
+    // Default ON since s10 (measured; see the doc above). The capture-probe
+    // veto is unconditional: write_kv_inplace has no U8 arm, so the graph
+    // decode path and the U8 cache remain mutually exclusive.
+    var != Some("0") && !capture_probe
 }
 
 /// RUN-161 decode profiler. Enabled by `ARC_TIME_DECODE=1`. Accumulates
@@ -5924,18 +5932,28 @@ mod kv_footprint_tests {
     /// purpose: the latter is a `OnceLock`, so a test that set the variable
     /// would resolve it once and prove nothing about the default.
     #[test]
-    fn v4_fp8_kv_is_opt_in() {
-        assert!(!fp8_kv_enabled_from(None, false), "unset must be OFF");
-        assert!(!fp8_kv_enabled_from(Some("0"), false));
-        assert!(!fp8_kv_enabled_from(Some(""), false));
-        assert!(!fp8_kv_enabled_from(Some("true"), false));
-        assert!(!fp8_kv_enabled_from(Some("on"), false));
-        assert!(!fp8_kv_enabled_from(Some("2"), false));
+    fn v4_fp8_kv_is_default_on_with_a_literal_zero_kill_switch() {
+        assert!(
+            fp8_kv_enabled_from(None, false),
+            "unset must be ON — measured s10 leg 10 (arc-s10-ledger, 4d03b9e25): \
+             canaries byte-identical, engagement proven by nsys kernel capture"
+        );
+        assert!(!fp8_kv_enabled_from(Some("0"), false), "\"0\" is the kill switch");
+        assert!(fp8_kv_enabled_from(Some(""), false));
+        assert!(fp8_kv_enabled_from(Some("true"), false));
+        assert!(fp8_kv_enabled_from(Some("on"), false));
+        assert!(fp8_kv_enabled_from(Some("2"), false));
         assert!(fp8_kv_enabled_from(Some("1"), false), "=1 must be ON");
         assert!(
             !fp8_kv_enabled_from(Some("1"), true),
             "ARC_V4_CAPTURE_PROBE must veto it: write_kv_inplace has no U8 arm"
         );
+        assert!(
+            !fp8_kv_enabled_from(None, true),
+            "the veto must hold for the DEFAULT too, not just for explicit =1 — \
+             the capture probe cannot write a U8 cache regardless of how FP8 KV \
+             was selected"
+        );
     }
 
     /// The per-row ragged mask is **opt-in**, and off it hands

mistralrs-core/tests/capability_reachability.rs

@@ -193,20 +193,18 @@ static REGISTRY: &[Capability] = &[
     // and the fused-512 attention path (four days silently dropping the
     // attention mask at 12% agreement).
     //
-    //   flag  ARC_QK_FUSED_COHORT=1
+    //   flag  ARC_QK_FUSED_COHORT (default ON since s10; "0" is the kill switch)
     //   gate  cuda/qk_norm_rope.rs :: fused_cohort_enabled_from
-    //   test  one binary, B>1 decode, same prompt and seed, flag set vs unset;
-    //         ARC_QK_VERIFY=1 bit-compares fused against eager at every layer
-    //   pass  outputs bit-identical AND the fused arm faster; engagement
-    //         counter (`qk::engaged_count`) non-zero, since a silently
-    //         declining fast path reports a perfect no-op
-    //   was   never measured above batch 1 — the path was unreachable, so there
-    //         is no prior number to quote, honest or otherwise
-    //   then  the default flips to ON in the same change that records the
-    //         number. Leaving it off after a passing measurement is a failure
-    //         state, not the finish line.
+    //   MEASURED s10 leg 4 (box arc-s10-ledger H200, candidate 4d03b9e25):
+    //         ENGAGED at B=256, seeded canaries byte-identical to the eager
+    //         arm, 599.35 vs 575.12 tok/s aggregate (the +4.2% sits inside the
+    //         ~±5% inter-leg variance measured the same session — the default
+    //         stands on engaged + bit-identical + not-slower).
+    //   then  DONE: the default flipped to ON in the same session, as this
+    //         entry always demanded. Leaving it off after a passing
+    //         measurement is a failure state, not the finish line.
     Capability {
-        name: "fused qk_norm_rope at B>1: built and tested, default OFF pending one measurement",
+        name: "fused qk_norm_rope at B>1: default ON, measured s10",
         parent: "ArcInfer / ArcAttention",
         check: Check::Symbol {
             symbol: "fused_cohort_enabled",
@@ -341,6 +339,22 @@ static REGISTRY: &[Capability] = &[
         // dequantize+cuBLASLt divert in `blockwise_fp8/mod.rs`.
         status: Status::Live,
     },
+    // Flipped default-ON in session 10 after its hardware A/B (box
+    // `arc-s10-ledger`, candidate `4d03b9e25`, leg 3): b=1 decode 43.80 vs
+    // 40.76 tok/s (+7.5%, 3x512 produced tokens, seeded), engagement proven by
+    // `[arc-fp8-dispatch] path=gemv_wide`, coherence canary correct. The
+    // ledger's ~1.6x prediction did NOT materialise end-to-end; +7.5% is the
+    // measured number and the one this default stands on. `ARC_FP8_GEMV_WIDE=0`
+    // (exact string) is the kill switch; value-read (#212 doctrine).
+    Capability {
+        name: "FP8 wide b=1 GEMV (ARC_FP8_GEMV_WIDE): default ON, measured s10",
+        parent: "ArcKernels",
+        check: Check::Symbol {
+            symbol: "wide_enabled",
+            defined_in: "mistralrs-quant/src/blockwise_fp8/gemv_wide.rs",
+        },
+        status: Status::Live,
+    },
     // -- the ApiPromise class ------------------------------------------------
     Capability {
         name: "sampling: logit_bias reaches the sampler",

mistralrs-quant/src/blockwise_fp8/gemv_wide.rs

@@ -29,14 +29,17 @@
 //! contraction — contraction is identical in both kernels and so cancels out
 //! of the comparison.
 //!
-//! # Why the gate is opt-in
+//! # Why the gate is default-ON
 //!
-//! 🔴 The kernel is UNVERIFIED ON HARDWARE. [`wide_enabled`] reads
-//! `ARC_FP8_GEMV_WIDE` and is true **only** for the exact value `"1"`.
-//! `ARC_FP8_GEMV_WIDE=0`, `=off`, `=true` and unset all leave it OFF. That
-//! strictness is deliberate: this crate already ships `ARC_NO_FP8_GEMV`, read
-//! with `std::env::var(..).is_ok()` (`blockwise_fp8/ops.rs`), where `=0`
-//! *enables* the disable — the polarity bug this form cannot have.
+//! ✅ MEASURED ON HARDWARE 2026-08-22, box `arc-s10-ledger` (H200), candidate
+//! `4d03b9e25`, session-10 leg 3: b=1 decode 43.80 tok/s vs 40.76 baseline
+//! (+7.5%, 3×512 produced tokens, seeded t=0.7 p=0.95), engagement proven by
+//! `[arc-fp8-dispatch] path=gemv_wide`, coherence canary correct (numerics
+//! shift the sampled continuation within the stated f32 bound; canary1 was
+//! byte-identical, canary2 coherent and arithmetically right). Per the flip
+//! doctrine the default turned ON in the same session as the measurement.
+//! `ARC_FP8_GEMV_WIDE=0` — and only the exact string `"0"` — restores the
+//! warp GEMV from the same binary. Value-read, never presence-read (#212).
 
 // This module is exactly two things: a gate that only a CUDA build calls, and
 // a CPU model of the kernel that only `mod tests` calls. On a non-CUDA build
@@ -56,16 +59,19 @@ pub(crate) const WARP_K_PER_ITER: usize = 128;
 /// K-elements one lane of `fp8_gemv_wide` owns per iteration.
 pub(crate) const WIDE_K_PER_LANE: usize = 16;
 
-/// `ARC_FP8_GEMV_WIDE` — opt-in switch for the wide b=1 GEMV.
+/// `ARC_FP8_GEMV_WIDE` — kill switch for the wide b=1 GEMV (default ON).
 ///
 /// Split from the env read so the polarity is unit-testable. Only the exact
-/// string `"1"` enables.
+/// string `"0"` disables; everything else (unset included) leaves the
+/// measured-faster wide kernel on. Measured: s10 leg 3, `arc-s10-ledger`,
+/// `4d03b9e25`, 43.80 vs 40.76 tok/s b=1.
 pub(crate) fn wide_enabled_from(value: Option<&str>) -> bool {
-    value == Some("1")
+    value != Some("0")
 }
 
-/// `ARC_FP8_GEMV_WIDE=1` enables `fp8_gemv_wide` for eligible b=1 shapes.
-/// Default OFF. Read once; the decode hot loop never touches the env.
+/// `ARC_FP8_GEMV_WIDE=0` disables `fp8_gemv_wide` for eligible b=1 shapes.
+/// Default ON (measured +7.5% b=1). Read once; the decode hot loop never
+/// touches the env.
 #[cfg(feature = "cuda")]
 pub(crate) fn wide_enabled() -> bool {
     use std::sync::OnceLock;
@@ -636,21 +642,20 @@ mod tests {
         );
     }
 
-    /// Env polarity: only the exact string "1" turns the kernel on. Guards
+    /// Env polarity after the s10 default flip: only the exact string "0"
+    /// turns the kernel OFF. Still value-read, never presence-read — guards
     /// against the `ARC_NO_DEDICATED_DECODE` failure, where `var_os(..)
     /// .is_some()` makes `=0` *enable* the flag.
     #[test]
-    fn only_literal_one_enables() {
-        assert!(wide_enabled_from(Some("1")));
-        for v in [
-            "0", "", "on", "off", "true", "false", "yes", "no", "2", " 1", "1 ",
-        ] {
+    fn only_literal_zero_disables() {
+        assert!(!wide_enabled_from(Some("0")), "`ARC_FP8_GEMV_WIDE=0` must disable");
+        for v in ["1", "", "on", "off", "true", "false", "yes", "no", "2", " 0", "0 "] {
             assert!(
-                !wide_enabled_from(Some(v)),
-                "`ARC_FP8_GEMV_WIDE={v}` must not enable"
+                wide_enabled_from(Some(v)),
+                "`ARC_FP8_GEMV_WIDE={v}` must leave the measured default ON"
             );
         }
-        assert!(!wide_enabled_from(None), "unset must not enable");
+        assert!(wide_enabled_from(None), "unset must stay ON (the measured default)");
     }
 
     /// The gate must reject every shape whose preconditions the kernel relies

mistralrs-quant/src/blockwise_fp8/mod.rs

@@ -686,10 +686,30 @@ pub fn blockwise_fp8_linear_b(
 /// **The measurement that settles this constant:** re-run the M sweep with the
 /// flatten fix present, three arms, on one binary —
 /// `ARC_FP8_CUBLAS_MIN_M`/`ARC_NO_FP8_WMMA` toggle all three with no rebuild —
-/// and set the default to the first M at which (a) beats `min((b), (c))`. Until
-/// then the default stays where it is, because moving it would be a *second*
-/// constant with no measurement behind it, which is the exact defect this doc
-/// exists to record.
+/// and set the default to the first M at which (a) beats `min((b), (c))`.
+///
+/// # ✅ THAT SWEEP HAS NOW RUN — s10, and 64 is the measured crossover
+///
+/// 2026-08-22, box `arc-s10-ledger` (H200), candidate `4d03b9e25`, V4-Flash,
+/// `arc-tools/fp8_threshold_sweep.sh`, LADDER {8, 64, 256}, three arms, one
+/// binary, flatten fix present, >=1000-token floor per row (M=8 rows VOID on
+/// the floor and make no claim):
+///
+/// ```text
+///     M       arm  tokens   wall_s     tok/s   status
+///    64     tiled    2811     12.6     223.1   CLEAN
+///    64      wmma    2811      8.6     328.2   CLEAN
+///    64  cublaslt    3072      5.6     552.4   CLEAN  <- (a) beats min((b),(c))
+///   256     tiled   12288     43.0     285.8   CLEAN
+///   256      wmma   11244     18.1     620.4   CLEAN
+///   256  cublaslt   12288     17.2     715.9   CLEAN  <- and again
+/// ```
+///
+/// First clean M where (a) beats min((b), (c)): **64**. The default moves to
+/// that measured point. M=8 stays with the native arms (its rows were void, so
+/// no claim is made about 5..63 — the conservative side of a void row is to
+/// leave it with the incumbent). `fp8_gemv_warp` still owns M <= 4; the M=1
+/// floor row from the earlier sweep (cuBLASLt 0.72x at M=1) still stands.
 ///
 /// `fp8_gemv_warp` still owns `M <= 4` and is not affected by this constant.
 ///
@@ -698,11 +718,11 @@ pub fn blockwise_fp8_linear_b(
 #[cfg(feature = "cuda")]
 fn arc_fp8_cublas_min_m() -> usize {
     use std::sync::OnceLock;
-    // Still 512. NOT because 512 is right -- the doc above shows it is a single
-    // prefill point -- but because the only sweep in hand measured the scalar
-    // kernel that #200 displaced, and lowering this to 5 on that evidence would
-    // make the tensor-core GEMM dead code at every M >= 5. Re-sweep three arms.
-    const DEFAULT_MIN_M: usize = 512;
+    // 64 = the first CLEAN M at which dequantize+cuBLASLt beats min(WMMA,
+    // tiled) in the s10 three-arm sweep (arc-s10-ledger, 4d03b9e25; rows in
+    // the doc above). M=8 rows were VOID on the token floor, so 5..63 stays
+    // with the native arms rather than being claimed by interpolation.
+    const DEFAULT_MIN_M: usize = 64;
     static CACHE: OnceLock<usize> = OnceLock::new();
     *CACHE.get_or_init(|| {
         std::env::var("ARC_FP8_CUBLAS_MIN_M")