feat(nvidia): the shipped GEMM becomes the first declared delegate - #652
Merged
Merged
Conversation
The delegation contract (#649/#650) had zero live users. Measured on the fleet's sm_120 before this change, all four NVIDIA matmul candidates returned None from `contract_for_candidate` while the Tier-3 hand-tuned lane won arbitration on tier priority alone -- no declared accuracy budget, no determinism claim, no architecture claim, no coverage claim. This declares the shipped `libtessera_nvidia_gemm` mma.sync GEMM as a `DelegatedCandidate`, so its tier and budget are derived from a contract the Target IR verifier also enforces rather than hand-set in Python. Four things the first real use forced, each measured rather than assumed: * **A delegate can keep its name.** `name` is a dispatch/cache key, not a claim: the autotune corpus and the E3 `force` hatch key on this exact string. Deriving it from `callee` would invalidate every persisted verdict on a C-symbol rename, with no error. * **A delegate may declare a contract family.** The shipped GEMM binds a different symbol per dtype (`..._f16` / `..._bf16`) and `callee` is identity, so one declared callee for two bound symbols is exactly the Python-vs-IR drift the contract exists to stop. Family members must agree on the fields that describe the delegate rather than one route. A per-dtype *tolerance* looked like a second motivation and measurement refuted it: bf16 and f16 agree within 25% at every K from 32 to 8192, because the oracle's reference rounds its operands to the storage dtype first, so the residual is f32 accumulation error and does not depend on storage width. * **The architecture claim was wrong in both directions.** `tessera_nvidia_gemm.cpp` NVRTC-compiles for the live device and the kernel needs only `mma.sync.aligned.m16n8k16` (sm_80+); the sm_120 cubin is an AOT fast path, not the envelope. The contract now says `sm_80+`; `mma_arch` stays `sm_120` as a footprint-model key, which is a different thing and is documented as such. * **A fixed absolute tolerance is the wrong shape for the claim.** Measured at M=N=256, absolute error grows about K^1.2 while relative error grows near sqrt(K): 5e-3 keeps only ~6x headroom at K=8192 and would be breached past K~65536 on a kernel that is not wrong. Both bounds are now declared, and the relative one carries large K. `determinism="deterministic"` is grounded, not assumed: the kernel gives one warp each 16x8 output tile and reduces K serially into four accumulator registers -- no atomics, no split-K, no cross-block reduction. Its `mma.sync.m16n8k16.row.col` fragment layout was checked against the PTX ISA canonical mapping for A, B and D before the claim. Also: device-resident latency for both GEMM lanes. Two of four candidates had none, so the only available comparison was host-dominated wall time -- the compiled Tile lane measures 2.99 ms of device work inside 34.0 ms of wall time at 2048^3. A Tier-3 delegate with no device timer can never be displaced by a faster compiled kernel, which is the whole of Decision #28. Device proof in tests/device/nvidia/test_shipped_gemm_delegate.py is owed on sm_120 and is not claimed here.
…t time it The device timer I added for `nvidia_mma_gemm_emitted` returned rc=5 on sm_120 every time, and the cause is a real divergence rather than a missing table entry. The launch bridge's `benchmarkTileGemm16` launches `gx = ceil(N/tileN)`, `gy = ceil(M/tileM)` -- x maps to N. The NVIDIA Tile lowering agrees (`NVIDIALowering.cpp`: `mt = blockY*16`, `nt = blockX*8`), which is why both Tile candidates time correctly through it and why their existing latencies are sound. `ptx_emit` uses the opposite convention (`mt = ctaid.x*16`, `nt = ctaid.y*8`), as does the shipped AOT kernel. So the NVIDIA backend carries two block-index conventions. Registering the emitted kernel's geometry in `tileLaunchConfig` would launch it transposed: at 512x512 that covers rows to 1024 and columns only to 256, leaving half the output unwritten while still reporting a plausible latency -- a number that is worse than no number. Swapping the tile dims is not a fix either; it lines up only when M == N. Rather than ship that, the timer is removed, the divergence is recorded where the code is, and the device test names the one candidate without a timer so an unexplained `None` stays a regression while this one stays a tracked gap. Decision #28's displacement test is unaffected: the delegate is measured against the two compiled Tile lanes. A unit test now pins `ptx_emit`'s axis mapping, so "fixing" one side of the divergence fails loudly instead of silently producing a mis-shaped grid.
The device timer's first real output contradicts the arbiter's default.
On sm_120, f16, square, device-resident timing with 0.000-0.008 ms
spreads across repeats:
shape shipped(T3) tile_shared(T2) faster max|err|
512^3 0.043 ms 0.059 ms delegate, 37% both 2.48e-05
1024^3 0.320 ms 0.312 ms compiled, 2.3% both 6.10e-05
2048^3 2.448 ms 2.051 ms compiled, 16.2% both 1.54e-04
Errors are equal at every shape, so Decision #28's in-budget half is
satisfied outright and the displacement condition holds at 1024^3 and
above -- while `arbitrate()` still returns the delegate, because tier
priority is the default.
The test I first wrote asserted "the delegate is measurably the fastest"
and checked only 512^3, the one shape where that is true. It passed, and
it would have reported green for a default that is 16% wrong at 2048^3.
It now asserts the crossover in both directions, and checks that the
faster compiled kernel is no less accurate -- otherwise "faster" is not a
displacement argument.
This is not an argument for deleting the delegate: it wins by 37% at
512^3, and a flat "measurement beats tier" switch would regress that.
It is an argument for shape-bucketed measured selection, which is what
Decision #28's lead-safety is for.
Recorded in the NVIDIA queue with both follow-ups: wiring measured
selection into the OP_MATMUL path, and the two block-index conventions
that keep the emitted lane unmeasurable.
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: e8382823c5
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| def accuracy_budget(self, region: Any) -> "tuple[float | None, float | None]": | ||
| """``(atol, rtol)`` the F4 oracle must hold this candidate to for | ||
| ``region``. Defaults to the flat class attributes. | ||
|
|
There was a problem hiding this comment.
This introduces a region-specific accuracy hook in the shared Candidate runtime contract and extends the shared delegated-candidate semantics, but the commit updates only the NVIDIA plan; a repo-wide check of the Apple, ROCm, and x86 plans found no sibling-backend outcome. Record each backend as follow-up required, parity validated, or not applicable so consumers of this shared contract do not drift.
AGENTS.md reference: AGENTS.md:L81-L85
Useful? React with 👍 / 👎.
Comment on lines
+247
to
+249
| small = _device_timings(region, *_operands(512, 512, 512, "float16")) | ||
| assert min(small, key=small.__getitem__) == SHIPPED, ( | ||
| f"the delegate no longer wins at 512^3: {small}") |
There was a problem hiding this comment.
When this device suite runs on any NVIDIA GPU where the MMA runtime and PTX bridge load, nvidia_mma_ptx_launch_available() returns true without checking the model or compute capability (tests/_support/nvidia.py, lines 35–42). This assertion hard-codes a performance ranking measured only on an RTX 5070/sm_120, so an sm_80/sm_90 host, another sm_120 GPU, or a different driver can fail despite correct kernels; gate on the exact measured device or restrict the test to timing validity and accuracy-budget invariants.
Useful? React with 👍 / 👎.
gstoner
added a commit
that referenced
this pull request
Aug 31, 2026
…wups docs+test: address the #652 review — backend sync records, device-gate the ranking
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
The delegation contract from #649/#650 had zero live users. Measured on the fleet's sm_120 before this branch, all four NVIDIA matmul candidates returned
Nonefromcontract_for_candidatewhile the Tier-3 hand-tuned lane won arbitration on tier priority alone — no declared accuracy budget, no determinism claim, no architecture claim, no coverage claim.This declares the shipped
libtessera_nvidia_gemmmma.sync GEMM as aDelegatedCandidate, so its tier and budget are derived from a contract the Target IR verifier also enforces rather than hand-set in Python.The finding
Giving the delegate a device timer immediately contradicted the arbiter's default. sm_120, f16, square, device-resident CUDA-event timing, spreads of 0.000–0.008 ms across repeats:
shipped(T3)tile_shared(T2)Errors are equal at every shape, so Decision #28's in-budget half is satisfied outright. The displacement condition holds at 1024³ and above, and
arbitrate()still returns the delegate.This was invisible before because end-to-end wall time ranks them the other way (9.4 ms vs 33.1 ms at 2048³) — it is host-dominated: the Tile lane spends 2.99 ms on device inside 34.0 ms of wall. The Tier-3 lane had no device timer at all, so a Tier-3 delegate could never honestly lose.
Not an argument for deleting the delegate — it wins by 37% at 512³. It is the concrete argument for shape-bucketed measured selection.
What first use forced
nameis a dispatch/cache key, not a claim; the autotune corpus and E3forcekey on it. Deriving it fromcalleewould invalidate persisted verdicts on a C-symbol rename, with no error.calleeis identity. A per-dtype tolerance looked like a second motivation and measurement refuted it: bf16 and f16 agree within 25% at every K from 32→8192, because the oracle's reference rounds operands to the storage dtype first, so the residual is f32 accumulation error and is independent of storage width.mma.sync.aligned.m16n8k16(sm_80+); the sm_120 cubin is an AOT fast path. Contract now sayssm_80+;mma_archstayssm_120as a footprint-model key, documented as a different thing.6× headroom at K=8192 and breaks past K65536 on a correct kernel. Both bounds declared; the relative one carries large K.determinism="deterministic"is grounded, not assumed: one warp per 16×8 output tile, serial K reduction into four accumulators, no atomics/split-K. Themma.sync.m16n8k16.row.colfragment layout was checked against the PTX ISA canonical mapping for A, B and D before the claim.A correction inside the branch
I added a device timer for
nvidia_mma_gemm_emittedtoo; it returned rc=5 every time. Cause is real: the NVIDIA backend carries two block-index conventions —ptx_emitand the shipped AOT kernel map x→M, whileNVIDIALowering.cppand the launch bridge'sbenchmarkTileGemm16map x→N. Registering its geometry would launch a transposed grid (at 512×512: rows to 1024, columns only to 256 — half the output unwritten, with a plausible latency). Removed rather than shipped; recorded where the code is; a unit test pins the mapping so "fixing" one side fails loudly. The two Tile lanes are unaffected and their existing latencies are sound.Verification
main's sources over the same file set: 21 failures on this branch vs 23 on main. The one failure unique to this branch istest_dk3_rocm_moe_transport_perf_baseline_is_bounded, a timing-bounded test that passes 5/5 in isolation and is order-dependent; a different timing test flipped the other way on main.test_apple_gpu_delta_erase_routing.py::test_jit_path_threads_erase_end_to_endfails identically onmain(proven by stashing).🤖 Generated with Claude Code