fix(ArcKernels): FP8 bank conflict + the dispatch threshold's provenance, rebased onto master

[heydryft]

Replaces #201, which was red for exactly one reason: its base was
release/openrouter-ready and the Base branch lane refuses any base that is
not master. Every other lane on #201 was green. This branch is cherry-picked
onto origin/master (6ffdac7ae at the time of writing), so
release/openrouter-ready's history — including 145bbbc0a, the
docs/engineering/OPENROUTER_READY.md commit that exists only because
gh pr create refused an empty branch — does not come along. That file is
absent from this branch; git diff --stat origin/master...HEAD is two files.

What landed, and what did not

#201 carried three commits. Two land; one was already fixed on master in a
better form, and cherry-picking it now would be a regression.

Landed — b5de0f1d4, the 4-way shared-memory bank conflict. Unchanged, and
its author's ⚠️ NEVER COMPILED / must-be-A/B'd label is kept verbatim. Verified
the arithmetic against the shipped constants rather than the commit message:
blockwise_fp8_gemm.cu:623-624 and :641-642 instantiate fp8_matmul_tiled
with TILE = 32, TILE_K = 32, so blockDim.x = 32 and a warp is 32
consecutive tx. The inner product at :110 reads s_weight[tx][k], stride
BLOCK_K + pad floats. At pad = 4: stride 36, bank = (4·tx + k) mod 32
takes 8 distinct values, 4 lanes each — a 4-way conflict on every FMA. At
pad = 1: stride 33, bank = (tx + k) mod 32, all 32 banks, conflict-free.
Output is bit-identical by construction (the k loop is bounded k < BLOCK_K,
so the pad columns are written and read by nobody), and shared memory drops
9216 B → 8448 B per block. The predicted 562 ms → ~134 ms stays UNVERIFIED:
there is no nvcc on this host and this change was made with zero GPU.

Landed with the constant reverted — 52a8750f5, the dispatch threshold.
The doc lands, the 512 → 5 does not. Reasons in the next section.

Dropped — 445b92063, the rank-3 flatten. Already on master in the general
form: UnquantLinear::forward flattens at
mistralrs-quant/src/unquantized/mod.rs:155-161, landed by #202
(611d236c2), whose own comment names 445b92063 and says it "subsumes the
caller-side workaround". Applying the caller-side version now is not merely
redundant — master's flatten is deliberately gated && !a.device().is_cpu(),
pinned by cpu_forward_still_takes_the_broadcast_path_bit_for_bit
(unquantized/mod.rs:857), because flattening changes CPU F16 accumulation
order and re-baselines every V4 batched-vs-solo tolerance at once (that comment
records 1.57e-2 measured against a 1e-2 tolerance on x86). 445b92063 is not
device-gated, so it would reintroduce the CPU flatten for blockwise-FP8 layers
only — where no test is watching.

The threshold: shipping 512, and what would settle it

ARC_FP8_CUBLAS_MIN_M stays at 512. Not because 512 is right — the doc this
PR lands says plainly that 512 was a single prefill point, and that
M = 5..511 fell through to a kernel whose own comment admits it has "no
tensor-core instruction anywhere in it". It stays at 512 because the sweep
behind 5 no longer measures the arm that runs.

52a8750f5 was authored 01:47 on 2026-08-21. At 07:04 the same day, #200
(9ee459191) merged a tensor-core blockwise-FP8 GEMM. fp8_blockwise_matmul_impl
now selects it by default for everything fp8_gemv_warp does not own
(use_wmma, ops.rs), and our shipped weight_block_size [128, 128] clears
fp8_wmma_eligible against that kernel's N_BLK = 64, K_BLK = 128
(128 % 64 == 0, 128 % 128 == 0). So the sweep's "native" arm — the scalar
fp8_matmul_tiled — is now reachable only with ARC_NO_FP8_WMMA=1, and
lowering this constant to 5 would route every M >= 5 into the dequantize path
and leave #200's kernel unreachable at every batch size that matters. Silently:
it compiles and the tests pass. That kernel's own docs already state the merge
rule — three arms, not two, and "do not merge a threshold on the two-arm
result".

Two more reasons the rows cannot be cashed as they stand: they predate the
flatten, so arm (a) was dispatching a B-way batched GEMV with m = 1 and
understates cuBLASLt; and arm (a) pays an uncached full-model dequantize_w()
on every forward (~12.7 GB/step) that arms (b) and (c) do not — nothing in this
PR caches it, so 5 would pay that at every size.

The measurement that settles the constant: re-run the M sweep with the
flatten 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. Moving 512 to 5 before that exists would be a second constant
with no measurement behind it, which is the defect the doc was written to
record. The sweep table, its clean-row criteria, and the M=1 → 0.72x GEMV
floor are all preserved in the doc so the re-run starts from data, not memory.

A conditional threshold (5 when WMMA is ineligible, 512 otherwise) was
considered and rejected: it would make ARC_NO_FP8_WMMA=1 flip two things at
once and destroy #200's control arm.

Worth a separate change: cuBLASLt block-scaled FP8 is reachable

CUBLASLT_MATMUL_MATRIX_SCALE_BLK128x128_32F exists at our pin — cudarc
0.19.4 (Cargo.lock:1333-1336, arriving via the candle git dep at
Cargo.toml:52-53), at
~/.cargo/registry/src/index.crates.io-*/cudarc-0.19.4/src/cublaslt/sys/mod.rs:747
(cfg cuda-12090/cuda-13000) and :759 (cfg cuda-13010/cuda-13020).
CUBLASLT_MATMUL_DESC_A_SCALE_MODE / _B_SCALE_MODE are at :618-619,
:655-656, :693-694 (first gated at cuda-12080). Candle enables
cuda-version-from-build-system, so the block-128 variant needs a CUDA 12.9+
toolkit at build time; 12.8 gets the scale-mode attributes but not
BLK128x128_32F. Our weight_block_size [128, 128] is an exact match for the
descriptor's geometry.

Our own wrapper only sets *_SCALE_POINTER today
(mistralrs-quant/src/cublaslt/matmul.rs:182-200), but the generic setter
set_matmul_desc_attribute (cudarc-0.19.4/src/cublaslt/result.rs:111) takes
any descriptor attribute, so no cudarc bump is needed. The caveat from the
earlier check holds: A must also be FP8 — there is no bf16×fp8 mixed-type
path — so this needs an activation quantiser before it is usable. If that is
built, the dequantize disappears entirely and the threshold question above
changes shape. Not built here; reporting reachability only.

[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        29464        20161         6264         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                212        47877            0        37224        10653
 |- BASH                  72         1655         1203          331          121
 |- C                      3           17           17            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                  67         2058         1723           77          258
 |- TOML                   6          207          164            0           43
 |- YAML                   5           41           36            5            0
 (Total)                            53727         4766        37766        11195
─────────────────────────────────────────────────────────────────────────────────
 Rust                    681       340749       293058        18073        29618
 |- Markdown             504        32196          471        27754         3971
 (Total)                           372945       293529        45827        33589
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
 Total                  1352       515093       362988        97876        54229
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━