Repository navigation
Conversation
…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>
…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.
56a95d8 to
ab21027
Compare
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.
…settle its arity for free
The probe as authored swept the two descriptor byte-offset fields with the B
tile fixed n-contiguous and `imm-trans-b` fixed at 0. If either of those two
fixed choices is the wrong one it can only return NO_ENCODING_MATCHED — a whole
box round-trip that teaches nothing. All four unknowns now go in the sweep:
trans_b {0,1} x B layout {n-contiguous, k-contiguous} x 10 leading-dim
candidates x 10 stride candidates = 400 wgmma issues, four launches.
Two other changes to stop a false PASS:
* POISON POOL. The B tile is 256 bytes; alone in shared memory a wrong stride
can walk off it into zeros and still reproduce the reference. The tile now
sits at the front of an 8 KB pool pre-filled with a large nonzero sentinel,
so any out-of-tile read perturbs the sum. The compare stays bit-exact.
* "could not run" outranks any verdict. A config that failed to launch, or whose
mma.sync control disagreed with the CPU reference, now forces rc=2 for the
whole probe rather than being averaged into a matches count (D18 rule 2).
ARITY IS NOT A HARDWARE QUESTION. wgmma's register-A form takes either 7 or 8
operands after the destination — whether an `imm-trans-a` precedes
`imm-trans-b` is exactly what PTX ISA 9.7.16.5.2 would say if its body were not
truncated in the published HTML (checked 2026-08-17; 9.7.16.5.1.2.2, the
descriptor bit table, is truncated the same way). But a wrong arity is a ptxas
ERROR, and ptxas runs on a GPU-less runner. So the fat-binary job now compiles
the probe under both arities and reports which one ptxas accepts, failing only
if NEITHER does. Arity is settled by the compiler, the descriptor by the
hardware, and neither is written from documentation.
`wave66_wgmma_desc_probe.sh` is the box runner. It reads its source out of the
object store with `git show` and never checks anything out: the box's working
tree is shared with a curve run in flight.
…dexed in device code `CAND[li]` with a runtime `li` inside the kernel odr-uses a namespace-scope constexpr array, and nvcc rejects that — there is no device copy, so it is "identifier CAND is undefined in device code". The probe had never been compiled, so nobody had hit it. Replaced with a `__device__ __host__` function over the same sequence (0 and the powers of two to 256), which is a constant expression on both sides and needs no `__constant__` symbol or `cudaMemcpyToSymbol`.
The header said Blackwell is "compile-checked separately" without saying where. The tcgen05 snippet lives inline in wave65_arctarget_gate.sh (sm_100a, labelled COMPILE-ONLY / NEVER EXECUTED), not in a standalone probe file — the original header named `arctarget_tcgen05_compile_probe.cu`, which does not exist.
…osed
The free ptxas lane answered the arity question on its first real run, and the
answer is legible even though the step failed:
WGMMA_HAS_TRANS_A=0 (7 operands):
error : Instruction 'wgmma.mma_async with floating point types'
not supported on .target 'sm_90'
WGMMA_HAS_TRANS_A=1 (8 operands):
error : Arguments mismatch for instruction 'wgmma.mma_async'
ptxas validated the 7-operand form's ARGUMENTS and rejected it only on target
grounds; it rejected the 8-operand form's arguments outright. So the register-A
form takes `d, a, b-desc, scale-d, imm-scale-a, imm-scale-b, imm-trans-b` — no
`imm-trans-a`. That is the thing PTX ISA 9.7.16.5.2 would have said if its body
were not truncated, and it cost no GPU time.
The same log exposed two real bugs, both of which would have burned a box slot:
1. THE CONTROL WAS MALFORMED. `mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16`
carries only three type qualifiers; ptxas: "Unexpected instruction types
specified for 'mma'". It needs four — dtype, atype, btype AND ctype —
`.f32.bf16.bf16.f32`. This is the CONTROL, the instruction that decides
whether any wgmma verdict can be believed, so a silently malformed one takes
the whole probe down with it.
2. `-arch=sm_90a` DID NOT PRODUCE sm_90a. nvcc accepted it without complaint
and emitted a `compute_90` intermediate, so ptxas rejected wgmma as "not
supported on .target 'sm_90'". Now built with an explicit
`-gencode arch=compute_90a,code=sm_90a`, and — because being wrong here is
silent — both the CI step and the box runner now emit PTX and read the
`.target` line back before believing anything.
That second one also invalidates a check we already shipped: the box runner and
the wave65 gate both grep `nvcc --list-gpu-code` for `sm_90a`. That listing
NEVER prints the arch-specific variants (CUDA 12.9 shows sm_90/sm_100/sm_103
and no `a` forms), so the check would report "toolkit cannot target sm_90a" on
a toolkit that can. Replaced with the direct test: ask the tool what it DID,
not what it claims to support.
b6e0955 to
1d39712
Compare
…mmit log The 7-operand result was in a commit message; a future reader looks at the file. Records the ptxas discrimination verbatim (7 operands -> arguments validated, target objected; 8 operands -> arguments rejected), why a register operand has no transpose immediate, and why WGMMA_HAS_TRANS_A stays as a build knob: it is what made the compiler answer, and it is how the answer gets re-checked on a toolkit we have not tried.
…'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>
|
ArcGate: HELD on #108, which is genuinely red ( Recording why this PR should land with #108 rather than be folded into it or dropped: #108 authored the wgmma descriptor probe but never ran it, and the probe as authored could only ever answer "no". It swept two byte-offset fields with the B tile pinned n-contiguous and This PR puts all four unknowns in the sweep ( Worth stating plainly since it will come up: the probe exists because the documentation cannot answer it. The published PTX ISA lists §9.7.16.5.1.2.2 "Matrix Descriptor Format" and §9.7.16.5.2 in its table of contents with the bodies truncated — the bit table and operand order are not obtainable from the document. And a wrong descriptor does not fault; it reads the wrong bytes and returns wrong numbers, silently, on the one kernel nobody else can adopt from us. Also: 1 check here — the comment bot only. No CI lane has ever run on this PR (stacked-PR trigger gap, now fixed on master). Any push starts all 16. Nothing deleted; branch untouched. |
|
Genuinely blocked on unmerged #108 — must stay stacked. Not rebased. Attempted Two hard dependencies on #108, both real:
Rebasing all 8 commits onto Correct sequencing: land #108 first, then this rebases onto |
Held: this one genuinely depends on unmerged #108. Rebasing onto master would be wrong.One of five PRs based on another PR's branch — here For the other four the fix was "rebase onto master". Not for this one.
Replaying all 8 commits onto master would resolve, but the result is wrong in a way CI would not catch: #109 would become a superset of #108, swallowing the entire So: land #108 first, then this rebases to its own 1-commit diff. The base stays as-is until then, and it is correctly red under the new guard in the meantime — which is the accurate signal, not a false one. The zero-CI symptom here was a consequence of a real dependency, not a mis-set base. |
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 📚 Stack order — this is the TOPThis PR used to be based on its parent PR's branch, which is exactly the pattern the Both are now retargeted at Merge the bottom PR FIRST, then this one. Until the bottom lands, this PR's diff against the integration branch will also contain the bottom's commits. |
1 participant
What this is
The wgmma descriptor probe, made able to actually answer — plus the half of the wgmma question that needs no GPU, moved onto a free runner.
Stacked on #108, which authored the probe but never ran it.
Why probe at all — verified, not assumed
wgmma.mma_async(sm_90a) takes its B operand through a 64-bit shared-memory matrix descriptor. I checked the published PTX ISA HTML on 2026-08-17: §9.7.16.5.1.2.2 "Matrix Descriptor Format" and §9.7.16.5.2 "wgmma.mma_async" are listed in the table of contents and their bodies are truncated. The bit table and the operand order are not obtainable from the document.That is the justification for probing, and it should be on the record so nobody later reads this as someone being cautious. A wrong descriptor does not fault — it reads the wrong bytes and returns wrong numbers, silently, on the one kernel nobody else can adopt from us.
The probe as authored could only ever say "no"
It swept the two byte-offset fields with the B tile pinned n-contiguous and
imm-trans-bpinned to 0. If either pin is the wrong choice, the only reachable outcome isNO_ENCODING_MATCHED— a negative that proves nothing, and is indistinguishable from "this encoding does not exist". That is a formality, not a check. Same family as everything else this week: a check that cannot fail proves nothing.All four unknowns now go in the sweep:
imm-trans-bs[k*N+n], k-contiguouss[n*K+k]400 wgmma issues, four launches.
Two more things that would have let a false PASS through
The poison pool. The B tile is 256 bytes. Alone in shared memory, a descriptor with a wrong stride can walk off it into zeros and still reproduce the reference. The tile now sits at the front of an 8 KB pool pre-filled with a large nonzero sentinel, so any out-of-tile read perturbs the sum. The comparison stays bit-exact — a tolerance is how a wrong descriptor sneaks through.
"Could not run" outranks any verdict. A config that failed to launch, or whose
mma.synccontrol disagreed with the CPU reference, now forcesrc=2for the whole probe instead of being averaged into a matches count. The control gate is unchanged and still load-bearing: two wrong runs that agree are not evidence.Arity: SETTLED — 7 operands, no GPU spent
The free ptxas lane answered it on its first real run. The step reported failure, but the failure is the result:
ptxas validated the 7-operand form's arguments and objected only to the target; it rejected the 8-operand form's arguments outright. So the register-A form is:
No
imm-trans-a— a register operand has no shared-memory layout to transpose. That is what §9.7.16.5.2 would state if its body were not truncated, obtained for zero GPU time.WGMMA_HAS_TRANS_Astays as a build knob: it is what made the compiler answer, and it is how the answer gets re-checked on a toolkit we have not tried.Two bugs that run exposed, each of which would have burned a box slot
1. The control was malformed.
mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16carries three type qualifiers; it needs four — dtype, atype, btype and ctype (.f32.bf16.bf16.f32). ptxas: "Unexpected instruction types specified for 'mma'". This is the control — the instruction that decides whether any wgmma verdict can be believed — so a silently malformed one takes the whole probe down with it, while looking like a descriptor result.2.
-arch=sm_90adid not produce sm_90a. nvcc accepted the flag with no diagnostic and emitted acompute_90intermediate, so ptxas rejected wgmma as "not supported on .target 'sm_90'" — indistinguishable from a hardware limitation. Now built with an explicit-gencode arch=compute_90a,code=sm_90a, and because being wrong here is silent, both the CI step and the box runner emit PTX and read the.targetline back before believing anything.That second one invalidated a check we had already shipped: the box runner grepped
nvcc --list-gpu-codeforsm_90a. That listing never prints theavariants — CUDA 12.9 showssm_90/sm_100/sm_103and no arch-specific forms — so it would have reported "toolkit cannot target sm_90a" on a toolkit that can. A capability listing was consulted where only the artefact could answer. Both are recorded as D18 instance 13.Arity is not a hardware question
The register-A form takes either 7 or 8 operands after the destination — whether an
imm-trans-aprecedesimm-trans-bis exactly what §9.7.16.5.2 would say if its body were not truncated.But a wrong arity is a ptxas error, and ptxas runs on a GPU-less runner:
-arch=sm_90acompiles device code to SASS at build time. So the ArcTarget fat-binary job now compiles the probe under both arities and reports which one ptxas accepts, failing only if neither does.Arity is settled by the compiler, the descriptor by the hardware, and neither is written from documentation. It also means a broken probe source is caught here instead of on a rented box.
What a result of AMBIGUOUS would mean — read it as data, not failure
At
m64n8k16the B tile is 16x8 bf16 = 256 B, which is exactly two 8x16-byte core matrices. Two core matrices give the descriptor only one thing to describe — the offset from the first to the second — so it is likely that only ONE of the leading-dim / stride fields binds and the other is a free variable. The probe would then reportPROBE_STATUS=AMBIGUOUSand print every match.That is still a useful run, and deliberately so: the printed match set shows which field is pinned and which is free, across all four (layout, trans_b) configs. It is not a "no". Binding both fields needs at least four core matrices, i.e. widening to
m64n16k16— a targeted follow-up that the match set makes small and specific, rather than another blind sweep.UNIQUE_ENCODING_IDENTIFIED(rc 0) is the best case;AMBIGUOUS(rc 1) is the expected case;NO_ENCODING_MATCHED(rc 1) is the one that would mean the assumed bit layout or accumulator fragment layout is wrong, and the partial-hit counts are printed to separate those two.Scope
Does not touch any kernel. The per-arch wgmma/tcgen05 paths are a separate PR gated on this probe's verdict — deliberately not bundled.
wave66_wgmma_desc_probe.shreads its source out of the object store withgit showand never checks anything out: the box's working tree is shared with a curve run in flight.