Land Wave 0 of the integrated compiler plan, plus W1.1 typing groundwork - #490
Merged
Merged
Conversation
Executes all ten W0 items from docs/audit/compiler/INTEGRATED_COMPILER_PLAN.md
("stop the bleeding": live defects, fail-open paths, inert machinery, false
documentation), then the two follow-ons W0 exposed and the first W1.1 slice.
W0 — governance and correctness
* W0.8: adopt Decisions #21a, #10a, #29, #30, #31, #32 in CLAUDE.md, each
with its originating defect. #29/#31 are drift-gated by the new
tests/unit/test_governance_declarations.py: every primitive_coverage axis
must name an existing consumer, and no two ODS files may declare the same
dialect. The two genuinely unconsumed axes are explicit ratchet waivers.
* W0.1: `manifold` is now a verified EBM_ManifoldAttr and Canonicalize fails
closed instead of silently defaulting to "euclidean" (a Euclidean step on
spherical state converges and reports a confidently wrong answer).
* W0.2: CheckpointInnerLoop out of the default EBM pipeline — its three
attributes have zero consumers tree-wide — with a #10a CHECK-NOT fixture.
* W0.4: jacrev records the forward pass once and reuses the tape, as its
docstring always claimed (measured: 1 evaluation instead of 4).
* W0.6: delete three duplicate dialect ODS files. The #31 gate found the
third on its first run: TileMemoryOps.td declared the same `tile` dialect
with contradictory mnemonics, tablegen'd but never included or registered.
* W0.7: split "[GA8 stub]" into honest [annotation-only] labels; remove the
false claim that GA8 passes gate on `canonical`.
* W0.3: EBM energies are traced through autodiff.tape when a cotangent path
exists, with central differences retained for raw-NumPy callbacks
(1 energy evaluation instead of 2·2^n on Cl(3,0)).
* W0.9: add a real MLIR parse + dialect-load + verifier harness to the Target
IR contract test, keeping the substring assertions as smoke coverage.
* W0.10: decided to BUILD tessera_x86 rather than grant Decision #19 a
carve-out, and built it.
Side effect: ts-ebm-opt never registered `arith`, so 6 of its 12 lit
fixtures could not parse — invisible because the EBM backend is OFF by
default. EBM lit is now 12/12.
W0.9 fallout — no Python-emitted Target IR was valid MLIR
Fixed four stacked defects, each invisible to substring assertions:
undialect-prefixed module attributes; an invented `<dialect>.func` container
no dialect defined; ops emitted with signatures their own ODS rejects; and
five op names emitted but never declared. A second, duplicate emitter family
in matmul_pipeline.py had the same defects. The gate now also parses every
committed golden, which is what caught the fifth defect.
tessera_x86 (Decision #19)
New hardware-free Target IR dialect with a real !tessera_x86.tile type,
registered in tessera-opt. Value-carrying AMX ops are separated from
directives, and abi_call models the C-shim boundary rather than hiding it.
Ships positive AND negative fixtures — the negative one proves the verifier
rejects a dot-product whose operands never came from a tile load.
Per project direction, AMX lowering is optional (expected supersession by
ACE); the ops remain as the IR-level contract.
cpu reference lane
Collapsed `tessera.cpu.<source-op>` (one name per Graph IR op, unenumerable
in ODS) to a single declared tessera.cpu.reference. Lossless: the verifier
already required a `source` attribute. No target is excluded from the
parse/verify gate any more except NVIDIA, whose dialect is off by default.
W1.1 groundwork
* docs/audit/compiler/W1_1_TYPING_INVENTORY.md — the producer/consumer
inventory W1.1 requires before tightening. Headline: W1.1 is a migration
already in progress, not a design job.
* tile.async_copy / tile.wait_async gain a verified typed-token form.
* The missing !tile.fragment + tile.mma fixtures now exist (positive plus
four negative cases) — the regression net the producer migration needs.
* Records the blocker that stops the migration: the typed mma form (exactly
3 raw operands) and warp-spec token sync (requires a token operand) are
mutually exclusive, verified empirically. That is why no C++ producer
emits the typed form.
Verification on Ubuntu/Strix Halo (gfx1151, LLVM 23): unit 14148 passed /
0 failed; lit main 288, EBM 12, Clifford 13; 24 generated docs in sync.
Fleet routing and the CLAUDE.md toolchain section retargeted to this box as
the primary for core compiler work.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
![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: 0003e9d7ad
ℹ️ 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".
…sync
CI lint
Two F401s orphaned by the W0.4 jacrev rewrite (`grad`,
`_normalize_argnums`). ruff is clean on python/tessera/, which is CI's scope.
P1 — jacrev raised on outputs the tape never produced
`jacrev(lambda x: x)` and `jacrev(lambda x: <constant>)` both raised. The
pre-W0.4 implementation wrapped fn in `sum(out * cotangent)` through ops.*,
which made the backward target tape-produced regardless of what fn returned;
reusing one tape removed that accidental shield. Both cases have well-defined
Jacobians and are now resolved structurally (identity / zero).
P1 — aliased cotangents were first-match-wins instead of summed
`Multivector.coefficients` returns a fresh whole-array view per access, so an
energy reading the state twice registers two tape ids, each holding half of
dE/dx. Returning the first match produced a gradient of 1 where the true
value is 2 — a silently wrong Langevin step, not an error. Now accumulates
across every alias. Does not reproduce with a plain ndarray (both operands
share one id), which is why the Multivector path is the regression test.
P1 — _copy_target_op dropped the typed-op fields
operand_types / result_type / prelude were not copied, so probe annotation
downgraded typed ops back to `() -> ()` while keeping their operand
references: the annotated module stopped parsing ("expected 3 operand types
but had 0"). Fixing that uncovered a second defect underneath — the
`*.profiler_probe` ops are emitted into every lane but declared by no
dialect, the same class W0.9 fixed for five other names. Declared them for
ROCm, Apple (cpu + gpu), and x86; `cpu` already had one. All ten
target x {plain, probe-annotated} combinations now parse and verify.
Cross-backend sync (AGENTS.md)
This PR changes shared Target IR, so all four backend plans now carry a
`TARGET-IR-CONFORMANCE-2026-08-02` entry with an architecture-specific
outcome. NVIDIA is explicitly **follow-up required, NOT validated**: its
dialect is off in the default build, so the new parse/verify gate SKIPS
rather than passes it, and tessera_nvidia.profiler_probe is likely to carry
the same undeclared-op defect. A green run here is not evidence for SM120.
Three regression tests added so none of these can silently return.
Verification: unit 14155 passed / 0 failed; lit main 288, EBM+Clifford 25;
24 generated docs in sync; ruff clean in CI scope.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
gstoner
pushed a commit
that referenced
this pull request
Aug 3, 2026
P1 — NVIDIA lowering still counted RAW operands
MMAOp::verify() now counts data operands, but NVIDIALowering.cpp compared
raw op->getNumOperands() against 3/5 and then indexed raw operands. So the
exact typed-plus-token form this change exists to unblock would have hit
emitError + signalPassFailure during SM120 lowering -- and operand 3 would
have been the async token rather than an NVFP4 scale. Both the count and the
indexing now use tessera::tile::dataOperands, matching the verifier and the
ROCm lowering.
Verified by BUILDING it rather than by inference. ROCm and NVIDIA cannot both
register in one tessera-opt ("the lean registration path has no arm for
them"), which is exactly why this went unnoticed: the default build has
NVIDIA off. Configured a second tree (build-nvidia) and compiled it.
That also closes a gap carried since PR #490: with tessera_nvidia actually
registered, the W0.9 parse/verify gate PASSES for sm90 / sm100 / sm120, plain
and probe-annotated. It had been SKIPPING every run. The contract test now
discovers either build tree, so NVIDIA is no longer silently unmeasured --
38 passed, zero skips.
P1 — cross-backend sync (AGENTS.md)
All four backend plans record TILE-MMA-DATA-OPERANDS-2026-08-03 with an
architecture-specific outcome: NVIDIA parity validated (and its prior
NOT-VALIDATED status closed), ROCm parity validated with no behavior change
(it composes the shared rule with its own token type rather than forking),
x86 not applicable with a reason (no tile.mma consumer; re-assess when the
x86vector lowering lands), Apple not applicable to the operand contract but
touched by the storage-dtype enforcement ordering.
The reviewer's second point was well aimed: the missing cross-backend
assessment is what allowed the incompatible NVIDIA raw-operand check to sit
unnoticed.
Verification: unit 14210 passed / 0 failed; lit main 288, EBM+Clifford 25;
24 generated docs in sync; ruff clean.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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.
Executes all ten W0 items from
INTEGRATED_COMPILER_PLAN.md— the "stop the bleeding" wave: live defects, fail-open paths, inert machinery, and false documentation — then the two follow-ons W0 exposed, and the first W1.1 slice.W0 was the only wave with no dependencies. The plan's risk table warns specifically against starting later ("Someone starts at W3"), so the ordering here is deliberate.
W0 — governance and correctness
CLAUDE.md, each recorded with its originating defect. #29/#31 are drift-gated by a newtest_governance_declarations.py.manifoldis a verifiedEBM_ManifoldAttr;Canonicalizefails closed instead of silently defaulting to"euclidean". A Euclidean step on spherical state doesn't diverge — it converges and reports a confidently wrong answer.CheckpointInnerLoopout of the default EBM pipeline (its three attributes have zero consumers tree-wide), with a Decision #10aCHECK-NOTfixture.jacrevrecords the forward pass once and reuses the tape, as its docstring already claimed. Measured: 1 evaluation instead of 4.[GA8 stub]split into honest[annotation-only]labels; removed the false claim that GA8 passes gate oncanonical.autodiff.tapewhen a cotangent path exists, central differences retained for raw-NumPy callbacks. 1 energy evaluation instead of 2·2ⁿ on Cl(3,0).tessera_x86rather than grant Decision #19 a carve-out — and built it.What the new gates found immediately
The value here is mostly in what the drift gates caught on their first run:
TileMemoryOps.tddeclared the sametiledialect as the production ODS with contradictory mnemonics (mma.tcgen05vs the livetcgen05.mma) — tablegen'd, but never#included and never registered.<dialect>.funccontainer no dialect defines; ops emitted with signatures their own ODS rejects; and five op names emitted but never declared. A second, duplicate emitter family inmatmul_pipeline.pyhad the same defects.ts-ebm-optnever registeredarith, so 6 of its 12 lit fixtures couldn't parse — invisible because the EBM backend is OFF by default. Now 12/12.tessera_x86(Decision #19)x86 was the one backend with no Target IR dialect. It now has one, with a real
!tessera_x86.tiletype, registered intessera-opt. Value-carrying AMX ops are separated from directives, andabi_callmodels the C-shim boundary instead of hiding it so Decision #28's arbiter can distinguish compiler-generated from delegated work. Ships positive and negative fixtures — the negative one proves the verifier rejects a dot-product whose operands never came from a tile load, which is precisely what a substring test cannot check.Per project direction, AMX lowering is optional (expected supersession by ACE); the ops stay as the IR-level contract. Attributed as owner direction rather than asserted as verified fact.
W1.1 groundwork
W1_1_TYPING_INVENTORY.mdis the producer/consumer inventory W1.1 requires before tightening anything. Headline: W1.1 is a migration already in progress, not a design job —MMAOp::verify()already enforces the typed contract and explicitly keeps a legacy branch.tile.async_copy/tile.wait_asyncgain a verified typed-token form, and the missing!tile.fragment+tile.mmafixtures now exist.The migration itself is not started, and the inventory records why: the typed
mmaform (exactly 3 raw operands) and warp-spec token sync (requires a token operand) are mutually exclusive — verified empirically. That is why no C++ producer emits the typed form. The proposed fix is noted but deliberately left for its own verification pass.Corrections to my own work, kept in the record
Two inventory claims were wrong and are corrected in place rather than quietly edited:
Tile_AsyncTokenTypeis used but never ODS-typed (not a #29 violation), and the "10 producers" are really 5 construction sites. Both were exposed by implementing rather than reading — and one wrong verifier rule was caught by six existing fixtures, which is the concrete argument for W1.1's own "inventory before tightening" precondition.Verification
Run on Ubuntu/Strix Halo (gfx1151, LLVM 23):
Fleet routing and the
CLAUDE.mdtoolchain section are retargeted to this box as primary for core compiler work. Also corrects a routing error: AMX proof was routed to "the Zen5 box", but AMX is Intel-only — no machine in the fleet can run the lane added in #489.🤖 Generated with Claude Code