nvidia: make kernel_call a real delegation contract; add inline_ptx - #649
Merged
Merged
Conversation
…_ptx Enabling step for the bootstrap prune. The 19 NVIDIA bootstrap packagers contain legitimate fast paths -- vendor libraries, hand-tuned kernels, inline PTX -- and pruning before they have a declared boundary to land on would delete capability. NVIDIA first because it has both the largest gap and working profiling tools; the arbiter is MEASURED, so a boundary on an unprofilable target is bookkeeping. kernel_call was a summary line and nothing else. It inherited the shared attr-dict, so `callee` -- the one fact naming what is delegated to -- rode as an unvalidated discardable attribute; an emitter could name any symbol, or none, and verify. The dialect header says why it existed: to keep Python-emitted IR parseable. It was a parse-compatibility stub for the very packagers being pruned, which is Decision #29's anti-pattern. Both pathways are declared, as two ops rather than one with a mode. kernel_call takes callee/arch/binding/provenance/accuracy; inline_ptx takes ptx/constraints/arch/accuracy/has_side_effects. Separate because the delegate differs in kind -- a binding resolved at link time versus text carried in the artifact -- and their failure modes differ: an empty callee is an unresolved symbol, an empty ptx body is a silently successful no-op. One op with a mode attribute needs a verifier that decides which half of its own attributes to trust. The attributes are the arbiter's inputs, which is what makes this real rather than decoration. `accuracy` is the budget half of "fastest in-budget candidate": tolerance_bounded must state a tolerance, reference_exact must not carry one. Semantic key, never defaults (#21a). Evidence on The-Super-Bear with the full driver: builds clean, positive fixture parses both ops, the new negative fixture rejects 7 cases, and the NVIDIA lit suite is 60/60. Also folds in the review P2 on #648: the lean driver's core spine is LINKED (TesseraPM PUBLIC-links TesseraIR/TesseraScheduleIR) but not REGISTERED. "Not linked" was an inaccurate build diagnosis and would misdirect anyone debugging a lean-driver failure. Not done and stated as such: nothing yet queries these attributes to score a candidate. The verifier is a real consumer, but arbiter integration is the next step, and the ROCm equivalent is still owed.
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: 06724a6b7b
ℹ️ 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".
Step 1 of the prune sequence. The contract landed in the previous commit verified but unread -- nothing scored a delegate, which was the stated gap. This closes it. `DelegatedCandidate` derives BOTH arbiter-relevant facts from the IR contract rather than accepting them at the registration site: tier comes from `provenance` and the F4 budget from `accuracy`/`tolerance`. A delegate therefore cannot claim in Python a budget it did not declare to the verifier. Everything else about arbitration is unchanged, which is the point -- a delegate is enumerated, F4-gated and selected by the same arbitrate() path as compiled candidates, so Decision #28 scores a hand-tuned kernel AGAINST compiled output rather than above it. Both provenances map to Tier.HAND_TUNED: a vendor library and an in-tree hand-tuned kernel differ in origin but not in what the arbiter must do with them. `provenance` is still carried because it is what distinguishes them in a dispatch log. `reference_exact` yields atol=None -- the oracle's default budget, the same standard compiled output is held to -- deliberately not 0.0, which would reject a correct candidate. An exact claim is "no worse than the reference", not "bit-identical in floating point". The Python validator is not a second contract. It rejects exactly what TesseraNVIDIADialect.cpp rejects, and test_delegate_contract.py asserts that case for case, including a count check so neither side can grow a case alone. Two enforcers of one contract is the shape that produced the Apple two-disconnected-compilers defect; the differential is what makes this a bridge rather than another instance of it. Fixture grew to 9 cases (added unknown_provenance, empty_arch) to match. Evidence: 19 Python tests, NVIDIA lit 60/60 on The-Super-Bear, ruff and mypy clean, 201 arbiter-adjacent tests pass.
Step 2 began by asking which NVIDIA packagers to migrate onto the new delegation boundary. The measurement says: none of them. nvidia_native.py has ZERO references to NVRTC, cuBLAS/cuDNN/CUTLASS, any .so, or raw device source. 13 of its 19 bootstrap packagers construct Tile IR and compile it through tessera-opt -- the MLIR pipeline already runs from Tile onward, and what bypasses it is Graph -> Schedule -> Tile. NVIDIA's real delegation surface is ptx_emit.py, emit/nvidia_cuda.py and runtime.py: different files, different work. Across all four backends: 24 of 34 bootstrap packagers are IR-constructing, 1 delegates, 1 both, 8 are wrappers. So the prune is overwhelmingly an ABSORPTION job, not a delegation-migration job. That corrects a claim I made when justifying the boundary work -- the boundary was still right to land first, but for the delegation surface that actually exists rather than for these packagers. The taxonomy is a new dashboard axis so it tracks rather than decays. Its detector is a regex, not a literal, because each backend spells the helper differently -- NVIDIA _compile_tile_ir, ROCm _compile_attention_tile_ir, x86 emit_matmul_tile_ir plus a direct tessera-opt call. Keyed to the first spelling it classified the other two backends as "other", which is a taxonomy that reports nothing; the first run said 14/2 and the corrected one says 24/1. Kind counts are sub-rows of the bootstrap total (they sum to 34), not of the gap count -- the first placement implied a relationship that did not exist. Tests assert every backend's spelling is detected, that no packager falls out of the accounting, and that the surface stays mostly IR-constructing, so a flip in its character is visible rather than silently re-scoping the prune.
Two live defects in the delegation contract as first shipped, found by stress-testing the design rather than by a failing test. Determinism was undeclarable. Tessera guarantees @jit(deterministic=True) and a split-K delegate accumulating with atomics is not reproducible run to run, so the arbiter could have selected one inside a deterministic region. Same shape as the Decision #5 scar: a guarantee defeated through a path nobody checked. `determinism` is now a required enum on both ops. The accuracy claim was absolute-only while Candidate already carried atol AND rtol. An absolute bound is meaningless without the result's magnitude -- 1e-6 is vacuous at 1e6 and unsatisfiable at 1e-9 -- so a delegate whose real claim was relative had to overclaim. tolerance_rel added; either or both now satisfy a bounded claim, and DelegatedCandidate carries both budgets. Decision #19 amended. Its stated justification was wrong twice over: "hardware-free" was never true (its own example, tessera_rocm.mfma, names AMD hardware) and lit-testability is not what the layer buys (NVVM and ROCDL are MLIR dialects; mlir-opt verifies them on any host). What it actually buys is CONTRACT CARRIAGE -- it is the last level where numeric_policy, layout, distribution and arbiter metadata still exist, and lowering straight to NVVM would drop them, which is exactly Decision #32's recorded scar. The corrected justification supplies a membership test: an op belongs in tessera_<backend> when it carries a Tessera contract upstream cannot express, and its description must name which. That is deliberately NOT a bar on duplication -- a Tessera FMA may resemble vector.fma and earn its place by carrying an accumulator contract or an avx10.2 feature gate. So operator expansion is expected, and Apple and x86 are named as the two that should grow. Also records the remaining design gaps found in the same pass, ordered by whether the design is wrong or merely incomplete: per-op accuracy budgets do not compose, fusion foreclosure is uncosted so the arbiter is biased toward delegates, kernel_call does not verify its callee ABI, delegates are unversioned, and the arbiter sits on the Python side of the prune. Evidence: NVIDIA lit 60/60 on The-Super-Bear with the new required attribute, 22 Python contract tests, mypy and ruff clean.
A sweep of all 32 decisions (plus sub-decisions) against the direction: MLIR/LLVM core, prune the Python bootstrap backend path, contract- carrying Target IR, measured three-tier arbiter. Most support it. Six did not, and one was an unresolved conflict. #28 vs #31 -- the one that mattered. #31 says one production lowering per boundary and delete the second; #28 keeps three tiers of kernels deliberately competing for the same op. Read literally, every Tier-3 candidate is a #31 violation, and the delegation contract sits exactly on the ambiguity. The bootstrap prune therefore had no principled stopping point: #31 could be cited to delete the whole Tier-3 population, which is the ceiling #28 exists to protect. Reconciled in both places -- #31 governs lowering PATHS (how IR descends a level), #28 governs implementation SELECTION (which kernel runs for one op at one level). The test is not how many kernels exist but how many authorities decide what the next level looks like. #1 was actively harmful, not merely stale: it named AMX as the only execution path (retired) and gated GPU work behind isa >= SM_90, which reads as excluding sm_120 -- the live NVIDIA lane. Applied literally it gates off working hardware. #11 keys the autotune cache on {op, shape, dtype, arch, layout, numeric_policy, movement} with nothing versioned. Under #28 a cached entry is a measurement, and a measurement is only valid for the code that produced it; a toolkit upgrade silently invalidates every entry without invalidating the cache. Same failure as the Krylov ratchet, latent in a database instead of a JSON file. #12's schema cannot say which route produced a latency, so three competing tiers are not comparable. Practice was already ahead of the rule -- record_sm120_packet.py stamps `route` -- so this is a schema gap, and the added field is additive. #26a's "revisit on architectural grounds" trigger arrived. The architectural gap is real (Apple's Target IR declares dispatch containers and no machine primitives, while apple_msl.py already models simdgroup_matrix) but it is answered by up-levelling the dialect, NOT by emitting AIR -- NVVM and ROCDL sit above LLVM IR too. That strengthens the deferral rather than reversing it. #29 gains the sequencing corollary that keeps operator expansion honest: add each op only when its producer and consumer land with it. Gates: docs lint, 36 governance/audit tests, 28 generated docs in sync.
This was referenced Aug 30, 2026
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.
Process note first: these commits were originally pushed to
fix/nvidia-lean-driver-claimafter #648 merged, so they were stranded on a closed PR and never reached main. Rebased onto main here. (Second time this session I continued work on an already-merged branch — same shape as #641.)Why this had to land before any pruning
The 19 NVIDIA bootstrap packagers contain legitimate fast paths — vendor library entries, hand-tuned kernels, inline PTX. Deleting them before those paths have a declared boundary to land on would be capability loss, which is exactly what Decision #31's ordering caveat exists to prevent.
NVIDIA over ROCm because it has both the largest gap (19 of 34 bootstrap packagers) and working profiling tools — the second reason matters more: Decision #28's arbiter is measured, so a delegation boundary on a target that can't be profiled is bookkeeping, not a candidate.
What
kernel_callwasA summary line and nothing else. It inherited
TesseraNVIDIA_Op's sharedattr-dict, socallee— the single fact naming what is delegated to — rode as an unvalidated discardable attribute. An emitter could name any symbol, or none, and still verify.The dialect header states why it existed: Python emitters "may add
tessera_nvidia.kernel_call", and registering it "keeps the emitted surface parseable". It was a parse-compatibility stub for the very packagers being pruned — Decision #29's anti-pattern.Both pathways, as two ops rather than one with a mode
kernel_callcallee,arch,binding∈ {cuda_kernel,c_abi},provenance∈ {vendor_library,handwritten_kernel},accuracyinline_ptxptx,constraints,arch,accuracy, optionalhas_side_effectsSeparate ops because the delegate differs in kind — a binding resolved at link/launch time versus text carried in the artifact — and their failure modes differ: an empty
calleeis an unresolved-symbol error; an emptyptxbody is a silently successful no-op. One op with a mode attribute would need a verifier that decides which half of its own attributes to trust, which is the shape that lets a malformed candidate through.What makes it real rather than decoration
The attributes are the arbiter's inputs.
accuracyis the budget half of "fastest in-budget candidate": atolerance_boundeddelegate must state itstolerance, andreference_exactmust not carry one — two contradictory claims leave a reader unable to tell which is honoured. Semantic key, never defaults (#21a).provenanceis what lets the arbiter tell delegated from compiler-generated work when scoring.Evidence (The-Super-Bear, full driver)
tessera-nvidia-optbuilds cleanbinding, empty constraints, empty ptxNot done, named rather than implied
Nothing yet queries these attributes to score a candidate. The verifier is a real consumer so Decision #29 is satisfied, but arbiter integration — reading
provenance/accuracyto admit or reject a Tier-3 candidate against compiled output — is the next step. Until it lands, this is a contract without a scorer. The ROCm equivalent is also still owed.Also folds in the #648 review P2: the lean driver's core spine is linked but not registered (
TesseraPMPUBLIC-linksTesseraIR/TesseraScheduleIR); "not linked" was an inaccurate build diagnosis that would misdirect a lean-driver debug.🤖 Generated with Claude Code