compiler: CuTe IR assessment — verified layout algebra, LAYOUT-ALG-1

[gstoner]

Assessment of NVIDIA/cutlass#3426 (CuTe IR MLIR dialect, merged 2026-08-03 at 3664fe2), its executable math contract, and the LAYOUT-ALG-1 work it opens.

What the PR under review actually is

475 files, ~125k lines, 374 of them tests. 63 ops, 8 attribute-backed types. Zero math ops — no tensor/MMA/copy/atom files at all. It is indexing arithmetic, which is why the title misleads and why it matters to us.

1. Mathematical verification

Every worked example re-derived independently and machine-checked — 25 tests, pure integer layout evaluation, no cutegen, no numpy.

The algebra is sound, including the case most likely to be wrong: logical_divide((6,8):(8,1), (3,4):(1,3)), where the inner map is the identity on [0,48) yet the result is not the source regrouped, because A(j) = 8l₀ + 24(l₁ mod 2) + (l₁ div 2) + 2l₂ is genuinely non-affine in the l₁ mode. A layout system carrying only flat extent/stride arrays cannot express that result at all.

Three documentation defects, all prose, none algebra — each carried as a negative fixture so a port cannot inherit it:

Defect	Detail
Composition-with-shape prose contradicts its own example	Tutorial says A ∘ make_layout(shape); that yields [0..7] where the stated result yields [0,1,4,5,8,9,12,13]. Correct reading is mode-wise truncation — the ODS is right, the tutorial is wrong
Product table row for logical_product	Reads verbatim as its own blocked_product row; the examples show them different
complement's "covers exactly [0,M)" guarantee	False for the no-cotarget form — bijects onto a strict superset. A port reading the prose as the contract would build a wrong verifier

2. Scoping: 63 ops, four primitives

composition, complement, coalesce, right_inverse — everything else is mode regrouping. Confirmed: all six product variants and all four divide variants are one leaf multiset in different bracketings.

3. The convergence finding

Five in-tree items independently ask for layout reasoning, each proposing a different weaker mechanism: FORGE W1 (a six-point lattice approximating what composition + divisibility decides exactly), FORGE W2 (residency = cosize), SparDA §III.3 (GQA-fold = group_modes/logical_divide, not expressible today — rearrange fails closed on exactly that spec), TileSight §3.2 (rasterization = composition(grid, raster_layout)), G1b (the shared butterfly/FFT pass needs a shared representation to be shared in).

Layout algebra is not a tenth substrate investment — it is the missing implementation under CORE_SUBSTRATE_VIEW S9, whose ⊑ operator that document already flags unowned.

Changes since first review — the three decisions

1. L0 decided: one C++ implementation, Python binds via ctypes

This reverses my first-draft recommendation, which is recorded as withdrawn rather than quietly edited. The decision is the stronger one against our own governance: one implementation satisfies Decision #31 by construction rather than by declared-oracle exemption, and it makes the MLIR consumers first-class, so FORGE's ⊑ verifier and the emitter index math call the same code.

The cost I flagged is folded into L1 acceptance rather than dismissed — coupling emit/ to a build artifact is the failure the Apple dylib already demonstrates, and layout algebra sits under every emitter:

• A1 — the binding fails closed with one named diagnostic and ships no fallback path, because a fallback is a second implementation in disguise (Decision #21a).
• A2 — the build dependency is in the ordinary target set with a binding-loads test, so breakage surfaces at test time with a fix instruction, not inside an emitter.

Sizing moves: L1 ~1w → ~2w (ABI, build wiring, A1/A2); L5 ~2-3w → ~1-2w, since there is no second implementation to differential-test.

2. LAYOUT-ALG-1 bound in INTEGRATED_COMPILER_PLAN.md §4

Ordered L1…L5, with L1 gated on L2 being committed (otherwise L1 is a Decision #29 violation by construction) and L5 sequenced after W1.1 step 4 per the same #31 ordering caveat W1_1_TYPING_DESIGN.md makes for #tile.mma_desc. The assessment keeps mathematical and acceptance authority; the integrated plan keeps order and promotion.

3. Rasterization claim corrected — but only the half that is stale

Verified by executing emit/nvidia_cuda.py::_raster_launch, not by reading imports:

row_major     int mt=blockIdx.x*16, nt=blockIdx.y*8,
column_major  _tsr_mt = blockIdx.x % ((M+15)/16); _tsr_nt = blockIdx.x / ((M+15)/16)
grouped_m     _tsr_raster_pp = 4 * ((N+7)/8);  … panel walk …

All four emitters consume tile_rasterization.py; the MLX swizzle_log heuristic was retired. So "no emitter consumes it" is closed.

TileSight's substantive conclusion survives and is preserved as such. raster_order is carried, not swept — row_major remains the production choice everywhere, and automatic enumeration is deliberately withheld pending an architecture-owned correlation/retain verdict, because ROCM-CALIB-1 established that an unvalidated locality metric must not change a production raster choice. The lever is now expressible on every backend and still unpulled; the blocker moved from codegen to measured device evidence. Correcting this to "the finding is stale" would have erased a live gap.

That rescopes L4 down (~2w → ~1-2w): the plumbing exists, so L4 is a consolidation onto shared algebra with bit-identical output for every reachable (raster_order, raster_group) as the gate — a stronger acceptance test than the row-major-only check first proposed — and the measured half stays out of scope.

Files

File	Change
docs/audit/compiler/CUTE_IR_ASSESSMENT.md	New — review, verification, L0…L5
tests/unit/test_layout_algebra_contracts.py	New — 25 tests, three doc defects as negative fixtures
docs/audit/compiler/INTEGRATED_COMPILER_PLAN.md	LAYOUT-ALG-1 bound in §4
docs/audit/compiler/TILESIGHT_ASSESSMENT.md	§3.2 dated correction; original bullets preserved
docs/audit/compiler/CORE_SUBSTRATE_VIEW.md	§1 demand-matrix row corrected
docs/audit/compiler/README.md	Catalog entry + routing row

Gates

Gate	Result
tests/unit/test_layout_algebra_contracts.py	25/25 pass
ruff check	clean
scripts/check_generated_docs.sh	ok, 25 docs in sync
pre-push: op-catalog/spec + registry + docs-freshness	ok
Full tests/unit -m "not slow" sweep	still running (~40 min in) — will report. Docs plus one hermetic test with no project imports; no production code

Mac / Homebrew python3 3.14.6. No device lanes — the substrate is host-free by construction, so nothing here needs the Strix Halo or NR2 Pro boxes.

🤖 Generated with Claude Code

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 6b0b3d6910

ℹ️ 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".

[gstoner]

P2 — Preserve mode nesting in the layout fixtures

Fixed in 78791e4. The finding is correct, and the harness was weaker than the assessment claimed it was.

Confirming the hole. Four of the six product variants were written as the same tuple:

"logical_product": ((3,4,2,5), (4,1,12,24)),
"zipped_product":  ((3,4,2,5), (4,1,12,24)),
"tiled_product":   ((3,4,2,5), (4,1,12,24)),
"flat_product":    ((3,4,2,5), (4,1,12,24)),

In the docs those are ((3,4),(2,5)), ((3,4),(2,5)), ((3,4),2,5) and (3,4,2,5) — three distinct structures collapsed by hand, with leaves() sorting mode order away on top.

One consequence worth naming beyond the report: assessment §2's scoping claim — "all six variants are one leaf multiset regrouped", which is what "63 ops, four primitives" rests on — was close to circular for those four. The test asserted what I had typed, not what the algebra does.

Fix, stronger than adding a nested type. Fixtures are now the documented layout strings verbatim — layout("((3,4),(2,5)):((4,1),(12,24))") — parsed into nested (shape, stride) trees. That preserves nesting and removes transcription as a source of error. The parser carries its own test (including profile-mismatch rejection), since the fixtures are only as trustworthy as it is.

Every operation is now asserted on two independent axes:

• function — image over [0, size), invariant under regrouping;
• structure — the exact nested shape/stride, which is what distinguishes the variants and what downstream slicing/partitioning consume.

Neither implies the other, so both are kept. The equivalence classes are pinned exactly: the six products are five distinct structures, because logical_product and zipped_product genuinely coincide for a plain-layout tiler — asserting "all six differ" would have been wrong, asserting nothing let the collapse through. The multiset test survives but now states that it is meaningful only because the structural test proves the six are not the same object.

Verified by mutation, not by assertion. Against the new harness:

Mutation	Result
implementation flattens every product result	3 failed
blocked_product returns logical_product's grouping	3 failed
zipped_divide returns logical_divide's grouping	3 failed
the depth-2 nested divide split is flattened	1 failed

Against the old harness, that same zipped_divide mutation passed 25/25 — and "flatten everything" was literally what the old file already encoded. That is the empirical proof the concern was real rather than hypothetical.

25 → 32 tests. Assessment §1/§2 updated to describe what the harness now proves, with the earlier weakness recorded rather than quietly overwritten, and LAYOUT-ALG-1's L1 acceptance in the integrated plan now requires the same two-axis check of the port — so an implementation returning the right elements under the wrong grouping must fail.