From f2f6aaa35919c8376c16187047c4aded53cb6535 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 15:36:47 +0100 Subject: [PATCH 1/2] fix(ArcKernels): record why FP8's dispatch threshold is 512, and why 512 stays Cherry-pick of 52a8750f5 (`fix/fp8-cublas-crossover`), resolved against `origin/master` @ 6ffdac7ae. The doc lands; the constant does not. `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 at 52a8750f5, 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.) 52a8750f5 concluded from those rows that the default should be 5. THAT CONCLUSION DOES NOT SURVIVE THE REBASE, so this commit keeps 512: * The sweep was taken at 01:47 on 2026-08-21. At 07:04 the same day, #200 (9ee459191) merged a tensor-core blockwise-FP8 GEMM, and `fp8_blockwise_matmul_impl` now selects it by default for everything `fp8_gemv_warp` does not own (`use_wmma`, ops.rs). Our shipped `weight_block_size` is [128, 128] and the kernel tiles N_BLK=64, K_BLK=128, so `fp8_wmma_eligible` passes: 128 % 64 == 0, 128 % 128 == 0. At M = 5..511 the native arm is no longer the scalar kernel the sweep measured. * Lowering the default to 5 would therefore route every M >= 5 into the dequantize path and make #200's kernel unreachable at every batch size that matters -- silently, because it compiles and the tests pass. That kernel's own docs already name the merge rule: three arms, not two, and "do not merge a threshold on the two-arm result". * The sweep also predates the rank-3 flatten in the next commit, so arm (a) was dispatching a B-way batched GEMV with m=1. Those rows understate cuBLASLt and have to be retaken regardless. * Arm (a) pays an uncached full-model `dequantize_w()` on every forward that arms (b) and (c) do not. Nothing here caches it, so the crossover this sweep found is the crossover of the uncached path. What settles the 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 take the first M at which dequantize+cuBLASLt beats the better of the two native kernels. Until that exists, moving 512 to 5 would be a second constant with no measurement behind it -- the exact defect this doc was written to record. Also recorded as a recurring failure, because 512 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. `ARC_FP8_CUBLAS_MIN_M` remains the kill switch in both directions. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- mistralrs-quant/src/blockwise_fp8/mod.rs | 103 +++++++++++++++++++++-- 1 file changed, 98 insertions(+), 5 deletions(-) diff --git a/mistralrs-quant/src/blockwise_fp8/mod.rs b/mistralrs-quant/src/blockwise_fp8/mod.rs index b5276bc5f..2cf3b9fa5 100644 --- a/mistralrs-quant/src/blockwise_fp8/mod.rs +++ b/mistralrs-quant/src/blockwise_fp8/mod.rs @@ -551,22 +551,115 @@ pub fn blockwise_fp8_linear_b( /// Parent system: ArcKernels. /// /// Row threshold at or above which a blockwise-FP8 linear prefers the -/// dequantize + cuBLASLt path over the scalar `fp8_matmul_tiled` kernel. +/// dequantize + cuBLASLt path over the native FP8 kernels in +/// `blockwise_fp8/ops.rs`. /// /// 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. +/// # 🔴 512 IS 512 BECAUSE 512 WAS THE ONLY POINT EVER MEASURED +/// +/// The previous revision of this 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. (That was the dispatch before #200; see below for what M = +/// 5..511 selects now.) +/// +/// 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 +/// ``` +/// +/// **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 any lowered value is >= 5, never +/// 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. +/// +/// # 🔴 AND IT IS STILL 512: THE SWEEP NO LONGER MEASURES THE ARM THAT RUNS +/// +/// That sweep was taken at 01:47 on 2026-08-21. At 07:04 the same day, #200 +/// (`9ee459191`) merged a **tensor-core** blockwise-FP8 GEMM, and +/// `fp8_blockwise_matmul_impl` now selects it by default for everything +/// `fp8_gemv_warp` does not own — see `use_wmma` in `blockwise_fp8/ops.rs`. Our +/// shipped `weight_block_size` is `[128, 128]` and the kernel's tiling is +/// `N_BLK = 64`, `K_BLK = 128`, so `128 % 64 == 0` and `128 % 128 == 0`: +/// `fp8_wmma_eligible` passes. **At M = 5..511 the native arm is no longer +/// `fp8_matmul_tiled`.** +/// +/// The rows above therefore compare arm (a) against an arm that is now only +/// reachable with `ARC_NO_FP8_WMMA=1`. Lowering this constant to 5 today would +/// route every M >= 5 into the dequantize path and leave #200's kernel +/// **unreachable at every batch size that matters** — the wired-but-dead +/// failure this repo already tracks, and silent, because it compiles and the +/// tests pass. The instruction is written into that kernel's own docs: sweep +/// three arms, not two, and *"do not merge a threshold on the two-arm result"*. +/// +/// ```text +/// (a) dequantize_w() + cuBLASLt <- what the sweep above measured +/// (b) native FP8 tensor-core GEMM <- #200, ARC_NO_FP8_WMMA unset +/// (c) native scalar fp8_matmul_tiled <- ARC_NO_FP8_WMMA=1 +/// ``` +/// +/// Two further reasons the sweep cannot be cashed as-is: +/// +/// * It ran **before** the rank-3 flatten in `forward` above. Arm (a) was +/// dispatching `batch_matmul` with `stride_b = 0` — a B-way batched GEMV with +/// m=1 — so those rows understate cuBLASLt and must be retaken regardless. +/// * Arm (a) pays a full-model `dequantize_w()` on **every** forward (~12.7 +/// GB/step, +8.48 GB resident BF16) that arms (b) and (c) do not. Cache that +/// and (a)'s crossover moves down; leave it uncached and (a) pays at every M. +/// Nothing here caches it. +/// +/// **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. +/// +/// `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; + // 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; static CACHE: OnceLock = OnceLock::new(); *CACHE.get_or_init(|| { From 9195d0e96a0cfdd74ad8ee869b4ea539181dfb77 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 2/2] =?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 aaba92fad..dad9e8330 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;