From 52a8750f5b23a972246df36c9ac8221952f43a89 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 01:47:54 +0100 Subject: [PATCH 1/3] fix(ArcKernels): FP8 dispatch threshold was 512 because 512 was the only point measured `arc_fp8_cublas_min_m` defaulted to 512, and the doc said why outright: "the default is set to that measured point rather than to an interpolated crossover". `fp8_gemv_warp` owns M <= 4, so M = 5..511 fell through to `fp8_matmul_tiled` -- the kernel whose own comment says there is "no tensor-core instruction anywhere in it". A profile of the shipped default caught it at 7,525 launches for 17.6% of GPU time in one decode window. Swept on an H200 against V4-Flash, one binary and one env toggle (`ARC_FP8_CUBLAS_MIN_M`), aggregate tok/s, CLEAN ROWS ONLY (>=95% achieved concurrency, <15% derived-vs-measured spread): M=1 34.92 vs 25.27 cuBLASLt 0.72x <- the GEMV floor is real M=8 49.25 vs 64.95 cuBLASLt 1.32x M=16 45.61 vs 52.55 cuBLASLt 1.15x M=128 114.21 vs 130.76 cuBLASLt 1.14x (M=32, M=64 and M=256 came back dirty at this generation length and are excluded from the decision rather than averaged in.) The crossover is not 512; it sits immediately above the GEMV's domain. Default is now 5, the first row `fp8_gemv_warp` does not own. M = 5..7 is unmeasured, bounded below by that kernel and above by the +32% at M=8. KEEP THE FLOOR: at M=1 cuBLASLt is 0.72x, so forcing the library path over the dedicated GEMV costs 28% of b=1 decode. That row is why this is 5 and not 1. Mechanism: the tiled kernel is scalar, so its cost is instruction-bound and grows with M while the memory controller idles; cuBLASLt dequantizes once and hands a tensor-core GEMM the whole tile. The counter shows it -- memory-controller utilisation moved 2-3% -> 6-19% across the sweep, the first intervention this session to move it at all. Recorded in the doc as a recurring failure, because this is the second such gate: `qtip::gather_policy`'s tile-fill predicate needed n >= 683 and kept the grouped GEMM unreachable in every decode step (worth 1.45x at B=128). Both were derived from a single working point and neither was ever swept. Fast path by default; `ARC_FP8_CUBLAS_MIN_M` remains the kill switch and a value above any real batch restores the previous always-native dispatch. Co-Authored-By: Claude Opus 5 (1M context) --- mistralrs-quant/src/blockwise_fp8/mod.rs | 60 ++++++++++++++++++++++-- mistralrs-quant/src/blockwise_fp8/ops.rs | 3 +- 2 files changed, 56 insertions(+), 7 deletions(-) diff --git a/mistralrs-quant/src/blockwise_fp8/mod.rs b/mistralrs-quant/src/blockwise_fp8/mod.rs index b7b900268..a291e0e11 100644 --- a/mistralrs-quant/src/blockwise_fp8/mod.rs +++ b/mistralrs-quant/src/blockwise_fp8/mod.rs @@ -555,18 +555,68 @@ pub fn blockwise_fp8_linear_b( /// Measured on an H200 against DeepSeek-V4-Flash at 512 prompt tokens: the /// cuBLASLt path is 27x faster on the MLA `q_proj`/`o_proj` projections /// (706.1 -> 26.2 ms/step) and 21x on the shared expert (253.2 -> 11.9), for -/// 1.513x on prefill TTFT end to end. The default is set to that measured -/// point rather than to an interpolated crossover. +/// 1.513x on prefill TTFT end to end. /// -/// The decode regime is untouched: it has its own dedicated `fp8_gemv_warp` -/// kernel for M <= 4 and never reaches this threshold. +/// # 🔴 THE DEFAULT WAS 512 BECAUSE 512 WAS THE ONLY POINT MEASURED +/// +/// The previous doc said so outright — *"the default is set to that measured +/// point rather than to an interpolated crossover"* — and 512 then became a +/// wall. `fp8_gemv_warp` owns `M <= 4`, so **`M = 5..511` fell through to +/// `fp8_matmul_tiled`**, the kernel whose own comment says there is *"no +/// tensor-core instruction anywhere in it"*. A profile of the shipped default +/// caught it running **7,525 times for 17.6% of GPU time** in one decode +/// window. +/// +/// Swept on an H200, V4-Flash, one binary and one env toggle, aggregate tok/s, +/// **clean rows only** (>=95% achieved concurrency, <15% derived-vs-measured +/// spread): +/// +/// ```text +/// M=1 34.92 vs 25.27 cuBLASLt 0.72x -- the GEMV floor is REAL +/// M=8 49.25 vs 64.95 cuBLASLt 1.32x +/// M=16 45.61 vs 52.55 cuBLASLt 1.15x +/// M=128 114.21 vs 130.76 cuBLASLt 1.14x +/// ``` +/// +/// So the crossover is not 512; it is immediately above the GEMV's domain. The +/// default is now **5** — the first row `fp8_gemv_warp` does not own. +/// `M = 5..7` is **unmeasured**, bounded below by the GEMV domain and above by +/// the +32% at M=8. +/// +/// **Keep the floor.** At M=1 cuBLASLt is *worse* (0.72x): the dedicated +/// `fp8_gemv_warp` kernel exists for a reason and forcing the library path over +/// it costs 28% of b=1 decode. That row is why this constant is 5 and not 1. +/// +/// **Mechanism, not just a number.** The tiled kernel is scalar, so its cost is +/// instruction-bound and grows with M while the memory controller idles. The +/// cuBLASLt path dequantizes once and hands a tensor-core GEMM the whole tile. +/// The signature is visible in the counter: memory-controller utilisation went +/// **2-3% -> 6-19%** across the sweep, the first intervention this session to +/// move it at all. A scalar kernel that keeps the card 100% busy and the memory +/// controller at 2% is not doing the work; it is doing arithmetic *about* the +/// work. +/// +/// # A recurring failure in this codebase: one measured point frozen as a gate +/// +/// This is the **second** gate found this way. `qtip::gather_policy`'s tile-fill +/// predicate required `6n/256 >= 16`, i.e. n >= 683, and so kept the amortizing +/// grouped GEMM unreachable in every decode step; removing it was worth 1.45x +/// at B=128. Both gates were derived from a single working point and neither +/// was ever swept. **If you are about to set a dispatch threshold from one +/// measurement, sweep it first and write down which rows were clean** — a +/// threshold is a claim about every value it excludes, not just the one you +/// tried. +/// +/// `fp8_gemv_warp` still owns `M <= 4` and is not affected by this constant. /// /// `ARC_FP8_CUBLAS_MIN_M` overrides; a value above any real batch restores /// the previous always-native dispatch. #[cfg(feature = "cuda")] fn arc_fp8_cublas_min_m() -> usize { use std::sync::OnceLock; - const DEFAULT_MIN_M: usize = 512; + // 5 = the first row `fp8_gemv_warp` does not own. Swept, clean rows only; + // M=1 measured 0.72x for cuBLASLt, which is why this is not 1. + const DEFAULT_MIN_M: usize = 5; 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 85eb12b23..0baa281e2 100644 --- a/mistralrs-quant/src/blockwise_fp8/ops.rs +++ b/mistralrs-quant/src/blockwise_fp8/ops.rs @@ -1034,8 +1034,7 @@ fn fp8_blockwise_matmul_impl( match input.dtype() { DType::F16 => { - let output = - crate::arc_outbuf::alloc_out_fully_written::(&dev, (m * n) as usize)?; + let output = crate::arc_outbuf::alloc_out_fully_written::(&dev, (m * n) as usize)?; let input_s = match &*input_storage { Storage::Cuda(cuda_storage) => cuda_storage.as_cuda_slice::()?, From 445b920630e91e266d66695c99b3e80feb8eedae Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 02:01:19 +0100 Subject: [PATCH 2/3] fix(ArcKernels): the cuBLASLt FP8 branch never flattened rank-3 to 2-D MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The native branch flattens `[B, T, hidden]` to `[B*T, hidden]` before its GEMM. The dequantize + cuBLASLt branch did not — it passed `x` through untouched. V4's activation is rank-3, and bias is `None` on every V4 linear, so `UnquantLinear::forward` took `w.broadcast_left(B)` into `cublaslt.batch_matmul` with `stride_b = 0`. At decode T=1, making it a **B-way batched GEMM with m=1**: a batched GEMV wearing a GEMM's name, and the one shape cuBLASLt has no advantage in. The documented 27x was measured at prefill, where `[1, 512, hidden]` collapses to batch=1 and the call is a real GEMM. This invalidated the threshold A/B rather than merely slowing it: lowering `ARC_FP8_CUBLAS_MIN_M` selected the batched-GEVM shape, so the sweep was comparing the tiled kernel against a degenerate call and would have supported "cuBLASLt loses at decode" — a conclusion about a shape the change was never meant to select. Flattening makes the two branches comparable, which is the premise the A/B rests on. Co-Authored-By: Claude Opus 5 (1M context) --- mistralrs-quant/src/blockwise_fp8/mod.rs | 28 ++++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/mistralrs-quant/src/blockwise_fp8/mod.rs b/mistralrs-quant/src/blockwise_fp8/mod.rs index a291e0e11..14e5be51f 100644 --- a/mistralrs-quant/src/blockwise_fp8/mod.rs +++ b/mistralrs-quant/src/blockwise_fp8/mod.rs @@ -142,6 +142,34 @@ impl QuantMethod for BlockwiseFP8Linear { weight, self.bias.clone(), )))?; + + // 🔴 FLATTEN TO 2-D FIRST — the native branch above already does this at + // the top of `forward`, and this branch did not. + // + // V4's activation is rank-3 `[B, T, hidden]`. Passed through untouched, + // `UnquantLinear::forward` takes `w.broadcast_left(B)` and, because bias + // is `None` on every V4 linear, dispatches `cublaslt.batch_matmul` with + // `stride_b = 0`. At decode T=1, so that is a **B-way batched GEMM with + // m=1** — a batched GEMV wearing a GEMM's name, which is the one shape + // cuBLASLt has no advantage in. The documented 27x was measured at + // prefill, where `[1, 512, hidden]` collapses to batch=1 and the call is + // a genuine `[512, hidden] x [hidden, N]` GEMM. + // + // Without this, lowering `ARC_FP8_CUBLAS_MIN_M` measures the batched-GEMV + // path and concludes "cuBLASLt loses at decode" — a wrong conclusion + // drawn from a shape the threshold change was never meant to select. + // Flattening makes the two branches comparable, which is the whole + // premise of the A/B. + let orig_dims = x.dims().to_vec(); + if orig_dims.len() > 2 { + let features = orig_dims[orig_dims.len() - 1]; + let rows: usize = orig_dims[..orig_dims.len() - 1].iter().product(); + let out = unquant.forward(&x.reshape((rows, features))?)?; + let out_features = out.dim(out.rank() - 1)?; + let mut new_dims = orig_dims[..orig_dims.len() - 1].to_vec(); + new_dims.push(out_features); + return out.reshape(new_dims); + } unquant.forward(x) } From b5de0f1d41ba46cca5aef552bbc78950f56c604d Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 03:19:02 +0100 Subject: [PATCH 3/3] =?UTF-8?q?perf(ArcKernels):=20UNVERIFIED=20=E2=80=94?= =?UTF-8?q?=20remove=20the=204-way=20shared-memory=20bank=20conflict?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ⚠️ NEVER COMPILED. The box went down (port 22 closed, 100% ICMP loss) before this could be built or measured. Committed so it is not lost; it must be built and A/B'd before any number is quoted from it. `fp8_matmul_tiled`'s inner product reads `s_weight[tx][k]` with `tx` varying fastest within a warp, so consecutive lanes sit `BLOCK_K + pad` floats apart. At the shipped BLOCK_K=32 (`TILE_K = 32`, blockwise_fp8_gemm.cu:381/399) a +4 pad gives stride 36; 36 mod 32 = 4, gcd(4,32) = 4, so the warp reaches only 8 of the 32 shared-memory banks — a 4-way conflict on every FMA of the k loop. A +1 pad gives stride 33, coprime with 32, so all 32 banks are hit. Output is bit-identical by construction: the k loop is bounded `k < BLOCK_K`, so the pad columns are written by nobody and read by nobody. The padding exists only to set the row stride. This does not change the kernel's ceiling. The inner loop still issues two shared-memory loads per single FMA with no register blocking, which bounds it near an eighth of even the scalar roofline. The real fix is a tensor-core blockwise-FP8 GEMM. Co-Authored-By: Claude Opus 5 (1M context) --- .../kernels/blockwise_fp8/blockwise_fp8_gemm.cu | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) diff --git a/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu b/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu index 838014466..def6b9066 100644 --- a/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu +++ b/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu @@ -51,8 +51,19 @@ __global__ void fp8_matmul_tiled(const T *__restrict__ input, T *__restrict__ output, int M, int N, int K, int scale_row_stride, int block_size_y, int block_size_x) { - __shared__ float s_input[BLOCK_M][BLOCK_K + 4]; - __shared__ float s_weight[BLOCK_N][BLOCK_K + 4]; + // Padding is +1, not +4. The inner product reads `s_weight[tx][k]` with `tx` + // varying fastest within a warp, so consecutive lanes are `BLOCK_K + pad` + // floats apart. At BLOCK_K=32 a +4 pad gives stride 36; 36 mod 32 = 4, so + // gcd(4,32)=4 and the warp lands on only 8 of the 32 shared-memory banks -- + // a 4-way conflict on every FMA of the k loop. A +1 pad gives stride 33, + // which is coprime with 32, so all 32 banks are hit and the access is + // conflict-free. + // + // Bit-identical output: the k loop is bounded by `k < BLOCK_K`, so columns + // BLOCK_K..BLOCK_K+pad-1 are written by nobody and read by nobody. The pad + // exists only to change the row stride. + __shared__ float s_input[BLOCK_M][BLOCK_K + 1]; + __shared__ float s_weight[BLOCK_N][BLOCK_K + 1]; const int bx = blockIdx.x; const int by = blockIdx.y;