diff --git a/mistralrs-core/src/cuda/qk_norm_rope.rs b/mistralrs-core/src/cuda/qk_norm_rope.rs index 013d7f41b..3f15456dc 100644 --- a/mistralrs-core/src/cuda/qk_norm_rope.rs +++ b/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" ); } } diff --git a/mistralrs-core/src/models/deepseek4.rs b/mistralrs-core/src/models/deepseek4.rs index 5b007f92b..64361e466 100644 --- a/mistralrs-core/src/models/deepseek4.rs +++ b/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 diff --git a/mistralrs-core/tests/capability_reachability.rs b/mistralrs-core/tests/capability_reachability.rs index 18fa40eeb..2ba4de18f 100644 --- a/mistralrs-core/tests/capability_reachability.rs +++ b/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", diff --git a/mistralrs-quant/src/blockwise_fp8/gemv_wide.rs b/mistralrs-quant/src/blockwise_fp8/gemv_wide.rs index e173f4f01..0f38c182f 100644 --- a/mistralrs-quant/src/blockwise_fp8/gemv_wide.rs +++ b/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 diff --git a/mistralrs-quant/src/blockwise_fp8/mod.rs b/mistralrs-quant/src/blockwise_fp8/mod.rs index 5ef1cecc3..897f1e068 100644 --- a/mistralrs-quant/src/blockwise_fp8/mod.rs +++ b/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 = OnceLock::new(); *CACHE.get_or_init(|| { std::env::var("ARC_FP8_CUBLAS_MIN_M") diff --git a/mistralrs-quant/src/blockwise_fp8/ops.rs b/mistralrs-quant/src/blockwise_fp8/ops.rs index 8d8e8c5c9..347eb8288 100644 --- a/mistralrs-quant/src/blockwise_fp8/ops.rs +++ b/mistralrs-quant/src/blockwise_fp8/ops.rs @@ -1069,11 +1069,12 @@ fn fp8_wmma_eligible(block_size_y: i32, block_size_x: i32) -> bool { /// (see [`gemv_wide::wide_eligible`]); any shape that fails one keeps /// `fp8_gemv_warp`, unchanged, with no error. /// -/// 🔴 Default OFF. `ARC_FP8_GEMV_WIDE=1` — and *only* the literal `"1"` — -/// turns it on. The kernel has never run on a GPU and is not bit-identical to -/// `fp8_gemv_warp` (it re-associates the f32 accumulation; the exact claim and -/// its CPU proof are in `blockwise_fp8/gemv_wide.rs`). It must be A/B'd on -/// hardware against the default before anyone considers flipping it. +/// ✅ Default ON since the s10 hardware A/B (box `arc-s10-ledger`, candidate +/// `4d03b9e25`, leg 3: 43.80 vs 40.76 tok/s b=1, engagement line +/// `path=gemv_wide`, coherent canary). `ARC_FP8_GEMV_WIDE=0` — and *only* the +/// literal `"0"` — restores `fp8_gemv_warp` from the same binary. It is not +/// bit-identical to `fp8_gemv_warp` (it re-associates the f32 accumulation; +/// the exact claim and its CPU proof are in `blockwise_fp8/gemv_wide.rs`). /// /// The pointers passed here are the ones the kernel will receive, so the /// 16-byte alignment a `uint4` load requires is checked on the real address @@ -1323,7 +1324,7 @@ fn fp8_blockwise_matmul_impl( k as usize, ); if let Some(scale_shift) = wide_shift { - // 🔴 UNVERIFIED ON HARDWARE. `ARC_FP8_GEMV_WIDE=1` only. + // ✅ Measured s10 leg 3 (arc-s10-ledger, 4d03b9e25): +7.5% b=1. Default ON; `ARC_FP8_GEMV_WIDE=0` restores warp. unsafe { ffi::launch_fp8_gemv_wide_f16( input_ptr as *const _, @@ -1427,7 +1428,7 @@ fn fp8_blockwise_matmul_impl( k as usize, ); if let Some(scale_shift) = wide_shift { - // 🔴 UNVERIFIED ON HARDWARE. `ARC_FP8_GEMV_WIDE=1` only. + // ✅ Measured s10 leg 3 (arc-s10-ledger, 4d03b9e25): +7.5% b=1. Default ON; `ARC_FP8_GEMV_WIDE=0` restores warp. unsafe { ffi::launch_fp8_gemv_wide_bf16( input_ptr as *const _,