From 589a4a7256aab0dda228117c5eeb18dbc1f5de67 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 1/6] 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 2cf3b9fa5..05ade1090 100644 --- a/mistralrs-quant/src/blockwise_fp8/mod.rs +++ b/mistralrs-quant/src/blockwise_fp8/mod.rs @@ -143,6 +143,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 f60b7e8a6f8c38f6e5227324d51df69eb4bcb9b5 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 19:43:40 +0100 Subject: [PATCH 2/6] perf(ArcKernels): pow2 shift fast path for the tiled FP8 GEMM's per-element scale divisions MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit get_scale in fp8_matmul_tiled performs TWO runtime signed integer divisions per weight element staged into shared memory (n / block_size_y, k / block_size_x). Both divisors are kernel arguments, so nvcc expands each into a ~15-20 instruction reciprocal-and-fixup sequence — paid once per element of every weight tile, ~33.9G times per B=256 decode step by the source-derived count. The shipped weight_block_size is [128, 128]: both powers of two. This adds a POW2_SCALE template arm: the host launcher tests pow2-ness once per launch, computes shift = log2(block_size) on the host, and the kernel indexes the scale grid with two SHFs instead. Non-power-of-two geometries keep the division path, unchanged. Bit-identical by construction: for the non-negative indices used here, n >> log2(d) == n / d for every input — same quotient, same scale word fetched, same arithmetic after it. Dispatching on POW2_SCALE can never change what an A/B leg measures, only how fast the scalar arm runs. The C ABI (launch_fp8_matmul_{f16,bf16}) is unchanged; no Rust edits needed. Compile-unverified locally (macOS has no nvcc); covered by the nvcc CI lane (.github/workflows/cuda_compile_check.yaml), which builds this TU for sm_80/sm_90 on every PR touching it. The bank-conflict half of this lane's brief (+4 -> +1 tile padding) was already landed by PR #215 and is an ancestor of this branch; the kernel's stride-33 comment documents the bank arithmetic. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --- .../blockwise_fp8/blockwise_fp8_gemm.cu | 118 ++++++++++++++---- 1 file changed, 93 insertions(+), 25 deletions(-) diff --git a/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu b/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu index dad9e8330..077b7c473 100644 --- a/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu +++ b/mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu @@ -40,17 +40,54 @@ __device__ __forceinline__ float get_scale(const float *__restrict__ scale, return __ldg(&scale[sr * scale_stride + sc]); } +// Shift variant of `get_scale`, for power-of-two scale-block sizes. +// +// `block_size_y`/`block_size_x` arrive as kernel ARGUMENTS, so `n / +// block_size_y` above is `int / int` by a non-constant divisor. NVIDIA GPUs +// have no integer-divide instruction; nvcc expands each division inline to a +// reciprocal-multiply-and-fixup sequence of ~15-20 instructions, and +// `fp8_matmul_tiled` pays TWO of them per weight element it stages into +// shared memory. The shipped `weight_block_size` is [128, 128] — both powers +// of two — so the host launcher computes `shift = log2(block_size)` once per +// launch and the kernel indexes with two SHFs instead. +// +// Bit-identical: for non-negative `n` and a power-of-two divisor, +// `n >> log2(d)` IS `n / d` — the same quotient for every input, no rounding +// involved anywhere. The two paths read the same scale word; only the +// instruction count changes. Non-power-of-two geometries keep the division +// path above. +__device__ __forceinline__ float get_scale_pow2(const float *__restrict__ scale, + int n, int k, int scale_stride, + int scale_shift_y, + int scale_shift_x) { + int sr = n >> scale_shift_y; + int sc = k >> scale_shift_x; + return __ldg(&scale[sr * scale_stride + sc]); +} + // ============================================================================ // FP8 Matmul Kernel // ============================================================================ -template +// `POW2_SCALE` selects, at compile time, how the per-element weight scale is +// indexed while staging the weight tile: +// * true -> `get_scale_pow2`: two SHFs, using `scale_shift_y`/`scale_shift_x` +// (host-computed `log2(block_size_y)`/`log2(block_size_x)`). +// * false -> `get_scale`: two runtime integer divisions (the general path, +// kept for non-power-of-two scale geometries). +// The two instantiations produce BIT-IDENTICAL output (`n >> log2(d) == n / d` +// for the non-negative indices used here), so dispatching on it can never +// change what a leg of an A/B measures — only how fast the scalar arm runs. +// Both `block_size_*` and `scale_shift_*` are always passed; each +// instantiation reads only its own pair. +template __global__ void fp8_matmul_tiled(const T *__restrict__ input, const __nv_fp8_e4m3 *__restrict__ weight, const float *__restrict__ weight_scale, T *__restrict__ output, int M, int N, int K, int scale_row_stride, int block_size_y, - int block_size_x) { + int block_size_x, int scale_shift_y, + int scale_shift_x) { // 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 @@ -106,8 +143,14 @@ __global__ void fp8_matmul_tiled(const T *__restrict__ input, if (gn < N && gk < K) { __nv_fp8_e4m3 w; w.__x = __ldg(reinterpret_cast(&weight[gn * K + gk])); - float s = get_scale(weight_scale, gn, gk, scale_row_stride, - block_size_y, block_size_x); + float s; + if constexpr (POW2_SCALE) { + s = get_scale_pow2(weight_scale, gn, gk, scale_row_stride, + scale_shift_y, scale_shift_x); + } else { + s = get_scale(weight_scale, gn, gk, scale_row_stride, block_size_y, + block_size_x); + } val = fp8_to_float(w) * s; } s_weight[ln][lk] = val; @@ -625,41 +668,66 @@ __global__ void fp8_moe_gemm(const T *__restrict__ input, // C API // ============================================================================ -extern "C" void launch_fp8_matmul_f16(const __half *input, - const __nv_fp8_e4m3 *weight, - const float *weight_scale, __half *output, - int M, int N, int K, int scale_row_stride, - int block_size_y, int block_size_x, - cudaStream_t stream) { +// Host-side pow2 test + exact log2 for the tiled kernel's scale fast path. +// Computed ONCE per launch, on the host, so the device never divides. +static inline bool fp8_is_pow2(int v) { return v > 0 && (v & (v - 1)) == 0; } +static inline int fp8_ilog2_exact(int v) { + int s = 0; + while ((1 << s) < v) + ++s; + return s; // caller guarantees fp8_is_pow2(v), so (1 << s) == v +} + +// Shared launcher: picks the POW2_SCALE instantiation when BOTH scale-block +// sizes are powers of two (the shipped geometry is [128, 128], so serving +// always takes it), otherwise the general division path. The two +// instantiations are bit-identical in output — see the kernel comment. +template +static void launch_fp8_matmul_impl(const T *input, const __nv_fp8_e4m3 *weight, + const float *weight_scale, T *output, int M, + int N, int K, int scale_row_stride, + int block_size_y, int block_size_x, + cudaStream_t stream) { constexpr int TILE = 32; constexpr int TILE_K = 32; dim3 block(TILE, TILE); dim3 grid(CEILDIV(N, TILE), CEILDIV(M, TILE)); - fp8_gemm::fp8_matmul_tiled - <<>>(input, weight, weight_scale, output, M, N, K, - scale_row_stride, block_size_y, - block_size_x); + if (fp8_is_pow2(block_size_y) && fp8_is_pow2(block_size_x)) { + fp8_gemm::fp8_matmul_tiled + <<>>(input, weight, weight_scale, output, M, N, + K, scale_row_stride, block_size_y, + block_size_x, fp8_ilog2_exact(block_size_y), + fp8_ilog2_exact(block_size_x)); + } else { + fp8_gemm::fp8_matmul_tiled + <<>>(input, weight, weight_scale, output, M, N, + K, scale_row_stride, block_size_y, + block_size_x, 0, 0); + } CUDA_CHECK(cudaGetLastError()); } +extern "C" void launch_fp8_matmul_f16(const __half *input, + const __nv_fp8_e4m3 *weight, + const float *weight_scale, __half *output, + int M, int N, int K, int scale_row_stride, + int block_size_y, int block_size_x, + cudaStream_t stream) { + launch_fp8_matmul_impl(input, weight, weight_scale, output, M, N, K, + scale_row_stride, block_size_y, block_size_x, + stream); +} + extern "C" void launch_fp8_matmul_bf16(const __nv_bfloat16 *input, const __nv_fp8_e4m3 *weight, const float *weight_scale, __nv_bfloat16 *output, int M, int N, int K, int scale_row_stride, int block_size_y, int block_size_x, cudaStream_t stream) { - constexpr int TILE = 32; - constexpr int TILE_K = 32; - - dim3 block(TILE, TILE); - dim3 grid(CEILDIV(N, TILE), CEILDIV(M, TILE)); - - fp8_gemm::fp8_matmul_tiled<__nv_bfloat16, TILE, TILE, TILE_K> - <<>>(input, weight, weight_scale, output, M, N, K, - scale_row_stride, block_size_y, - block_size_x); - CUDA_CHECK(cudaGetLastError()); + launch_fp8_matmul_impl<__nv_bfloat16>(input, weight, weight_scale, output, M, + N, K, scale_row_stride, block_size_y, + block_size_x, stream); } // Warp-per-row GEMV launcher for the decode regime (M <= 4). Row-grouping From 4c7a61ced9a1369ebed253744e2290ea85a84571 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 19:43:51 +0100 Subject: [PATCH 3/6] =?UTF-8?q?feat(ArcLab):=20ARC=5FLOG=5FFP8=5FDISPATCH?= =?UTF-8?q?=20=E2=80=94=20engagement=20lines=20proving=20WHICH=20FP8=20ker?= =?UTF-8?q?nel=20served=20a=20forward?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Every threshold sweep on this lane needs an engagement assertion: a server-log line proving which kernel actually ran, per leg. None existed — a leg could silently measure the wrong arm (the ARC_NO_DEDICATED_DECODE incident produced two void A/Bs exactly this way, both arms running the same code). ARC_LOG_FP8_DISPATCH=1 (read by VALUE via env_flag_is_set, latched once per process) makes each dispatch path print ONE line per process: [arc-fp8-dispatch] path= first_shape=mM_nN_kK Sites: the native dispatch chain in fp8_blockwise_matmul_impl (both dtype arms, mirroring the if/else exactly) and the ARC_FP8_CUBLAS_MIN_M divert to dequantize+cuBLASLt in BlockwiseFP8Linear::forward. Cost when unset: one latched bool read. When set: one relaxed atomic swap per call — no locks, so enabling it does not perturb the rates a sweep records. Registered in mistralrs-core/tests/capability_reachability.rs (Status::Live); the registry run passes (11/11). The Rust is cfg(cuda) — compile-unverified locally on macOS; covered by the cuda-typecheck job of the nvcc CI lane. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --- .../tests/capability_reachability.rs | 15 +++ mistralrs-quant/src/blockwise_fp8/mod.rs | 14 +++ mistralrs-quant/src/blockwise_fp8/ops.rs | 111 +++++++++++++++++- 3 files changed, 136 insertions(+), 4 deletions(-) diff --git a/mistralrs-core/tests/capability_reachability.rs b/mistralrs-core/tests/capability_reachability.rs index c005e0153..7e1d0d7a7 100644 --- a/mistralrs-core/tests/capability_reachability.rs +++ b/mistralrs-core/tests/capability_reachability.rs @@ -326,6 +326,21 @@ static REGISTRY: &[Capability] = &[ // D2H that also disqualifies all rows from the GPU sampler fast path. status: Status::Live, }, + Capability { + name: "FP8 dispatch engagement log (ARC_LOG_FP8_DISPATCH)", + parent: "ArcLab", + check: Check::Symbol { + symbol: "log_fp8_dispatch", + defined_in: "mistralrs-quant/src/blockwise_fp8/ops.rs", + }, + // The engagement line the FP8 threshold sweep + // (`arc-tools/fp8_threshold_sweep.sh`) greps per leg to prove WHICH + // kernel actually served it. Value-read (`env_flag_is_set`), latched + // once per process, one line per path per process. Referenced from + // both native dispatch arms in `ops.rs` and from the + // dequantize+cuBLASLt divert in `blockwise_fp8/mod.rs`. + status: Status::Live, + }, // -- the ApiPromise class ------------------------------------------------ Capability { name: "sampling: logit_bias reaches the sampler", diff --git a/mistralrs-quant/src/blockwise_fp8/mod.rs b/mistralrs-quant/src/blockwise_fp8/mod.rs index 05ade1090..5ef1cecc3 100644 --- a/mistralrs-quant/src/blockwise_fp8/mod.rs +++ b/mistralrs-quant/src/blockwise_fp8/mod.rs @@ -95,6 +95,20 @@ impl QuantMethod for BlockwiseFP8Linear { // A/B runs from a single binary. let m_rows: usize = x.dims().iter().rev().skip(1).product(); let prefer_cublas = m_rows >= arc_fp8_cublas_min_m(); + if prefer_cublas + && matches!(x.device(), candle_core::Device::Cuda(_)) + && ffi::HAVE_BLOCKWISE_GEMM_KERNELS + { + // Engagement line for the threshold sweep: this forward was + // DIVERTED from the native FP8 kernels to dequantize+cuBLASLt + // by `ARC_FP8_CUBLAS_MIN_M`. `ARC_LOG_FP8_DISPATCH=1` only. + ops::log_fp8_dispatch( + ops::Fp8DispatchPath::DequantCublaslt, + m_rows, + self.weight.dims().first().copied().unwrap_or(0), + x.dims().last().copied().unwrap_or(0), + ); + } if !prefer_cublas && matches!(x.device(), candle_core::Device::Cuda(_)) && ffi::HAVE_BLOCKWISE_GEMM_KERNELS diff --git a/mistralrs-quant/src/blockwise_fp8/ops.rs b/mistralrs-quant/src/blockwise_fp8/ops.rs index 75a1d1f42..8d8e8c5c9 100644 --- a/mistralrs-quant/src/blockwise_fp8/ops.rs +++ b/mistralrs-quant/src/blockwise_fp8/ops.rs @@ -1008,8 +1008,7 @@ fn fp8_wmma_enabled() -> bool { // Under the presence test this replaced, it silently selected the control // arm — which is the arm this kernel's own docs tell you to run first, so // an operator following them would have recorded the control twice. - static ENABLED: LazyLock = - LazyLock::new(|| !crate::env_flag_is_set("ARC_NO_FP8_WMMA")); + static ENABLED: LazyLock = LazyLock::new(|| !crate::env_flag_is_set("ARC_NO_FP8_WMMA")); *ENABLED } @@ -1099,6 +1098,81 @@ fn wide_gemv_shift( ) } +/// Parent system: ArcLab. +/// +/// Which blockwise-FP8 kernel a forward actually selected. Used only by +/// [`log_fp8_dispatch`]; the discriminants index its once-per-process latch. +#[cfg(feature = "cuda")] +#[derive(Clone, Copy)] +pub(crate) enum Fp8DispatchPath { + GemvWide = 0, + GemvWarp = 1, + Wmma = 2, + Tiled = 3, + DequantCublaslt = 4, +} + +#[cfg(feature = "cuda")] +impl Fp8DispatchPath { + fn name(self) -> &'static str { + match self { + Fp8DispatchPath::GemvWide => "gemv_wide", + Fp8DispatchPath::GemvWarp => "gemv_warp", + Fp8DispatchPath::Wmma => "wmma", + Fp8DispatchPath::Tiled => "tiled", + Fp8DispatchPath::DequantCublaslt => "dequant_cublaslt", + } + } +} + +/// `ARC_LOG_FP8_DISPATCH=1` — engagement logging for the blockwise-FP8 +/// dispatch. Read by VALUE via `env_flag_is_set` (`=0` keeps it off), latched +/// once per process so the decode hot loop never touches the env. +/// +/// Registered in `mistralrs-core/tests/capability_reachability.rs`. +#[cfg(feature = "cuda")] +fn fp8_dispatch_log_enabled() -> bool { + use std::sync::LazyLock; + static ENABLED: LazyLock = + LazyLock::new(|| crate::env_flag_is_set("ARC_LOG_FP8_DISPATCH")); + *ENABLED +} + +/// Parent system: ArcLab. +/// +/// Print ONE line per dispatch path per process naming the kernel that +/// actually ran — the engagement assertion every threshold sweep must grep +/// for. A sweep leg that cannot prove WHICH kernel served it is measuring an +/// assumption (`ARC_NO_DEDICATED_DECODE` produced exactly that null result: +/// both arms of two A/Bs ran the same code and nobody could tell). +/// +/// Cost when `ARC_LOG_FP8_DISPATCH` is unset (the default): one latched bool +/// read. When set: one relaxed atomic swap per call — no locks, no +/// allocation — so enabling it does not perturb the rates the sweep records. +/// The shape printed is the FIRST shape that reached the path, which is +/// enough to identify the projection; per-call logging would be ~430 lines +/// per V4 step at B=256 and is deliberately not offered. +#[cfg(feature = "cuda")] +pub(crate) fn log_fp8_dispatch(path: Fp8DispatchPath, m: usize, n: usize, k: usize) { + use std::sync::atomic::{AtomicBool, Ordering}; + static LOGGED: [AtomicBool; 5] = [ + AtomicBool::new(false), + AtomicBool::new(false), + AtomicBool::new(false), + AtomicBool::new(false), + AtomicBool::new(false), + ]; + if !fp8_dispatch_log_enabled() { + return; + } + if !LOGGED[path as usize].swap(true, Ordering::Relaxed) { + eprintln!( + "[arc-fp8-dispatch] path={} first_shape=m{m}_n{n}_k{k}", + path.name() + ); + } +} + /// FP8 blockwise matmul. /// Computes output = input @ weight.T where weight is FP8 blockwise quantized. /// - input: [M, K] in fp16/bf16 @@ -1221,8 +1295,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::()?, @@ -1234,6 +1307,21 @@ fn fp8_blockwise_matmul_impl( let (input_ptr, _input_guard) = slice_ptr(input_s, input_l.start_offset()); let wide_shift = wide_gemv_shift(use_gemv, k, block_size_x, weight_ptr, input_ptr); + // Mirrors the if/else chain below exactly; keep the two in step. + log_fp8_dispatch( + if wide_shift.is_some() { + Fp8DispatchPath::GemvWide + } else if use_gemv { + Fp8DispatchPath::GemvWarp + } else if use_wmma { + Fp8DispatchPath::Wmma + } else { + Fp8DispatchPath::Tiled + }, + m as usize, + n as usize, + k as usize, + ); if let Some(scale_shift) = wide_shift { // 🔴 UNVERIFIED ON HARDWARE. `ARC_FP8_GEMV_WIDE=1` only. unsafe { @@ -1323,6 +1411,21 @@ fn fp8_blockwise_matmul_impl( let (input_ptr, _input_guard) = slice_ptr(input_s, input_l.start_offset()); let wide_shift = wide_gemv_shift(use_gemv, k, block_size_x, weight_ptr, input_ptr); + // Mirrors the if/else chain below exactly; keep the two in step. + log_fp8_dispatch( + if wide_shift.is_some() { + Fp8DispatchPath::GemvWide + } else if use_gemv { + Fp8DispatchPath::GemvWarp + } else if use_wmma { + Fp8DispatchPath::Wmma + } else { + Fp8DispatchPath::Tiled + }, + m as usize, + n as usize, + k as usize, + ); if let Some(scale_shift) = wide_shift { // 🔴 UNVERIFIED ON HARDWARE. `ARC_FP8_GEMV_WIDE=1` only. unsafe { From 6b951cf900702f866d57a6c6f5acc1fc5c6dbc8b Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 19:44:05 +0100 Subject: [PATCH 4/6] =?UTF-8?q?feat(ArcLab):=20three-arm=20FP8=20threshold?= =?UTF-8?q?=20sweep=20harness=20=E2=80=94=20the=20measurement=20allowed=20?= =?UTF-8?q?to=20move=20ARC=5FFP8=5FCUBLAS=5FMIN=5FM?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit arc_fp8_cublas_min_m's doc and the WMMA kernel's doc both forbid merging a threshold from a two-arm result: the only sweep ever taken predates both the WMMA GEMM (#200) and the rank-3 flatten, so it compared cuBLASLt-as-batched- GEMV against the scalar kernel only. This script is the three-arm re-run they demand, on ONE binary, env toggles only, fresh server per leg (the gates are OnceLock-latched, so re-exporting into a live server would silently measure the previous leg). Design: the threshold candidates {5, 8, 64, 256, 512} are swept as the OFFERED DECODE BATCH M (a grid that varies only the env at one fixed batch routes every MIN_M <= batch leg to the same kernel and cannot name a crossover in M), and each ladder point runs all three arms: (a) cublaslt: ARC_FP8_CUBLAS_MIN_M=5 (b) wmma: ARC_FP8_CUBLAS_MIN_M=1000000, ARC_NO_FP8_WMMA=0 (c) tiled: ARC_FP8_CUBLAS_MIN_M=1000000, ARC_NO_FP8_WMMA=1 A leg counts only if ALL of: 1. ENGAGEMENT — its expected [arc-fp8-dispatch] path= line present in the server log and both rivals ABSENT (a WMMA leg whose log shows tiled is eligibility silently failing: VOID, loudly); 2. FLOOR — summed usage.completion_tokens >= MIN_TOKENS_FLOOR (default 1000); the rate divides by tokens the server SAYS it generated, never by requested max_tokens; 3. CANARY — greedy fixed-prompt stream vs the baseline leg (tiled @ M=5); cross-arm divergence is expected numerics and is REPORTED with its first index; a degenerate stream (<5 tokens / no finish_reason) voids the leg. The summary table names, per M, cublaslt vs min(wmma, tiled) and the FIRST clean M where cuBLASLt wins — the value ARC_FP8_CUBLAS_MIN_M may move to. No clean crossover => the default stays 512. Script only — NOT run here (no GPU; D14). House lock/preflight/provenance discipline copied from arcspec_token_identity_b8.sh. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --- arc-tools/fp8_threshold_sweep.sh | 412 +++++++++++++++++++++++++++++++ 1 file changed, 412 insertions(+) create mode 100755 arc-tools/fp8_threshold_sweep.sh diff --git a/arc-tools/fp8_threshold_sweep.sh b/arc-tools/fp8_threshold_sweep.sh new file mode 100755 index 000000000..05883dd13 --- /dev/null +++ b/arc-tools/fp8_threshold_sweep.sh @@ -0,0 +1,412 @@ +#!/bin/bash +# Parent system: ArcKernels + ArcQuant/Legacy (blockwise FP8). +# +# THREE-ARM FP8 DENSE-GEMM THRESHOLD SWEEP — the measurement that is allowed +# to move `ARC_FP8_CUBLAS_MIN_M` off 512. +# +# ───────────────────────────────────────────────────────────────────────────── +# WHY THREE ARMS, AND WHY A TWO-ARM RESULT MUST NOT BE MERGED +# +# `arc_fp8_cublas_min_m`'s own doc (blockwise_fp8/mod.rs) and the WMMA +# kernel's doc (blockwise_fp8/ops.rs, `fp8_wmma_enabled`) both record the +# standing instruction: the only sweep ever taken compared (a) dequantize + +# cuBLASLt against (c) the scalar tiled kernel, BEFORE #200 landed the +# (b) tensor-core WMMA GEMM — and before the rank-3 flatten, so arm (a) was a +# B-way batched GEMV wearing a GEMM's name. Freezing a threshold from that +# sweep would leave (b) dead code at every M that matters, silently. This +# script is the three-arm re-run those docs demand: +# +# (a) dequantize_w() + cuBLASLt ARC_FP8_CUBLAS_MIN_M=5 +# (b) native FP8 tensor-core WMMA ARC_FP8_CUBLAS_MIN_M=1000000, ARC_NO_FP8_WMMA=0 +# (c) native scalar fp8_matmul_tiled ARC_FP8_CUBLAS_MIN_M=1000000, ARC_NO_FP8_WMMA=1 +# +# ONE binary. Every leg differs only in environment; a leg that needed a +# rebuild would be measuring a different artifact. The env gates are latched +# per process (OnceLock/LazyLock), so every leg gets a FRESH server process — +# re-exporting into a running server changes nothing and would silently +# measure the previous leg's setting. +# +# ───────────────────────────────────────────────────────────────────────────── +# THE LADDER: M VARIES BY OFFERED CONCURRENCY, NOT ONLY BY THRESHOLD +# +# The deliverable is "the first M at which (a) beats min((b), (c))". A grid +# that sweeps only ARC_FP8_CUBLAS_MIN_M at one fixed batch cannot name that M: +# every leg with MIN_M <= batch routes the same decode shape to the same +# kernel and the legs replicate each other. So the threshold candidates +# {5, 8, 64, 256, 512} are swept as the OFFERED DECODE BATCH — M concurrent +# greedy completions, T=1 per step, so the dense projections see m_rows = M — +# and each ladder point runs all three arms. 15 legs. +# +# ───────────────────────────────────────────────────────────────────────────── +# WHAT MAKES A LEG COUNT (all three, or the leg is VOID) +# +# 1. ENGAGEMENT, asserted per leg. Every leg runs with ARC_LOG_FP8_DISPATCH=1; +# the dispatch sites print `[arc-fp8-dispatch] path=` once per path +# per process. The leg's expected kernel must be PRESENT in the server log +# and both rival kernels ABSENT. A leg that cannot prove which kernel +# served it is measuring an assumption (`ARC_NO_DEDICATED_DECODE`'s two +# void A/Bs are the record of what that costs). In particular: if the WMMA +# arm's log shows `path=tiled`, WMMA eligibility silently failed and the +# arm is the control wearing the treatment's name — VOID, loudly. +# 2. A MINIMUM-TOKEN FLOOR. Rates on fewer than MIN_TOKENS_FLOOR generated +# tokens are warm-up noise. The rate divides by SUMMED +# `usage.completion_tokens` — the tokens the server says it generated — +# NEVER by requested max_tokens (a leg that stalls and returns short would +# otherwise report a healthy rate for work it did not do). +# 3. A TOKEN CANARY vs the baseline leg (arm (c) at M=5 — the shipped scalar +# kernel, nearest to b=1). Greedy, fixed seed, fixed prompt. The arms are +# DIFFERENT NUMERICS by design (f32 scalar vs tensor-core promotion vs +# dequant-BF16 GEMM), so divergence at some index is expected and REPORTED, +# not failed; what hard-fails a leg is a degenerate stream (< 5 tokens or +# no finish_reason) — a failure to generate, which no rate can offset. +# +# The final table names, per ladder M, the three clean rates and min((b),(c)), +# and then the FIRST M at which (a) beats that min — the value +# ARC_FP8_CUBLAS_MIN_M should default to. If any arm at that M is void, the +# crossover is NOT ESTABLISHED and the default stays 512. +# +# USAGE (on the GPU box; do not run without a GPU — D14) +# setsid nohup bash arc-tools/fp8_threshold_sweep.sh /root/logs/fp8_sweep.log 2>&1 & +# cat /root/logs/fp8sweep/summary.txt +# +set -u + +ARCENV="${ARCENV:-/root/arcenv.sh}" +WT="${WT:-/root/arc}" +BIN="${BIN:-$WT/target/release/mistralrs}" +MODEL="${MODEL:-/root/models/v4-src}" +UQFF="${UQFF:-}" # optional: --from-uqff when set +ARCH="${ARCH:-deepseekv4}" +PREFLIGHT="${PREFLIGHT:-/root/arc-tools/gpu_box_preflight.sh}" +OUT="${OUT:-/root/logs/fp8sweep}" +STATUS="${STATUS:-/root/logs/fp8_sweep.status}" +PORT="${PORT:-1241}" +SEED="${SEED:-1234}" +MAX_TOKENS="${MAX_TOKENS:-48}" # per request; the floor is on the SUM +MIN_TOKENS_FLOOR="${MIN_TOKENS_FLOOR:-1000}" +MAX_SEQS="${MAX_SEQS:-512}" # must be >= the largest ladder point +LADDER="${LADDER:-5 8 64 256 512}" # offered decode batch = candidate M +LOCK="${LOCK:-/root/locks/gpu.lock}" +LOCK_WAIT="${LOCK_WAIT:-2400}" +OWNER="fp8sweep" + +mkdir -p "$OUT" "$(dirname "$STATUS")" "$(dirname "$LOCK")" +: > "$STATUS" +say() { echo "[$(date -u +%H:%M:%SZ)] $*" | tee -a "$STATUS"; } +die() { say "ABORT: $*"; echo "ABORT: $*" >> "$OUT/summary.txt"; exit 2; } + +say "fp8 three-arm threshold sweep — start" + +# ── GPU lock with OWNER-MATCHED RELEASE and evidence-based stale reclaim +# (same discipline as arcspec_token_identity_b8.sh; see the incident notes +# there for why release refuses locks this process does not own). +release_lock() { + if [ -f "$LOCK" ] && grep -q "^$OWNER pid=$$ " "$LOCK" 2>/dev/null; then + rm -f "$LOCK" + fi +} +STALE_AFTER="${STALE_AFTER:-120}" +lock_is_stale() { + [ -f "$LOCK" ] || return 1 + local age holder pid used procs + age=$(( $(date +%s) - $(stat -c %Y "$LOCK" 2>/dev/null || echo 0) )) + [ "$age" -ge "$STALE_AFTER" ] || return 1 + holder=$(cat "$LOCK" 2>/dev/null) + pid=$(printf '%s' "$holder" | sed -n 's/.*pid=\([0-9]*\).*/\1/p') + if [ -n "$holder" ] && [ -n "$pid" ] && kill -0 "$pid" 2>/dev/null; then + return 1 + fi + used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits 2>/dev/null | head -1) + procs=$(nvidia-smi --query-compute-apps=pid --format=csv,noheader 2>/dev/null | wc -l) + [ "${used:-1}" -lt 1024 ] && [ "${procs:-1}" -eq 0 ] +} +acquire_lock() { + local waited=0 + while :; do + if ( set -o noclobber; echo "$OWNER pid=$$ started=$(date -u +%FT%TZ)" > "$LOCK" ) 2>/dev/null; then + trap 'release_lock' EXIT INT TERM + say "lock: acquired $LOCK" + return 0 + fi + if lock_is_stale; then + say "lock: RECLAIMING stale lock (age>=${STALE_AFTER}s, owner absent or dead, GPU idle)" + rm -f "$LOCK" + continue + fi + if [ "$waited" -ge "$LOCK_WAIT" ]; then + say "lock: HELD BY -> $(cat "$LOCK" 2>/dev/null)" + return 1 + fi + [ "$waited" = 0 ] && say "lock: waiting for -> $(cat "$LOCK" 2>/dev/null)" + sleep 15; waited=$((waited + 15)) + done +} +acquire_lock || die "GPU lock held by another chain for >${LOCK_WAIT}s" + +[ -f "$ARCENV" ] && . "$ARCENV" +if [ -x "$PREFLIGHT" ]; then + "$PREFLIGHT" > "$OUT/preflight.log" 2>&1 || die "preflight refused the box" + say "preflight OK" +else + say "WARN: no preflight at $PREFLIGHT — continuing without it" +fi +[ -x "$BIN" ] || die "$BIN not found — build once, then sweep (one binary, no rebuild)" +SHA=$(git -C "$WT" rev-parse HEAD 2>/dev/null) +say "binary ref: $SHA" + +# ── The prober: one canary + (M-1) filler prompts, all concurrent, all +# greedy. Non-streaming so `usage.completion_tokens` is authoritative. +cat > "$OUT/probe.py" <<'PYEOF' +"""Drive M concurrent greedy completions; report summed completion_tokens +(from usage — NEVER from requested max_tokens), wall seconds over the +concurrent burst, and the canary text (request 0).""" +import argparse, json, sys, threading, time, urllib.request + +CANARY = "List the first twelve prime numbers, separated by commas, then stop." +FILLERS = [ + "Name the seven continents, one per line.", + "Count from one to twenty in words, separated by spaces.", + "Write the first eight letters of the alphabet, separated by dashes.", + "Name eight colours of the visible spectrum, comma separated.", + "List eight consecutive even numbers starting at two.", + "Name the days of the week in order, one per line.", + "List the first eight powers of two, comma separated.", +] + +def one(port, prompt, max_tokens, seed, timeout): + body = json.dumps({ + "model": "default", "prompt": prompt, "max_tokens": max_tokens, + "temperature": 0, "seed": seed, "stream": False, + }).encode() + req = urllib.request.Request( + f"http://127.0.0.1:{port}/v1/completions", data=body, + headers={"Content-Type": "application/json"}) + with urllib.request.urlopen(req, timeout=timeout) as r: + d = json.load(r) + ch = (d.get("choices") or [{}])[0] + usage = d.get("usage") or {} + return { + "text": ch.get("text", ""), + "finish_reason": ch.get("finish_reason"), + "completion_tokens": usage.get("completion_tokens", 0), + } + +def burst(port, m, max_tokens, seed, timeout): + prompts = [CANARY] + [FILLERS[i % len(FILLERS)] for i in range(m - 1)] + out = [None] * m + errs = [] + def w(i): + try: + out[i] = one(port, prompts[i], max_tokens, seed, timeout) + except Exception as e: # noqa: BLE001 — a leg must record, not crash + errs.append(f"req[{i}]: {e}") + out[i] = {"text": "", "finish_reason": None, "completion_tokens": 0} + t0 = time.monotonic() + ts = [threading.Thread(target=w, args=(i,)) for i in range(m)] + for t in ts: t.start() + for t in ts: t.join() + wall = time.monotonic() - t0 + return out, wall, errs + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--port", type=int, required=True) + ap.add_argument("--m", type=int, required=True) + ap.add_argument("--max-tokens", type=int, required=True) + ap.add_argument("--seed", type=int, required=True) + ap.add_argument("--timeout", type=int, default=1800) + ap.add_argument("--warmup", action="store_true") + a = ap.parse_args() + + if a.warmup: + # Small unmeasured round: load caches, emit the first-dispatch log + # lines, absorb ramp. Its tokens are never counted. + burst(a.port, min(a.m, 8), 16, a.seed, a.timeout) + json.dump({"warmup": True}, sys.stdout) + return + + out, wall, errs = burst(a.port, a.m, a.max_tokens, a.seed, a.timeout) + total = sum(r["completion_tokens"] for r in out) + json.dump({ + "m": a.m, + "wall_s": wall, + "completion_tokens_sum": total, # from usage — the only numerator + "tok_per_s": (total / wall) if wall > 0 else 0.0, + "canary_text": out[0]["text"], + "canary_finish": out[0]["finish_reason"], + "canary_tokens": out[0]["completion_tokens"], + "errors": errs, + "short_streams": sum(1 for r in out if r["completion_tokens"] < 5), + }, sys.stdout) + +main() +PYEOF + +# ── One leg = one fresh server process + warm-up + measured burst. +# $1 arm name, $2 ladder M, then env pairs for the arm. +run_leg() { + local arm="$1" m="$2"; shift 2 + local leg="${arm}.M${m}" + local slog="$OUT/server.$leg.log" + say "leg $leg — starting server ($*)" + + local uqff_args=() + [ -n "$UQFF" ] && uqff_args=(--from-uqff "$UQFF") + # shellcheck disable=SC2086 + env RUST_LOG=info ARC_LOG_FP8_DISPATCH=1 "$@" \ + "$BIN" serve -p "$PORT" -m "$MODEL" -a "$ARCH" "${uqff_args[@]}" \ + --max-seqs "$MAX_SEQS" --prefix-cache-n 0 --seed "$SEED" \ + < /dev/null > "$slog" 2>&1 & + local spid=$! + local up=0 + for _ in $(seq 1 450); do + curl -s -m 5 "http://127.0.0.1:$PORT/v1/models" >/dev/null 2>&1 && { up=1; break; } + kill -0 "$spid" 2>/dev/null || break + sleep 2 + done + [ "$up" = "1" ] || { say "leg $leg — server never came up"; tail -20 "$slog" | tee -a "$STATUS"; return 1; } + + # Provenance: the running server must be the code we think it is. + local logged + logged=$(grep -m1 -o 'git revision: [0-9a-f]*' "$slog" | awk '{print $3}') + if [ -n "$SHA" ] && [ -n "$logged" ] && [ "$logged" != "$SHA" ]; then + say "leg $leg — PROVENANCE FAILED: server $logged != worktree $SHA" + kill -TERM "$spid" 2>/dev/null || true + return 1 + fi + + say "leg $leg — warm-up (unmeasured)" + python3 "$OUT/probe.py" --port "$PORT" --m "$m" --max-tokens "$MAX_TOKENS" \ + --seed "$SEED" --warmup > /dev/null || { say "leg $leg — warm-up failed"; kill -TERM "$spid" 2>/dev/null; return 1; } + + say "leg $leg — measured burst: M=$m x max_tokens=$MAX_TOKENS" + python3 "$OUT/probe.py" --port "$PORT" --m "$m" --max-tokens "$MAX_TOKENS" \ + --seed "$SEED" > "$OUT/$leg.json" || { say "leg $leg — probe failed"; kill -TERM "$spid" 2>/dev/null; return 1; } + + kill -TERM "$spid" 2>/dev/null || true + for _ in $(seq 1 30); do kill -0 "$spid" 2>/dev/null || break; sleep 2; done + kill -KILL "$spid" 2>/dev/null || true + sleep 5 + + # ENGAGEMENT: expected kernel present, both rivals absent, in the server log. + local want forbid1 forbid2 + case "$arm" in + cublaslt) want="dequant_cublaslt"; forbid1="wmma"; forbid2="tiled" ;; + wmma) want="wmma"; forbid1="dequant_cublaslt"; forbid2="tiled" ;; + tiled) want="tiled"; forbid1="wmma"; forbid2="dequant_cublaslt" ;; + *) say "leg $leg — unknown arm '$arm'"; return 1 ;; + esac + local ok=1 + grep -q "\[arc-fp8-dispatch\] path=$want " "$slog" || ok=0 + grep -q "\[arc-fp8-dispatch\] path=$forbid1 " "$slog" && ok=0 + grep -q "\[arc-fp8-dispatch\] path=$forbid2 " "$slog" && ok=0 + { + echo "leg=$leg" + grep '\[arc-fp8-dispatch\]' "$slog" | sed 's/^/ /' + echo " engagement=$([ "$ok" = 1 ] && echo PASS || echo FAIL)" + } >> "$OUT/engagement.txt" + if [ "$ok" != 1 ]; then + say "leg $leg — ENGAGEMENT FAILED (want $want, forbid $forbid1/$forbid2) — leg VOID" + echo '{"void": "engagement"}' > "$OUT/$leg.void" + fi + say "leg $leg — done" + return 0 +} + +# ── The 15 legs. Baseline (canary reference) FIRST: tiled @ smallest M. +MIN_M_NATIVE=1000000 # keeps every real m on the native path in arms (b),(c) +for m in $LADDER; do + run_leg tiled "$m" ARC_FP8_CUBLAS_MIN_M=$MIN_M_NATIVE ARC_NO_FP8_WMMA=1 \ + || die "tiled M=$m leg failed" +done +for m in $LADDER; do + run_leg wmma "$m" ARC_FP8_CUBLAS_MIN_M=$MIN_M_NATIVE ARC_NO_FP8_WMMA=0 \ + || die "wmma M=$m leg failed" +done +for m in $LADDER; do + run_leg cublaslt "$m" ARC_FP8_CUBLAS_MIN_M=5 \ + || die "cublaslt M=$m leg failed" +done + +# ── Summary: floors, canaries, the crossover table, the verdict. +python3 - "$OUT" "$MIN_TOKENS_FLOOR" $LADDER <<'PYEOF' | tee "$OUT/summary.txt" +import json, os, sys + +out, floor = sys.argv[1], int(sys.argv[2]) +ladder = [int(x) for x in sys.argv[3:]] +ARMS = ["tiled", "wmma", "cublaslt"] + +def load(arm, m): + leg = f"{arm}.M{m}" + if os.path.exists(os.path.join(out, leg + ".void")): + return None, "VOID(engagement)" + p = os.path.join(out, leg + ".json") + if not os.path.exists(p): + return None, "MISSING" + with open(p) as f: + d = json.load(f) + if d.get("errors"): + return d, f"ERRORS({len(d['errors'])})" + if d["completion_tokens_sum"] < floor: + return d, f"VOID(floor {d['completion_tokens_sum']}<{floor})" + if d["canary_tokens"] < 5 or not d.get("canary_finish"): + return d, "VOID(degenerate canary)" + return d, "CLEAN" + +base, base_status = load("tiled", ladder[0]) +base_text = base["canary_text"] if base else "" + +def canary(d): + if d is None or base is None: + return "-" + a, b = base_text, d["canary_text"] + if a == b: + return "identical" + n = min(len(a), len(b)) + i = next((j for j in range(n) if a[j] != b[j]), n) + return f"diverges@char{i}" + +print("=== FP8 three-arm threshold sweep ===") +print(f"floor: rate counted only when summed usage.completion_tokens >= {floor}") +print(f"canary baseline: tiled @ M={ladder[0]} [{base_status}]") +print() +print(f"{'M':>5} {'arm':>9} {'tokens':>7} {'wall_s':>8} {'tok/s':>9} {'status':>26} canary_vs_baseline") +rates = {} +for m in ladder: + for arm in ARMS: + d, status = load(arm, m) + rates[(arm, m)] = d["tok_per_s"] if (d and status == "CLEAN") else None + if d: + print(f"{m:>5} {arm:>9} {d['completion_tokens_sum']:>7} {d['wall_s']:>8.1f} " + f"{d['tok_per_s']:>9.1f} {status:>26} {canary(d)}") + else: + print(f"{m:>5} {arm:>9} {'-':>7} {'-':>8} {'-':>9} {status:>26} -") +print() +print("=== CROSSOVER (the number ARC_FP8_CUBLAS_MIN_M may move to) ===") +crossover = None +for m in ladder: + a = rates[("cublaslt", m)] + b = rates[("wmma", m)] + c = rates[("tiled", m)] + if a is None or b is None or c is None: + print(f"M={m:>4}: NOT ESTABLISHED — a leg is void; no claim about this M") + continue + native = min(b, c) + beat = a > native + print(f"M={m:>4}: cublaslt {a:.1f} vs min(wmma {b:.1f}, tiled {c:.1f}) = {native:.1f}" + f" -> {'cuBLASLt WINS' if beat else 'native wins'}") + if beat and crossover is None: + crossover = m +print() +if crossover is not None: + print(f"VERDICT: first clean M where cuBLASLt beats min(WMMA, tiled): {crossover}.") + print(f" Recommended default: ARC_FP8_CUBLAS_MIN_M={crossover} — merge WITH these") + print(" rows quoted in the PR (a threshold is a claim about every value it") + print(" excludes; these are the rows that back the claim).") +else: + print("VERDICT: no clean crossover on this ladder — the default STAYS 512.") + print(" Do not move the threshold from a partial table; re-run the void legs.") +PYEOF + +say "DONE — $OUT/summary.txt (engagement detail: $OUT/engagement.txt)" From 919bb361d8159884f9001d5177b06cc9e3847065 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 19:44:19 +0100 Subject: [PATCH 5/6] =?UTF-8?q?ci(ArcGate):=20free=20SASS=20probe=20for=20?= =?UTF-8?q?the=20FP8=20dense-GEMM=20lane=20=E2=80=94=20LDS=20width=20censu?= =?UTF-8?q?s=20+=20WMMA=20HMMA=20presence?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two questions this settles with nvcc + cuobjdump alone, no GPU: 1. Whether the 4-way bank-conflict cost model for fp8_matmul_tiled describes the compiled artifact: the model is a claim about SCALAR 32-bit LDS in the inner product, and if ptxas had vectorized those shared loads to LDS.128 its arithmetic would be about instructions that do not exist. The probe prints a per-kernel LDS/LDS.64/LDS.128 census. (Note: the +1 pad that kills the conflict also forecloses LDS.128 on tile rows — scalar LDS is the EXPECTED reading on current source.) 2. Whether the never-executed tensor-core GEMM (blockwise_fp8_gemm_wmma.cu) actually contains tensor-core instructions: HARD FAIL if the TU has zero HMMA/HGMMA — the three-arm sweep would otherwise compare cuBLASLt against two scalar kernels while calling one of them WMMA. Teeth: exactly 4 fp8_matmul_tiled kernels asserted (2 dtypes x 2 POW2_SCALE instantiations — a dropped instantiation makes the census about a kernel not in the build), and every tiled kernel must show >=1 LDS or the extractor lost the function body. Wired into the nvcc CI lane as a step beside the QTIP spill gate, so it runs for sm_80 and sm_90 on every Rust/kernel PR; also runnable standalone: arc-tools/fp8_gemm_sass_check.sh sm_90 Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --- .github/workflows/cuda_compile_check.yaml | 10 ++ arc-tools/fp8_gemm_sass_check.sh | 155 ++++++++++++++++++++++ 2 files changed, 165 insertions(+) create mode 100755 arc-tools/fp8_gemm_sass_check.sh diff --git a/.github/workflows/cuda_compile_check.yaml b/.github/workflows/cuda_compile_check.yaml index 2aa62730b..2bd1c949c 100644 --- a/.github/workflows/cuda_compile_check.yaml +++ b/.github/workflows/cuda_compile_check.yaml @@ -305,6 +305,16 @@ jobs: - name: QTIP beam spill gate (LOCAL:0, all geometries) run: bash arc-tools/qtip_beam_res_usage_check.sh sm_${{ matrix.compute_cap }} + # Same free-instrument class as the QTIP gate above, for the blockwise- + # FP8 dense-GEMM lane: asserts all 4 fp8_matmul_tiled instantiations + # exist (2 dtypes x 2 POW2_SCALE scale-index arms), prints the LDS width + # census that decides whether the bank-conflict cost model describes the + # compiled kernel, and HARD-FAILS if the never-executed WMMA GEMM + # contains no tensor-core (HMMA/HGMMA) instructions. nvcc + cuobjdump + # only; no GPU. + - name: FP8 GEMM SASS probe (LDS census + WMMA HMMA presence) + run: bash arc-tools/fp8_gemm_sass_check.sh sm_${{ matrix.compute_cap }} + - name: Compile QTIP CUDA kernels (mistralrs-quant) run: cargo build -p mistralrs-quant --lib --features cuda diff --git a/arc-tools/fp8_gemm_sass_check.sh b/arc-tools/fp8_gemm_sass_check.sh new file mode 100755 index 000000000..6e918ad40 --- /dev/null +++ b/arc-tools/fp8_gemm_sass_check.sh @@ -0,0 +1,155 @@ +#!/usr/bin/env bash +# Parent system: ArcKernels + ArcGate. +# +# FREE SASS PROBE for the blockwise-FP8 dense-GEMM lane. NO GPU REQUIRED: +# nvcc + cuobjdump only, so it runs on the GPU-less nvcc CI lane +# (.github/workflows/cuda_compile_check.yaml) and on any dev box with the +# toolkit. +# +# WHAT IT SETTLES +# --------------- +# 1. `fp8_matmul_tiled` (blockwise_fp8_gemm.cu) — HOW ptxas actually loads the +# shared tiles in the inner product `acc += s_input[ty][k] * s_weight[tx][k]`. +# The source-derived cost model for this kernel (shared traffic 4.13x over +# supply x a 4-way bank conflict from the old +4 tile padding) is a claim +# about SCALAR 32-bit LDS. If ptxas had already vectorized those loads to +# LDS.128, the conflict arithmetic would be about instructions that do not +# exist and the model's 7% agreement would be a coincidence. This probe +# prints the per-kernel LDS width census so that question is settled by the +# compiled artifact, not by reading C++. +# NOTE the padding interaction: the +1 pad (stride 33 floats, PR #215) +# that kills the 4-way bank conflict ALSO forecloses LDS.128 on tile rows +# (rows are only 4-byte aligned at stride 33). Scalar LDS here is +# therefore the EXPECTED outcome on current source, not a defect — the +# probe exists so a future retune that trades padding for vector width +# does it with the census in front of them. +# 2. `blockwise_fp8_gemm_wmma.cu` — whether the "tensor-core" GEMM actually +# contains tensor-core instructions (HMMA/HGMMA). That kernel has NEVER +# executed on any GPU; if ptxas lowered it without a single matrix-multiply +# instruction, every derivation in its header is void. HARD FAIL if absent. +# +# TEETH (a parser that matches nothing must not pass) +# --------------------------------------------------- +# * Kernel-count assertion: the tiled TU must yield exactly 4 +# `fp8_matmul_tiled` kernels (2 dtypes x 2 POW2_SCALE instantiations — +# see the launcher in blockwise_fp8_gemm.cu). A count drift means an +# instantiation was dropped and the census is about a kernel not in the +# build. +# * Every tiled kernel must show at least one LDS — a tiled shared-memory GEMM +# with zero shared loads means the SASS extractor lost the function body. +# +# Usage: arc-tools/fp8_gemm_sass_check.sh [sm_90] +set -euo pipefail + +ARCH="${1:-sm_90}" +REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +KDIR="${REPO_ROOT}/mistralrs-quant/kernels/blockwise_fp8" +OUT="$(mktemp -d)" +trap 'rm -rf "${OUT}"' EXIT + +# Mirror mistralrs-quant/build.rs exactly: a different flag set is a different +# instruction selection, and therefore a probe of a kernel we do not ship. +NVCC_ARGS=( + -std=c++17 + -O3 + -U__CUDA_NO_HALF_OPERATORS__ + -U__CUDA_NO_HALF_CONVERSIONS__ + -U__CUDA_NO_HALF2_OPERATORS__ + -U__CUDA_NO_BFLOAT16_CONVERSIONS__ + --expt-relaxed-constexpr + --expt-extended-lambda + --use_fast_math +) + +echo "== blockwise-FP8 GEMM SASS probe ==" +nvcc --version | tail -1 +echo "arch: ${ARCH}" +echo + +# --------------------------------------------------------------------------- +# Per-function SASS census. cuobjdump -sass emits one "Function : " +# header per kernel; slice on those and count instruction mnemonics inside +# each slice with awk. +# --------------------------------------------------------------------------- +census() { # $1 = cubin, $2 = function-name substring, $3 = extended regex to count + cuobjdump -sass "$1" | awk -v want="$2" -v pat="$3" ' + /Function :/ { fname = $NF; infn = index(fname, want) > 0; if (infn) { order[++n] = fname; count[fname] = 0 } } + infn && $0 ~ pat { count[fname]++ } + END { for (i = 1; i <= n; i++) printf "%s %d\n", order[i], count[order[i]] } + ' +} + +# --------------------------------------------------------------------------- +# 1. The scalar tiled GEMM: LDS width census. +# --------------------------------------------------------------------------- +echo "--- blockwise_fp8_gemm.cu : fp8_matmul_tiled ---" +nvcc -cubin -arch="${ARCH}" "${NVCC_ARGS[@]}" \ + -o "${OUT}/gemm.cubin" "${KDIR}/blockwise_fp8_gemm.cu" + +if command -v c++filt >/dev/null 2>&1; then + cuobjdump -sass "${OUT}/gemm.cubin" | grep -E 'Function :' | c++filt | grep fp8_matmul_tiled || true +fi + +# LDS.128 / LDS.64 are prefixes of the wide forms; plain "LDS " (or LDS.U) +# without a width suffix is the scalar 32-bit load. +lds_any="$(census "${OUT}/gemm.cubin" fp8_matmul_tiled '\|\')" +lds_64="$(census "${OUT}/gemm.cubin" fp8_matmul_tiled '\')" + +echo "kernel LDS(all) LDS.128 LDS.64" +fail=0 +n_tiled=0 +while read -r name total; do + [ -z "${name}" ] && continue + n_tiled=$((n_tiled + 1)) + w128="$(awk -v k="${name}" '$1 == k { print $2 }' <<<"${lds_128}")" + w64="$(awk -v k="${name}" '$1 == k { print $2 }' <<<"${lds_64}")" + printf '%-44s %8s %7s %6s\n' "${name:0:44}" "${total}" "${w128:-0}" "${w64:-0}" + if [ "${total}" -eq 0 ]; then + echo "FAIL: ${name} has ZERO shared-memory loads — the SASS extractor lost" + echo " the function body; every census line above is unverified." + fail=1 + fi +done <<<"${lds_any}" + +if [ "${n_tiled}" -ne 4 ]; then + echo "FAIL: expected 4 fp8_matmul_tiled kernels (2 dtypes x 2 POW2_SCALE" + echo " instantiations), found ${n_tiled}. An instantiation was dropped;" + echo " the census is about kernels that are not in the build." + fail=1 +fi + +total_128="$(awk '{ s += $2 } END { print s + 0 }' <<<"${lds_128}")" +if [ "${total_128}" -eq 0 ]; then + echo "VERDICT[tiled]: scalar LDS only, no LDS.128 — the shared loads are" + echo " 32-bit, so the bank-conflict arithmetic (stride mod 32) is about" + echo " instructions that really exist, and the +1-pad fix is live." +else + echo "VERDICT[tiled]: ptxas emitted LDS.128 — the scalar-LDS conflict model" + echo " does NOT describe this build; re-derive before quoting it." +fi +echo + +# --------------------------------------------------------------------------- +# 2. The WMMA GEMM: tensor-core instructions must exist. +# --------------------------------------------------------------------------- +echo "--- blockwise_fp8_gemm_wmma.cu : HMMA presence ---" +nvcc -cubin -arch="${ARCH}" "${NVCC_ARGS[@]}" \ + -o "${OUT}/wmma.cubin" "${KDIR}/blockwise_fp8_gemm_wmma.cu" + +hmma="$(cuobjdump -sass "${OUT}/wmma.cubin" | grep -cE '\<(HMMA|HGMMA)\>' || true)" +echo "HMMA/HGMMA instructions in TU: ${hmma}" +if [ "${hmma}" -eq 0 ]; then + echo "FAIL: the tensor-core GEMM contains no HMMA/HGMMA. It has never run on" + echo " hardware; if it also has no tensor-core instructions, every" + echo " derivation in its header is void and the three-arm sweep would" + echo " be comparing cuBLASLt against two scalar kernels." + fail=1 +fi +echo + +if [ "${fail}" -ne 0 ]; then + echo "RESULT: FAILED — see above." + exit 1 +fi +echo "RESULT: census complete (${ARCH}); 4 tiled kernels present, WMMA carries tensor-core instructions." From 4c9c07e5b7b16375057860287c529ec41e5332fb Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 20:21:36 +0100 Subject: [PATCH 6/6] =?UTF-8?q?fix(ArcGate):=20SASS=20probe=20=E2=80=94=20?= =?UTF-8?q?replace=20GNU-awk-unsupported=20\<=20\>=20ERE=20escapes=20with?= =?UTF-8?q?=20cuobjdump=20-fun=20extraction?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The probe's first CI run went red on itself, honestly: GNU awk treats \< \> (and warns on \.) as plain characters in ERE, so the awk body-splitter's counting regexes matched NOTHING, the per-function census returned zero LDS for all 4 fp8_matmul_tiled instantiations, and the fail-on-zero guard refused to trust its own zero — which is exactly what it exists for. This revision removes regex body-splitting entirely: the mangled kernel names come from the Function listing (awk '{print $NF}', no regex), and `cuobjdump -sass -fun ` extracts each function's SASS directly. All counting is grep -cE with POSIX classes and bracket expressions only — no \< \> anywhere. Also, per review: * The two failure modes are now DISTINGUISHED: FAIL[extractor] (no instruction lines came back — the census is unverified) vs FAIL[zero-LDS] (body extracted, genuinely no shared loads — the conflict model has no subject). Body presence is judged on /**/ instruction lines, so the two cannot be confused. * The VERDICT[tiled] line prints ONLY when the per-function census is complete and non-empty; otherwise it prints NOT ESTABLISHED. A TU-level cross-check line (whole cubin, includes the GEMV/MoE kernels) is always printed, labeled informational. * HMMA counting moved off the \< \> GNUism onto the same POSIX pattern. Verified locally against a stub cuobjdump with a realistic SASS fixture: happy path (4 kernels, LDSM excluded, LDS/LDS.64/LDS.128 counted per function, scalar arithmetic correct, exit 0) plus three negative controls (broken -fun -> FAIL[extractor] + verdict NOT ESTABLISHED; LDS-free body -> FAIL[zero-LDS]; HMMA-stripped WMMA TU -> hard fail), all exit 1. The real cuobjdump run is CI's — if -fun ever returns nothing there, the probe fails loudly rather than lying. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --- arc-tools/fp8_gemm_sass_check.sh | 124 ++++++++++++++++++++++--------- 1 file changed, 89 insertions(+), 35 deletions(-) diff --git a/arc-tools/fp8_gemm_sass_check.sh b/arc-tools/fp8_gemm_sass_check.sh index 6e918ad40..48634f447 100755 --- a/arc-tools/fp8_gemm_sass_check.sh +++ b/arc-tools/fp8_gemm_sass_check.sh @@ -67,17 +67,33 @@ echo "arch: ${ARCH}" echo # --------------------------------------------------------------------------- -# Per-function SASS census. cuobjdump -sass emits one "Function : " -# header per kernel; slice on those and count instruction mnemonics inside -# each slice with awk. +# Per-function SASS extraction. +# +# PORTABILITY LESSON (this script's first CI run went red on itself): GNU awk +# does not honor \< \> word-boundary escapes in ERE — it warns "escape +# sequence treated as plain character" and matches nothing — so the original +# awk body-splitter returned zero LDS for every function and the fail-on-zero +# guard below correctly refused to trust it. This revision does no regex body +# splitting at all: cuobjdump lists the mangled kernel names, and +# `cuobjdump -sass -fun ` extracts exactly one function's SASS. All +# counting is grep -cE with POSIX classes and bracket expressions only — +# no \< \> anywhere. # --------------------------------------------------------------------------- -census() { # $1 = cubin, $2 = function-name substring, $3 = extended regex to count - cuobjdump -sass "$1" | awk -v want="$2" -v pat="$3" ' - /Function :/ { fname = $NF; infn = index(fname, want) > 0; if (infn) { order[++n] = fname; count[fname] = 0 } } - infn && $0 ~ pat { count[fname]++ } - END { for (i = 1; i <= n; i++) printf "%s %d\n", order[i], count[order[i]] } - ' +fn_names() { # $1 = cubin, $2 = fixed-string name filter + cuobjdump -sass "$1" | grep 'Function : ' | awk '{print $NF}' | grep -F "$2" || true +} +count() { # $1 = file, $2 = POSIX ERE + grep -cE "$2" "$1" || true } +# SASS instruction lines carry a /**/ comment; body-vs-empty is +# judged on those, so "extractor returned nothing" and "zero LDS" cannot be +# confused with each other. +RE_INSN='/[*][0-9a-fA-F]+[*]/' +# LDSM ('M' after LDS) must not count; the char after LDS is '.' or space. +RE_LDS_ALL='[[:space:]]LDS[[:space:].]' +RE_LDS_128='[[:space:]]LDS[.](U[.])?128[[:space:]]' +RE_LDS_64='[[:space:]]LDS[.](U[.])?64[[:space:]]' +RE_MMA='[[:space:]](HMMA|HGMMA)[[:space:].]' # --------------------------------------------------------------------------- # 1. The scalar tiled GEMM: LDS width census. @@ -87,30 +103,48 @@ nvcc -cubin -arch="${ARCH}" "${NVCC_ARGS[@]}" \ -o "${OUT}/gemm.cubin" "${KDIR}/blockwise_fp8_gemm.cu" if command -v c++filt >/dev/null 2>&1; then - cuobjdump -sass "${OUT}/gemm.cubin" | grep -E 'Function :' | c++filt | grep fp8_matmul_tiled || true + cuobjdump -sass "${OUT}/gemm.cubin" | grep 'Function : ' | c++filt | grep fp8_matmul_tiled || true fi -# LDS.128 / LDS.64 are prefixes of the wide forms; plain "LDS " (or LDS.U) -# without a width suffix is the scalar 32-bit load. -lds_any="$(census "${OUT}/gemm.cubin" fp8_matmul_tiled '\|\')" -lds_64="$(census "${OUT}/gemm.cubin" fp8_matmul_tiled '\')" - -echo "kernel LDS(all) LDS.128 LDS.64" fail=0 +extract_ok=1 n_tiled=0 -while read -r name total; do - [ -z "${name}" ] && continue +sum128=0 +names="$(fn_names "${OUT}/gemm.cubin" fp8_matmul_tiled)" +if [ -z "${names}" ]; then + echo "FAIL[extractor]: no fp8_matmul_tiled function names in the cubin's" + echo " SASS listing — the name pass found nothing, so there is no census." + fail=1 + extract_ok=0 +fi + +echo "kernel insns LDS(all) LDS.128 LDS.64 scalar" +for name in ${names}; do n_tiled=$((n_tiled + 1)) - w128="$(awk -v k="${name}" '$1 == k { print $2 }' <<<"${lds_128}")" - w64="$(awk -v k="${name}" '$1 == k { print $2 }' <<<"${lds_64}")" - printf '%-44s %8s %7s %6s\n' "${name:0:44}" "${total}" "${w128:-0}" "${w64:-0}" - if [ "${total}" -eq 0 ]; then - echo "FAIL: ${name} has ZERO shared-memory loads — the SASS extractor lost" - echo " the function body; every census line above is unverified." + sfile="${OUT}/tiled_${n_tiled}.sass" + cuobjdump -sass -fun "${name}" "${OUT}/gemm.cubin" > "${sfile}" 2>/dev/null || true + insns="$(count "${sfile}" "${RE_INSN}")" + lds_all="$(count "${sfile}" "${RE_LDS_ALL}")" + w128="$(count "${sfile}" "${RE_LDS_128}")" + w64="$(count "${sfile}" "${RE_LDS_64}")" + scalar=$((lds_all - w128 - w64)) + sum128=$((sum128 + w128)) + printf '%-48s %6s %8s %7s %6s %6s\n' "${name:0:48}" "${insns}" "${lds_all}" "${w128}" "${w64}" "${scalar}" + if [ "${insns}" -eq 0 ]; then + echo "FAIL[extractor]: ${name}: 'cuobjdump -sass -fun' returned no" + echo " instruction lines. This is an EXTRACTION failure, not a zero-LDS" + echo " kernel — this function's census row is unverified." + fail=1 + extract_ok=0 + elif [ "${lds_all}" -eq 0 ]; then + echo "FAIL[zero-LDS]: ${name}: body extracted (${insns} instructions) but" + echo " it contains no shared-memory loads at all. A shared-tile GEMM" + echo " without LDS means the wrong function was extracted or the kernel" + echo " no longer stages tiles — either way the conflict model has no" + echo " subject here." fail=1 fi -done <<<"${lds_any}" +done if [ "${n_tiled}" -ne 4 ]; then echo "FAIL: expected 4 fp8_matmul_tiled kernels (2 dtypes x 2 POW2_SCALE" @@ -119,14 +153,32 @@ if [ "${n_tiled}" -ne 4 ]; then fail=1 fi -total_128="$(awk '{ s += $2 } END { print s + 0 }' <<<"${lds_128}")" -if [ "${total_128}" -eq 0 ]; then - echo "VERDICT[tiled]: scalar LDS only, no LDS.128 — the shared loads are" - echo " 32-bit, so the bank-conflict arithmetic (stride mod 32) is about" - echo " instructions that really exist, and the +1-pad fix is live." +# TU-level cross-check: computed on the WHOLE cubin dump, independent of the +# per-function extractor, so it survives an extractor failure and gives the +# reader a signal either way. Informational — the per-function census is the +# only thing the VERDICT may rest on (the TU includes the GEMV/MoE kernels). +cuobjdump -sass "${OUT}/gemm.cubin" > "${OUT}/gemm.sass" 2>/dev/null || true +echo "TU-level cross-check (whole cubin, informational):" \ + "LDS(all)=$(count "${OUT}/gemm.sass" "${RE_LDS_ALL}")" \ + "LDS.128=$(count "${OUT}/gemm.sass" "${RE_LDS_128}")" \ + "LDS.64=$(count "${OUT}/gemm.sass" "${RE_LDS_64}")" + +# The VERDICT prints ONLY on a complete, non-empty per-function census. +# Printing it off TU-level counts (or off an empty extraction) would be a +# lying instrument: confident prose about numbers that were never measured. +if [ "${extract_ok}" -eq 1 ] && [ "${n_tiled}" -eq 4 ] && [ "${fail}" -eq 0 ]; then + if [ "${sum128}" -eq 0 ]; then + echo "VERDICT[tiled]: scalar LDS only, no LDS.128 — the shared loads are" + echo " 32-bit, so the bank-conflict arithmetic (stride mod 32) is about" + echo " instructions that really exist, and the +1-pad fix is live." + else + echo "VERDICT[tiled]: ptxas emitted LDS.128 — the scalar-LDS conflict model" + echo " does NOT describe this build; re-derive before quoting it." + fi else - echo "VERDICT[tiled]: ptxas emitted LDS.128 — the scalar-LDS conflict model" - echo " does NOT describe this build; re-derive before quoting it." + echo "VERDICT[tiled]: NOT ESTABLISHED — the per-function census is empty or" + echo " partial (see FAILs above). Do not quote any LDS conclusion from this" + echo " run; the TU-level line above is a cross-check, not a verdict." fi echo @@ -137,7 +189,8 @@ echo "--- blockwise_fp8_gemm_wmma.cu : HMMA presence ---" nvcc -cubin -arch="${ARCH}" "${NVCC_ARGS[@]}" \ -o "${OUT}/wmma.cubin" "${KDIR}/blockwise_fp8_gemm_wmma.cu" -hmma="$(cuobjdump -sass "${OUT}/wmma.cubin" | grep -cE '\<(HMMA|HGMMA)\>' || true)" +cuobjdump -sass "${OUT}/wmma.cubin" > "${OUT}/wmma.sass" 2>/dev/null || true +hmma="$(count "${OUT}/wmma.sass" "${RE_MMA}")" echo "HMMA/HGMMA instructions in TU: ${hmma}" if [ "${hmma}" -eq 0 ]; then echo "FAIL: the tensor-core GEMM contains no HMMA/HGMMA. It has never run on" @@ -152,4 +205,5 @@ if [ "${fail}" -ne 0 ]; then echo "RESULT: FAILED — see above." exit 1 fi -echo "RESULT: census complete (${ARCH}); 4 tiled kernels present, WMMA carries tensor-core instructions." +echo "RESULT: census complete (${ARCH}); 4 tiled kernels present (per-function" +echo " extraction verified non-empty), WMMA TU carries ${hmma} HMMA/HGMMA."