Metal TensorOps Primitive Spec

Status: 📋 Proposed (design only; no implementation yet) Scope: Close the gap between Iron’s existing coop_tile_* cooperative-tensor primitive and the Metal 4 / MetalPerformancePrimitives (MPP) TensorOps surface, so the 15 hand-written butter_*.metal kernels Butter still ships become emitted Iron kernels — and so PSOCache.isMppKernel’s cross-repo name sniff retires. Out of scope: rewriting the Op::CoopTile* IR family, the _nax kernels, or the CUDA/SPIR-V mirrors — this spec extends them. Also out of scope: model loading, graph execution, tokenization. Iron is an optimized-kernel generator, not an inference engine. References: WWDC26 session 330, “Optimize custom machine learning operations with Metal tensors” — https://developer.apple.com/videos/play/wwdc2026/330/ · Crosley, “Metal for Machine Learning in 2026” — https://blakecrosley.com/blog/metal-machine-learning-2026 · Rigel: Reverse-Engineering the Metal 4.1 Tensor Compute Path on the Apple M4 Max GPU, arXiv:2606.12765 · BaseRT: Advancing Best-in-Class LLM Inference with Apple M5 Neural Accelerators, arXiv:2607.19438 · ggml-org/llama.cpp#16634 · Butter planning/known-issues.md “Kernel provenance” item 3 (2026-09-07) and planning/gpu-saturation-and-bandwidth-spec.md §S10.

Read ../developing.md “Kernel-authoring hazards” first. Every kernel this spec moves is a mode = Reduction cooperative-tensor kernel; a wrong dispatch geometry there hard-freezes the machine, and an empty body ships all-zeros through xcrun metal, the smoke build, and MSL snapshots.

---

1. Motivation

Butter tracks 15 hand-written .metal files under Sources/IronSwift/Resources/kernels/ — 4 050 lines of MPP cooperative-tensor and simdgroup_matrix GEMM that Iron does not emit. Butter’s own audit (planning/known-issues.md item 3, 2026-09-07) files them as blocked:

iron has no MPP / TensorOps primitive yet; until it does, these cannot be emitted. Metal 4 TensorOps support in iron is the prerequisite […] — file as an iron spec item.

That premise is wrong, and the correction is the whole point of this spec. Iron has had a first-class TensorOps primitive since the Op::CoopTile* family landed. Two of the fifteen — butter_nvfp4_moe_gather_bm16_mpp_xrows_{bf16,f32} — literally begin // Generated by Iron and carry SSA artifacts (v25, __ml_sub_offset) in their bodies. They are Iron’s own output, forked by a hand patch, with a comment that predicts its own loss:

NOTE: if this file is ever regenerated via make regenerate-kernels (Iron DSL, sibling ../iron repo), this hand patch will be silently dropped; port it to the Iron source template too.

So the blocker is coverage and maturity of an existing primitive, not its absence. Sorting the 15 against what Iron ships today:

pile	count	what it actually needs
A — already Iron output, forked	2	Re-land one hand patch into the Iron source. No new primitive.
B — expressible on today’s coop_tile_* / simdgroup_elem_*	8	Port work only. No codegen change.
C — needs new DSL surface	5	Real gaps: native bfloat operands, relaxed_precision, register-direct fragment fill, dual destination tiles, mem_none.

Eight of the fifteen are pure dtype twins (four {bf16, f32} pairs) — roughly 1 500 of the 4 050 lines are erased by a dtype axis alone. Ten of the fifteen collapse into three parametrised definitions (§5.6).

The cost of leaving them hand-written is not just duplication. Every one is outside the manifest, the PSO smoke test, and the empty-body MSL detector; the staging-bias hand patch in pile A is one make regenerate-kernels away from silent deletion; and PSOCache.isMppKernel decides live-compilation by substring-matching _mpp_ or bgemm in the kernel name, which already misfires on four kernels that contain no MPP at all.

---

2. Goals and non-goals

Goals

• Emit all 15 from wh-iron-std, verified bit-exact against the archived hand-written kernel before the hand-written file is deleted.
• Replace the _mpp_ / bgemm name sniff with per-kernel manifest metadata derived from Kernel::requires_cooperative_tensors(), which already computes exactly the right predicate.
• Settle, by on-device probe rather than by comment, whether matmul2d accepts native bfloat cooperative tensors — the question that gates the T23 staging-bias fold, coop_stage(), and pile C.
• Document coop_tile_* and simdgroup_* in ../STYLE_GUIDE.md §10, which today lists neither.
• Add a ## COOPERATIVE TENSOR INVARIANTS convention alongside ## DISPATCH INVARIANTS for the geometry rules a matmul2d kernel cannot self-check.

Non-goals

• A parallel Op::TensorOp* IR family. One primitive; new fields and sibling ops on Op::CoopTile*.
• Native quantized MTLTensor parameters with E8M0 scale planes (§4.6). That changes the kernel ABI from buffers to tensors and needs its own spec.
• Widening the _nax / _mpp twin split any further. It should shrink.
• Any change to Butter beyond the repin and the PSOCache swap.

---

3. Current state

3.1 What Iron already has

#[kernel] DSL      coop_tile_setup / _zero / _load_a / _load_b / _run / _store_c
(wh-iron-macros)   / _capacity / _coord / _get ; coop_stage(T)
                   crates/wh-iron-macros/src/kernel/body.rs:661-674, 2058-2350
        │
        ▼
IR                 Op::CoopTile*  — crates/wh-iron-core/src/ir/op.rs:1184-1301
(wh-iron-core)     CoopTileScope::{SimdGroup,Threadgroup}   — op.rs:407-419
                   CoopTileAccMode::{Overwrite,MultiplyAccumulate} — op.rs:421-428
                   Kernel::requires_cooperative_tensors()  — ir/kernel.rs:255-273
        │
        ├── MSL    msl/emit_block.rs:618-800  → mpp::tensor_ops::matmul2d
        │          msl/features.rs:63-67,206-236 → KernelFeatures::needs_mpp
        │          msl/mod.rs:326-336, 557-566   → header + body #if guard
        ├── CUDA   cuda/mod.rs:634-707          → nvcuda::wmma
        └── SPIRV  spirv/mod.rs:592-711         → cooperative matrix

Forty-five source files call coop_tile_setup. The library already covers most of Butter’s shapes: gemm/block_scaled_qmm_nax.rs carries a 28-format variant axis that includes nvfp4; moe/moe_mpp_bm16_bk32_block_scaled.rs emits iron_nvfp4_moe_gather_qmm_bm16_bk32_mpp; moe/moe_mpp_bm64_bk64_nvfp4.rs is the BK=64 NVFP4 gather; gemm/steel/steel_gemm_gather_nax.rs is the gather GEMM. The canonical minimal reference is probes/mpp_matmul_probe.rs — a hand-built Kernel IR (not #[kernel]) doing exactly Setup → Zero → LoadA → LoadB → Run → StoreC.

A second, older MMA surface also exists and matters here: Op::SimdgroupAlloc / SimdgroupElemLoad / SimdgroupElemStore / SimdgroupLoad / SimdgroupMatMul (op.rs:978-1027, DSL spellings simdgroup_alloc::<T,M,N>(), simdgroup_elem_{load,store}, simdgroup_matmul). It needs no OS gate — Apple7+.

coop_tile_setup already carries direct_inputs plus a_is_tg/a_ei/a_eo and b_is_tg/b_ei/b_eo, so device-address-space operands through tensor_inline already work (emit_block.rs:715-755 emits const_cast<device T*>(ptr)). That is most of the register-direct idea already present.

3.2 The 15 Butter kernels

All paths relative to ~/Development/personal/ai/butter/Sources/IronSwift/Resources/kernels/. Every one is live-compiled: eleven match _mpp_, four match bgemm. Flags reach the dispatch sites through ButterEnv[ButterFlag], and every MPP site is additionally gated on Ops.mppAutoCapable (macOS ≥ 26 and supportsFamily(.apple9), Sources/ButterSwift/Ops/Ops.swift:120-136).

Pile A — already Iron-emitted, forked

kernel	lines	descriptor / geometry	dispatch site	flag
butter_nvfp4_moe_gather_bm16_mpp_xrows_bf16	596	matmul2d_descriptor(16, 32, 16, ta=f, tb=t, tc=f, multiply_accumulate), execution_simdgroup; BM16/BN32/BK16, TPG 32, 1 SG; grid (nOut/32, ⌈mTotal/16⌉, 1); <half,half,float>	Ops/OpsNVFP4.swift:679, 729	BUTTER_NVFP4_MPP_XROWS (off)
butter_nvfp4_moe_gather_bm16_mpp_xrows_f32	582	same, <float,float,float>	Ops/OpsNVFP4.swift:680, 730	same

Non-expressible today: nothing. The delta from moe/moe_mpp_bm16_bk32_block_scaled.rs’s family is the x_rows A-row indirection plus, in the bf16 file only, the T23 stage-bias fold (kNvfp4StageBias = 2^-14 folded into both staged operands, kNvfp4StageComp = 2^28 applied once at the store). The f32 file has no kNvfp4StageBias at all — it does not need one.

These two also carry the only real pre-Metal-4 fallback in the set:

#else
    // Pre-Metal-4 stub: mpp::tensor_ops is unavailable.
#endif

which is exactly what msl/mod.rs:557-566 emits.

Pile B — expressible on today’s DSL surface

kernel	lines	primitive	geometry	dispatch site	flag
butter_nvfp4_gather_qmm_rhs_nax_mpp_bf16	238	matmul2d(32,32,32, tb=t, mul_acc), <half,half,float>	BM=BN=64, BK=32, TPG 128, 4 SG (2×2); grid (nOut/64, ⌈mTotal/64⌉, 1); nOut%64==0, kIn%32==0, blockSize==16	Ops/OpsNVFP4Gather.swift:664, 776	BUTTER_NVFP4_GATHER_NAX (off)
butter_nvfp4_gather_qmm_rhs_nax_mpp_f32	180	same	same	OpsNVFP4Gather.swift:665, 777	same
butter_nvfp4_gather_qmm_rhs_nax64_mpp_bf16	272	same descriptor, two 64×32 half-blocks staged per K-tile	BM=BN=BK=64, TPG 128; kIn%64==0	OpsNVFP4Gather.swift:499, 556	BUTTER_NVFP4_GATHER_NAX64, BUTTER_NVFP4_GATHER_BK64 (off)
butter_nvfp4_gather_qmm_rhs_nax64_mpp_f32	207	same	same	OpsNVFP4Gather.swift:500, 557	same
butter_nvfp4_gather_steel_bgemm_bf16	202	simdgroup_matrix<bfloat,8,8>, no MPP	wn=2, subM=subN=32, bm=bn=64, 4 accs, TPG 128; ragged expert-run walk	OpsNVFP4Gather.swift:126, 377	BUTTER_NVFP4_GATHER_STEEL (off)
butter_nvfp4_gather_steel_bgemm_dense_bf16	197	same, single expert	same, 9 buffers (no indices)	OpsNVFP4Gather.swift:72	same
butter_nvfp4_gather_steel_bm16_bgemm_bf16	203	same	wn=4, subM=subN=16, bm=16, bn=64, 2 accs	OpsNVFP4Gather.swift:172, 210	..._STEEL_BM, ..._THRESHOLD
butter_nvfp4_gather_steel_bm16_bn128_bgemm_bf16	175	same	wn=4, subM=16, subN=32, bm=16, bn=128, 2 accs	OpsNVFP4Gather.swift:309	same

The four steel_* kernels are the surprise. Their includes are <metal_stdlib> + <metal_simdgroup_matrix> — no MetalPerformancePrimitives.h, no mpp::, no __METAL_VERSION__ guard. They are legacy simdgroup_matrix MMA with both operands’ thread_elements() filled straight from device memory, zero threadgroup memory, zero barriers. Iron’s Op::SimdgroupElem{Load,Store} + SimdgroupMatMul express exactly this. They are in the “15 MPP kernels” bucket only because they were deliberately named *bgemm* to opt into live compilation — the dense kernel says so:

Named with a bgemm substring so PSOCache live-compiles it from this source against the running OS’s Metal toolchain (same mechanism the NAX/NAX64 MPP kernels use).

They also stage into bfloat rather than half and therefore carry no T23 fold — “that overflow class structurally cannot occur (empirically verified by the magnitude-ramp test, not merely assumed)”. They need no Ops.mppAutoCapable gate either, which is a genuine portability advantage.

Pile C — needs new DSL surface

kernel	lines	what Iron cannot express	dispatch site	flag
butter_dense_gemm_nax2_mpp_bf16	335	See §3.3.	Models/Text/LagunaMppGemm.swift:72, dispatchNax2; TPG 256	BUTTER_LAGUNA_NAX2_GEMM — default ON
butter_dense_gemm_mpp_bf16	208	matmul2d(32,32,32, tb=t, mul_acc), half-staged, BM=BN=64, BK=32, TPG 128, half As[2048]/half Bs[2048]/float OutScratch[4096] = 24 KiB. Structurally iron_gemm_q8_mpp with a dense bf16 weight; blocked only on the bf16 question and the T23 fold.	LagunaMppGemm.swift:65, 270; grid ((n+63)/64, (m+63)/64, 1)	BUTTER_LAGUNA_MPP_GEMM — default ON
butter_nvfp4_moe_gateup_bm16_mpp_bf16	194	Two destination cooperative tensors live at once (ct_cg, ct_cu) fused gate+up over one shared X load. coop_tile_setup names exactly one C tile. Store extent is transposed relative to the tile (extents<int,32,16> for a 16×32 tile). 14 buffers — the widest in the set. TPG 32, no simdgroup_index_in_threadgroup.	Ops/OpsNVFP4.swift:579; grid (nOut/32, ⌈mTotal/16⌉, 1)	BUTTER_NVFP4_MPP_FUSED_GATEUP (off)
butter_nvfp4_moe_gateup_bm16_mpp_f32	206	same, plus native <float,float,float> MPP operands with threadgroup float staging — not a dtype swap of the bf16 file’s I/O	Ops/OpsNVFP4.swift:578	same
butter_nvfp4_moe_gather_eg_mpp_bf16	255	matmul2d(32,32,32), BM=BN=64, TPG 128 — the matmul half is pile B. What is new is the precomputed tile plan (tile_expert/tile_row_start/tile_row_count; no m_total param) built by the butter_nvfp4_moe_gather_eg_tileplan_u32 inline-MSL prepass. Iron’s moe/moe_tile_plan_builder*.rs family already covers that shape.	Ops/OpsNVFP4ExpertGather.swift:337, 599, 668; grid (nOut/64, plan.maxTiles, 1)	BUTTER_NVFP4_MPP_GATHER — default ON

Three of the fifteen are behind default-ON flags. Those three carry production risk on any migration and should land last within their phase, each with a Laguna Paris coherence run.

3.3 nax2 in detail — the one genuine outlier

butter_dense_gemm_nax2_mpp_bf16.metal is the only register-direct neural-accelerator-structured kernel in the set. BM=64, BN=128, BK=256, WM=2 × WN=4 = 8 simdgroups / 256 threads, zero threadgroup memory. It needs six things Iron does not have:

1. The 6-arg descriptor overload.

   constexpr auto desc = mpp::tensor_ops::matmul2d_descriptor(
       /*M=*/16, /*N=*/32, /*K=*/16,
       /*transpose_left=*/false, /*transpose_right=*/true,
       /*relaxed_precision=*/false,
       mpp::tensor_ops::matmul2d_descriptor::mode::multiply_accumulate);

   Iron emits only the 4-arg + mode form (emit_block.rs:659-663).
2. Native <bfloat, bfloat, float> cooperative tensors — and specifically on the element-wise-write path, not the ct.load() threadgroup-tensor path. The kernel’s own header records that distinction. Iron’s coop_stage(T) forces bf16 → half unconditionally.
3. Per-lane fragment writes — ct_a[i] = a[i], ct_b[8 + i] = b1[i], fed by a frag_load helper that reads A/B straight from device memory. Iron has the read side (CoopTileGet / CoopTileCoord / CoopTileCapacity) but no write side.
4. Paired-N issue — one 16×16 A fragment against a pair of adjacent B fragments per call, halving issue count vs a square walk.
5. threadgroup_barrier(mem_flags::mem_none) as a pure execution regrouper. Iron has simdgroup_barrier_mem_none() but no threadgroup form.
6. volatile int compiler_barrier; as a register-pressure fence, plus #pragma clang loop unroll(disable) on both K loops — Iron emits unroll(full) keyed off CoopTileCapacity but has no way to disable.

It also does a threadgroup-ID swizzle with early exit ((tgid.y << swizzleLog) + (tgid.x & ((1u << swizzleLog) - 1u))), masks nothing at all, and enforces m%64 == 0 && n%128 == 0 && k%256 == 0 purely Swift-side in LagunaMppGemm.nax2Eligible.

3.4 What none of the 15 uses

reduce_rows, map_iterator, is_compatible_as_left_input, get_left_input_cooperative_tensor(source), tensor_handle kernel parameters, slice(), dynamic_length_v, and quantized MTLTensor scale planes. Every one hand-decodes NVFP4 — E2M1 nibbles through a select ladder to {0,0.5,1,1.5,2,3,4,6}, E4M3 block scales through a 5-op branchless carry-sign-fold — into threadgroup memory or registers. Those TensorOps features matter for fused FlashAttention (gpu-saturation-and-bandwidth-spec.md §S2/§S10) and for a macOS 27 native-quantized path (§4.6); they are not on the critical path for these 15.

---

4. Design

4.1 The seam

Everything below is an extension of Op::CoopTile* — new fields on CoopTileSetup, new sibling ops in the same family, the same needs_mpp feature flag, the same #if __METAL_VERSION__ >= 400 guard, the same CUDA/SPIR-V lowering obligations. iron_std’s de-facto rule that NAX kernels assert !any(Op::InlineMsl) (present in ~20 files) holds: none of this is a raw-MSL escape hatch.

Adding an op means, per ../developing.md: variant in op.rs with flag attributes → fmt_ir arm → parser in body.rs’s match name → emit arm in emit_block.rs (+ cuda/hip/spirv) → KernelFeatures case if it needs a header → any pass that pattern-matches Op::CoopTile* (today: licm.rs, remap.rs) → requires_cooperative_tensors() → MSL snapshot fixture.

4.2 Staging dtype — retire the silent bf16 downgrade

// today — gemm/quantized_nax.rs and ~20 siblings
coop_tile_setup("gemm", 32, 32, 32, coop_stage(T), "accumulate", "simdgroup", f32, ...);
//                                  ^^^^^^^^^^^^^ bf16 → half, unconditionally

// proposed — coop_stage keeps working; a kernel may opt into native bf16
coop_tile_setup("gemm", 32, 32, 32, coop_elem(T), "accumulate", "simdgroup", f32, ...);

coop_elem(D) resolves to D unchanged, resolved per-instantiation in dtype_from_expr_arg exactly as coop_stage is (body.rs:2065-2096).

Do not decide this from comments. Iron asserts Apple’s matmul2d mishandles bfloat cooperative tensors; Butter ships nax2 with native bfloat and claims it removes the T23 fold. Extend probes/mpp_matmul_probe.rs with a <bfloat, bfloat, float> case on both the load() and the element-write path, and have tests/mpp_matmul_probe.rs assert on-device. coop_stage() stays for kernels with a measured reason to downgrade.

4.3 Descriptor extensions

Two new fields on Op::CoopTileSetup:

• relaxed_precision: bool (default false) — selects the 6-arg matmul2d_descriptor overload.
• n_dest: u32 (default 1) — number of destination cooperative tensors.

coop_tile_setup is already an 18-argument positional call with silent .unwrap_or(0) / .unwrap_or(false) defaults, and arg 6 of coop_tile_load_* is overloaded (bool → direct, otherwise → offset). A typo’d bool becomes a zeroed tile dimension with no diagnostic. Adding two more positional args without fixing this is a mistake. Migrate the intrinsic to a struct-literal form in the same change:

coop_tile_setup("gemm", CoopTile {
    m: 16, n: 32, k: 16,
    elem: coop_elem(T), acc: f32,
    acc_mode: "accumulate", scope: "simdgroup",
    tb: true,
    relaxed_precision: false,
    n_dest: 2,
});

with the positional form kept as a deprecated shim for one release so the ~45 existing call sites migrate incrementally.

4.4 Multiple destination tiles

coop_tile_zero("gu", 0u32);          coop_tile_zero("gu", 1u32);
coop_tile_load_a("gu", "Xs", ...);   // one shared A load
coop_tile_load_b("gu", "Wg", ...);   coop_tile_run_into("gu", 0u32);
coop_tile_load_b("gu", "Wu", ...);   coop_tile_run_into("gu", 1u32);
coop_tile_store_c("gu", 0u32, "OutG", ...);
coop_tile_store_c("gu", 1u32, "OutU", ...);

Emits {name}_ct_c0, {name}_ct_c1, … Retires butter_nvfp4_moe_gateup_bm16_mpp_{bf16,f32} and is the cheapest real win in the set: one X load, two weight matmuls. Note the store extent may be transposed relative to the tile — the destination cooperative tensor’s store layout is N-major, and CoopTileStoreC’s ei/eo must be allowed to express that.

n_dest’s generated guard (§4.7 covers the mechanism) validates three things a hand-written kernel currently gets right by inspection: destination indices are in range and non-duplicate, every destination is zeroed or otherwise initialized before its first run_into, and each destination’s store layout — the N-major transpose above is per-destination, not shared — matches what the emitted coop_tile_store_c call expects. Fusing two destinations is the cheap win this section describes, but it also doubles live accumulators and adds a second store’s worth of shared-storage pressure at once; the guard must size the fused case against the §4.7 32 KiB threadgroup budget, not reuse the single-destination sizing it replaces.

4.5 Register-direct fragment fill

The read side exists; add the mirror, and generalise the accessors off C-only:

// Op::CoopTileSet   { name, side: CoopTileSide::{A,B,C}, dest: u32, idx, data }
// Op::CoopTileCapacity / CoopTileCoord / CoopTileGet gain `side` + `dest`
let cap = coop_tile_capacity("gemm", "a");
for i in range(0u32, cap, 1u32) {           // emitter MUST attach unroll(full)
    let r = coop_tile_coord("gemm", "a", i, 0u32);
    let c = coop_tile_coord("gemm", "a", i, 1u32);
    coop_tile_set("gemm", "a", i, load(x[(m_base + r) * k_in + kb + c]));
}

The per-lane layout of a cooperative tensor is implementation-defined; the only correct way to recover which element idx is remains get_multidimensional_index. Butter’s nax2 and all four steel_* kernels hardcode the map instead (fm = (qid & 4) | ((lane >> 1) & 3), fn = ((qid & 2) | (lane & 1)) * 4) — that is a reverse-engineered hardware constant, and porting it verbatim would bake an undocumented layout into the DSL. Use the accessor and let the probe in probes/mma_layout_probe.rs prove equivalence.

emit_block.rs already keys #pragma clang loop unroll(full) off CoopTileCapacity in its Op::Loop arm — documented as load-bearing (an unpragma’d attempt measured +22.8 % isolated occupancy but a wired regression). Extend that keying to the new ops, and add an unroll(disable) counterpart for nax2’s outer K loops.

4.6 Deferred: native quantized tensors

macOS/iOS 26 TensorOps quantization covers int4/int8 only; fp4, fp8, int2 and the E8M0 block-scale auxiliary plane arrive in macOS/iOS 27 (WWDC26 s330; Crosley 2026). NVFP4’s E4M3-per-block × global-f32 scaling is not the MX E8M0 shape, so even on 27 the mapping is not free. When 27 is the floor, the surface is a new param kind:

#[kernel(mode = Reduction)]
pub fn iron_mxfp4_qmm_native<T>(
    x: Tensor<T>,
    #[quantized(fmt = mxfp4, scales = e8m0, block = 32)] w: Tensor<u32>,
    mut out: Tensor<T>,
) { ... }

emitting tensor_blockwise<tensor_plane_scales, device metal_fp8_ue8m0_format, 32, 1> and a tensor_handle-backed metal::tensor parameter, with the host binding an MTLTensor carrying an MTLTensorAuxiliaryPlaneDescriptor. That changes ParamManifest and every Swift wrapper — separate spec.

4.7 Cooperative-tensor dispatch invariants

Reduction kernels already carry ## DISPATCH INVARIANTS (../developing.md:222-235). TensorOps adds constraints the kernel cannot self-check. Add a ## COOPERATIVE TENSOR INVARIANTS sub-block and require it on every kernel this spec touches:

• TPG = 32 × (simdgroups in the execution scope). Mismatch is silent wrong output, not a fault.
• At least one of the descriptor’s M / N / K is a multiple of 32, and at least one of M / N a multiple of 16 — otherwise a static_assert fires. sdpa/steel_attn/steel_attention_nax.rs works around this today with BD=32 and a hand-rolled D-chunk loop; no pass validates it.
• Fragment granularity is 8×8, per-lane capacity 2·(M/8)·(N/8) for a single simdgroup (Rigel §4) — M and N multiples of 8, and never an assumed layout.
• Device-direct tensor_inline byte-stride, not a universal rule. Where a kernel’s coop_tile_load_* reads an operand straight out of device memory — no threadgroup staging in between — the inner stride must be 128-byte aligned (Rigel): K a multiple of 32 for f32, 64 for f16/bf16, 128 for fp8. Violations surface as misleading SFINAE errors.
• The threadgroup ct.load path is exempt from that stride figure. gemm/gemm_q8_mpp.rs:44-90 stages Q8_0-dequantized operands into threadgroup_alloc("Xs"/"Ws", 2048, coop_stage(T)) — 64×32 rows at half precision, a 64-byte row — then runs coop_tile_load_a/b against 32×32 views of that buffer; k_in % 32 == 0 bounds the K-loop step, not a device byte-stride. The emitter controls the staged layout, so what constrains this path is fragment extent (the granularity bullet above — each loaded view a whole multiple of the 32×32/16×16 fragment shape), not the 128-byte device-stride rule. Do not require the direct-input guard here; it would reject the shipped staging every _mpp_ gather/gemm kernel uses today. Emit the direct-input stride guard only on the operand path it actually governs, and gate its codegen on the dtype (f32/f16/bf16/fp8 pick different multiples) and the toolchain probe from §4.8 — not unconditionally.
• Accumulator is at least fp32 regardless of operand dtype (Rigel).
• Threadgroup operands need the bank skew the library already uses — BK + 4 (row stride 36), head_dim + 8, or the multiple-of-8 leading dim (stride 40) that steel_gemm_splitk_nax_bt_direct_32x64_pipelined_ld40 uses for aligned shared views.
• Threadgroup budget is 32 KiB including MPP’s own internal allocations. Butter’s nax64 kernels deliberately alias Xs/Ws onto OutScratch’s backing store because separate declarations hit the ceiling exactly, and the symptom was per-layer NaN starting at layer 38 of 39 — invisible to every unit test. A generated kernel that allocates naively is silently wrong, not merely slower. moe_gather_qmm_expert_mpp.rs shows the alternative: drop OutScratch entirely via the device-direct masked readout, 20 KiB → 4 KiB, 1 → 3 threadgroups per core.

Generate the guard, don’t just document it. The ## COOPERATIVE TENSOR INVARIANTS block above is prose for a reviewer; it is not what stops a bad dispatch. Every constraint in this section — threadgroup geometry, the aligned-axis contract (the “Asymmetric bounds contracts” risk in §6), the relevant quantization block-size constraints, and this section’s own scratch-budget requirement — must be generated, not hand-written, from the same resolved variant that produces the shader body and the buffer signature: one resolution, three consumers. This lands with the kernel, not after it — §5’s phase-completion criteria are amended so a ported kernel carrying this block ships its generated guard in the same commit as the shader and the signature, not as a follow-up.

Reject an unsupported configuration before it is encoded into the dispatch — never by launching it to find out. ../developing.md’s kernel-authoring hazards are exactly this failure mode: a bad dispatch geometry hard-freezes the machine, it does not raise. The generated guard is a precondition evaluated ahead of encoding, not a comment a reviewer can miss.

Preserve old symbols, buffer/constexpr ordering, and default emissions — including the §4.3 positional shim — while adding generated guards. New Metal-only functionality must resolve to either a correct CUDA/SPIR-V lowering or an explicit, generated “unsupported on this backend” result; existing CUDA/SPIR-V paths must not acquire changed semantics by inheriting a new default that was meant for the Metal guard.

This is a distinct resolution from §4.9’s. Variant resolution here picks the (M, N, K, dtype, …) row and runs at codegen time, feeding shader, signature, and guard together; §4.9’s compilation-policy resolution picks live-compile vs. AOT and runs at PSO-prepare time. Keep the two separate — only the codegen-time one is “the same resolved variant” this section means.

4.8 Availability gating

axis	today	needed
MSL version	#if defined(__METAL_VERSION__) && __METAL_VERSION__ >= 400 around header (msl/mod.rs:326-336) and body (:557-566), #else empty stub so the metallib still links	keep
Feature tier inside Metal 4	none	macOS 26 (int4/int8) vs 27 (fp4/fp8/int2/E8M0) differ. Needs a KernelFeatures::mpp_tier and a guard emitted only where §4.6 surface is used. The macro is unconfirmed — see §6.
Runtime hardware	Context::chip_family() >= 10 in Rust; Ops.mppAutoCapable in Swift; Kernel::requires_cooperative_tensors() drives iron test skips	surface in the manifest (§4.9) so Butter stops re-deriving it per call site

is_unsupported_cooperative_tensor_toolchain (wh-iron/src/runner/harness.rs:824-834) already classifies deferred-static-alloca and get_{left,right}_input_cooperative_tensor PSO-build failures as toolchain gaps. New surface must extend that matcher. Also note skip_unless_apple10 is duplicated verbatim across many tests/*.rs rather than living in tests/common/mod.rs — fold it there while touching these files.

4.9 The header-skew problem, and the recommendation

Why Butter live-compiles. Sources/IronSwift/PSOCache.swift:109-118:

The MetalPerformancePrimitives header inlines cooperative-tensor type IDs + descriptor layouts that drift between SDK versions — when SDK macOS != runtime macOS, an offline-compiled metallib’s baked-in MPP types disagree with the device runtime’s, producing bit-deterministic wrong output (e.g. cos 0.816 vs 0.999 oracle).

Measured twice. PSOCache.swift:160-175 cites Iron’s own moe_bm64_ragged_correctness: the metallib bm64 gives sumAbs 20718 vs the runtime-compiled 24583 on identical inputs. llama.cpp hit the same class independently — macOS 26.0’s shipped MPP framework lacked the bfloat tensor support the docs claimed, and an M2 Ultra compiled bfloat kernels that then segfaulted with a zero-thread-capacity pipeline descriptor (#16634). The failure mode is deterministic wrong numbers, not a crash.

Butter’s mitigation is isMppKernel(name) → liveCompileMppFunction(name): locate <name>.metal beside kernels.metallib, makeLibrary(source:), with an AOT fallback when the live compile itself fails (the named real case: "unsupported deferred-static-alloca-size" on the bm8 matmul2d on macOS 26). The detector is:

return name.contains("_mpp_") || name.contains("bgemm")

with a BUTTER_PSO_LIVE_COMPILE_MPP=0 escape hatch “in case a future kernel name accidentally matches”. It already misfires on the four steel_* kernels. And Iron encodes the dependency from its side too — gemm/gemm_q8_mpp.rs:16-18 says “Name contains _mpp_ so Iron’s PSOCache LIVE-COMPILES it”. A substring is load-bearing across a repo boundary.

Options. (1) Keep the name match — zero work, already broken. (2) Always ship MSL beside the metallib — already true (--emit all writes both, and KernelManifest::source points at the .metal); a precondition, not a fix. (3) Stamp the build SDK in the manifest and live-compile everything on a mismatch — correct but recompiles ~650 kernels to protect 15. (4) Per-kernel capability metadata. ✅ Recommended.

Iron already computes the exact predicate. Kernel::requires_cooperative_tensors() (ir/kernel.rs:255-273) returns true for any Op::CoopTile* and for Op::InlineMsl whose source contains "mpp::", walking nested blocks correctly, with a unit test pinning all four cases. It drives iron test skips and never reaches the manifest. Plumb it:

crates/wh-iron-codegen/src/emit.rs

#[derive(Serialize)]
pub struct KernelManifest {
    pub name: String,
    pub source: String,
    pub kernel_mode: String,
    pub params: Vec<ParamManifest>,
    pub constexprs: Vec<ConstExprManifest>,
    /// Capabilities the emitted MSL depends on: "mpp", "simdgroup_matrix", "bf16".
    #[serde(default, skip_serializing_if = "Vec::is_empty")]
    pub requires: Vec<String>,
    /// True when this kernel's MSL resolves framework headers whose ABI is not
    /// stable across SDK/runtime skew (today: MetalPerformancePrimitives
    /// cooperative tensors). The host MUST compile `source` at run time rather
    /// than load the function from the metallib.
    #[serde(default, skip_serializing_if = "std::ops::Not::not")]
    pub live_compile: bool,
    /// Minimum Metal language version (e.g. 400) and GPU family (e.g. 10).
    #[serde(default, skip_serializing_if = "Option::is_none")]
    pub min_metal_version: Option<u32>,
    #[serde(default, skip_serializing_if = "Option::is_none")]
    pub min_gpu_family: Option<u32>,
}

live_compile defaults to kernel.requires_cooperative_tensors(). Provide a #[kernel(live_compile = false)] opt-out for a kernel proven skew-safe — an opt-in would recreate today’s failure mode, where safe behaviour depends on someone remembering. (Precedent exists for setting such flags post-hoc on the IR instead — bfloat_reinterpret_cast, cuda_max_registers, wants_indirect_variant all work that way — but an attribute is the better choice here because the value is a property of the kernel’s source, not of a tuning decision.)

Butter then replaces isMppKernel’s substring match with a manifest lookup, and the four steel_* kernels stop being live-compiled for no reason. Additionally stamp the manifest with the build-time SDK version as a diagnostic, so “SDK 26.1, runtime 26.0” is legible in a bug report rather than inferred from a cosine.

Resolution, caching, and invalidation. The schema above must not collapse three separate claims into one bool. requires_cooperative_tensors() is a source requirement — cooperative tensors appear in the body, full stop; it cannot by itself say AOT or live compilation is numerically safe on a given SDK/runtime pair. live_compile is compilation policy: true is the requirement-driven MUST, live-compile until proven otherwise. false is not the requirement lapsing — it is an evidenced exception, settable only where a qualification (a probe run for the specific SDK/runtime pair, per §4.2’s mpp_matmul_probe pattern) has passed; no probe, no false. Symmetrically, when live compilation itself fails (the named "unsupported deferred-static-alloca-size" case), falling back to the metallib’s AOT function is a known-working fallback preserved through migration, not a claim that the AOT artifact is qualified — an unqualified AOT PSO does not become safe merely because the live compile that would have replaced it failed to build.

Resolve the policy once, when a PSO is created or prepared, never per dispatch, and cache the resolved decision. Key the cache on: the shader source after specialization (a variants(...) row is a distinct source), the effective manifest entry (live_compile, min_metal_version, min_gpu_family), a fingerprint of the compilation-policy / qualification table in effect, the compiler options passed to makeLibrary, and the device/SDK/runtime identity triple (§4.8). Invalidate the cached decision whenever any key component changes — including a policy or qualification update landing against byte-identical shader source: a kernel unqualified on SDK 26.0 that becomes qualified on 26.1 must not keep serving the 26.0 decision because the source hash alone didn’t move.

Keep manifest parsing, source hashing, and path resolution out of the warm dispatch path; they belong at PSO creation/preparation, where the cache above already amortizes them. The 2026-09-07 kernel-swap cuda-policy-host-optimization campaign (evidence retained on the perf/atlas-kernel-gauntlet branch, not on dev) is host evidence for this exact mistake on the CUDA side — eligible-policy cost fell ~53% once the repeated host-side path work was removed from the per-call check. That is host-bookkeeping evidence, explicitly not a model-speed claim.

Runtime overrides (BUTTER_PSO_LIVE_COMPILE_MPP=0 and any future kin) are a separate, genuinely dynamic mechanism layered after resolution — they do not enter the cache key. Document their own invalidation: process-lifetime scope, effective on the next PSO preparation rather than retroactively on an already-cached decision.

---

5. Implementation phases

Each phase is independently shippable and ends with make fmt-check && make clippy && make test green, the empty-body MSL detector clean over make emit-all OUT=/tmp/iron-smoke, and — where kernels move — make bench-vv rows in the PR body. Where a ported kernel carries a ## COOPERATIVE TENSOR INVARIANTS block (§4.7), its generated per-variant dispatch guard lands in the same commit as the shader and buffer signature — the guard is not a follow-up phase.

Phase 0 — settle native bfloat (S, ~1 day)

Extend probes/mpp_matmul_probe.rs with a <bfloat, bfloat, float> case on both the ct.load() and the element-write path; assert in tests/mpp_matmul_probe.rs; run on M4 and M5. Decides whether the T23 fold is a permanent tax or a workaround with an expiry date. Gates phases 2, 4, 6, 8. No kernels move.

Phase 1 — manifest capability metadata (S, ~1–2 days)

Status (2026-09-08): iron half landed — manifest v2 (requires, live_compile, min_metal_version, min_gpu_family, toolchain stamp) + IronKernels capability sets + #[kernel(live_compile = false)] opt-out + registry gate. Butter swap pending.

§4.9. Add a kernel_registry_consistency assertion that every requires_cooperative_tensors() kernel manifests live_compile: true. Butter repins and swaps isMppKernel for the lookup. Retires the bgemm misfire immediately. No kernels move. Acceptance criteria also cover §4.9’s resolution/caching contract: the policy resolves and caches at PSO creation/preparation, not per dispatch, and invalidates on any key change.

Manifest lookup miss. build_manifest (crates/wh-iron-codegen/src/emit.rs:146-160) enumerates only Iron Kernels; none of the fifteen hand-written MPP sources carry a manifest entry until their own phase lands (§5.6), so Butter’s new lookup can miss on any of them through phase 8. A miss must fall back to today’s routing — the isMppKernel substring match — not a rejection and not AOT selection; the latter reintroduces the §4.9 skew failure this phase exists to close. This fallback is explicitly transitional, and distinct from an opt-out: an absent name predates the manifest schema and keeps substring-routed live-compile, while live_compile: false on a present entry is the only way to turn cooperative-tensor handling off for a kernel. Acceptance criteria for this phase add a fixture proving the fallback: an unmigrated, default-ON MPP kernel (iron_gemm_q8_mpp, absent from the manifest at this phase) dispatched through the new lookup path, asserting it still live-compiles instead of falling through to AOT.

Phase 2 — pile A re-land (S, ~2 days; highest leverage)

Port the T23 stage-bias fold (or whatever survives phase 0) plus the x_rows A-row indirection into moe/moe_mpp_bm16_bk32_block_scaled.rs. Delete butter_nvfp4_moe_gather_bm16_mpp_xrows_{bf16,f32}.metal. 2 of 15 gone, no new primitive.

Phase 3 — the steel_* family (M, ~1 week)

Port the four simdgroup_matrix gather BGEMMs onto simdgroup_alloc / simdgroup_elem_{load,store} / simdgroup_matmul. One variant family collapses all four:

#[kernel(mode = Reduction, variants(
    (WN, SUBM, SUBN, BM, BN, RAGGED) = [
        (2, 32, 32,  64,  64, 1u32),   // steel_bgemm
        (2, 32, 32,  64,  64, 0u32),   // steel_bgemm_dense
        (4, 16, 16,  16,  64, 1u32),   // steel_bm16_bgemm
        (4, 16, 32,  16, 128, 1u32),   // steel_bm16_bn128_bgemm
    ], suffix = "bm{BM}_bn{BN}"))]

Carry the T34 loop order. These kernels nest tj outer / ti inner with a simdgroup_matrix<float,8,8> accs[4] array so each B element is dequantized once per (tj, K-step) rather than once per (ti, tj, K-step) — a measured 3.4–4.0× against the natural order. A generator that emits ti-outer silently loses it. 4 of 15 gone; no MPP involved, so this phase can land ahead of the rest.

Phase 4 — pile B nax / nax64 (M, ~1 week)

Reconcile butter_nvfp4_gather_qmm_rhs_nax{,64}_mpp_{bf16,f32} against moe/moe_mpp_bm64.rs and moe/moe_mpp_bm64_bk64_nvfp4.rs. Iron’s BK64 file documents a still-open quality death (“dual half-slab BK=64 staging empty-generated on full Laguna Paris”); Butter’s BK64 kernel is a shape that works. Fold the working staging — including the Xs/Ws-onto-OutScratch aliasing (§4.7) — back into Iron and expose {bf16, f32} on the existing dtype axis. 10 of 15.

Phase 5 — dual destination tiles (M, ~4 days codegen + 2 days port)

§4.3, §4.4. Retires butter_nvfp4_moe_gateup_bm16_mpp_{bf16,f32}. 12 of 15.

Phase 6 — butter_dense_gemm_mpp_bf16 (S, ~2 days, after phase 0)

A variant row on the iron_gemm_q8_mpp family with a dense bf16 weight. Default-ON flag → ship with a coherence run. 13 of 15.

Phase 7 — EG tile-plan gather (M, ~1 week)

Port butter_nvfp4_moe_gather_eg_tileplan_u32 onto moe/moe_tile_plan_builder*.rs, then the matmul half. Worth noting the payoff beyond one kernel: with a host-built plan, the while (sub < 64u) ragged expert-run scan — the single most awkward construct in piles A and B, present in eight of the fifteen — disappears from the DSL entirely. Default-ON flag. 14 of 15.

Phase 8 — nax2 register-direct (L, ~2–3 weeks, after phase 0)

§3.3 and §4.2/§4.3/§4.5. The largest lift and the one that most changes emit_block.rs. Default-ON flag; land last. 15 of 15 — PSOCache.isMppKernel and liveCompileMppFunction delete, replaced by the phase-1 manifest lookup, and Resources/kernels/butter_*.metal empties.

5.6 Consolidation banked

template	covers	axis
A — staged 64×64 matmul2d gather	6 files (dense_gemm_mpp, nax ×2, nax64 ×2, eg)	dtype, BK ∈ {32,64}, row-source ∈ {run-scan, tile-plan}, fold on/off
B — BM16 single-simdgroup matmul2d	4 files (gateup ×2, xrows ×2)	dtype, n_dest ∈ {1,2}
C — simdgroup_matrix steel	4 files	(WN, SUBM, SUBN, BM, BN, RAGGED)
nax2	1 file	stands alone

5.7 Testing

1. Bit-exact vs the archived hand-written kernel. A fixture compiling both the archived butter_*.metal and the emitted kernel, dispatching on identical buffers, asserting byte equality. Stronger than a cosine oracle and the only thing that catches a silently-dropped hand patch — pile A’s exact failure mode.
2. GPU correctness vs a naive CPU oracle in the same commit (crates/wh-iron-std/tests/<kernel>_gpu_correctness.rs), behind the shared skip_unless_apple10 gate. Non-negotiable: empty-body MSL passes xcrun metal, the smoke build, and MSL snapshots, and still ships all-zeros.
3. MSL snapshots following tests/steel_msl_snapshots.rs — the model for pinning the emit path of kernels whose GPU tests cannot run on CI hardware. Refresh with cargo insta test --accept -p wh-iron-std --test <name>.
4. make bench-vv parity before any hand-written kernel is deleted. A generated kernel that is correct but slower is not a migration. Follow the docs/benchmarks/nax-*/README.md evidence format: predeclared gate, paired rounds, forward+reverse arm order, F64 oracle, mutation negative, frozen hashes, losses reported alongside wins.
5. Butter-side Laguna Paris coherence after each phase, and mandatorily for the three default-ON flags. Pile A and Iron’s own BK64 history are both cases where every unit test passed and the multi-layer residual still died.

Every new test file needs a row in docs/test-inventory-status.tsv plus a regenerated docs/test-inventory.md (make check-test-inventory).

---

6. Risks and open questions

Risks

risk	why it bites a code generator
Silent-correctness hazards dominate	The nax64 threadgroup aliasing and the T23 fold are load-bearing, invisible to unit tests, and not expressible as “emit a matmul”. §4.7 exists to make them reviewable.
Loop-order dependence	Phase 3’s tj-outer/ti-inner nesting is worth 3.4–4.0×. Emit the natural order and you silently lose it.
Undocumented fragment layouts	nax2 and the steel_* kernels hardcode the lane→(row,col) map. §4.5 routes around it via get_multidimensional_index; if that path is slower, it becomes a real decision rather than an oversight.
deferred-static-alloca-size	A generated kernel that changes tile sizes can trip the macOS 26 toolchain limit that already forces Butter’s AOT fallback on the bm8 matmul2d.
Asymmetric bounds contracts	Nearly every kernel masks M but not N or K, enforced only by a Swift-side eligible predicate; nax2 masks nothing. The DSL needs a contract-aligned-axis declaration that generates the Swift guard in lockstep, or the two drift.
Per-pack scale indexing	The scale index resolves per pack ((k0 + pack_in_row * 8) / block_size), not per tile. A per-tile shortcut compiles and is wrong only at block_size seams.
Three default-ON flags	BUTTER_LAGUNA_MPP_GEMM, BUTTER_LAGUNA_NAX2_GEMM, BUTTER_NVFP4_MPP_GATHER — phases 6, 7 and 8 touch production paths directly.

Open questions — for Tom

1. Native bfloat. gemm/quantized_nax.rs:18-20 and ~20 siblings assert Apple’s matmul2d mishandles bfloat cooperative tensors; nax2 uses <bfloat, bfloat, float> and claims it removes the T23 fold. Which OS / silicon / SDK produced each observation? Is the Iron comment stale, is nax2 M5-only, or is the distinction exactly load-vs-element-write? Answered (M5 Max, macOS 27.0) — probes/mpp_matmul_probe.rs’s four (F16, BF16) × (TensorLoad, ElementWrite) variants all pass their GPU correctness test, including a half-range regime (±2^17 mixed with ±2^-20) that a silent bf16→half downgrade cannot survive without producing non-finite output — both bf16 variants stayed finite and within 2e-3 relative tolerance of the f64 CPU oracle. On this machine, native <bfloat, bfloat, float> matmul2d works correctly on both the tensor-load and the element-write path — no T23 fold needed here. M4 (Apple9) is still pending; whether the Iron comment is M4-specific or simply stale remains open until that run lands.
2. relaxed_precision. nax2 passes false. Measured, or safe default? Does it change NAX routing?
3. _nax vs _mpp. gemm/block_scaled_qmm_nax.rs:8-13 says the twins are byte-identical and co-exist purely for consumer dispatch tables; KERNEL_AUDIT.md:51 gives an intended MXU-fallback-vs-M4-only distinction the code does not encode. Can the duplication retire here?
4. BK=64. moe/moe_mpp_bm64_bk64_nvfp4.rs documents an open quality death on the dual half-slab; Butter’s nax64 ships the shape. Is Butter’s staging the fix, or a different algorithm?
5. Paired-N. Is nax2’s (16, 32, 16) paired-fragment issue a general win, or specific to the K=2048 projections?
6. steel_gemm_splitk_nax_bt_direct_32x64_pipelined_ld40’s stride-40 shared rows vs the BK + 4 = 36 skew everywhere else — is one of them the convention this spec should standardise on? Answered — no. Do not standardise LD40 across backends: the measured win is inline-PTX-specific, Metal screening lost to the held stride-36 pipeline, and WMMA did not share it. Pitch, staging, and split stay per-variant with backend-specific qualification.

Open questions — for Eric

7. Ordering vs ek/model-api-consistency. gpu-saturation-and-bandwidth-spec.md says every item there is a post-clean-up task and that the branch moves several cited call sites. Does phase 1 (manifest-only, no kernel motion) count as in-scope now?
8. macOS floor. Does §4.6 justify a macOS 27 floor for a future kernel tier, or must every kernel keep a 26 path?
9. Archiving. Keep the hand-written .metal under crates/wh-iron-std/tests/fixtures/ as bit-exactness oracles after deletion, or delete outright and trust the CPU oracle? Answered — keep them. The archived implementation stays as a compatibility oracle, alongside the independent high-precision oracle from §5.7 — not a replacement for it.
10. live_compile scope. Default-on for every cooperative-tensor kernel is ~45 Iron kernels each paying a runtime makeLibrary(source:). Acceptable, or narrow it once a skew-detection probe exists?
11. Split the steel_* phase out? Phase 3 needs no TensorOps work at all and could land as its own PR ahead of this spec.

Needs external confirmation before implementation

12. The macOS 27 feature-tier preprocessor macro (§4.8). This spec uses __METAL_VERSION__ >= 410 as a placeholder; Apple has not documented a feature macro for the fp4/fp8/int2/E8M0 tier. Confirm against the shipping SDK headers.
13. execution_simdgroups<N> (plural). WWDC26 s330 distinguishes execution_simdgroup (singular — each simdgroup owns complete rows, which is what makes single-pass softmax work) from execution_simdgroups<N>. Iron’s CoopTileScope::Threadgroup maps to metal::execution_threadgroup, which may not be the same thing. Confirm before any FlashAttention work depends on it.

---

7. Why this earns its place

• Deletes 4 050 lines of Metal that sits outside Iron’s manifest, PSO smoke test, empty-body detector, GPU-correctness gate, and MSL snapshots.
• Removes a cross-repo substring dependency (_mpp_ / bgemm) that is load-bearing for numerical correctness and already misfires on four kernels.
• Closes a silent-data-loss path: pile A’s hand patch is one make regenerate-kernels from disappearing, and its absence is non-finite output at real activation magnitudes.
• Resolves a documented contradiction about native bfloat that splits Iron’s DSL from Butter’s newest kernel, with a measurable answer either way.
• Follows ../developing.md verbatim: “Improve the compiler, don’t hand-write MSL. If the DSL can’t express a pattern, extend the codegen. Don’t bypass it.”
• It is step (d) of five in Butter’s kernel-provenance plan and the prerequisite for the ~90-kernel inline-MSL pass that is step (e).

---

8. References

source	what it establishes here
WWDC26 s330, “Optimize custom ML operations with Metal tensors” (link)	tensor_handle / tensor_inline, matmul2d_descriptor, cooperative tensors, reduce_rows, map_iterator, is_compatible_as_left_input, quantized tensor types + E8M0 auxiliary plane, slice() on multi-plane tensors
Crosley, “Metal for ML in 2026” (link)	OS feature split (26: int4/int8; 27: fp4/fp8/int2/E8M0); M5 in-core NAX; get_left_input_cooperative_tensor replacing the macOS 26 threadgroup round-trip
Rigel, arXiv:2606.12765	8×8 fragment tiling, per-lane capacity 2·(M/8)·(N/8), the 128-byte tensor_inline inner-stride rule, ≥fp32 accumulation — and that fp8 matmul2d on M4 Max runs at ~0.94× fp16 (emulated, not accelerated)
BaseRT, arXiv:2607.19438	matmul2d at simdgroup/threadgroup scope drives the M5 NAX; operands read straight from device memory; prefill ≤6.4× llama.cpp / 3.9× MLX, decode 1.02–1.75×; -std=metal4.0 required
ggml-org/llama.cpp#16634	Independent confirmation of SDK/runtime MPP header skew (§4.9)
liuliu/example_matmul_metal4	“at least one of M or N a multiple of 16, else static_assert”

Iron: ../developing.md (hazards, ## DISPATCH INVARIANTS, the empty-body MSL detector) · ../testing.md (test layers) · ../STYLE_GUIDE.md §10 (which §2 proposes to extend) · KERNEL_AUDIT.md · KERNEL_CONSOLIDATION_PLAN.md.

Butter: planning/known-issues.md “Kernel provenance” (2026-09-07) item 3 · planning/gpu-saturation-and-bandwidth-spec.md §S10 · Sources/IronSwift/PSOCache.swift:96-236.