feat(arctarget): make the GPU arch matrix real — multi-arch cubins, asserted with cuobjdump (D16/D18) by heydryft · Pull Request #108 · aeonmindai/arc

heydryft · 2026-08-17T18:07:51Z

The result

ArcTarget: verified cubins sm_80,sm_90a,sm_100a,sm_103a in libmistralrsquant.a

Four architectures, confirmed present by cuobjdump in the crate that carries the moat kernels. That is D16 satisfied and demonstrated for the first time. Until today every Arc binary was single-arch SASS for whatever card happened to build it, while the release matrix, install.sh and the docs all advertised Ada/Hopper/Blackwell. The machinery in this PR works, and the line above is the evidence.

What is still open, stated plainly

The same CI run failed on the auxiliary crate:

CompilationFailed { path: "src/cuda/gemv_bf16.cu", message: "nvcc error:\n\n" }

An empty message — with Segmentation fault (core dumped) two lines earlier in the raw log. nvcc crashed and cudaforge surfaced it as a blank compile error (now recorded as D18 instance 13: a misattributed signal, and the first one in a dependency rather than our code). So this is not "gemv_bf16.cu doesn't compile"; it is "the toolchain died while holding that file", which is a different failure with a different owner.

Two things were added to answer it rather than route around it:

a per-arch nvcc probe that compiles every kernel standalone, once per architecture, printing a line per (file, arch) pair. It cannot fail the job — it exists so the gate is interpretable instead of blank.
cargo build -j 1, because the failing run compiled this crate and mistralrs-quant concurrently, each spawning parallel nvcc across four architectures on a 4-core runner. Serialising removes memory pressure as a variable, so the next run distinguishes contention from a genuine toolkit limit.

What was deliberately not done: the gate's subject was not switched to mistralrs-quant, the crate that passes. The workflow chose arc-cuda-graph for speed (small kernel set), so swapping it would have looked like a sensible optimisation while being exactly what this PR exists to stop — weakening a check until it goes green. If the probe shows a real toolkit limit on one architecture, the honest gate is the arches this toolkit can actually produce, named out loud, with the crash recorded — a scoping result, not a reason to drop an architecture from the matrix.

The problem, measured

On a live H200 on 2026-08-17:

PASS  multi-arch compile
cubin arches in archive: sm_90a
FAIL  D16 unmet: archive is missing sm_100a (has: sm_90a)

ARC_CUDA_ARCHS=90,100,103 compiled clean and emitted only sm_90a. Nothing failed. The compile passing was read as the architectures being present — so every Blackwell claim Arc makes (the release matrix, install.sh's arch detection, D16 itself) was unverified, and on that box false. This is D18 applied to the arch matrix: the absence of a signal read as a specific signal.

Root cause: cudaforge derives exactly one -gencode per file from CUDA_COMPUTE_CAP or the build box's nvidia-smi (cudaforge-0.1.5/src/builder.rs:401-405), and no build.rs in the workspace passed an arch flag of its own. Every Arc binary is single-arch SASS for whatever card happened to be in the build machine. arc-cuda-graph had no arch handling at all, while src/cuda/sampling_kernel.cu:25 names sm_89/sm_90/sm_100 as its targets.

Fixing this first is the precondition for every other arch workstream. #99's SM90+ grouped-GEMM path and #94's Hopper/Blackwell TurboQuant kernels both compile today into archives that may contain no cubin for the architecture they were written for.

What this changes

arc-target/ — one dependency-free crate, used both as a [build-dependencies] entry and at runtime, so the build matrix and the dispatch cannot disagree:

CudaArch — parsing/emission including the a (architecture-specific) and f (family-specific) suffixes. A bare capability ≥ 90 selects a, matching cudaforge's own suffixing and intent: wgmma/tcgen05 exist only in the arch-specific variants, so a plain cubin would silently be the weaker binary.
ArcTier — the three specialisations. Consumer Blackwell (sm_120) is deliberately not folded into Blackwell: it has no tcgen05, and calling it Blackwell would recreate in the type system exactly the over-claim being removed.
Cubin coverage rules: an sm_100a cubin does not serve an sm_103 device, and vice versa.
cuobjdump --list-elf parsing, with the H200 failure as a regression test.

Build scripts (mistralrs-quant, arc-cuda-graph) — ARC_CUDA_ARCHS makes the arch list authoritative (first target becomes cudaforge's, the rest are appended as -gencode; nvcc documents --generate-code as repeatable, which is what produces the fat binary). Then the archive is inspected and the build fails if a requested architecture is absent. Left unset, the build is byte-for-byte what it is today — each extra arch is a full recompile of every .cu, which is not worth paying on a rental that runs one card.

Runtime (mistralrs_quant::arc_target) — device capability → tier, checked against the architectures the build observed itself producing (not the ones it was asked for). One line, once per capability, on every outcome. A device with no cubin is named before the driver's cudaErrorNoKernelImageForDevice; and a dispatch that silently takes the Ampere path on Hopper is no longer indistinguishable from one that took the Hopper path.

CI — a free Blackwell gate. nvcc emits Blackwell cubins without a Blackwell, so a GPU-less ubuntu-latest job builds sm_80 + sm_90a + sm_100a + sm_103a and confirms the cubins with an independent cuobjdump call. It carries a negative control: with CUOBJDUMP=/bin/false the build must refuse, because a green gate that cannot fail proves nothing.

Three-state observation, not two. cuobjdump could not be run is neither a pass nor a failure. Requested archs + unverifiable ⇒ hard failure (a claim we cannot back). No archs requested + unverifiable ⇒ the runtime reports Unknown, never "not covered".

The wgmma descriptor probe

arc-tools/probe/arctarget_wgmma_desc_probe.cu determines the shared-memory matrix descriptor encoding empirically. wgmma (sm_90a) and tcgen05 (sm_100a+) are mutually exclusive and both take their B operand through a 64-bit descriptor whose fields are specified semantically rather than literally — the PTX ISA section is truncated in the published HTML. A wrong descriptor returns wrong numbers, not an error, on the moat kernel. So the probe builds a tile whose product is known, sweeps 100 candidate field encodings, and reports which reproduce the reference exactly; an mma.sync control must pass first, so two wrong runs cannot agree into a pass. Exit codes separate "answered negatively" from "could not run".

arc-tools/wave65_arctarget_gate.sh runs the whole thing on a box, including asking nvcc which predefined macro distinguishes sm_90a from plain sm_90 — the guard every future Hopper path needs, whose name is not in the public docs.

Measured vs compile-only

Nothing in this PR is a performance claim. No benchmark is reported and none was run.
Blackwell is compile-only, by construction. We rent SM90. sm_100a/sm_103a cubins are proven to exist; no line here says they were executed.
The H200 gate above is the measurement that motivated the change; the on-box gate re-runs it against this branch.

Coordination

Does not touch qtip_grouped_gemm.cu / src/qtip/grouped.rs (perf(qtip): arch-specialise the trellis grouped GEMM (D16) + close the silent-failure holes (D18) #99) or turbo_paged_attention.* / mistralrs-paged-attn/build.rs (feat(turboquant): CUDA kernels at head_dim 64/128/256/512, Hopper + Blackwell #94).
feat(turboquant): CUDA kernels at head_dim 64/128/256/512, Hopper + Blackwell #94 added an inline ARC_CUDA_ARCHS block to mistralrs-paged-attn/build.rs. It appends to the autodetected arch rather than replacing it (so an H200 build of 90,100,103 emits sm_90a twice), and it has no verification. On merge it should be replaced with the two arc_target::build::* calls used here — that also brings paged-attn under the same gate.
Also corrects mistralrs-quant/build.rs's SageAttention comment, which advertised SM89/SM90/SM100 kernels that are never compiled behind cfgs no source file reads.

…sserted `ARC_CUDA_ARCHS=90,100,103` compiled clean on a live H200 and produced an archive whose only cubin was `sm_90a`. Nothing failed. The compile passing was read as the architectures being present, so every Blackwell claim Arc makes — the release matrix, install.sh's arch detection, D16 itself — was unverified and, on that box, false. Root cause: `cudaforge` derives exactly ONE `-gencode` per file from `CUDA_COMPUTE_CAP` or the build box's `nvidia-smi` (0.1.5 builder.rs:401-405). No build.rs in the workspace passed an arch flag of its own, so the binary was always single-arch SASS for whatever card was in the machine. ArcTarget is now a system rather than a place in the taxonomy: * `arc-target/` — one dependency-free crate, used by build scripts AND at runtime, owning arch parsing (`a`/`f` suffixes), the three tiers (Ampere/Hopper/Blackwell), cubin coverage rules, and `cuobjdump` parsing. 17 unit tests, including the measured failure as a regression case. * `mistralrs-quant` + `arc-cuda-graph` build scripts: `ARC_CUDA_ARCHS` makes the list authoritative (first target is cudaforge's, the rest appended as `-gencode`), then **`cuobjdump --list-elf` is run over the produced archive and the build FAILS if a requested arch is absent**. Unset keeps today's single-arch autodetect unchanged. * Runtime (`mistralrs_quant::arc_target`): device capability -> tier, checked against the architectures the build OBSERVED itself producing. One line, once per capability, on every outcome — a device with no cubin is named before the driver's `cudaErrorNoKernelImageForDevice`, and a silent Ampere fallback on Hopper is no longer indistinguishable from a Hopper run (D18). * CI: a GPU-less job builds sm_80+sm_90a+sm_100a+sm_103a and confirms the cubins independently, with a negative control that proves the check can still say no. nvcc needs no Blackwell to emit Blackwell. * `arc-tools/probe/arctarget_wgmma_desc_probe.cu`: the wgmma shared-memory matrix descriptor is determined EMPIRICALLY — the encoding is swept and the hardware selects, because a wrong descriptor returns wrong numbers rather than an error and the PTX ISA section is not literal. Gated behind an mma.sync control so two wrong runs cannot agree into a pass. Also corrects `build.rs`'s SageAttention comment, which advertised SM89/SM90/SM100 kernels that are never compiled and cfgs no source reads. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

…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.

…n arch `mistralrs-paged-attn/build.rs` appended one `-gencode` per entry in `ARC_CUDA_ARCHS` on top of the one cudaforge derives itself, and applied the same `a` suffix for cap >= 90 that cudaforge's `GpuArch::auto_suffix` does. Verified against cudaforge 0.1.5 rather than assumed: compute_cap.rs auto_suffix(90).to_gencode_arg() -> "-gencode=arch=compute_90a,code=sm_90a" builder.rs:401 let gencode_arg = gpu_arch.to_gencode_arg(); builder.rs:405 command.arg(&gencode_arg)... builder.rs:412 for arg in &self.extra_args { command.arg(arg); } cudaforge emits its own gencode first and then appends `extra_args` verbatim, so on an H200 building `ARC_CUDA_ARCHS=90,100,103` nvcc received `-gencode=arch=compute_90a,code=sm_90a` twice, byte-identical. Fixed by hoisting the existing `get_compute_cap()` binding above the loop (it is what decides which requested arches are genuinely extra) and skipping the one cudaforge already covers. The later duplicate binding is deleted; there is now one. This is the dedupe half of #108's `arc_target::build::split_primary()`, done locally because `arc-target` does not exist on this branch. The other half — `verify_and_export()`, the `cuobjdump` gate that fails a build whose archive is missing a requested arch — genuinely needs that crate and is left to #108, which should replace this block with the two calls. Until it lands, `arc-tools/wave64_v4_turboquant_kv_gate.sh` STEP 2 asserts the arch list externally with `cuobjdump --list-elf`, so a missing arch is caught by the gate even though it is not yet caught by the build. Not verified here: whether nvcc rejects or tolerates the duplicate flag. macOS cannot run the cuda build path. Emitting it once is correct either way, which is why this is worth doing without waiting on that answer.

…NDENCIES The new job failed on its first run, before reaching ArcTarget's gate at all: `candle-kernels`' build script detects its architecture by shelling out to `nvidia-smi`, which does not exist on a GPU-less runner, so the build died with ComputeCapDetectionFailed while compiling a dependency. Every other CUDA job in this file already sets CUDA_COMPUTE_CAP (the nvcc matrix at job level, and the workspace `cargo check`). This job is new and missed it — which is exactly why the miss surfaced only when the job first ran. It does not weaken the gate. mistralrs-quant and arc-cuda-graph take their architectures from `split_primary` via `compute_cap_arch()`, which overrides cudaforge's env lookup outright, and the verification step asks `cuobjdump` what is actually in the archive regardless.

github-actions · 2026-08-17T19:31:24Z

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                189        37703            0        28931         8772
 |- BASH                  70         1618         1189          313          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)                            43427         4663        29455         9309
─────────────────────────────────────────────────────────────────────────────────
 Rust                    662       307804       266560        13893        27351
 |- Markdown             477        22853          471        19669         2713
 (Total)                           330657       267031        33562        30064
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━
 Total                  1277       451971       330166        73659        48146
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━

…cc against itself First real run of this job failed with: CompilationFailed { path: "src/cuda/gemv_bf16.cu", message: "nvcc error:\n\n" } An EMPTY error message — which is this PR's own subject reappearing in its own CI. The real stderr, two lines up, was `Segmentation fault (core dumped)`: nvcc crashed, and cudaforge surfaced that as a blank compile error. Same run also printed, for the crate that carries the moat kernels: ArcTarget: verified cubins sm_80,sm_90a,sm_100a,sm_103a in libmistralrsquant.a So the multi-arch machinery and its cuobjdump verification demonstrably work across all four architectures. What is unproven is arc-cuda-graph specifically. Two changes, neither of which weakens the gate: * A per-arch nvcc probe compiles every kernel in the crate standalone, once per architecture, printing a line for each (file, arch) pair. It cannot fail the job — it exists so the gate below is interpretable instead of blank. If all four arches compile a file standalone, the failure was contention, not the toolkit; if one arch fails everywhere, that is a toolkit limit and belongs in a different bucket than an Arc bug, exactly as the wave65 gate already separates them. * `cargo build -j 1`. The failing run built this crate and mistralrs-quant concurrently, each spawning parallel nvcc across four architectures on a 4-core runner. Serialising removes memory pressure as a variable.

…downgrades `-arch=sm_90a` is accepted by nvcc without any diagnostic and then emits a `compute_90` intermediate — proved the same day by the wgmma probe, where ptxas rejected the instruction as "not supported on .target 'sm_90'" over a file named `...compute_90.ptx`. So the diagnostic added one commit ago would have printed `OK sm_90a` for a file it had actually compiled as sm_90: a false OK, produced by the step added to make a false failure interpretable. Verification code is not exempt, and this is the third time today that rule has been the one that mattered. Now uses the same explicit `-gencode arch=compute_90a,code=sm_90a` form that `CudaArch::gencode()` already emits — which is why the real fat-binary build produced genuine sm_90a/sm_100a/sm_103a cubins while this probe would not have.

…rol was malformed Four corrections, all established on hardware, none of which change the multi-arch machinery itself (which is now verified: `verified cubins sm_80,sm_90a,sm_100a,sm_103a in libmistralrsquant.a`). 1. `nvcc --list-gpu-code` NEVER prints the `a` variants — CUDA 12.9 lists sm_90 / sm_100 / sm_103 and no arch-specific forms. The gate grepped it for `sm_90a` and so reported "toolkit cannot target sm_90a" on a toolkit that can: a capability listing consulted instead of the artefact, where the answer was never going to be. Replaced with a compile-and-read-back: `-gencode arch=compute_90a,code=sm_90a -ptx`, then assert the `.target` line. Believe the artefact, not the advertisement. 2. `-arch=sm_90a` is ACCEPTED by nvcc with no diagnostic and produces a `compute_90` intermediate, so ptxas rejects wgmma as "not supported on .target 'sm_90'" — which reads exactly like a hardware limitation rather than a flag silently not honoured (D18 #13; every dependency that reports on our behalf is an unaudited narrator). Every `-arch=` in the gate and the probe's build line is now `-gencode`, and the readback in (1) detects the non-honoured case explicitly rather than inferring it. 3. The probe's mma.sync CONTROL was malformed: `mma` takes FOUR type qualifiers (dtype, atype, btype, ctype) and this had three — ptxas: "Unexpected instruction types specified for 'mma'". The control is the instruction that decides whether any wgmma verdict is believable, so a silently malformed control takes the whole probe down AND the failure presents as a descriptor result. Now `.f32.bf16.bf16.f32`, verbatim from qtip_grouped_gemm.cu:132. 4. wgmma register-A arity is settled at 7 operands, no `imm-trans-a`, decided without a GPU by making ptxas discriminate: the 7-operand form had its ARGUMENTS validated and only the target objected, while the 8-operand form was rejected outright. Recorded in the probe so it is not re-derived on a rented box. (The probe already used the 7-operand form; it is now documented as established rather than assumed.) Also removes the same bad oracle from build.rs's failure message, which had been telling readers to check `nvcc --list-gpu-code`, and annotates the informational call in CI so nobody gates on it later. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

…'s file #109 (`perf/wgmma-descriptor-probe`) is based on THIS branch and rewrites `arctarget_wgmma_desc_probe.cu` wholesale (+270/-108: 400-variant sweep, poison pool, `cand()`, fixed `mma.sync` ctype, explicit `-gencode`). My annotation commit touched the same hunks and flipped #109 to mergeable=false. Restored byte-identical to d2d6317, the commit #109 diffed against, so its diff applies cleanly again. Nothing is lost: the ctype fix and the `-gencode` build line both exist on that branch already, and the two findings worth keeping — the settled 7-operand register-A arity (established by making ptxas discriminate: 7 operands had its ARGUMENTS validated with only the target objecting, 8 were rejected outright) and the "-arch=sm_90a is accepted without diagnostic and downgrades to compute_90" note — are being routed to that chain rather than re-derived on a rented box. The gate script keeps its own `-gencode` invocation of the probe; that file is mine and is unaffected. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

heydryft · 2026-08-18T08:58:01Z

ArcGate: NEEDS-OWNER — genuinely red, and the red is not yours.

The failing lane is ArcTarget fat binary (sm_80 + sm_90a + sm_100a + sm_103a). Every other lane is green, and this PR has no CI complete either (it predates that required job), so it needs an absorb regardless.

Not closing and not ranking it down. The headline in the body is real and is the first time it has been true in this repo:

ArcTarget: verified cubins sm_80,sm_90a,sm_100a,sm_103a in libmistralrsquant.a

Four architectures, confirmed present by cuobjdump, in the crate that carries the moat kernels. Until this, every Arc binary was single-arch SASS for whatever card happened to build it while the release matrix and install.sh advertised Ada/Hopper/Blackwell.

The failure is CompilationFailed { path: "src/cuda/gemv_bf16.cu", message: "nvcc error:\n\n" } — an empty message with Segmentation fault (core dumped) two lines earlier. As the body says, that is not "gemv_bf16.cu does not compile"; it is the toolchain died while holding that file, surfaced blank by cudaforge. Different failure, different owner. The per-arch nvcc probe and -j 1 added here are the right response — making the gate interpretable rather than routing around it.

What it needs: absorb master (which now has #124's rewritten grouped-GEMM — expect interaction), then re-run and read the per-arch probe output to see whether the segfault survives serialisation. If it does, it is an nvcc bug to pin to a version, not a code fix.

⚠️ #109 is stacked on this one and inherits the red. #109 is the probe made able to answer — it widens the sweep to all four unknowns (imm-trans-b, B smem layout, leading-dim offset, stride offset), which matters because the probe as authored here could only ever return NO_ENCODING_MATCHED, a negative that proves nothing. Worth landing together once this is green.

Nothing deleted; branch untouched.

heydryft · 2026-08-20T11:03:31Z

Triage verdict: REBASE, not close. Still unique and the conflict is trivial.

Evidence against current master: there is no arc-target/ crate (only arc-cuda-graph), and ARC_CUDA_ARCHS appears only in memory/mission/KERNEL_RULES.md — i.e. D16 dual-arch is documented doctrine with no implementation. Every binary we ship is still single-arch.

Conflict: 1 file, Cargo.lock only. Mechanical.

#109 (perf/wgmma-descriptor-probe) stacks on this branch and is CLEAN+MERGEABLE, so this must land first; the pair should go together. Then #99, whose body names #109's descriptor probe as its unblock.

heydryft · 2026-08-21T15:36:44Z

Retargeted at the integration branch

Base changed: master → release/openrouter-ready.

The queue is being restructured to the shape the owner asked for: one PR open against master (#194), with everything else merging into a single integration branch. Agents branch off release/openrouter-ready and PR back into it; #194 is the single gate from there to master.

Two things had to land on master first for this to be workable, and both have:

The base-branch CI lane used to hard-fail any PR not targeting master (ci(ArcGate): let the base-branch lane permit the single integration branch — 7 PRs are red for addressing, not quality #216). Seven PRs were red for their addressing, not their quality. The lane now accepts master or release/openrouter-ready — an exact-match allowlist, nothing else — and its intent is intact: release/openrouter-ready reaches master only through Arc → OpenRouter-ready: the single integration PR (everything else merges into this branch) #194, whose own base is master, so nothing enters master without a full run whose base is master.
release/openrouter-ready was re-cut from current master. It had diverged badly — merging it as-is would have reverted TCFRAG (perf(ArcQuant/ArcKernels): TCFRAG-2B — put the qtip2b trellis GEMV on the tensor cores #203) and re-broken the flag-polarity work (fix(ArcGate): read boolean ARC_* flags by value, so =0 means off #212). It is now master plus only the work it uniquely carried.

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 release/openrouter-ready and resolve conflicts against it rather than against master. CI lanes will now actually run and report on this PR instead of failing at the base check.

📚 Stack order — this is the BOTTOM

This PR and its dependent were both retargeted at release/openrouter-ready, so they are now siblings rather than a stack. That makes the merge order load-bearing and no longer enforced by the base ref:

Merge this PR FIRST. Its dependent carries changes that assume this one is already in the branch, and merging them out of order will produce conflicts or silently drop this PR's changes.

🔴 `needs-work` — this one is red for a REAL reason

Retargeting fixes the addressing problem, not this PR's own problem. Labelled needs-work so it is not mistaken for "just needs a rebase" when the queue is triaged.

This PR has a real fat-binary failure in the multi-arch cubin work. The base change will make the lanes run against the integration branch, but the fat-binary failure is a genuine defect in the change itself and will persist until it is fixed. Do not merge it on the strength of a green base-branch check alone.

Note its dependent #109 (the wgmma descriptor probe) is gated on this — see the stack-order note above.

…n arch `mistralrs-paged-attn/build.rs` appended one `-gencode` per entry in `ARC_CUDA_ARCHS` on top of the one cudaforge derives itself, and applied the same `a` suffix for cap >= 90 that cudaforge's `GpuArch::auto_suffix` does. Verified against cudaforge 0.1.5 rather than assumed: compute_cap.rs auto_suffix(90).to_gencode_arg() -> "-gencode=arch=compute_90a,code=sm_90a" builder.rs:401 let gencode_arg = gpu_arch.to_gencode_arg(); builder.rs:405 command.arg(&gencode_arg)... builder.rs:412 for arg in &self.extra_args { command.arg(arg); } cudaforge emits its own gencode first and then appends `extra_args` verbatim, so on an H200 building `ARC_CUDA_ARCHS=90,100,103` nvcc received `-gencode=arch=compute_90a,code=sm_90a` twice, byte-identical. Fixed by hoisting the existing `get_compute_cap()` binding above the loop (it is what decides which requested arches are genuinely extra) and skipping the one cudaforge already covers. The later duplicate binding is deleted; there is now one. This is the dedupe half of #108's `arc_target::build::split_primary()`, done locally because `arc-target` does not exist on this branch. The other half — `verify_and_export()`, the `cuobjdump` gate that fails a build whose archive is missing a requested arch — genuinely needs that crate and is left to #108, which should replace this block with the two calls. Until it lands, `arc-tools/wave64_v4_turboquant_kv_gate.sh` STEP 2 asserts the arch list externally with `cuobjdump --list-elf`, so a missing arch is caught by the gate even though it is not yet caught by the build. Not verified here: whether nvcc rejects or tolerates the duplicate flag. macOS cannot run the cuda build path. Emitting it once is correct either way, which is why this is worth doing without waiting on that answer.

…lackwell (#94) * feat(turboquant): instantiate the CUDA kernels at head_dim 64/128/256/512 TurboQuant's paged kernels were hardcoded to head_dim 128 — `constexpr int HS = 128`, a single `SGN[128]` sign table, `rotate128`, and six `extern "C"` launchers each opening with `if (hs != 128) return;`. That early return is a silent no-op: the output buffer is left uninitialized, so a mismatched head dim produces garbage *fast* rather than failing. Kernels - `tq_attn`, `tq_attn_clocked`, `tq_cache_k`, `tq_cache_v` and the two clocked cache variants now take HS as a template parameter. NT stays 128 for every HS, so the warp-reduction topology is unchanged; wider heads give each thread DPT = HS/NT output dims instead of adding threads. At HS=128 the emitted work is the same as before. - The per-lane packed-K gather is now one vector load (`tq_load_klane`) instead of BPL scalar byte loads. The x=16 cache swizzle keeps each lane's run contiguous and naturally aligned, which `packed_k_bytes_are_swizzle_aligned` asserts. - Every `if (hs != 128) return;` is replaced by a dispatch over the instantiated widths. D16 dual-arch - `TQ_PREFETCH` software-pipelines the K gather on SM90/SM100/SM103 and runs flat below Hopper — same arithmetic, scheduled for the part it runs on. - Launches now opt in to the arch's real dynamic shared-memory budget. The logits array scales with context length and CUDA caps dynamic shared memory at 48 KB without an explicit opt-in, so contexts past ~12k tokens could not launch at all on any arch. - `ARC_CUDA_ARCHS=90,100,103` emits cubins for Hopper *and* Blackwell rather than only the build box's GPU. Tables - Codebooks and sign diagonals are dimension-dependent, so 512 needs its own. All 20 tables are emitted by `emit_turboquant_cuda_tables`, which calls the same `generate_signs`/`get_codebook` the Rust compressor uses — correct by construction, not by transcription. The generator reproduces the shipped d=128 table byte-for-byte. - `turboquant::cuda_tables` re-derives every table at test time and diffs the checked-in header, and pins the Rust head-dim list against both the kernel's dispatch switch and the FFI wrapper's list. Nothing previously asserted that the tree's four copies of the sign table agreed with the generator or with each other. Gates - `TURBOQUANT_HEAD_DIM = 128` exact-match becomes membership in `TURBOQUANT_CUDA_HEAD_DIMS`; K and V are compressed independently and no longer have to be equal widths. Not measured on a GPU: the kernels are written and gated but this box has no nvcc, so nothing here is a claim about hardware behaviour. * test(gpu): wave63 gate — dual-arch compile proof for the head_dim 512 kernels Proves the four head-dim instantiations compile for sm_90a AND sm_100a/sm_103a (D16), not just the build box's GPU, and that the shipping decode path did not regress. Deliberately contains no quality A/B: TurboQuant quality is settled by prior measurement and re-proving it would spend GPU budget on a known number. * fix(build): watch turbo_paged_attention.cu for changes Every other .cu in this crate had a rerun-if-changed line; this one did not. Emitting any rerun-if-changed disables cargo's watch-the-whole-package default, so edits to the largest CUDA file in the crate were relying entirely on cudaforge's object cache to trigger a rebuild. * fix(turboquant): 64-bit cache offsets, per-device smem opt-in, clocked-kernel parity Three problems from review of the head_dim generalization. 1. 32-bit overflow on cache byte offsets. `pb*kbs`, `pb*vbs`, `bi*cbs` and `bi*vbs` were all `int`. At head_dim 512 with nkvh=8 and block_size=32, `kbs` is 65536 B, so `pb*kbs` wraps once the K cache passes 2 GB — an entirely reachable allocation on an H200, and 4x sooner than at 128. Now computed in `long long` at all eight sites. 2. The shared-memory opt-in cached a process-wide `granted` flag per kernel instantiation. The requestable maximum differs by architecture, so one flag cannot be right across a mixed fleet, and a failed request left `granted` at 0 and fell through to a launch that then failed with no diagnostic. The cache is gone; the call is cheap and now always attempted. 3. `tq_attn_clocked` kept the scalar byte gather while `tq_attn` moved to a vector load plus the arch-gated pipeline. Identical results, but the phase-2 stamps would have described a kernel nobody launches — at head_dim 512 that is 8 scalar loads measured against the 1 vector load actually executed. The loop now mirrors `tq_attn`. Verified without nvcc by type-checking the kernel bodies against a CUDA shim (clang++ -Wall -Wextra -Wshadow, both __CUDA_ARCH__ branches): clean. Still not run on a GPU. * fix(gate): D18 exit codes, branch-independent preflight, source-landmine guard The gate exited 1 when gpu_box_preflight.sh was absent. That read as a code failure for a couple of minutes when it was really 'this machine cannot answer'. Per D18 rule 2: 0 pass, 1 genuine failure (the only signal to act on), 2 environment could not answer. Nine conditions now exit 2 — missing/failing/incomplete preflight, no repo, no nvcc, unreachable remote, un-checkout-able branch, no archive to inspect, no cuobjdump. Four exit 1, all genuine: nvcc compile failure, a native build failure, a test failure, and the archive missing an arch it claims to carry. The preflight is now looked up at ARC_PREFLIGHT, then /root/arc-tools, then /usr/local/lib/arc, and only last inside the repo. It lives on wave61/box-preflight-shared-prefix, so a repo-relative path vanishes the moment a runner checks out any other branch — which is exactly how this gate came to refuse to run. Guarded the sourced-preflight landmine directly. gpu_box_preflight.sh ends on [ "$_arc_pf_sourced" = "1" ] || exit 1 so when SOURCED on a *failing* box that test succeeds, short-circuits the '|| exit 1', and becomes the last command — 'source pf || handler' returns 0 and the handler never fires. The gate ignores the return value entirely and reads _ARC_PF_FAILED, treating an unset flag as 'did not reach a verdict' rather than as a pass. Also reordered so the sm_90a+sm_100a+sm_103a compile is STEP 1 and needs no full workspace build, and it now asserts on cuobjdump output that both cubins are actually in the archive — a build that merely succeeds is not the claim. Exit paths verified locally against a stub reproducing the real preflight's sourced-exit structure: all four environment cases return 2. * fix(paged): keep the TurboQuant DEFAULT at head_dim 128; widen only what is asked for Widening the acceptance gate from head_dim == 128 to every instantiated width was correct for what a user may REQUEST, but PagedCacheType:: TurboQuant carries #[default] (cache_engine.rs:22), so it also silently moved standard-layout models at head_dim 64/256/512 onto kernels that have never executed — and silently dropped prefix caching with them, since supports_prefix_cache() is false for every TurboQuant variant. Two regressions, neither visible at the call site, for users who never asked for TurboQuant at all. This repo has paid for that exact shape once: FP8 KV shipped default-on and unmeasured (wave43-BU) and every V4 request died. #98 cites that incident by name as its own reason not to default itself on; this takes the same choice. resolve_for_model already distinguished explicit from ambient-default in order to pick error-vs-fallback, so the fix is to let that same 'forced' flag also pick WHICH head-dim set applies: compiled -> widens what you may ask for (TURBOQUANT_CUDA_HEAD_DIMS) measured -> widens what you get unasked (TURBOQUANT_DEFAULT_HEAD_DIMS) TURBOQUANT_DEFAULT_HEAD_DIMS is [128] and moves when a hardware gate passes, not when a kernel compiles. Nothing about the kernels changes: the six 'if (hs != 128) return;' no-ops (D18 instance 3) and the 48 KB dynamic smem cap fix are untouched, and explicit --pa-cache-type turboquant still reaches every instantiated width including V4's 512. The unsupported-geometry message now separates 'no kernel exists' from 'a kernel exists but is unmeasured, so the default will not choose it for you' — the second tells the user to opt in rather than to give up. Tests: the old fixture asserted the regression (its expected-to-fall-back list moved [64,96,192,256] -> [96,192,320,1024] precisely because 64 and 256 stopped falling back). Replaced with an explicit-path test over every instantiated width, plus two pins on the default path. Both pins were mutation-checked by reverting the const to the full set: both fail, and the first revision of one passed vacuously because widening the const emptied its loop, so it now asserts the loop is non-empty before iterating. * docs: correct the public record for the widened TurboQuant kernels This PR branched before #101 ("correct the public record") landed, so it had never seen that text. Merging master in makes four README claims false — three falsified by this PR's own kernels, one that was already wrong when it was written. Falsified by this PR: - L31 "the paged kernel exists at head_dim 128 only" - L31 "there is no kernel at head_dim 512, so DeepSeek V4 cannot use it" DeepSeek V4 reports KvCacheLayout::Standard at head_dim 512 (normal_loaders.rs), so the 512 instantiation genuinely reaches it. - L141 an explicit `--pa-cache-type turboquant` off-128 "is a hard error" It is now accepted at any instantiated width; the hard error moved out to the uninstantiated set. Already wrong before this PR: - L141 "TurboQuant is not the default anywhere" `defaults::PAGED_CACHE_TYPE` is `PagedCacheType::TurboQuant` and the CLI's `--pa-cache-type` has no clap default, so leaving it unset on CUDA gives a standard-layout head_dim-128 model TurboQuant KV with no flag — and silently drops prefix caching, which no TurboQuant variant supports. That was true on master; the narrowing in a3a5faa bounds it to 128 but does not remove it. Stating the opposite understated the risk to anyone reading the README to decide whether to opt out, so it is corrected in place rather than quietly reworded, and the Rust example's comment now says `Auto` opts *out* rather than implying TurboQuant is off until asked for. Also refreshes the `--pa-cache-type` help text, which still described the explicit path as 128-or-error. No behaviour change: documentation and one clap doc comment. * fix(build): stop emitting a duplicate -gencode for the build box's own arch `mistralrs-paged-attn/build.rs` appended one `-gencode` per entry in `ARC_CUDA_ARCHS` on top of the one cudaforge derives itself, and applied the same `a` suffix for cap >= 90 that cudaforge's `GpuArch::auto_suffix` does. Verified against cudaforge 0.1.5 rather than assumed: compute_cap.rs auto_suffix(90).to_gencode_arg() -> "-gencode=arch=compute_90a,code=sm_90a" builder.rs:401 let gencode_arg = gpu_arch.to_gencode_arg(); builder.rs:405 command.arg(&gencode_arg)... builder.rs:412 for arg in &self.extra_args { command.arg(arg); } cudaforge emits its own gencode first and then appends `extra_args` verbatim, so on an H200 building `ARC_CUDA_ARCHS=90,100,103` nvcc received `-gencode=arch=compute_90a,code=sm_90a` twice, byte-identical. Fixed by hoisting the existing `get_compute_cap()` binding above the loop (it is what decides which requested arches are genuinely extra) and skipping the one cudaforge already covers. The later duplicate binding is deleted; there is now one. This is the dedupe half of #108's `arc_target::build::split_primary()`, done locally because `arc-target` does not exist on this branch. The other half — `verify_and_export()`, the `cuobjdump` gate that fails a build whose archive is missing a requested arch — genuinely needs that crate and is left to #108, which should replace this block with the two calls. Until it lands, `arc-tools/wave64_v4_turboquant_kv_gate.sh` STEP 2 asserts the arch list externally with `cuobjdump --list-elf`, so a missing arch is caught by the gate even though it is not yet caught by the build. Not verified here: whether nvcc rejects or tolerates the duplicate flag. macOS cannot run the cuda build path. Emitting it once is correct either way, which is why this is worth doing without waiting on that answer.

heydryft mentioned this pull request Aug 17, 2026

probe(wgmma): make the descriptor probe able to answer, and settle its arity on a free runner #109

Open

heydryft mentioned this pull request Aug 17, 2026

feat(turboquant): CUDA kernels at head_dim 64/128/256/512, Hopper + Blackwell #94

Merged

This was referenced Aug 17, 2026

perf(qtip): arch-specialise the trellis grouped GEMM (D16) + close the silent-failure holes (D18) #99

Open

ci: run the full lane on stacked PRs, and refuse to report green when it did not (D18 #11) #106

Merged

heydryft and others added 4 commits August 17, 2026 20:55

heydryft mentioned this pull request Aug 21, 2026

Arc → OpenRouter-ready: the single integration PR (everything else merges into this branch) #194

Draft

heydryft changed the base branch from master to release/openrouter-ready August 21, 2026 15:36

heydryft added the needs-work Red for a real defect, not for addressing — needs code changes, not just a rebase label Aug 21, 2026

Navigation Menu

feat(arctarget): make the GPU arch matrix real — multi-arch cubins, asserted with cuobjdump (D16/D18) - #108

feat(arctarget): make the GPU arch matrix real — multi-arch cubins, asserted with cuobjdump (D16/D18)#108
heydryft wants to merge 6 commits into
release/openrouter-readyfrom
arctarget/multiarch-specialization

heydryft commented Aug 17, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Aug 17, 2026 •

edited

Loading

Uh oh!

heydryft commented Aug 18, 2026

Uh oh!

heydryft commented Aug 20, 2026

Uh oh!

heydryft commented Aug 21, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

heydryft commented Aug 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

The result

What is still open, stated plainly

The problem, measured

What this changes

The wgmma descriptor probe

Measured vs compile-only

Coordination

Uh oh!

github-actions Bot commented Aug 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

heydryft commented Aug 18, 2026

Uh oh!

heydryft commented Aug 20, 2026

Uh oh!

heydryft commented Aug 21, 2026

Retargeted at the integration branch

📚 Stack order — this is the BOTTOM

🔴 needs-work — this one is red for a REAL reason

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

heydryft commented Aug 17, 2026 •

edited

Loading

github-actions Bot commented Aug 17, 2026 •

edited

Loading

🔴 `needs-work` — this one is red for a REAL reason