Conversation
…nly point measured
`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, 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.)
The crossover is not 512; it sits immediately above the GEMV's domain. Default
is now 5, the first row `fp8_gemv_warp` does not own. M = 5..7 is unmeasured,
bounded below by that kernel and above by the +32% at M=8.
KEEP THE FLOOR: at M=1 cuBLASLt is 0.72x, so forcing the library path over the
dedicated GEMV costs 28% of b=1 decode. That row is why this is 5 and not 1.
Mechanism: the tiled kernel is scalar, so its cost is instruction-bound and
grows with M while the memory controller idles; cuBLASLt dequantizes once and
hands a tensor-core GEMM the whole tile. The counter shows it --
memory-controller utilisation moved 2-3% -> 6-19% across the sweep, the first
intervention this session to move it at all.
Recorded in the doc as a recurring failure, because this 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.
Fast path by default; `ARC_FP8_CUBLAS_MIN_M` remains the kill switch and a
value above any real batch restores the previous always-native dispatch.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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) <noreply@anthropic.com>
…nflict⚠️ 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) <noreply@anthropic.com>
Code Metrics Report━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Language Files Lines Code Comments Blanks ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ C Header 5 305 210 52 43 CSS 2 1181 1036 34 111 CUDA 78 27867 19381 5568 2918 Dockerfile 1 39 22 8 9 JavaScript 16 3546 2676 482 388 Jinja2 7 694 656 5 33 JSON 74 4600 4597 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 10258 6849 2734 675 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 204 45330 0 35183 10147 |- BASH 72 1655 1203 331 121 |- C 3 17 17 0 0 |- CUDA 2 84 56 16 12 |- JSON 18 708 708 0 0 |- PowerShell 1 1 1 0 0 |- Python 23 1008 787 113 108 |- Rust 66 2051 1716 77 258 |- TOML 6 207 164 0 43 |- YAML 5 41 36 5 0 (Total) 51102 4688 35725 10689 ───────────────────────────────────────────────────────────────────────────────── Rust 678 335044 288629 17172 29243 |- Markdown 497 30361 471 26183 3707 (Total) 365405 289100 43355 32950 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Total 1337 502989 357396 92632 52961 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
…kernel dead code The forward() dispatch chooses between the native FP8 path and dequantize_w() + cuBLASLt BEFORE reaching the WMMA kernel. Master's threshold is 512 so B=256 reaches it; PR #201 lowers it to 5, which would leave this kernel unreachable at every batch size that matters -- silently, since it still compiles and still passes its tests. The two are answers to the same question asked before and after the premise changed: cuBLASLt won the M=8..128 sweep because it had tensor cores and the native path did not. This kernel removes that asymmetry without paying the ~12.7 GB/step dequantize or the +8.48 GB of resident BF16 weights. Records the three-arm sweep that has to replace the two-arm one. Nothing has run; the expectation that (b) beats (a) is a derivation from bytes moved. UNVERIFIED ON HARDWARE -- never run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
…per 128-K (#200) * perf(ArcKernels): blockwise-FP8 GEMM on tensor cores, scale promoted per 128-K UNVERIFIED ON HARDWARE -- never run. No GPU was available for this wave. Nothing in this commit has executed. Every number below is a DERIVATION from published machine limits, not a measurement. Do not quote any of it as a result, and do not claim a speedup. WHAT `fp8_matmul_tiled` is 66% of a B=256 decode step -- 524 ms of 794. Its inner loop is `acc += s_input[ty][k] * s_weight[tx][k]`: two shared-memory float loads to feed one FMA, no register blocking, no tensor-core instruction anywhere in it. Shared memory delivers 32 floats/clk/SM against 128 FP32 lanes, so that loop cannot exceed ~1/4 of the FP32 CUDA-core rate however it is tuned. It is a structural ceiling, not a tuning problem. This adds `fp8_matmul_wmma` next to it: a real tensor-core GEMM that keeps the weights in FP8 and applies the `[N/128, K/128]` block scale by promoting an FP32 accumulator at every scale-block boundary along K -- DeepGEMM's structure. `K_BLK == block_size_x`, so one K-tile is exactly one scale block and the scale is a single scalar per tile. Numerically this is better than what it replaces, not a trade. The scalar kernel computes `f32(act) * f32(w * scale)` and rounds every product. Here the FP8 -> fp16/bf16 weight conversion is exact (e4m3's 3 mantissa bits and 2^-9..2^8 exponent range are representable in both), the tensor core forms each product at full width into f32, and the scale is applied once per 128 K-elements rather than once per element. DERIVED COST (a derivation, NOT a measurement) V4 at B=256: 7 sites/layer x 43 layers = 4.238 G params, so FLOP = 2 * 256 * 4.238e9 = 2.170 TFLOP. H200 FP32 non-tensor 67 TFLOP/s -> 32.4 ms (the scalar kernel's own bound; it achieves 6.8%) H200 BF16/FP16 tensor 989.5 TFLOP/s -> 2.19 ms <-- this kernel's bound H200 FP8 tensor 1979 TFLOP/s -> 1.10 ms HBM 4.8 TB/s over 4.238 GB -> 0.88 ms Compute-bound at a derived 2.19 ms floor against a measured 524 ms today. A WMMA GEMM of this shape typically realises 50-70% of peak, which would put it at a derived 3.1-4.4 ms. Arithmetic only; none of it has run. The remaining 2x to 1.10 ms is NOT reachable by tuning this kernel: `mma` with e4m3 operands needs BOTH operands in FP8, hence FP8 activations with their own per-token 128 scales. Named, not built. WHY NOT cuBLASLt, AND WHY NOT wgmma cuBLASLt does support exactly this layout -- `CUBLASLT_MATMUL_DESC_B_SCALE_MODE = CUBLASLT_MATMUL_MATRIX_SCALE_BLK128x128_32F` -- and the docs put it at compute capability 9.0, i.e. Hopper, our target. But the enum does not exist before CUDA 12.9 (absent in 12.4.1, absent in 12.8.0, present in 12.9.0) and .github/workflows/cuda_compile_check.yaml pins 12.4.1. It also wants FP8 activations with VEC128 scales and TN layout, so it is not the one-attribute change it first looks like. Checked before writing the kernel, as instructed. wgmma is reachable -- cudaforge auto-suffixes sm_90 to sm_90a -- and worth maybe another 1.3-1.5x. Not used, because this file cannot be run before it is committed and a hand-rolled wgmma descriptor that is subtly wrong returns plausible logits rather than an error. `nvcuda::wmma` fixes the fragment layout in the compiler, compiles for both sm_80 and sm_90a, and has a working precedent in this tree (kernels/mxfp4/mxfp4_gemm_wmma.cu). The wgmma rung belongs on a box that can run the A/B this commit sets up. KILL SWITCH Default-on per house fast-path-default policy, so the first box A/Bs both arms from one binary with no rebuild: (default) -> tensor-core WMMA GEMM ARC_NO_FP8_WMMA=1 -> the scalar `fp8_matmul_tiled`, i.e. exactly the behaviour of the commit before this one Run the control arm first. NO M THRESHOLD, DELIBERATELY This codebase has frozen a dispatch threshold from a single measured point twice -- `ARC_FP8_CUBLAS_MIN_M = 512` (512 was the only M measured, so M = 5..511 fell through to the scalar kernel) and `qtip::gather_policy`'s `n >= 683`. A threshold is a claim about every value it excludes. I have zero measured points, so inventing one here would be strictly worse than either. Sweep on hardware first, then add one if the sweep shows one. WHAT WAS ACTUALLY VERIFIED * `cargo check -p mistralrs-quant` (no cuda) green. * `cargo test -p mistralrs-quant --test cuda_kernel_build_guard` -- 5/5, including `expected_kernel_count_matches_disk`. * EXPECTED_KERNEL_COUNT 41 -> 43 (kernel + its cc<8.0 link stub); the dummy-stub count in that file's own note updated 5 -> 6. * nvcc for sm_80 and sm_90 via the free no-GPU GitHub Actions gate. Nothing else. In particular: no kernel executed, no output compared against a reference, no timing taken. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcKernels): WMMA ldm must be a multiple of 32 B, not 16 -- the C++ guide understates the ISA UNVERIFIED ON HARDWARE -- never run. Two NVIDIA documents govern `wmma::load_matrix_sync`'s leading dimension and they DISAGREE. The laxer one is the one people quote, and the first version of this kernel followed it -- which compiles, and is wrong. CUDA C++ Programming Guide 12.4 §7.24.1: "mptr must be a 256-bit aligned pointer ... and [ldm] must be a multiple of 8 for __half element type or multiple of 4 for float element type. (i.e., multiple of 16 bytes in both cases)." PTX ISA 12.4 §9.7.13.3.2, working our exact shape through: "The starting address of each instance of the leading dimension (row or column) must be aligned with the size of the corresponding fragment in bytes." ... for `wmma.load.a.sync.aligned.row.m16n16k16.f16` the fragment is 32 B (eight `.f16x2` elements), so "p is a multiple of 32" and "2*s is a multiple of 32" i.e. ldm must be a multiple of SIXTEEN __half elements, double the guide's number, and the base pointer must be 32 B aligned, not 16. The +8 pad gave ldm = 136 elements = 272 B. 272 % 16 == 0 satisfies the guide; 272 % 32 == 16 violates the ISA. Worst of both worlds, because nothing would have complained until the numbers came out subtly wrong on a rented box. * SMEM_PAD 8 -> 16, so ldm = 144 elements = 288 B = 9 * 32 B. Satisfies both documents, for f16 and bf16 alike. Asserted at compile time now, with the ISA rule as the assertion message so the next person does not "optimise" the pad back down to the guide's number. * The shared tile was `__align__(16)`; every fragment pointer is `base + k*32`, so a 16 B aligned base put all of them on 16 B boundaries. Now `__align__(128)`. * Shared memory 34,816 -> 36,864 B, still under the 48 KB no-opt-in limit (asserted). Also renames N_SUBTILES -> N_SUB_TILES; the `typos` CI lane reads the concatenation as a misspelling of SUBTITLES. Found by verifying the WMMA contract against the shipped CUDA 12.4 headers and both specs rather than against memory. Nothing here has run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * docs(ArcKernels): flag that lowering ARC_FP8_CUBLAS_MIN_M makes this kernel dead code The forward() dispatch chooses between the native FP8 path and dequantize_w() + cuBLASLt BEFORE reaching the WMMA kernel. Master's threshold is 512 so B=256 reaches it; PR #201 lowers it to 5, which would leave this kernel unreachable at every batch size that matters -- silently, since it still compiles and still passes its tests. The two are answers to the same question asked before and after the premise changed: cuBLASLt won the M=8..128 sweep because it had tensor cores and the native path did not. This kernel removes that asymmetry without paying the ~12.7 GB/step dequantize or the +8.48 GB of resident BF16 weights. Records the three-arm sweep that has to replace the two-arm one. Nothing has run; the expectation that (b) beats (a) is a derivation from bytes moved. UNVERIFIED ON HARDWARE -- never run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * docs(ArcKernels): name the occupancy limit to check first, before anyone measures At V4's B=256 shapes the grid is ~128 blocks against 132 SMs -- one wave, one block per SM, 8 warps of a possible 64. That is the most likely reason the kernel lands short of its 2.19 ms derived bound, and it is where a first profile should point. Lists the three knobs (smaller tiles, cp.async double buffering, split-K) and deliberately guesses none of them, for the same reason the dispatch threshold is left unset. UNVERIFIED ON HARDWARE -- never run. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Contributor
Author
|
Superseded by #215, which is the same work cherry-picked onto This PR was red for exactly one reason — the Cherry-picked rather than retargeted, deliberately: this branch carries What changed in the move:
Closing so the fixes are not orphaned a second time. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Measured on an H200 against V4-Flash, on top of the MoE crossover fix (#197).
1. The threshold was 512 because 512 was the only point measured
arc_fp8_cublas_min_mdefaulted to 512, and the doc said why outright: "thedefault is set to that measured point rather than to an interpolated crossover."
fp8_gemv_warpownsM <= 4, so M = 5..511 fell through tofp8_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.
Default is now 5, the first row
fp8_gemv_warpdoes not own.M = 5..7isunmeasured. Fast by default;
ARC_FP8_CUBLAS_MIN_Mremains the kill switch.2. The cuBLASLt branch never flattened rank-3 to 2-D — and that invalidated the A/B
The native branch flattens
[B,T,hidden]to[B*T,hidden]; the dequantize+cuBLASLtbranch did not. V4's activation is rank-3 and bias is
Noneon every V4 linear, soUnquantLinear::forwardtookw.broadcast_left(B)intocublaslt.batch_matmulwithstride_b = 0. At decodeT=1that is a B-way batched GEMM with m=1 — a batchedGEMV wearing a GEMM's name, the one shape cuBLASLt has no advantage in. The documented
27× was prefill, where
[1,512,hidden]collapses to batch=1 and the call is a real GEMM.Worth +10.4% on its own at B=256 (156.85 → 173.15).
3. Measured
Post-flatten, one binary, one toggle (
ARC_FP8_CUBLAS_MIN_M=999999= kill switch):vs the old 512 default at B=256: 139.33 → 173.15 = 1.24×.
b=1 is unchanged (34.98 vs 34.86). Both arms take
fp8_gemv_warpatM=1, so thefloor is untouched by this constant — which is the check that mattered.
The clean confirmation for the constant did not survive. A re-run at NTOK=160
(which previously gave 96–97% fill) completed both arms — 8 rows recorded — and the box
went down (port 22 closed, 100% ICMP loss) before the log could be read. It did not
come back over ~10 minutes of retries.
So the shipped
5currently rests on the dirty B=64/B=256 rows above plus thepre-flatten clean rows, and those pre-flatten rows compared against the degenerate
batched-GEVM shape described in §2. Every arm measured, clean or dirty, favours
cuBLASLt above the floor, and none is a clean post-flatten row. Re-run
arc-tools/batch_sweep.pyat NTOK=160 before quoting a ratio from this.4. Two follow-ups this work turned up
The real fix is a descriptor, not a kernel. CUDA 13.0's cuBLASLt already exposes
exactly what blockwise FP8 needs:
BLK128x128_32Fis precisely the[N/128, K/128]gridweight_scale_invalreadycarries. That means native block-scaled FP8 on tensor cores with no dequantize at
all — one descriptor attribute plus the existing scale tensor, not a kernel rewrite.
512 was never a kernel property. It is the amortisation point of
dequantize_w(),which runs uncached on every forward (~12.7 GB/step). Cache that and the correct
threshold is 5 regardless.
5. A recurring failure, recorded in the doc
This is the second gate found this way.
qtip::gather_policy's tile-fill predicaterequired
6n/256 >= 16, i.e.n >= 683, keeping the amortising grouped GEMM unreachablein every decode step — worth 1.45× at B=128 when removed. Both gates were derived from a
single working point and neither was ever swept. The doc now says: if you are about to
set a dispatch threshold from one measurement, sweep it first and write down which rows
were clean.
The last commit (
b5de0f1d4, bank-conflict padding+4→+1, bit-identical byconstruction) is labelled UNVERIFIED — never compiled, for the same box outage.