Apple GPU Tier-3: conv2d via MPSGraph convolution2D by gstoner · Pull Request #21 · gstoner/tessera

gstoner · 2026-05-30T12:41:14Z

Completes the last open item in the Apple GPU Tier-2/Tier-3 plan: a 2-D
convolution runtime lane built on MPSGraph's
convolution2DWithSourceTensor:weightsTensor:descriptor:.

What

tessera_apple_gpu_conv2d_{f32,f16} (NHWC source [N,H,W,Cin], HWIO
weights [kH,kW,Cin/groups,Cout], optional bias [Cout]): full
stride / padding / dilation / groups, fp32 internal accumulation, cached
MPSGraph + RAII buffer-pool acquires for every allocation.
bf16 runs via a host fp32 round-trip (MPSGraph has no native bf16 conv).
Non-Apple reference fallback (reference_conv2d_f32 + stub parity).
runtime.py dispatcher + ctypes wrappers; tessera.conv2d wired into the
metadata op-dispatch loop, the _APPLE_GPU_CONV_OPS envelope, and driver.py
gating so a single conv2d plan reports execution_mode="metal_runtime".

Layout

Matches the existing eager tessera.ops.conv2d (NHWC / HWIO) exactly, so the
GPU path is a drop-in for the numpy reference — asserted by
test_conv2d_matches_reference_ops_conv2d.

Tests

tests/unit/test_apple_gpu_conv2d.py — 13 tests, validated on Apple
Silicon: f32 across 7 shape/stride/pad/dilation/groups/depthwise/rect cases vs
numpy (rtol 1e-4), native f16, bf16 round-trip, symbol export, envelope
membership, and agreement with the eager reference.

Verification (local, Apple Silicon)

conv2d suite: 13/13 pass
apple_gpu + driver + runtime sweep: 542 passed, 3 skipped
buffer-pool RAII gate + runtime ABI audit: pass (dashboard regenerated, 109 to 113 symbols)
mypy ratchet: conv2d code clean (only the pre-existing environmental torch-import error remains)

CI on this repo is uniformly red on main (Python 3.8 to 3.11 matrix, missing
optional deps) — same state PRs #17 to #20 merged through. The local signal above
is green.

🤖 Generated with Claude Code

Adds a 2-D convolution runtime lane (NHWC source, HWIO weights) built on MPSGraph's convolution2DWithSourceTensor:weightsTensor:descriptor:, completing the last open Tier-3 item in the plan. - apple_gpu_runtime.mm: mpsg_run_conv2d (cached graph, RAII buffer-pool acquires for source/weights/optional-bias, fp32 internal accumulation) + tessera_apple_gpu_conv2d_{f32,f16} + conv2d_out_h/out_w helpers + reference_conv2d_f32 fallback. Full stride/pad/dilation/groups, optional bias. - apple_gpu_runtime_stub.cpp: non-Apple parity (f32 reference conv; f16 zero stub so python upcasts on fallback). - runtime.py: _apple_gpu_conv2d_{f32,f16} ctypes wrappers + a _apple_gpu_dispatch_conv2d dispatcher (f32/f16 native, bf16 host round-trip, None-fallback for other dtypes); tessera.conv2d added to _APPLE_GPU_CONV_OPS + the runtime envelope + the metadata op-dispatch loop. - driver.py: tessera.conv2d gated into _APPLE_GPU_CONV_OPS / runtime envelope so a single conv2d plan reports execution_mode="metal_runtime". - tests/unit/test_apple_gpu_conv2d.py: 13 tests — f32 (7 shape/stride/pad/ dilation/groups/depthwise/rect cases) vs numpy, native f16, bf16 round-trip, symbol export, envelope membership, agreement with eager ops.conv2d. - docs: plan conv2d row marked done; runtime_abi dashboard regenerated (109 -> 113 symbols). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: b318ae1442

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

chatgpt-codex-connector · 2026-05-30T12:44:11Z

 })

+# 2026-05-30 — Tier-3 conv2d via the MPSGraph convolution2D node (NHWC/HWIO).
+_APPLE_GPU_CONV_OPS: frozenset[str] = frozenset({"tessera.conv2d"})


Add canonical conv2d_nhwc to the Apple GPU envelope

The compiler catalog canonicalizes user-facing ops.conv2d to tessera.conv2d_nhwc (op_catalog.py maps conv2d that way and aliases legacy tessera.conv2d to it), and existing IR/tests use tessera.conv2d_nhwc. With the envelope containing only tessera.conv2d, a normal Apple-GPU conv2d plan is still rejected by _is_apple_gpu_mps_executable and never reaches the new dispatcher; only uncanonicalized legacy names do. Please include the canonical op name here and in the runtime dispatch set.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-05-30T12:44:11Z

+  if (outH > 0 && outW > 0)
+    std::memset(O, 0, static_cast<std::size_t>(N) * outH * outW * Cout * 2);


Return a real f16 fallback instead of zeroing outputs

On non-Darwin _load_apple_gpu_runtime() compiles this stub, and _apple_gpu_dispatch_conv2d treats the exported f16 symbol as a valid implementation, so every f16 conv2d run through the stub returns all zeros instead of falling back to a reference result. This also happens on Apple if the Metal path is unavailable; I confirmed the new tests/unit/test_apple_gpu_conv2d.py::test_conv2d_f16_native fails on Linux with an all-zero actual output.

Useful? React with 👍 / 👎.

@jit

A plain Python `for` loop in @jit(target="apple_gpu") now lowers to the GraphFn tessera.control_for path and executes on Apple GPU, with no explicit jit_fori_loop call: @jit(target="apple_gpu") def f(x, w): for _ in range(N): x = ts.ops.silu(ts.ops.matmul(x, w)) return x - New python/tessera/compiler/graphfn_bridge.py — IR-to-IR translation: detect_loop_fn reads the @jit graph_ir op-list (loop body inline between tessera.scf.for.{begin,end} markers), recovers the single tensor carry structurally (the one arg both read in and re-bound by the body), and build_graphfn replays the body ops through the GraphFn builder into for_loop -> run_via_target_ir. Reuses the entire G-A/G-B/G-C machinery; no new C ABI / ODS. - jit.py: JitFn.__init__ detects at decoration (best-effort, cached per arg-shape/dtype); __call__ dispatches the matched loop before the existing apple_gpu branch. - Dispatch policy (Decision #21): auto-route any matching single-carry bounded loop; if the shape matches but a body op has no GraphFn builder (sqrt/conv/ einsum/...) raise a stable diagnostic naming the op + target -- never a silent host-Python fallback. Non-matching functions (multi-carry, no-carry, dynamic trip) keep the existing @jit path. - v1: loop is the whole function over its args, single tensor carry, static trip, f32. bf16 / control_if / control_while / scan are the remaining close-out phases. +9 tests (tests/unit/test_jit_apple_gpu_loop_bridge.py). Broad jit/compiler sweep (843) + production lane (315) green; mypy clean host + linux; test_coverage census regenerated for the new test file. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…safety fix Generative round-trip test (gen op-chain -> render DSL -> parse via lower_text_to_graph_ir -> assert recovered op names) plus a malformed-input fuzz that asserts the parser only ever raises a NAMED diagnostic (FrontendSyntaxError/FrontendSemanticError), never an uncaught crash. The fuzz immediately found a real crash-safety bug: parse_module()'s loop guard 'while _peek_text() != "}"' is also true at EOF, so a module missing its closing '}' reached an 'assert token is not None' (a bare AssertionError crash) instead of a named parser error. Fixed to emit FrontendSyntaxError 'unexpected end of input: expected }' (Decision #21 / crash-safety contract). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

The Multivector front-end allow-list is {Cl(3,0), Cl(1,3)} but v1 ships Apple-GPU kernels only for Cl(3,0) (cl30). A @clifford_jit callable invoked with a spacetime Cl(1,3) Multivector previously routed silently to the numpy reference inside tessera.ga.* — the silent-fallback anti-pattern (Decision #21). Add a call-time signature gate (mirroring the existing decoration-time dtype!='f32' gate): _require_cl30_args refuses any non-Cl(3,0) Multivector arg with CLIFFORD_UNSUPPORTED_SIGNATURE, naming the offending signature and pointing to the plain tessera.ga.* numpy lane for non-Cl(3,0) algebras. Wired into both the IR-compiled and lazy (first-call-compile) call boundaries — the lazy gate fires before the numpy trace. New diagnostic code ConstrainedDiagnosticCode.CLIFFORD_UNSUPPORTED_SIGNATURE. The plain tessera.ga.* lane is unaffected (numpy Cl(1,3) still works); only the GPU-plan decorator gates. Tests: tests/unit/test_clifford_jit_signature_gate.py (Cl(3,0) runs; Cl(1,3) single-op + multi-op plans gated; plain ga lane unaffected). 165 clifford_jit/ dialect regression tests green; mypy host+linux + ruff clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Makes the grouped_layout / scale_layout contract load-bearing instead of audit-decoration, then adds a genuinely new capability. Rung A — runtime gate (real correctness value): * grouped_layout.APPLE_SUPPORTED_GROUPED_KINDS = {dense, contiguous}; grouped_kind_unsupported_message() (Decision #21 — names op + target + kind); validate_grouped_alignment() (each group's M a multiple of the declared alignment). * Both the eager grouped_gemm(kind=, alignment=) and the Apple runtime dispatch now READ the contract: dispatch by kind, reject masked / k_grouped with a clear diagnostic (instead of silently computing contiguous), and enforce a declared alignment. Default (contiguous, no alignment) is the existing ragged path — unchanged. Rung B — quantized grouped GEMM (new capability, not plumbing): * grouped_layout.apply_quant_for_grouped() — dependency-injected quantize-then-dequantize of x + per-expert w per the canonical scale layout (fp8_e4m3/e5m2 per-tensor, nvfp4 1x16 block). * grouped_gemm(..., quant="fp8_e4m3"|"nvfp4"|...) and the runtime dispatch (quant kwarg) now run a correct dequant-on-host quantized grouped GEMM, with the f32 fused MSL kernel doing the matmul on Apple. fp8 rel err ~0.037, nvfp4 ~0.131 vs the f32 grouped GEMM. Tests: tests/unit/test_grouped_gemm_contract.py grows to 35 — alignment validation, eager+runtime kind rejection (Decision-#21 message asserts target + kind), and the quantized-path oracle (eager == runtime, within precision budget, unsupported dtype rejected). spec_sync / mypy / ruff / generated-docs clean. Next: Step 3 (thread the contract through Graph→Schedule→Tile IR). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

@wg

The fused-kernel perf rung for the local MoE block: tessera_apple_gpu_moe_swiglu_f32 collapses the 3 grouped-GEMM + silu_mul dispatches of the composed path into a single MSL kernel — the grouped analog of the dense swiglu_f32 kernel. One thread per token t with expert e=Eids[t]: gate = x[t]@wg[e] ; up = x[t]@wu[e] ; hidden = silu(gate)*up ; O[t] = hidden@Wd[e] folding routing in (per-token expert id) and removing per-expert dispatch overhead. Wg/Wu (E,K,H), Wd (E,H,Kout); per-row stack buffers cap H,Kout ≤ 256 (the C symbol early-returns past that → CPU reference / composed fallback). - C++: dispatch_moe_swiglu_msl + reference_moe_swiglu_f32 + extern C symbol in apple_gpu_runtime.mm; non-Darwin reference parity in apple_gpu_runtime_stub.cpp. - Bridge: _apple_gpu_backend.gpu_moe_swiglu_block + ctypes signature; _SENTINEL_SYMBOL bumped to moe_swiglu_f32 so a stale prebuilt dylib recompiles. - Runtime: _apple_gpu_dispatch_moe_swiglu_block takes the fused fast path for f32 / no-quant / H,Kout ≤ 256, else the composed lanes (quant keeps exact per-GEMM scale semantics the single kernel can't express). kind/alignment contract enforced on both paths (masked/k_grouped rejected, Decision #21). 5 new tests (fused vs f64 reference ~6e-5, fused-fast-path vs composed rel<1e-5, ABI-symbol probe, large-H fallback). runtime_abi dashboard regenerated (+1 symbol). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Completes the Decision #19 follow-on started in the flagship slice. C++ (Tile→Apple fusion passes): - 3 more chain passes (matmul→softmax, matmul→gelu, matmul→rmsnorm) now consume an upstream tessera.fusion.intent (source="descriptor") and fall back to structural re-discovery (source="rediscovered"), with a Decision-#21 warning on descriptor/IR disagreement — same template as matmul→softmax→matmul. - 3 composite passes (swiglu / mla_decode / native_sparse_attn) emit the descriptor with source="composite_op" (they lower a pre-fused op, which is itself the descriptor — no chain re-discovery). All 7 fused calls now carry tessera.fusion.kernel + tessera.fusion.source, so the fusion decision is first-class/auditable in Target IR. Python (emit-half): - canonical_compile.stamp_fusion_intents(module) tags the terminal op of each recognized linear chain with tessera.fusion.intent (from the canonical _KNOWN_FUSION_CHAINS), so the frontend produces descriptor-annotated IR that the C++ passes consume. Idempotent. Built clean; full tessera-ir lit 119 pass / 0 fail. Tests: tests/tessera-ir/phase8/apple_gpu_fusion_descriptor.mlir (extended to gelu + rmsnorm) and tests/unit/test_fusion_intent_emitter.py (6, incl. an emit↔consume contract guard tying the Python intents to the C++ consumers). mypy clean. COMPILER_AUDIT updated. Remaining: auto-wire stamp_fusion_intents into a Target-IR lowering path. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

* feat(apple): Target IR fusion descriptor — emit + consume (Decision #19 slice 1) The Apple Target IR fusion passes re-discover the matmul→softmax[→matmul] / matmul→gelu / ... chains the canonical compile already recognized (COMPILER_AUDIT "fusion intent is too late"). This lands the flagship slice that makes the fusion decision first-class in Target IR. MatmulSoftmaxMatmulFusionToAppleGPU.cpp now: - EMITS the descriptor on the fused call — tessera.fusion.kernel = "matmul_softmax_matmul" + tessera.fusion.source = "descriptor" | "rediscovered" — so the fusion decision (which kernel, and whether the compiler's intent drove it vs. structural re-discovery) is auditable in the IR (Decision #19). - CONSUMES an upstream tessera.fusion.intent on the tail op: when present the fusion is descriptor-driven (source="descriptor"); absent it, the structural walk re-discovers it (source="rediscovered", back-compat). Both fuse to the same kernel. - On a descriptor/IR disagreement (intent set but structure doesn't match) emits a Decision-#21 warning naming the op instead of silently falling back. Built clean; full tessera-ir lit 119 pass / 0 fail (existing fusion fixtures unaffected by the new call attrs). Tests: tests/tessera-ir/phase8/apple_gpu_fusion_descriptor.mlir (lit) + tests/unit/test_apple_fusion_descriptor.py (2). test_coverage regenerated; COMPILER_AUDIT updated. Follow-on: same template for the other 6 Apple fusion passes + a Python emitter stamping tessera.fusion.intent from the canonical _KNOWN_FUSION_CHAINS. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(apple): fusion descriptor across all 7 passes + Python emit-half Completes the Decision #19 follow-on started in the flagship slice. C++ (Tile→Apple fusion passes): - 3 more chain passes (matmul→softmax, matmul→gelu, matmul→rmsnorm) now consume an upstream tessera.fusion.intent (source="descriptor") and fall back to structural re-discovery (source="rediscovered"), with a Decision-#21 warning on descriptor/IR disagreement — same template as matmul→softmax→matmul. - 3 composite passes (swiglu / mla_decode / native_sparse_attn) emit the descriptor with source="composite_op" (they lower a pre-fused op, which is itself the descriptor — no chain re-discovery). All 7 fused calls now carry tessera.fusion.kernel + tessera.fusion.source, so the fusion decision is first-class/auditable in Target IR. Python (emit-half): - canonical_compile.stamp_fusion_intents(module) tags the terminal op of each recognized linear chain with tessera.fusion.intent (from the canonical _KNOWN_FUSION_CHAINS), so the frontend produces descriptor-annotated IR that the C++ passes consume. Idempotent. Built clean; full tessera-ir lit 119 pass / 0 fail. Tests: tests/tessera-ir/phase8/apple_gpu_fusion_descriptor.mlir (extended to gelu + rmsnorm) and tests/unit/test_fusion_intent_emitter.py (6, incl. an emit↔consume contract guard tying the Python intents to the C++ consumers). mypy clean. COMPILER_AUDIT updated. Remaining: auto-wire stamp_fusion_intents into a Target-IR lowering path. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(apple): auto-wire fusion-intent emitter into the compile path Closes the Decision #19 loop end-to-end: the frontend now auto-produces descriptor-annotated Graph IR for Apple targets, which the Target IR fusion passes consume (source="descriptor"). - driver.compile_graph_module calls stamp_fusion_intents(module) before rendering the Graph IR, gated to apple_gpu / apple_cpu (the only Target IR consumers today; the descriptor is backend-agnostic, extends when others consume it). Lazy import avoids the canonical_compile ↔ driver cycle. - Fix: stamp the intent into the op's MLIR `attrs` field, NOT `kwargs`. kwargs are forwarded as the op's real call arguments in the reference/runtime execution path, so a descriptor there leaked into the numpy op call (gelu(**kwargs) got an unexpected 'tessera.fusion.intent'). attrs renders only into the MLIR text the C++ passes read. Verified: apple_gpu/apple_cpu compiles carry tessera.fusion.intent on chain terminals; cpu is gated (unstamped); the previously-broken apple_cpu transformer execution test passes; 116 apple/canonical/strict-dispatch tests green; mypy clean; drift in sync. emitter tests updated to assert via attrs + that the intent never leaks into kwargs. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…avor emission Implements the "adopt now" AMD-ecosystem patterns (rocWMMA / hipBLASLt / AITER / AMD Gluon) as hardware-free IR/metadata/dispatch surfaces, wires them into the audit registry + dashboards, and lands the honest B4 slice (arch-keyed FP8 flavor emission) so the FNUZ-vs-OCP lit fixtures can run. New modules + tests (134 unit tests): - compiler/rocm_mma.py (A1) — unified MFMA(CDNA)/WMMA(RDNA) descriptor: shape is the anchor, A/B operand layout + k_width derived; stable diagnostics for FP8-on-gfx1151, FP8-on-gfx90a, fp32-on-RDNA. - compiler/tile_layout.py (A2) — BlockedLayout/SliceLayout/LinearLayout (bit-basis → free reshape/permute) + costed convert_layout. - compiler/epilogue.py (A3) — hipBLASLt bit-flag Epilogue + EpilogueSpec + autodiff bridge (backward_epilogue / requires_aux; *_AUX pre-activation). - compiler/tuned_dispatch.py (A4/A5) — CSV tuned-config DB keyed on the problem signature (never solidx), de-dup keeps min latency, two-tier override, correctness-gated tune(), untuned worklist artifact. Extended: - rocm_target.py (A6) — arch-keyed FP8 semantics (fnuz/ocp/none) + fp8_dtype_flavor + profile.fp8_semantics. - grouped_layout.py (A6/A7) — hipBLASLt scale-mode -> ScaleLayout; classify_gemm_dispatch (batched vs grouped/device-resident). Registry wiring: - primitive_coverage.py — attaches metadata.rocm_mma (A1) + metadata.epilogue (A3) to GEMM-family / fused-epilogue ops. - backend_manifest.py — typed BackendKernelEntry.mma_descriptor field (ROCm GEMM entries), validated + serialized. - gpu_target_map.py — rocm_target_map.md dashboard gains a per-arch FP8 numeric-semantics table; generated docs regenerated (drift gate clean). B4 (honest slice): - TileToROCM.cpp — arch pass option + arch-keyed fp8_flavor attribute on tessera_rocm.mfma (FNUZ vs OCP), hard error for no-FP8 arch (Decision #21); tessera-rocm-opt rebuilt against MLIR 22.1.6. - lit fixtures fp8_flavor_arch_keyed.mlir / fp8_unsupported_arch.mlir (verified via FileCheck) + test_rocm_fp8_cpp_python_consistency.py (single-source enforcer: C++ emission == rocm_target.fp8_dtype_flavor). Docs: ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md (survey + ranked patterns + landed status); ROCM_AUDIT.md pointer. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…rectness Tested fixes (each with a regression test): - vjp_pow: evaluate log only on the positive subset (no divide-by-zero/invalid RuntimeWarning + masked NaN on non-positive bases). - Parameter.grad setter + tape._accumulate_param_grad: accumulate gradients in >=fp32 regardless of the parameter's storage dtype (fp16 lost precision). Removed the now-dead _TESSERA_TO_NUMPY_DTYPE map. - DistributedArray.parts(): shards now carry a replicated spec, not the parent's partitioned spec (avoids double-partition on a second parts()). - plan_all_to_all: derive source ranks from the np.array_split contiguous layout (remainder spread across the first ranks), not floor division onto the last. - TesseraRuntime: an explicitly requested library (arg or TESSERA_RUNTIME_LIB) that fails to load now raises instead of silently using the mock backend; also catch symbol-binding (AttributeError), not just dlopen (OSError). - TesseraRuntime._telemetry_events: bounded deque, not an unbounded list. - autotune._measure_gemm_wall_clock: label the result "wall_clock_reference" (it times a numpy fp32 reference GEMM, not the backend kernel) — Decision #21. - TileIRLoweringPass: emit a stable [TILE_IR_LOWERING] diagnostic + fail if a tessera.flash_attn/matmul survives unlowered (greedy rewrite returns success on zero matches) — Decision #21. Forward guard; positive lit fixtures pass. - dflash_speculative_verify: document that distribution preservation is exact only for untruncated supports (top-k/top-p renormalize differently). Guards kept for findings that were false positives on inspection: relu/clip cotangent dtype (numpy already promotes); cosine/nt_xent VJPs are finite (not NaN) at zero vectors. Verified: 235 affected tests pass; mypy ratchet 0; ruff clean; tessera-opt rebuilds and phase3 lit fixtures pass. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…erence" test_public_autotune_on_device_cpu_uses_wall_clock_measurement asserted the old method="on_device" label (and "wall-clock measurement" reason) that the Decision-#21 honesty fix renamed to "wall_clock_reference" (it times a numpy fp32 reference GEMM, not the backend kernel). Fixes the CI unit-job failure on a20099a. Missed in the original regression sweep — this test lives in test_profiling_autotuning_foundation.py, not the bayesian/loop autotune files. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…/qkv_projection/factorized_matmul/einsum on gfx1151 Closes the **matmul-family chains** group of the ROCm op surface. All five ops execute on the SAME compiler-generated WMMA GEMM kernel (the rocm_compiled spine), reshaped/batched/split in the runtime — the matmul analog of how flash_attn GQA/MQA reuse the FA kernel, so no new MLIR pass is needed: - batched_gemm — loop the gemm over leading batch dims (np.matmul semantics) - linear_general — axis=-1 reshape [...,K]→[M,K] + gemm + optional bias - qkv_projection — packed x@W_qkv projection (the 3-way split is a host view) - factorized_matmul — GPU matmul + exact host rank-r SVD-truncate epilogue - einsum — single-contraction two-operand specs → (batched) gemm via canonicalize+transpose; other specs emit a stable "unsupported" diagnostic (Decision #21) Shared lane rocm_matmul_family_compiled (one executor + execution_matrix row + KNOWN_EXECUTORS; five backend_manifest entries + fixtures). f16/bf16 storage, f32 accumulate (WMMA on gfx1151 has no f32 storage path). Validated on gfx1151 vs numpy across dtype × shape incl. multi-batch. Dashboards regenerated. Stacked on rocm/silu-mul-alibi (#125) — both touch the same shared files; rebases clean onto main once #125 merges. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…/qkv_projection/factorized_matmul/einsum on gfx1151 (#129) All five matmul-family ops execute on the SAME compiler-generated WMMA GEMM kernel (the rocm_compiled spine), reshaped/batched/split in the runtime — no new MLIR pass: - batched_gemm — loop the gemm over leading batch dims (np.matmul semantics) - linear_general — axis=-1 reshape [...,K]→[M,K] + gemm + optional bias - qkv_projection — packed x@W_qkv projection (the 3-way split is a host view) - factorized_matmul — GPU matmul + exact host rank-r SVD-truncate epilogue - einsum — single-contraction two-operand specs → (batched) gemm; other specs emit a stable "unsupported" diagnostic (#21) Shared lane rocm_matmul_family_compiled; GEMM-family compiled entries keep the unified MMA descriptor. f16/bf16, f32 accumulate. Validated on gfx1151 vs numpy. (Rebuilt clean onto post-#125 main; supersedes the original #126 branch which GitHub closed when its stacked base branch was deleted during the merge.) Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…roup_advantages (S11, was 0/0) (#166) PR-C of the loss campaign. The S11 RL losses were reference-only on both devices; this gives x86 a real lane. New x86_rl_loss_compiled lane. Kernel (avx512_policy_loss_f32.cpp, new file) — per-element surrogate over (logp_new, logp_old, advantages), ratio=exp(ln−lo): - ppo: −min(ratio·adv, clip(ratio,1−ε,1+ε)·adv) - cispo: −(min(ratio, ε_high)·adv·logp_new) exp reuses the Cephes core. grpo == ppo on provided advantages. normalize_group_advantages — (r−mean)/sqrt(var+eps) over the group axis — runs on the existing AVX-512 layer_norm kernel (transpose-to-last → kernel → back). Reduction none/mean/sum on the reduce kernel. The optional KL-penalty / entropy-bonus / masked-reduce add-ons are NOT in the fused lane — they emit a stable diagnostic (Decision #21), never a silent wrong result. Validation: - Standalone C++ test_policy_loss.cpp — ALL PASSED at 2e-5. - tests/unit/test_x86_rl_loss_compiled.py — 14 passed vs the tessera.rl reference (ppo/cispo/grpo × none/mean/sum; normalize over axes 1/-1/0; KL-term diagnostic; reject). Wiring: runtime `_X86_POLICY_OPS` + `_execute_x86_compiled_rl_loss` + `_x86_group_normalize` + symbol binding + executor table; execution_matrix catalog + row; backend_manifest `_X86_KERNELS` (fused, 4 ops) + `_NUMERICAL_FIXTURES`; dashboards regenerated. x86 lane count 83 → 87. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…e/asym + ppo/cispo/grpo/normalize) (#171) Second ROCm-mirror PR — closes the cross-device pair for the S11 binary losses and RL losses (x86 landed in #165/#166). Two new dedicated kernels + the norm lane for advantage normalization. Kernels (new passes, ROCDL-lowered, exp/log1p via math->rocdl): - GenerateROCMBinaryLossKernel: tessera_rocm.binary_loss — bce (kind 0) / asymmetric_bce (kind 1, pos/neg weights), stable softplus form. New rocm_binary_loss_compiled lane. - GenerateROCMPolicyLossKernel: tessera_rocm.policy_loss — ppo (kind 0) / cispo (kind 1) surrogate, ratio=exp(ln-lo), clip attr. New rocm_rl_loss_compiled lane. grpo == ppo on advantages; normalize_group_advantages routes through the rocm norm (layer_norm) lane over the group axis; KL/entropy/mask add-ons diagnose out (Decision #21). Runtime: `_rocm_launch_nary_elementwise` helper (shared K-input launch) + `_execute_rocm_compiled_binary_loss` / `_execute_rocm_compiled_rl_loss` + `_rocm_group_normalize` + executor table; pred/target broadcast for binary. Validation: - tessera-opt codegen + ROCDL lowering verified for both kinds of each. - test_rocm_binary_loss_compiled.py + test_rocm_rl_loss_compiled.py — 16 passed on gfx1151 vs tessera.losses/tessera.rl at 2e-5 (norm 2e-4) + codegen gates. Wiring: execution_matrix catalog + 2 rows; backend_manifest _ROCM_COMPILED (2 + 4 ops) + _NUMERICAL_FIXTURES; dashboards regenerated. bce/asym + ppo/cispo/ grpo/normalize now run on BOTH gfx1151 and AVX-512. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…wo (Spectral PR2) (#176) Spectral PR2 — the x86 power-of-two FFT core, consuming the PR1 SpectralPlan. New x86_fft_compiled lane (fft / ifft / rfft / irfft). Kernel (avx512_fft_f32.cpp, new file) — tessera_x86_fft_c2c_f32: in-place iterative Cooley-Tukey DIT C2C over a batch of power-of-two rows (interleaved complex64). Per-stage twiddles gathered to a contiguous table; the butterfly inner loop runs 16 complex/iteration via deinterleave/interleave permutes + an FMA complex multiply (vr=br*tr−bi*ti, vi=br*ti+bi*tr), scalar tail for the small early stages. Forward e^{−2πi}, inverse e^{+} UNNORMALIZED (the plan applies the scale). NaN/inf flow through. Runtime x86_fft_compiled lane: - fft/ifft — C2C; moveaxis→last, batched rows, plan.scale (ifft 1/N backward). - rfft — real→complex C2C, take [0, n/2]; irfft — Hermitian-reconstruct the full spectrum, inverse C2C, real part, 1/N. axis-generic via moveaxis. - SpectralPlan owns strategy + normalization; non-power-of-two lengths emit a stable diagnostic (Decision #21) until the Bluestein/DFT path (PR3). Validation: - Standalone C++ test_fft.cpp — ALL PASSED vs a naive DFT (fwd+inv, n=2..4096, batched; SIMD path exercised for n>=32). - tests/unit/test_x86_fft_compiled.py — 18 passed vs np.fft (fft/ifft over shapes + inner axis; rfft/irfft round-trip; non-pow2 diagnostic; reject). Wiring: runtime `_x86_fft_c2c_rows` + `_execute_x86_compiled_fft` + symbol binding + executor table; execution_matrix catalog + row; backend_manifest `_X86_KERNELS` (4 ops) + `_NUMERICAL_FIXTURES`; dashboards regenerated. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…ane (PR #205 review) group_norm/instance_norm accept optional weight/bias operands in the catalog, but the x86/rocm norm-compose device lane (like the rmsnorm/layer_norm device ops) is UNWEIGHTED — it read only operand[0] and silently dropped any weight/bias, so launch() reported success with an output missing the affine. Now the lane rejects an artifact carrying >1 operand with a stable diagnostic naming the op + operand count (Decision #21), directing the affine to compose separately via ops.mul/ops.add. Test: instance_norm with a weight operand is rejected (ok=False, "UNWEIGHTED" in reason). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

… (P5) (#205) * feat(s-series): group/instance/weight norm device lanes on x86 + ROCm (P5) The three remaining P5 normalization ops as device-backed composed lanes on both native targets — no new kernels: - group_norm / instance_norm: a row-wise mean/var normalize IS the layer_norm kernel applied to a reshaped [rows, cols] view. instance_norm reshapes to (N*C, spatial); group_norm to (N*G, (C/g)*spatial); the device layer_norm normalizes the rows; host reshapes back. UNWEIGHTED (like rmsnorm/layer_norm — the affine composes separately through ops.mul/ops.add). - weight_norm: w / sqrt(Σ_{¬axis} w² + eps) — the device reduce lane (sum over the last axis after moving `axis` to front + flatten) gives the per-axis sum-of-squares; host does the sqrt + divide. The same device-heavy / host-light split LAMB/Muon use. - runtime.py: `_norm_compose_compute` + `_execute_{x86,rocm}_compiled_normcompose` composing on `_device_layernorm_rows` / `_device_reduce_sum_rows` (axis=-1); registered in `_executor_table`. - backend_manifest.py / execution_matrix.py: fused x86 + compiled rocm entries, rows, path descriptions, fixture map. - tests: test_{x86,rocm}_normcompose_compiled.py — instance/group (various num_groups) + weight_norm (axes 0/1/2/-1) vs nn.functional + bad-divisor rejection. Both executed on hardware (AVX-512 + gfx1151): 17 passed, group/instance maxerr 2.4e-7, weight_norm 3e-8. (rmsnorm_safe already shipped; complex arithmetic + conformal geometry are a follow-up.) Dashboard: the three ops backend_kernel reference -> partial. Drift gate clean; ruff + mypy clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * fix(runtime): reject affine operands in the unweighted norm-compose lane (PR #205 review) group_norm/instance_norm accept optional weight/bias operands in the catalog, but the x86/rocm norm-compose device lane (like the rmsnorm/layer_norm device ops) is UNWEIGHTED — it read only operand[0] and silently dropped any weight/bias, so launch() reported success with an output missing the affine. Now the lane rejects an artifact carrying >1 operand with a stable diagnostic naming the op + operand count (Decision #21), directing the affine to compose separately via ops.mul/ops.add. Test: instance_norm with a weight operand is rejected (ok=False, "UNWEIGHTED" in reason). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * fix(docs): regenerate test_coverage after the affine-rejection test (PR #205 CI) The norm-compose affine-rejection test added an instance_norm reference; the test_coverage drift gate (and its unit mirror) needs the regenerated count. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

* feat(s-series): x86 AVX-512 flash_attn forward lane (P10) The AVX-512 partner to the shipped ROCm WMMA flash_attn — closing the attention x86 gap (P10 of S_SERIES_GAP_CLOSURE_PLAN). FA-style streaming / online softmax: for each query row we sweep the keys once keeping a running max + denominator + rescaled accumulator, so the S×S score matrix is never materialized (O(d) state per query). The two hot inner loops — the QKᵀ dot and the acc += p·V_j update — run as AVX-512 FMA over 16 lanes + a scalar tail (head dim need not be a multiple of 16). f32 throughout (softmax is f32 on every backend). Core MHA path: scale + causal (causal mask aligned to the sequence tail, matching the dense reference), cross-attention (Sq != Sk). GQA/MQA, sliding-window, logit-softcap, attn_bias and dropout are the ROCm lane's extras — the x86 lane rejects them with a stable diagnostic (Decision #21) rather than silently mis-running. Reachable via compiler_path="x86_flash_attn_compiled". Matches the dense attention reference; validated on AVX-512 (6 tests incl. batched heads, causal, cross-attention, GQA-rejection). Manifest flips flash_attn x86 → fused with the numerical fixture; test_no_x86_fused_attention updated to assert the new lane. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * fix(s-series): P10 review — reject attn_bias operand + multi_head_attention Two P1 review fixes on the x86 flash_attn lane: - A 4th operand is the additive attn_bias / mask (Graph IR + the Apple path normalize flash_attn(..., attn_bias=…) into a 4th operand). The lane now rejects >3 operands with a stable diagnostic instead of silently running un-biased attention. - Drop tessera.multi_head_attention from the accepted ops: its contract is rank-3 [B, S, H*D] + num_heads (needs a split to [B, H, S, D] this kernel does not do), so num_heads>1 would mis-run as single-head attention over H*D. The lane now accepts only tessera.flash_attn ([..., S, D]). Added tests locking both rejections. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * chore(s-series): regen test_coverage for P10 flash_attn test count The P10 review-fix commit added 2 tests (attn_bias + multi_head_attention rejection) without regenerating the drift-gated test_coverage dashboard, so the flash_attn direct-test-reference count drifted. Regenerate. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…ified_jit prose (#404) Three related pieces of cleanup on top of the Apple GPU C-ABI registry (#403): 1. Reconcile the ABI registry with the runtime (the registry net now PASSES on-device). The registry listed two device-limit probe symbols the runtime never exported, so the strict net (test_dylib_exports_resolve_every_registry _symbol) failed. Add them to apple_gpu_runtime.mm (+ the non-Darwin stub): - tessera_apple_gpu_max_threadgroup_memory_length ([device maxThreadgroupMemoryLength]; 0 = "use static floor") - tessera_apple_gpu_family_integer (raw MTLGPUFamilyApple* value, e.g. Apple7 == 1007; -1 sentinel) Enum values + selectors grounded in the on-machine SDK headers (Decision #27). apple_target.probe_apple_runtime_limits already binds these defensively. 2. Fix 4 pre-existing stale tests (all test-side, not product regressions): - test_apple_gpu_tiny_decode..._kv_cache: diagnostic wording unified to "KV-cache mutation ..." in c53010c; Decision #21 contract still holds. - test_apple_gpu_simple_moe...: tessera.moe gained a native Apple GPU compute lane (_APPLE_GPU_MOE_COMPUTE_OPS) → now metal_runtime; assert that + numerical proof. Renamed to ..._runs_metal_runtime. - test_apple_gpu_multi_op_with_non_gpu_op_stays_metal_artifact: moe is no longer a non-lane op; swap to tessera.flip (+ lane_for self-guard) to keep the conservative-residency-gate guard meaningful. - test_compile_loads_real_metal_package (pk1): c53010c's blanket "compiled"→"device_verified_jit" rename clobbered a repr assertion + docstrings; revert to "compiled". 3. Repo-wide follow-on to (2): c53010c intentionally renamed the *status token* compiled→device_verified_jit but over-reached into English prose. Revert every word-usage back to "compiled" (diagnostic message strings, NVRTC-/ HIPRTC-compiled, emit/* local variable names, docstrings/comments) while preserving all genuine status-token references (quoted "device_verified_jit" values, backtick doc-refs, and status-name prose like `native/ device_verified_jit`, `= device_verified_jit`, `fused (x86) / device_verified_jit (rocm)`). Each reversion git-verified against the pre-c53010c image. Generated dashboards regenerated (new ABI symbols, reworded notes, flip test count); drift gate clean (22 in sync). Apple suite: 2060 passed, 3 skipped. The emit/* variable renames are covered by test_kernel_cache / test_dynamic_shape_emit / test_spectral_candidates / test_tpp_candidates. Retest target: CUDA (sm_120) + ROCm (gfx1151) — this branch touches the nvidia/rocm emit + manifest prose and needs the ROCm-enabled tessera-opt those boxes have (the local Mac build is CPU+Apple only). Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

… proof Two classes of defect where a build or a generator silently produced something that looked correct. Both are now structurally prevented. tessera-opt: the lean artifact driver re-derived its intent from (ROCM || NVIDIA) && !CORE_TESSERA_IR. CMake clears CORE_TESSERA_IR for any NVIDIA build without CUDA, so configuring NVIDIA alongside the Apple backend linked TesseraApple, defined TESSERA_HAVE_APPLE_BACKEND, and compiled out every one of its registration blocks with no diagnostic. Same for Solvers, Neighbors, TPP, scaling-resilience, and both FA-4 dialects. Leanness is now one named CMake intent plus a feature ledger, and combining it with any other feature is a configure error that names the conflict and both fixes. --tessera-build-info reports profile and features, so telling two binaries apart no longer means diffing --help. The Apple value-lane envelope moved into the backend beside the lowering that consumes it. The emit pipelines share one spine, register only with the core IR, and fail through the registry error handler -- the PassPipelineRegistration<> wrapper takes a void builder, so a failure there would have installed a silently empty pipeline, worse than the abort it replaced (Decision #21). TSOL: the drift gate compares the dashboard against render_dashboard(), so a constant baked into the renderer is self-consistent and invisible to it. Three stale claims survived that way -- a 432-entry registry that is really 482, a line-number citation pointing at an unrelated table, and a hardcoded "zero". All are now derived, and the tests gate the class rather than the instances. The regeneration instruction now names the generator that writes both the .md and the .csv. Apple packet, from PR review: the void ..._f32 ABIs fall through to a numerically-identical CPU reference, so an oracle match proved the math and not the placement. Both entry points gained status-bearing twins following the documented TILE-1 precedent, in the .mm and the non-Darwin stub, and the recorder refuses to seal a fixture whose placement is not positively proven at both the fixture and timing shapes. Host identity is pinned per lane so an M3/M4 host cannot seal an apple7 packet, and source_fingerprint now hashes the runtime source instead of the toolchain digest. The earlier device_event diagnosis was wrong: the Metal command-buffer timer works once dispatch telemetry is enabled. The real gap is the MPSGraph matmul route, which has no device timer, and required_timing_domains is report-wide. Corrected in APPLE-DEVICE-EVENT-1. Tests now declare which tessera-opt passes they drive and skip with the binary's build profile when absent, instead of failing on "Unknown command line argument" -- which reads as a broken test rather than a build-selection problem. The pipeline drift gate learned the registerPassPipeline spelling. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Routing Apple execution through the generic synth->compile->cache loop (kernel_cache.build) was supposed to be a call-site swap. It is not, and the reason is worth recording: AppleMSLEmitter.emit() only ever produced the canonical *scalar* matmul-epilogue body, while run_fused_region prefers the coopmat (simdgroup_matrix) kernel for any eligible region at f16/f32/bf16. Switching production to build() first would therefore have silently downgraded every matrix-unit kernel to the scalar one. Both bodies compute the same numbers, so the F4 oracle would not have caught it — it would have shown up only as a large, unexplained throughput regression. That is precisely the "shared infra must never cap the lead backend's ceiling" rule in Decision #28. So the emitter learns the variants first: * `select_fused_variant(region, dtype)` is now the single rule. `emit()` resolves AUTO through it and `run_fused_region` consults the same function instead of re-deriving the predicate inline, so the kernel the generic loop emits is the kernel the launch path would run. (Same one-predicate-two- callers discipline as the reduce placement fix.) * An explicit `variant=` pins the body. That is the hook a measured arbiter needs to emit every candidate and time them, rather than inheriting this preference order as if it were a decision. * Pinning a variant a region cannot express (coopmat with a reduction, residual, or prologue) raises EmitError rather than quietly returning the scalar body under a coopmat label (Decision #21). Verified: build() now yields `synth_matmul_epi_coopmat` with a simdgroup body for an eligible region, dims sharing a bucket reuse one cache entry, and a dtype variant does not alias. Launch-path numerics unchanged on Metal (f32 max err 0.0, f16 6e-5). Sweep unchanged at 47 pre-existing failures, +3 new passes. test_apple_emitter_wraps_matmul_epilogue_byte_identical encoded the old contract ("the emitter yields the canonical scalar form"). It now pins SCALAR to keep testing byte-identical passthrough, and the AUTO behaviour it used to assert is covered by its own test. Still open before production routes through build(): run_fused_region does synthesis and dispatch in one call, so using a pre-built KernelSource needs those split. That is the next slice, not a rename. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Track 2a. `apple_gpu` is compile-on-launch — its compile_fn returns None and Metal builds the kernel inside run_* via newLibraryWithSource:. Right default for a JIT, wrong one for measuring: every distinct kernel pays a front-end compile on first launch and nothing survives the process. `apple_gpu_air` is a second registered target emitting the *same* MSL and compiling it ahead of time: MSL --`xcrun metal -c`--> .air --`xcrun metallib`--> .metallib Verified on this host with the synthesizer's real coopmat output: 7296-byte .air, 7390-byte .metallib, `deferred=False` with the artifact on disk, while `apple_gpu` still defers and keys separately. A separate target rather than a flag, so AOT-vs-JIT is a measured arbiter candidate per (op, shape-bucket, dtype, target) per Decision #28 — both can be built and timed — instead of a build-time switch nobody revisits. The emitter delegates to AppleMSLEmitter rather than duplicating the synthesis dispatch. A second copy would drift, and then "AOT vs JIT" would be comparing two different kernels while claiming to compare compile strategies. Without the Metal toolchain the lane raises MetalToolchainError with the xcode-select / downloadComponent commands, and never falls back to the JIT path — an AOT measurement that was quietly a JIT one is worse than no measurement. Compile failures name the failing stage (Decision #21). Artifacts are content-addressed on (source, entry), so identical source compiles once per machine; TESSERA_APPLE_AIR_CACHE relocates the cache. Grounding for the harder question (Decision #26a): the .air is LLVM bitcode — magic dec0170b, `target triple = "air64_v28-apple-macosx26.0.0"`, and our pinned LLVM 23 llvm-dis reads it. A test asserts the bitcode magic so a toolchain change that stopped producing it is caught here rather than by whoever later attempts direct AIR emission. Nothing in this lane depends on that; it only bypasses the MSL front end. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…order An assessment of TileSight (arXiv:2607.22432) surfaced that Tessera's hardware-free analytical cost model was a mock: schedule_planner's latency estimate has no memory term, autotune_v2's mock latency is a hand-drawn bowl with its minimum placed by hand, and the target profiles carried capability data but no performance data at all. Theory §4 step 3's "without silicon, score by the Tier 2 cost model" therefore ranked nothing. This lands the two prerequisites. The cost model itself is still open. compiler/target_perf.py — per-device peaks, DRAM bandwidth, LLC size and SMEM/CU keyed by canonical normalize_target() ids. Per-device rather than per-arch because nvidia_sm120 covers parts more than 2x apart. Two honesty rules are gated by tests: provenance is per field (MEASURED / DERIVED / SPEC) because a real row mixes all three, and a value we do not have is absent rather than estimated — accessors return None and require() raises a diagnostic naming the gap (Decision #21). compiler/tile_rasterization.py — block swizzle as a first-class knob. A rasterization order is a permutation of block ids, not a change of arithmetic, so it has a total hardware-free oracle: is_bijection() enumerates the grid and proves every tile is hit exactly once. Previously only Apple had such a knob, via an MLX-inherited hardcoded heuristic. Docs record the finding and its verdict in TILESIGHT_ASSESSMENT.md, with the Theory §4 caveat stating plainly that measured arbitration is the only load-bearing scorer until the estimator is replaced. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Extends the two survey documents with six more ROCm projects. Documentation only. Compiler survey gains section 4.7 on rocisa, TensileLite's nanobind assembly generator — the same Python-driving-C++ shape we have. Three findings worth copying: IR nodes carry a mandatory clone() deep-copy contract; exporting a vector to Python is a copy, so elements are mutable through their shared_ptr but cannot be assigned or replaced; and import raises if any C++ source is newer than the built extension. That last one is added to the take list — we lost time this session to a tessera-opt binary that silently did not match its sources. Patterns doc gains four project briefs and a rocWMMA re-read: rocFFT has the best cache design in the ecosystem. The kernel name is the cache key, with every differentiating parameter encoded into it, so profiler output and cache identity are the same string and the cache needs no schema update when a new parameter appears. Three further key fields guard staleness — architecture, HIP version, and generator version. A read-only system cache ships with the library alongside a read-write user cache, the shipped one populated at build time by a helper that shares the generator but is not installed. AOT and JIT are one path with a policy knob rather than two lanes. Also records that hipRTC holds process-wide locks, so parallel compilation needs a helper process. rocPRIM turns tuning output into generated headers, and its fallback_config is a typed fallback ladder: an untuned type inherits the config of a representative matched on size range and floating-pointness. That is dtype bucketing, the same move Decision #28 makes for shapes. rocRAND is the one with a direct bearing on us. Under dynamic ordering it picks launch geometry per device, and AMD states plainly that the number of generators and the sequence of generated numbers can vary as a result. So reproducibility versus performance is a named opt-in mode, not an emergent property. Worth confirming the same holds for Decision #18: if a tuned launch configuration ever fed an RNG offset scheme, autotuning would silently change numerical output. rocALUTION is included as a contrast, not a pattern. It selects execution location at run time via RTTI and silently migrates an object back to the host when a routine is unavailable on the accelerator. That is the opposite of Decision #21, which requires a diagnostic naming the op and target. Both are defensible for their audience; the contrast is worth recording because silent host migration is how a performance cliff hides. rocWMMA re-read adds that collaborative fragments are a movement concept and are explicitly unsupported in MMA functions, that partial and oversized tiles became the library's problem in 2.0.0, and that the wavefront-centric contract is undefined behaviour rather than a hint. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Correct, and it was a rule I have been applying elsewhere and missed here: I widened a SHARED IR op's accepted form inside one backend. Measured both halves of the report. `ViewOp::verify` accepted any count >= 3 for a pointer-backed view -- so a 4-operand view was legal and meaningless -- while `NVIDIALowering`'s materializer requires exactly 3. The same valid Tile IR therefore lowered on ROCm and was rejected on NVIDIA, with the shared verifier declining to say which form was real. The contract now lives in `ViewOp::verify`: exactly 3 `(base, rowOrigin, colOrigin)`, or 5 with `(rowBound, colBound)`, and anything else is TILE_VIEW_POINTER_ARITY. The ODS documents what the operands mean and that a backend which cannot mask must SAY so rather than ignore the bounds -- dropping them reads past the edge of the matrix. NVIDIA now emits NVFRAGMENT_BOUNDED_VIEW_UNSUPPORTED naming op and target (Decision #21) rather than folding the case into the generic arity message, which would have read as "malformed IR" for IR that is well-formed and merely unsupported there. ── One thing I could not verify, stated plainly ── The NVIDIA diagnostic is written and compiles; it has NOT been executed. The NVIDIA dialect is off by default in this build, and neither `--tessera-lower-to-gpu` nor `--tessera-nvidia-pipeline-sm120` reached the materializer with a bounded view on this host -- no diagnostic, no error. Verifying it needs `-DTESSERA_ENABLE_CUDA=ON`, and the numeric half needs an sm_120 host. Recorded as the NVIDIA sibling outcome rather than implied to work. The shared-verifier half IS verified: `tile_view_pointer_arity.mlir` rejects the 4- and 6-operand forms with the named code. 303 lit, 14442 unit, mypy 0, ruff clean, docs in sync. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Rows 28-31 were opened as shared-contract adoption items and had sat `active` since 2026-07-27. Each is now assessed against source rather than left open by default. Two are satisfied by work that has since landed; two are not, and are narrowed rather than closed. Row 28 (APPLE-ATTN-BWD-2) — CLOSED. E2E-REAL-5B consumes the shared backward artifact and VERIFIES split_count == 2 and reduction_order == (0, 1) rather than assuming them, selecting the MSL split route as the only faithful mapping of an ascending fixed-order reduction. Four exact-device configurations match an independent float64 VJP; repeats are bit-identical. No AMD schedule transferred. Row 29 (APPLE-ATTN-BWD-3) — CLOSED. Apple declares `apple7_recompute` and the ODS verifier enforces it on both schedule.attention and schedule.attention_backward, so Apple cannot silently inherit x86's save_lse or the gfx1151 threshold. The row's condition for retaining recompute is now measured rather than absent: APPLE-ATTN-BWD-PERF-1 answers the row-statistics cost by computing them once per launch in the declared row_prepass instead of saving across the fwd/bwd boundary. Row 30 (APPLE-ATTN-MODIFIERS-1) — ACTIVE, narrowed. Apple expresses causal, softcap, additive bias and MHA/GQA/MQA, and rejects the rest closed. Two specific gaps keep it open: windows are EXCLUDED rather than expressed (the MSL non-causal window is a symmetric half-window, not the shared window_left/window_right semantics, so admitting it would compute a different mask than requested), and rejections name their reason in message text but are not registered diagnostics in the Decision #21 sense. Row 31 (APPLE-STATEFUL-TRANSPORT-1) — ACTIVE, unchanged. No Apple consumer of the generalized target-keyed resident ABI schema exists — searched, not assumed. Apple is aligned with the retirement half (APPLE-PIPE-1 already rejects name-based #tile.buffer_ref) and retains its proven session-private ReplaySSM lifecycle, but Metal threadgroup scheduling against the generalized schema, MoE launch-workspace ownership and rank/device topology binding are untouched. Docs only; no code, IR, ABI or evidence claim changes, so sibling backends are not applicable and their queues are unchanged. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…r backends CI (Linux) failed 18 of the new tests. Root cause is NOT the operand-ordering fix: on non-Darwin hosts `_load_apple_gpu_runtime` compiles `runtime/apple_gpu_runtime_stub.cpp`, whose binary switch implements opcodes 0-8 and whose `default:` arm assigns `out[i] = x`. So `mod`(9), `floor_div`(10), the six comparisons(11-16), and the logical/bitwise ops(17-22) silently return the LEFT operand there instead of computing anything — including for calls that carry no `scalar_side` at all, which behave identically on main. The Mac loads the real MPSGraph symbol, so the first run could not see it. Consumer tests now run through BOTH dispatcher lanes via a `lane` fixture: * `host_reference` — forces the symbol lookup to miss, which is the dispatcher's own documented fallback. Deterministic on every host, so operand ordering stays covered in CI. The swap happens before the lane split, so this exercises exactly the code under test. * `live_kernel` — Darwin-gated, keeping real Metal coverage where the opcode table is fully implemented. Gated rather than probed because probing would mean asserting the very thing these tests assert; the skip reason names the stub gap so it is not mistaken for absent coverage. Mac: 123 passed (both lanes). Simulated non-Darwin: 75 passed, 48 skipped, 0 failed. The stub defect is left for its own change — it is pre-existing, independent of operand ordering, and a Decision #21 violation in its own right (a lowering the backend cannot carry must diagnose, never silently return an operand). Also records the AGENTS.md cross-backend assessment this PR owed, under sync key `SCALAR-SIDE-ORDERING-2026-08-19`: * apple — parity validated on Metal; follow-up required for the portable stub. * nvidia — not applicable; no NVIDIA path consumes the `scalar` kwarg (verified by an exhaustive sweep for `get("scalar"`/`["scalar"]`/`get("other"`). * rocm — not applicable; `_execute_rocm_compiled_binary` binds both operands positionally and raises when fewer are present, so the lifted-scalar form cannot reach the gfx1151 lane. Fails closed by construction. * x86 — not applicable; `_execute_x86_compiled_binary` raises the same way. No device evidence is produced or claimed for nvidia, rocm, or x86. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

`applies_to(region)` is shape-blind by construction: a region carries structure (epilogue chain, dtype, transpose flags) and the dimensions arrive with the operands. So an aligned-only lane could not decline a ragged shape at the applicability gate, and the F4 oracle could not cover for it -- its probe shape is fixed and its verdict is cached under a key with no shape in it. The lane declined inside `run` instead, by returning the numpy reference, after it had already won. Two harms, reproduced against the real NvidiaMmaGemmEmittedCandidate: - Starvation. It won on tier at a ragged shape and handed back numpy while a lower-tier lane that could serve the shape went untried -- the failure RocmWmmaGemmCandidate.available was hardened against on the availability axis (PR #289 review), one axis over. - A fabricated measurement. `_measure` timed the decline and stored 0.00525 ms of numpy under the kernel's name against a real 0.00196 ms rival. That number is not inert: with the backstop disabled the record comes back `separated: True, margin 0.59`, so #663's separation machinery certifies a 2.4x loss for a kernel that never ran. The execution tag already said so. The D3 arbiter log has described this as "a silent degrade ... an unsupported shape" since it was written, and nothing read it. - `Candidate.applies_to_inputs(region, *inputs)`, additive, default True, fail-OPEN on absent/malformed operands (an operand error must still raise through `run`, not be silently excluded -- Decision #21). - `candidate.live_candidates`, one statement of "who is racing", replacing the copy `arbitrate`, `measured_arbitrate` and `corpus_winner` each kept. - `_measure` reads the tag: a reference decline lands in `unmeasured`, not in `candidates`. Fail-CLOSED backstop for lanes that never adopt the hook. - Producers: the NVIDIA emitted GEMM (aligned-only, as its own docstring and device timer already said) and ROCm flash-attention (head_dim % 16, Tier-3 on the one AMD device that executes -- the worse-placed of the two). All four backends assessed under APPLIES-TO-SHAPE-BLIND-2026-09-01: NVIDIA and ROCm follow-up required (device proof owed), Apple and x86 not applicable with reasons. Mutation-verified: four mutations, each killing only its own tests. Device evidence is NOT claimed -- this Mac has no CUDA or ROCm, so harm 2 was reproduced under a simulated device and labelled. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

…ed lane Two follow-on findings from the same code path. Buckets are coarse -- bucket_key maps both (24,12,20) and (32,16,32) to (32,16,32) -- so a ragged workload really does read the aligned workload's corpus row, and run_arbitrated passes a corpus hint to arbitrate as `force`, which restricts to that one name. Making arbitrate shape-aware WITHOUT corpus_winner therefore converts the silent degrade into an ArbiterError. Verified by running the decoupled state, not inferred. corpus_winner withholds the hint because its own `live` set excludes the lane; that coupling is now pinned by a regression test instead of left to be rediscovered by whoever next touches one of the two. And the diagnostic that error carries named the wrong gate: one message, "not available", covered not-registered / wrong-region / unavailable-here alike, and once a shape axis existed it was actively wrong for the commonest case -- the lane IS available, on a host that has it, for a shape it cannot serve. It now names which of the four gates rejected it (Decision #21). test_nvidia_e3_contract asserts the reason rather than the generic string, which is a stronger check than it had. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

chatgpt-codex-connector Bot reviewed May 30, 2026

View reviewed changes

gstoner merged commit e81f2bb into main May 30, 2026
8 of 26 checks passed

gstoner deleted the apple-gpu-conv2d-mpsgraph branch May 30, 2026 12:48

gstoner mentioned this pull request May 30, 2026

Apple GPU Tier-3: conv3d via im2col + GPU batched matmul #22

Merged

gstoner mentioned this pull request Jun 11, 2026

Apple Target IR fusion descriptors — emit + consume (Decision #19) #63

Merged

gstoner mentioned this pull request Jun 25, 2026

feat(rocm): matmul-family compiled lane (batched_gemm/linear_general/qkv_projection/factorized_matmul/einsum) — gfx1151 #126

Closed

gstoner mentioned this pull request Jun 27, 2026

feat(x86): AVX-512 RL policy-loss lane — ppo/cispo/grpo + normalize_group_advantages (S11) #166

Merged

gstoner mentioned this pull request Jun 28, 2026

feat(rocm): binary-loss + rl policy-loss lanes (ROCm mirror) #171

Merged

gstoner mentioned this pull request Jun 28, 2026

feat(x86): AVX-512 radix-2 FFT lane — fft/ifft/rfft/irfft, power-of-two (Spectral PR2) #176

Merged

gstoner mentioned this pull request Jun 29, 2026

feat(s-series): x86 AVX-512 flash_attn forward lane (P10) #213

Merged

gstoner mentioned this pull request Jun 29, 2026

feat(control-flow): CF0 — control-flow contract inventory + unsupported-target diagnostic #225

Merged

gstoner mentioned this pull request Jul 29, 2026

Add calibrated target performance parameters and block rasterization order #475

Merged

gstoner mentioned this pull request Aug 3, 2026

PB-4: compose the TSOL spectral family over Tessera's own FFT; fix a silent wrong answer at non-power-of-two sizes #495

Merged

gstoner mentioned this pull request Aug 19, 2026

graph_ir/runtime: honour scalar_side so a left-side scalar literal lowers correctly #589

Merged

gstoner mentioned this pull request Sep 1, 2026

fix: let a candidate decline a workload its kernel cannot serve #672

Merged

Navigation Menu

Apple GPU Tier-3: conv2d via MPSGraph convolution2D - #21

Apple GPU Tier-3: conv2d via MPSGraph convolution2D#21
gstoner merged 1 commit into
mainfrom
apple-gpu-conv2d-mpsgraph

gstoner commented May 30, 2026

Uh oh!

chatgpt-codex-connector Bot left a comment

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

		if (outH > 0 && outW > 0)
		std::memset(O, 0, static_cast<std::size_t>(N) * outH * outW * Cout * 2);

Conversation

gstoner commented May 30, 2026

What

Layout

Tests

Verification (local, Apple Silicon)

Uh oh!

chatgpt-codex-connector Bot left a comment

Choose a reason for hiding this comment

💡 Codex Review

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant