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; 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(|| {