perf(ArcKernels/ArcQuant): FP8 dense-GEMM lane — pow2 scale shifts, dispatch engagement log, three-arm sweep harness, SASS probe

[heydryft]

FP8 dense-GEMM lane: scale-division fix, dispatch engagement log, three-arm sweep harness, free SASS probe

Parent systems: ArcKernels + ArcQuant/Legacy (blockwise FP8), ArcLab, ArcGate. Targets release/openrouter-ready.

At B=256 decode, fp8_matmul_tiled serves all 7 dense projections per layer plus the lm_head — ~50% of the step in the last profile. This PR ships everything that can be shipped for that lane without a GPU, plus the instruments the box A/B needs, and refutes one premise of the brief that was already fixed upstream.

What changed

1. Cherry-pick 445b92063 (rank-3→2-D flatten before the cuBLASLt branch). It was NOT an ancestor of this base — it lives only on fix/fp8-cublas-crossover (PR fix(ArcKernels): FP8 threshold was 512 because 512 was the only point measured; cuBLASLt branch never flattened to 2-D #201). Without it, lowering ARC_FP8_CUBLAS_MIN_M on V4 decode selects a B-way batched GEMM with m=1 and the threshold A/B returns a confident false negative. Applied cleanly.

2. Pow2 shift fast path for the tiled GEMM's scale indexing (mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu). get_scale performed two runtime integer divisions per weight element staged into shared memory; the divisors are kernel arguments, so each expands to a ~15–20-instruction sequence. New POW2_SCALE template arm: the host launcher tests pow2-ness once per launch and passes log2(block_size_{y,x}); the kernel uses two shifts. The shipped geometry [128, 128] always takes it; non-pow2 keeps the division path. Bit-identical by construction: n >> log2(d) == n / d for the non-negative indices used — same quotient, same scale word, same arithmetic after it. C ABI unchanged; no Rust edits needed for this item.

3. ARC_LOG_FP8_DISPATCH=1 engagement lines (ArcLab). Each FP8 dispatch path prints ONE line per process: [arc-fp8-dispatch] path=<gemv_wide|gemv_warp|wmma|tiled|dequant_cublaslt> first_shape=mM_nN_kK. Value-read via env_flag_is_set, OnceLock/LazyLock-latched, one relaxed atomic swap per call when on — it cannot perturb the rates a sweep records. Registered in mistralrs-core/tests/capability_reachability.rs (run: 11/11 pass).

4. arc-tools/fp8_threshold_sweep.sh — the three-arm sweep the arc_fp8_cublas_min_m and fp8_wmma_enabled docs demand before the threshold moves. One binary, env toggles only, fresh server per leg (gates are latched per process). 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 arms (a) cuBLASLt (MIN_M=5), (b) WMMA (MIN_M=1000000, ARC_NO_FP8_WMMA=0), (c) tiled (MIN_M=1000000, ARC_NO_FP8_WMMA=1). Per leg: engagement assertion (expected path= present, both rivals absent — a WMMA leg whose log shows tiled is eligibility silently failing and is VOID, loudly), a minimum-token floor (rate divides by summed usage.completion_tokens, never by requested tokens), and a greedy canary vs the baseline leg (tiled @ M=5; cross-arm divergence is expected numerics and is reported with its first index — a degenerate stream voids the leg). The summary names the first clean M where (a) beats min((b), (c)); no clean crossover ⇒ the default stays 512. Not run here — no GPU (D14). No perf claim is made anywhere in this PR.

5. arc-tools/fp8_gemm_sass_check.sh + a step in the nvcc CI lane (.github/workflows/cuda_compile_check.yaml, beside the QTIP spill gate, sm_80 + sm_90): compiles the tiled TU and prints a per-kernel LDS/LDS.64/LDS.128 census (settles whether the 4-way-conflict model's scalar-LDS premise describes the compiled artifact — note the +1 pad that kills the conflict also forecloses LDS.128 on tile rows, so scalar LDS is the expected reading), asserts all 4 fp8_matmul_tiled instantiations exist, and HARD-FAILS if the never-executed WMMA GEMM contains zero HMMA/HGMMA.

Premises verified / refuted against this base

• Refuted: "shared tiles declared [32][BLOCK_K + 4]" — the bank-conflict fix (+4 → +1, stride 33, with the bank arithmetic commented in-kernel) already landed via PR fix(ArcKernels): FP8 bank conflict + the dispatch threshold's provenance, rebased onto master #215 (250dd6d7f), which is an ancestor of release/openrouter-ready. Nothing to do; this PR does not touch the padding.
• Kernel moved to mistralrs-quant/kernels/blockwise_fp8/blockwise_fp8_gemm.cu; inner loop at :121 (pre-edit), get_scale divisions at :38–39 — confirmed.
• Dispatch: fp8_gemv_max_m (default 4) at ops.rs:914–928; use_wmma at ops.rs:1197–1200; ARC_NO_FP8_WMMA value-read at ops.rs:1012; arc_fp8_cublas_min_m (default 512) at mod.rs:657–671 — confirmed (some line numbers had drifted from the brief).
• WMMA kernel header says UNVERIFIED ON HARDWARE — never run: confirmed. This PR does not modify that kernel; the box A/B decides its fate.
• 445b92063 not an ancestor: confirmed, cherry-picked (item 1).

Follow-up note: per-tile scale staging in shared memory (deliberately NOT done)

With TILE=32 and any block_size_y % 32 == 0 && block_size_x % 32 == 0 geometry (the shipped [128,128] qualifies), a weight tile spans exactly one scale block, so the whole tile's scale is a single scalar. Staging it would need: a per-k_tile __shared__ float s_tile_scale written by one thread, plus a __syncthreads() between that write and the load loop that reads it — a third barrier per k-tile that does not exist today (the current two sit after the load loops and after the FMA loop). After the shift fix, the per-element cost that staging would remove is a same-address __ldg (an L1 broadcast) — trading that for an extra barrier per tile is plausibly a wash or a regression, so it is written up here instead of shipped. If a box profile still shows the scale fetch after this PR, the contained shape is: compute s_tile_scale for tile t+1 during tile t (prefetch into the existing end-of-iteration barrier), uniformity-gated on the host like POW2_SCALE.

CI / verification

• cargo check -p mistralrs-quant: pass. Scoped clippy lane (arc-bench arc-engine arc-cuda-graph arc-cli mistralrs-quant --tests --examples -- -D warnings): pass. cargo test -p mistralrs-core --test capability_reachability: 11/11.
• The .cu edit and all cfg(cuda) Rust are compile-unverified locally (macOS, no nvcc); the nvcc lane in cuda_compile_check.yaml builds both for sm_80/sm_90 on this PR, and the new SASS step additionally asserts the 4 tiled instantiations.
• Both scripts bash -n clean; embedded Python parses.

🤖 Generated with Claude Code

https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf

[github-actions]

Code Metrics Report

━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
 Language              Files        Lines         Code     Comments       Blanks
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
 C Header                  5          305          210           52           43
 CSS                       2         1181         1036           34          111
 CUDA                     81        29475        20161         6275         3039
 Dockerfile                1           39           22            8            9
 JavaScript               16         3546         2676          482          388
 Jinja2                    7          694          656            5           33
 JSON                     75         4896         4893            0            3
 Makefile                  1            6            5            0            1
 Metal Shading Lan|       33        12224         9431         1142         1651
 PowerShell                1          300          227           30           43
 Python                  147        15371        12677          824         1870
 Shell                    42        10304         6858         2769          677
 Plain Text                4         3801            0         2479         1322
 TOML                     33         1498         1294           54          150
 YAML                      3           25           23            2            0
─────────────────────────────────────────────────────────────────────────────────
 HTML                      4         2687         2604           43           40
 |- CSS                    2          543          479           37           27
 |- JavaScript             1         1233         1215           12            6
 (Total)                             4463         4298           92           73
─────────────────────────────────────────────────────────────────────────────────
 Jupyter Notebooks         4          122           83           23           16
 |- Markdown               1           60           30           22            8
 |- Python                 1          122          113            1            8
 (Total)                              304          226           46           32
─────────────────────────────────────────────────────────────────────────────────
 Markdown                213        48932            0        38081        10851
 |- BASH                  72         1655         1203          331          121
 |- C                      4           19           19            0            0
 |- CUDA                   2           84           56           16           12
 |- JSON                  19          779          779            0            0
 |- PowerShell             1            1            1            0            0
 |- Python                23         1008          787          113          108
 |- Rust                  68         2063         1727           78          258
 |- TOML                   6          207          164            0           43
 |- YAML                   5           41           36            5            0
 (Total)                            54789         4772        38624        11393
─────────────────────────────────────────────────────────────────────────────────
 Rust                    681       340988       293252        18096        29640
 |- Markdown             504        32562          471        28063         4028
 (Total)                           373550       293723        46159        33668
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
 Total                  1353       516771       363188        99077        54506
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━