Repository navigation
perf(qtip): arch-specialise the trellis grouped GEMM (D16) + close the silent-failure holes (D18) - #99
perf(qtip): arch-specialise the trellis grouped GEMM (D16) + close the silent-failure holes (D18)#99heydryft wants to merge 6 commits into
Conversation
…lumbing (D18)
Two kernels behind one dispatch, picked by the device's RUNTIME compute
capability:
sm_80/sm_89 qtip2b_grouped_gemm_kernel_sm80 — unchanged, byte-identical.
This is the path with 5/5 hardware parity; it is not touched.
sm_90/sm_100 qtip2b_grouped_gemm_kernel_wg — 64-pair m-tile, contiguous-run
trellis decode into a shared-memory weight tile, dynamic shared
memory claimed at the arch's real budget.
below sm_80 the launcher returns cudaErrorNotSupported rather than
launching a kernel whose body is compiled out.
The two levers, both aimed at the MEASURED ceiling (the gen-2 GEMV sweep
measured 98/98 variants and concluded the limit is per-symbol trellis decode,
not memory latency):
1. m-tile 16 -> 64. Total decode work is (number of m-tiles) x (expert weight
count): every m-tile re-reads and re-decodes its expert's whole matrix. At
64 routed pairs per expert that is 4 m-tiles at TILE_M=16 and 1 at 64 —
4x less decode AND 4x less HBM traffic for the same output. Below 16
pairs/expert both schedules issue one tile, so this is weakly better
everywhere. This was already in the file's own tuning notes, unexecuted.
2. Contiguous-run decode. q2b_state_from_window is ~11 ALU ops and the Ampere
kernel pays it per weight, because an mma B fragment gives each thread runs
of length 2. Each thread now owns a run of 32 symbols of one weight row,
seeds it once with the window identity and steps the trellis recurrence
(~2 ops). Runs stay independent — no cross-thread state chain, no warm-up
replay. Mirrored and property-tested on CPU by grouped::run_states_2b.
The trellis decode stays fused inside the multiply: packed 2-bit bytes are
what cross HBM, and decode lands in a single 64x64 shared-memory tile that is
overwritten every k-chunk. Full-size weights never materialise. Routing weights
through shared memory is also what the arch-native MMAs require — wgmma's B
operand must be in shared memory and tcgen05 takes no register operands — so
the producer stage is already the right shape for them.
D18 (silent success), on a path that had no launch-status check anywhere:
- every launcher returns cudaError_t and ends with cudaGetLastError(); the host
bails instead of handing back its alloc_zeros buffer as a valid all-zero MoE
output.
- `if (grid <= 0) return;` is now cudaErrorInvalidConfiguration.
- pre-Ampere devices get cudaErrorNotSupported, not an empty-body launch.
- the m-tile schedule is queried from the device and cross-checked against the
Rust mirror, because route and GEMM disagreeing on tile_m mis-bins the tile
map and yields wrong numbers rather than an error.
- qtip_grouped_curve asked the device for its m-tile instead of assuming the
Ampere constant, which on an H200 would have bailed 4x early and overcounted
tiles 4x in its traffic model.
- const-assert ties GROUPED_DECODE_RUN to (tile_n * tile_k) / threads so a
retuned tile cannot leave the kernel decoding a partial tile.
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 72 24328 17592 4018 2718 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 143 14830 12217 797 1816 Shell 23 5756 3964 1412 380 Plain Text 4 3801 0 2479 1322 TOML 33 1485 1292 43 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 188 37465 0 28743 8722 |- BASH 69 1614 1187 311 116 |- C 2 12 12 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 65 2048 1713 77 258 |- TOML 6 207 164 0 43 |- YAML 4 38 33 5 0 (Total) 43185 4661 29265 9259 ───────────────────────────────────────────────────────────────────────────────── Rust 662 306982 265873 13815 27294 |- Markdown 477 22543 471 19392 2680 (Total) 329525 266344 33207 29974 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Total 1276 450597 329477 73114 48006 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
The status check meant for `grouped_dtype!` landed in `dequant_dtype!`
instead. Both macros end with the same
`unsafe { $launch(...); } drop(out_guard); wrap_cuda_slice(...)` shape, and
the patch that introduced the check matched on that shape and took the FIRST
occurrence, which is the dequantize path ~800 lines earlier.
Two consequences, both real:
* `dequant_dtype!` got `let rc = unsafe { $launch(...) }` around a launcher
that returns `()`, so `if rc != 0` was `()` vs integer — E0308 x3 (bf16,
f16, f32).
* the bail message referenced `cc_major`, `cc_minor`, `n_pairs` and `tile_m`,
which are locals of `grouped_gemm_2b_cuda` and do not exist in the
dequantize function — E0425 x12. `n_rows` and `num_symbols` DO exist there,
which is exactly why those two names are absent from the error list.
* and the grouped GEMM itself — the launch this whole change is about — ended
up with NO status check at all. The D18 fix was itself silently not applied:
the patch asserted only that its anchor text was PRESENT, not that it was
UNIQUE, so "the anchor matched" was read as "it matched where I meant".
Same mechanical shape as the bug it was fixing.
Fixed: `dequant_dtype!` restored verbatim, the status check moved to
`grouped_dtype!`. Implicit `{ident}` format capture is kept — it resolves
correctly inside a `macro_rules!` body, including through `bail!`'s own macro
layer (verified with a standalone repro); the earlier hygiene theory was wrong
and the real cause was only the misplacement.
Also:
* `arc-tools/cuda_compile_check.sh` said "cuda kernels did not compile" for a
failure of `cargo build --features cuda`, which does TWO things: nvcc on the
.cu files, then rustc on the cuda-gated Rust. It sent the operator to the
nvcc output when all 33 kernels had compiled and the Rust had not. The
message now names both stages and says how to tell them apart, and keeps 40
lines of tail instead of 5 so the rustc errors are actually visible.
* `ffi.rs`: the new query fn had been inserted between the routing launcher's
doc comment and its signature. Comment restored to its declaration.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The grouped-curve harness queries the device's arch-dependent m-tile (D16) rather than assuming the Ampere constant, but the query fn lived in qtip::mod without a lib.rs re-export, so the cuda build could not see it and the nvcc lanes failed to compile the example. Export it under the same cuda cfg as its FFI, and fully-qualify the non-cuda constant so neither arm warns.
…ss (D18 #12) The dispatch this PR adds picks the warpgroup path from the DEVICE's compute capability. But `qtip2b_grouped_gemm_kernel_wg`'s body is `#if __CUDA_ARCH__ >= 900`, which is a property of the BINARY. Those two can disagree, and when they do nothing says so: a build whose archive carries only sm_80 SASS, run on an H200, JITs from `compute_80` PTX — PTX in which the warpgroup body was already compiled away. `cc_major` is 9, so dispatch takes the wide branch, launches a kernel that writes nothing, and `cudaGetLastError()` returns success. The caller receives its `alloc_zeros` buffer as a shape-correct, error-free, ALL-ZERO MoE layer. This is the same failure this PR already closes one level down — a discarded launch status handing back the zeroed buffer — reappearing one level up, and it is the exact shape hit eleven times today: the absence of a signal read as a specific signal. ARC_CUDA_ARCHS (PR #108) makes the sm_90 cubin exist, so the chain closes this in practice. That is not good enough. A guarantee that depends on another PR having landed, on nobody building without the env var, and on no branch being cut from an intermediate state is a sequencing hope, not a gate. So: `qgw_arch_witness_kernel` is a real kernel carrying the SAME `__CUDA_ARCH__` guard as the one it vouches for, compiled in the same TU with the same arch flags. It is launched once per process (function-local static, thread-safe initialization — this is reachable from every model-parallel worker at once) and reports what the running binary actually contains: 900 if the warpgroup body survived compilation, otherwise the arch it was compiled for. It does NOT default on failure. A witness that could not be taken is a nonzero cudaError_t and the caller refuses — "could not answer" is not "answered yes", which is the same rule the probe harness follows. Checked in two places on purpose: `qtip2b_grouped_query_schedule` reports it so the Rust caller can name the real cause in words ("this binary carries no SM90 device code … would have produced an all-zero MoE layer with no error"), and the launcher macro refuses on its own so the C ABI is safe for any other caller. `qtip_grouped_tile_m()` checks it too — a harness reading the SM90 m-tile off a binary that cannot run the SM90 kernel is modelling a fiction.
…nd must not cache a non-answer Two ways the witness could itself become the silent-failure it exists to catch. 1. GRAPH CAPTURE. The witness launches `<<<1,1>>>` on the default stream. If that stream is capturing, the launch is APPENDED TO A GRAPH rather than executed, and the readback returns the memset zero — reported as "this binary carries no SM90 device code" and refusing a path that works perfectly. A false negative is still a wrong answer. It now checks `cudaStreamIsCapturing` on the same stream the launch uses and declines to answer rather than answering wrongly. (Capture also makes the API return `cudaErrorStreamCaptureImplicit` in some arrangements; that is a non-success status and takes the same path.) 2. CACHING A NON-ANSWER. The result was a `static const` initialized once, so the first attempt's outcome was permanent. A witness that could not be taken during capture would then poison every later call and keep the path refused long after the condition cleared. Only a SUCCESSFUL witness is cached now; anything else is retried. Two threads racing the first call may both take the witness. That is benign and deliberate — the probe is idempotent and its result is a deterministic property of the binary, so both writes store identical bytes. A lock on a path that runs per MoE layer would cost more than the duplicate.
The gate defaulted every artefact to a shared, un-namespaced path: status.txt, the per-arch compile logs, the parity log, both curve logs and the baseline worktree. Two gates running concurrently on the same box overwrite each other's status file — and reading one agent's status as another's is a failure that has already cost a round trip on this project. Everything now lives under GATE_DIR (default /tmp/arc-gate-wave64-keystone), overridable, with STATUS defaulting inside it. This is the same convention the wave65 ArcTarget gate already follows; #99's gate predates it and did not. No behavioural change to what the gate measures.
|
ArcGate triage. NEEDS-OWNER — and explicitly NOT superseded by #124. I checked that first, because "the other QTIP grouped-GEMM PR landed" is the obvious wrong conclusion here. They are disjoint work under one filename. No commit is shared:
#124's own record states it does NOT satisfy D16 — SM90 measured, SM100 never compile-verified — and its author declined that credit explicitly. So this PR is the only dual-arch work in the tree. Closing it as superseded would have deleted the one thing that would earn D16, which is why I did not. Why it is not landing right now. #124 merged first (it had a real 17/17 They are not import collisions. They are the arch-dispatch block ( I aborted the merge rather than resolve it, and left this branch untouched on the remote. Reconciling an arch-dispatch fork with a rewritten decode loop is a change to the kernel's meaning, not a textual merge, and getting it wrong here is exactly the failure mode the moat cannot afford: a wrong resolution does not fault — it computes wrong numbers on the one kernel nobody else can adopt from us. What it needs from its owner:
Not stale, not superseded, not closeable. Parked on an owner. |
|
Triage verdict: REBASE, not close. Checked against current master specifically because this touches what is now the single largest cost in the decode step. Nothing of its two core changes has landed. Master's grouped GEMM did evolve — One conflict must be resolved deliberately, not mechanically. #167 (now on master) writes the amortization gate Conflicts: 4 files — Not merging it now: the m-tile change is a numeric claim and there is no GPU to re-measure it on. #167's own finding — the kernel sits ~60x above its bandwidth bound even at 230% fill — argues the m-tile change should be measured on prefill shapes rather than assumed. Land it when a box exists. The D18 error-check half is independently safe and could be split out. Productive ordering: #108 -> #109 -> #99. #109's descriptor probe is precisely the unblock this PR's body names for |
…192) #167 is the only one of the seven PRs merged in the 2026-08-20 queue sweep that changes behaviour for a flagless user, and it landed on a green-CI gate that cannot see performance. Recording it in the register so it stays a visible open question rather than a merged assumption. The gate moved the 9-682 token band from grouped GEMM to gather-GEMV -- batched decode at B>=9 and every prefill chunk under 683 tokens, which is where chunked prefill lives. The path it now prefers carries 3.15x redundant reads, so it may be slower exactly where it fires. Names the A/B that settles it (B=16/32/64, chunks 128/256/512) and the coupling with the open #99, whose GROUPED_TILE_M -> 64 on SM90+ moves the threshold ~4x and which must read grouped_tile_m_for_cc or the gate and kernel will disagree on tile size. Also records the caveat covering the whole sweep: every number attached to the seven merged PRs was measured on that branch's own base, not on current master, and none may be quoted until re-measured. Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Retargeted at the integration branchBase changed: The queue is being restructured to the shape the owner asked for: one PR open against Two things had to land on
This PR was not closed and is not considered stale. An audit of the queue found the overwhelming majority of it to be real work that was never merged, not noise. What you need to do: rebase onto |
Rebase attempt onto
|
1 participant
What this is
The trellis grouped GEMM is the one kernel nobody can adopt from us: FlashInfer's grouped GEMM is stock CUTLASS with a closed dtype enum, no weight-decode hook (
gemm/group_gemm.cuh:63,67) andElementA = ElementB(group_gemm_sm90.cuh:65-66). Trellis is a state machine, not a dtype. This PR takes it from an Ampere baseline to an arch-dispatched pair of kernels, and closes the silent-failure holes on the path.First, the framing correction
The 22%-of-roofline baseline this work was scoped against does not exist as a valid measurement.
FACTS.md:800-812retracts the 63.5 ms grouped-GEMM microbench (wave35-BM) for two independent defects — an E=64 fixture when V4-Flash has 256 experts, and aLazyLockcap trap that made the grouped-vs-GEMV switch a silent no-op, so every prior "gemv" row re-measured the grouped kernel.FACTS.md:877-879still derived "~22% of roofline" from that retracted 63.5 ms. It should not be quoted, including as a baseline to beat. (Now corrected inFACTS.md.)And the roofline is the wrong ceiling anyway. The project has already measured the real one: the gen-2 GEMV sweep ran 98/98 variants and found every latency-hiding axis negative — split-K, staged width, warp specialization −38% — concluding "the cp.async premise is disproven, not unrealized… the ceiling is per-symbol trellis decode serialization — decode FEWER SYMBOLS, do not hide latency" (
FACTS.md:815-829). That applies to this kernel too. It is ALU-issue-bound on decode. Faster tensor-core instructions address the multiply, which is not the bottleneck. Both changes below target decode work instead.What runs today vs what is dead
moe/experts.rs:604/608/638→lib.rs:1404→bitshift.rs:1824→:1588→cuda_ops.rs:1837. No flag required.cuda_ops.rs:1452→qtip_bitshift_tune.cu, baked variants 21/6 — the two V4-Flash expert shapes).qtip_bitshift_tune2.cu(gen-2, 98 variants) has no production caller: nothing intune.rs:71-83is ≥44, so it is reachable only viaARC_QTIP_GEMV_VARIANT. Dead in production.qtip2b_beam.cu,qtip_beam.cu,qtip_quantize.cuare bake-time only.qtip_gather_gemv.cu/qtip_gemv.cuserve the olderqtip2rung, not the shippedqtip2bartifact.The two changes
1. m-tile 16 → 64 on SM90+. Total decode work is
(number of m-tiles) × (expert weight count)— every m-tile re-reads and re-decodes its expert's whole matrix. At 64 routed pairs per expert that is 4 m-tiles at 16 and 1 at 64: 4× less decode and 4× less HBM traffic for the same output. Below 16 pairs/expert both schedules issue one tile, so it is weakly better everywhere. This was in the file's own tuning notes and had never been executed. 64 is alsowgmma's andtcgen05's M.2. Contiguous-run decode.
q2b_state_from_windowis ~11 ALU ops and the Ampere kernel pays it per weight, because an mma B fragment hands each thread runs of length 2. Each thread now owns a run of 32 symbols of one weight row, seeds it once with the window identity and steps the original trellis recurrence (~2 ops) — state reconstruction drops from ~11 ops/weight to ~2.3. Runs stay independent, so there is still no cross-thread state chain and no warm-up replay.The moat is intact. Packed 2-bit bytes are what cross HBM. Decode lands in a single 64×64 tile of shared memory that is overwritten every k-chunk; full-size weights never materialise. Routing weights through shared memory is also exactly what the arch-native MMAs demand —
wgmma's B operand must be in shared memory, andtcgen05takes no register operands at all — so this producer stage is already their required shape.D16 status — partial, and here is precisely what is missing
mma.sync.m16n8k16(native)sm_90a)mma.sync, notwgmmasm_100a)mma.sync, nottcgen05What is genuinely arch-specialised: the tile schedule, the decode structure, and the shared-memory budget (
cudaFuncSetAttribute— without it a block is capped at 48 KB regardless of the 227 KB that cc 9.0 and 10.x allow, and 8.9 allows only 99 KB).mma.sync.m16n8k16.bf16is documented "requires sm_80 or higher" under the onion-layer model, so it is valid on sm_90a and sm_100a — but it is not native there, and D16 nameswgmma/tcgen05specifically. I am not claiming otherwise.Why they are not in this PR.
wgmmais sm_90a-exclusive (PTX §9.7.16.5.2 "Requires sm_90a"; a-suffix targets "do not follow the onion layer model" and "cannot be run on later generation devices") andtcgen05is sm_100a+, so they are two separate kernels, not one. Both hinge on the shared-memory matrix descriptor: a 64-bit value pinning swizzle mode, leading-dimension byte offset and stride byte offset to a core-matrix layout the PTX doc specifies semantically rather than as a numeric table — and the two descriptor layouts differ (wgmma: swizzle in bits 62-63; tcgen05: bits 61-63, different encoding, extra fixed fields). A wrong descriptor produces wrong numbers, not an error. Writing that blind, with no GPU in this session, on the one kernel that is the moat, is how D18 gets its ninth instance.The unblock is cheap and specific: a ~30-line descriptor probe on any sm_90a box — build a known tile, run one
wgmma, diff against this kernel'smma.syncresult. Minutes of GPU time, and it converts the swap from guesswork into a checked change. The repo already vendors amake_smem_desc+wgmma_m64n{64,128}k16helper (src/sage_cuda/kernels/wgmma.cuh); per the PTX doc, adding bf16 is a one-token change (.f16.f16→.bf16.bf16) with the identical shape set, A-fragment layout and descriptors.D18 — this path had no launch-status check anywhere
grep -rn "cudaGetLastError\|cudaPeekAtLastError" mistralrs-quant/kernels/qtip/returned zero hits. The GEMV path hascheck_gather_gemv_pairsprecisely because a discarded launch status hands the caller thealloc_zerosbuffer — a silently all-zero MoE layer, shape-correct and error-free. The grouped path had no analogue. Fixed:cudaError_tand ends withcudaGetLastError(); the host bails with the code and the shape.if (grid <= 0) return;→cudaErrorInvalidConfiguration.cudaErrorNotSupported, instead of launching a kernel whose body is#if-compiled out (an empty kernel returns success and writes nothing).tile_mmis-bins the tile map and yields wrong numbers, not an error.qtip_grouped_curvenow asks the device for its m-tile instead of assuming the Ampere constant; on an H200 that guess would have bailed 4× early and overcounted tiles 4× in its traffic model.const _: () = assert!tiesGROUPED_DECODE_RUNto(tile_n × tile_k) / threads, so retuning one tile constant stops the build rather than leaving the kernel decoding a partial tile.I also found and fixed a real race in my own new kernel before it left the branch: the next chunk's
cp.asyncstaging could overwrite the buffer the current chunk's mma was still reading, with no barrier between them. Silent corruption, no error. There is now a trailing__syncthreads()with a comment saying why.A validation blind spot worth naming
The first push failed the CUDA lane with 15 Rust errors while 33/33
.cufiles compiled. My report said "Rust side: check + scoped clippy green" — true, and structurally incapable of catching this: on macOS there is no CUDA toolkit, so#[cfg(feature = "cuda")]modules (qtip/cuda_ops.rs,qtip/ffi.rs) are never type-checked at all. "cargo check green" on a Mac does not mean "compiles with cuda", and nobody should read it that way. The gate that catches it is free and already exists:cuda_compile_check.yamlrunscargo build -p mistralrs-quant --lib --features cudaon a GPU-less runner. It should be the first thing consulted on any PR touching cuda-gated Rust — it costs nothing and needs no rental.Two things fixed as a result:
arc-tools/cuda_compile_check.shblamed the wrong stage.cargo build --features cudadoes two things — nvcc on the.cufiles, then rustc on the cuda-gated Rust — and the script reported "mistralrs-quant cuda kernels did not compile" for a pure Rust failure, sending the operator into the nvcc output for two round trips. It now names both stages, says how to tell them apart, and keeps 40 lines of tail instead of 5 so the rustc errors are visible.grouped_dtype!landed indequant_dtype!~800 lines earlier: both macros end with the sameunsafe { $launch(...); } drop(out_guard); wrap_cuda_slice(...), and my patch asserted its anchor was present, not unique, so it took the first match.dequant_dtype!'s launcher returns()→ E0308 ×3; the bail message namedcc_major/cc_minor/n_pairs/tile_m, which do not exist in that function → E0425 ×12 (n_rowsandnum_symbolsdo exist there, which is exactly why they are absent from the error list). And the grouped GEMM — the launch this whole PR is about — ended up with no status check at all. "The anchor matched" was read as "it matched where I meant": the same mechanical shape as the bug being fixed, in the fix for it.Verification status — read this literally
cargo check, scoped clippy lane green, 257/257mistralrs-quantlib tests pass. Four new tests cover the run-decode identity, the warm-up region, the arch schedule selector and the grid bound; 4/4 mutations caught (off-by-one symbol, wrong state mask, off-by-one seed, shift-by-1). D14: these are unit tests, not a kernel measurement.cuda_compile_check.yamlcompiles all 33.cufiles for sm_80 and sm_90 and type-checks the cuda-gated Rust. The coordinator's H200 run independently confirmed 33/33 kernels compiled, so the new SM90 kernel is nvcc-clean; the failure was Rust-only and is fixed above.arc-tools/wave64_keystone_gate.shadds sm_100a and reports an explicitENV-CANNOT-ANSWERline when the toolkit predates it. Blackwell is a compile target only for us either way — we cannot rent one.arc-tools/wave64_keystone_gate.shanswers all three, in order: stage 1 compiles for sm_80/sm_90a/sm_100a (nvcc cross-compiles — any box with CUDA ≥12.8 does this, an A30 at ~$0.40/hr is enough); stage 2 runs the grouped parity tests on real silicon; stage 3 A/Bsqtip_grouped_curveagainst the baseline on the same box. Three exit codes — 0 pass, 1 genuine failure, 2 environment-could-not-answer — and a terminal line for every outcome including timeout, so silence cannot read as success. It asserts the parity run actually executed tests and that the curve produced data rows, because both have reported green over nothing before.The arch witness (D18 #12) — added on review
The dispatch above picks the warpgroup path from the device's compute capability, but
qtip2b_grouped_gemm_kernel_wg's body is#if __CUDA_ARCH__ >= 900, which is a property of the binary. Those two can disagree, and when they did, nothing said so:That is this PR's own launch-status hole reappearing one level up. #108 (
ARC_CUDA_ARCHS) makes the sm_90 cubin exist, so the chain closes it in practice — but a guarantee that depends on another PR having landed, on nobody building without the env var, and on no branch being cut from an intermediate state is a sequencing hope, not a gate.qgw_arch_witness_kernelis a real kernel carrying the same__CUDA_ARCH__guard as the one it vouches for, in the same TU with the same arch flags. Launched once per process, it reports what the running binary actually contains. Checked in three places:qtip2b_grouped_query_schedule(so the Rust caller names the cause in words), the launcher macro (so the C ABI is safe for any other caller), andqtip_grouped_tile_m()(a harness reading the SM90 m-tile off a binary that cannot run the SM90 kernel is modelling a fiction).Two ways the witness could have become the bug it exists to catch, both closed:
cudaStreamIsCapturingon the same stream the launch uses and declines to answer rather than answering wrongly.static constinitialized once, so a witness that could not be taken would poison every later call. Only a successful witness is cached; anything else is retried.It never defaults on failure — "could not answer" is not "answered yes".
Noticed, not shipped
dequant_dtype!(the qtip2b dequantize path) has the same D18 exposure the grouped path had — its launchers returnvoid, so a failed launch hands back thealloc_zerosbuffer. The launchers are actually inqtip_bitshift.cu:527-541(a singleQ2B_DEQUANT_LAUNCHERmacro), notqtip_dequantize.cu— so the fix is smaller than assumed:void->int, onecudaGetLastError(), an empty-grid guard, threeffi.rsdecls and anrccheck in the macro. Still a different subsystem from this PR's subject. Worth a separate change.Scope honesty
The grouped GEMM is not the current bottleneck and this PR does not make it one. At 111.69 tok/s aggregate we are at low single-digit percent of H200 bandwidth, and the measured gaps are dominated by host overhead and the
xshistory, not kernel math. This is the durable moat and the right long-term build; it is not today's limiting factor, and nobody should describe it as one.🤖 Generated with Claude Code