Apple GPU MLA: weight-absorption decode (the bandwidth win) + KV-cache shape by gstoner · Pull Request #25 · gstoner/tessera

gstoner · 2026-05-30T13:37:25Z

The real MLA bandwidth win — weight absorption. DeepSeek's up-projection
weights absorb into the query/output so attention runs directly against the
cached compressed latent (shared across all heads); per-head K/V are never
materialized.

q_abs  = q_nope @ Wukᵀ
s_nope = q_abs @ c_kvᵀ ;  s_rope = rope(q_rope) @ rope(k_rope)ᵀ
attn   = softmax((s_nope + s_rope)·scale)
ctx    = attn @ c_kv ;  O = ctx @ Wuv

The win

The KV cache stores only c_kv [Skv, Dl] + the shared k_rope [Skv, dr]
(shared across all H heads) instead of per-head K/V. For DeepSeek-V2 dims
(H=128, dn=128, dr=64, dv=128, Dl=512) that's 2304 vs 20480 bytes/token —
~8.9× smaller (reported by the benchmark's cache_ratio_vs_explicit).
Mathematically identical to the explicit-K decoupled-RoPE path.

Correctness anchor

The strongest check: derive explicit per-head K_nope = c_kv @ Wuk and
V = c_kv @ Wuv, run the already-merged explicit-K kernel, and require the
absorbed kernel to match it bit-for-bit (up to fp tolerance). Absorption is
exact, so this passes for both rotation styles.

KV-cache integration

test_absorb_incremental_kv_cache grows the cache (c_kv + k_rope) one token
at a time and decodes at each step — proving the latent cache is all you need
to store. The kernel's inputs are the cache contents.

Surface

tessera_apple_gpu_mla_absorb_decode_f32 (.mm) — one cached MPSGraph
(absorb-K matmul → two-term score → softmax → attn·c_kv → absorb-V), RAII
buffer-pool acquires, host reference fallback + non-Apple stub. Host applies
switchable RoPE (interleaved/half) and tiles shared operands to the B·H batch
(compute-only — the cache stays small).
runtime.py: _apple_gpu_mla_absorb_decode_f32 ctypes wrapper + numpy-friendly
_apple_gpu_mla_absorb_decode(...) dispatcher.

Tests + benchmark

tests/unit/test_apple_gpu_mla_weight_absorb.py — 14 tests: vs numpy
(4 shapes × both styles), exact match vs the explicit-K kernel, styles-differ,
incremental KV-cache decode, cache-size win, symbol export.
benchmarks/apple_gpu/benchmark_mla_absorb.py — absorbed vs explicit-K, reports
cache_bytes_per_token + cache_ratio_vs_explicit.

Note on latency: at Sq=1 decode, absorbed compute is higher than explicit
(absorption trades compute for memory) — the win is cache bandwidth at long
context, captured by the cache-ratio field, not raw compute latency on
already-resident data.

Verification (local, Apple Silicon)

weight-absorption suite: 14/14; full MLA + apple_gpu + buffer-pool + ABI sweep: 510 passed
buffer-pool RAII gate + ABI audit: pass (dashboard 117 → 118 symbols)
benchmark runs on-device (8.9× cache ratio at DeepSeek-V2 dims)
mypy ratchet: clean (only the pre-existing environmental torch-import error)

Follow-on

Paged-cache (KVCacheHandle) wiring for production serving.

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

🤖 Generated with Claude Code

…e shape Adds weight-absorbed MLA decode — DeepSeek's up-projection weights absorb into the query/output so attention runs directly against the cached compressed latent (shared across all heads); per-head K/V are never materialized. q_abs = q_nope @ Wukᵀ s_nope = q_abs @ c_kvᵀ ; s_rope = rope(q_rope) @ rope(k_rope)ᵀ attn = softmax((s_nope + s_rope)·scale) ctx = attn @ c_kv ; O = ctx @ Wuv The KV cache stores only c_kv [Skv,Dl] + shared k_rope [Skv,dr] instead of per-head K/V (~8.9x smaller for DeepSeek-V2 dims). Mathematically identical to the explicit-K decoupled-RoPE path. - apple_gpu_runtime.mm: tessera_apple_gpu_mla_absorb_decode_f32 + mpsg_run_mla_absorb_f32 (one cached MPSGraph: absorb-K matmul -> two-term score -> softmax -> attn@c_kv -> absorb-V; RAII buffer-pool acquires) + reference_mla_absorb_f32 fallback. Host applies switchable RoPE and tiles the shared operands to the B*H batch (compute-only; the cache stays small). - apple_gpu_runtime_stub.cpp: non-Apple reference. - runtime.py: _apple_gpu_mla_absorb_decode_f32 ctypes wrapper + numpy-friendly _apple_gpu_mla_absorb_decode dispatcher. - tests/unit/test_apple_gpu_mla_weight_absorb.py: 14 tests — vs numpy (4 shapes x both rotation styles), exact match against the explicit-K kernel (absorption is exact), styles-differ, incremental KV-cache decode (grow c_kv/k_rope token by token), cache-size win, symbol export. - benchmarks/apple_gpu/benchmark_mla_absorb.py: absorbed vs explicit-K, reports cache_bytes_per_token + cache_ratio_vs_explicit. - docs: plan item 5 (c) marked done; paged-cache (KVCacheHandle) wiring noted as the remaining follow-on. runtime_abi dashboard regenerated (117 -> 118 symbols). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

+          double acc = 0;
+          for (int32_t j = 0; j < Skv; ++j) {
+            double cv = 0;
+            for (int32_t l = 0; l < Dl; ++l) cv += ckvb[j * Dl + l] * wuvb[l * dv + d];


chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 22b0b9bceb

ℹ️ 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-30T13:39:16Z

+    sym = getattr(runtime, "tessera_apple_gpu_mla_absorb_decode_f32", None)
+    if sym is None:
+        return None


Require the new MLA absorb symbol before accepting a cached runtime

When TESSERA_APPLE_GPU_RUNTIME_LIB or an existing build/.../libTesseraAppleRuntime was built before this commit, _load_apple_gpu_runtime() still accepts it because the prebuilt-library probe list was not extended for tessera_apple_gpu_mla_absorb_decode_f32. In that stale-runtime scenario this lookup returns None, causing _apple_gpu_mla_absorb_decode() to return None (and the new benchmark to time a no-op) instead of falling through to rebuild/load a runtime that contains the new symbol.

Useful? React with 👍 / 👎.

…essment Investigated the P2 "long-tail closure" item. Finding: it's already closed for everything with a provable rule — transpose_rule 0 partial; batching_rule 4; sharding_rule 39 — and ALL 43 residual partials sit in genuinely distributed-mesh-gated categories (reasoning-model fused attention, spectral distributed-FFT, distributed linalg, moe all-to-all transport, ebm, state-space, sparse). The _SHARDING_RULE_BY_CATEGORY classifier marks each with a documented "known but mesh-aware" reason; the standard attention family was already proven complete via test_attention_sharding_mock_mesh.py. These are Phase-G-gated by design, not bookkeeping debt. Flipping them without real mock-mesh proofs would be the audit-inflation Decision #25 forbids. Recorded in MASTER_AUDIT P2 so the item reads honestly: closable closure is done; the rest is correctly gated pending genuine per-variant proofs or Phase-G hardware. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Trim CLAUDE.md from 828 lines / ~14k words to 343 lines / ~3k words: - Remove the trailing ~5k-word sprint changelog and all hard-coded counts, deferring to docs/audit/MASTER_AUDIT.md + generated dashboards per Decisions #25/#26. Full prior file preserved verbatim at docs/audit/roadmap/archive/CLAUDE_MD_FULL_2026-06-13.md. - Keep all 27 "Do Not Revisit" decisions, design contracts, build/test commands, and source-location maps (condensed). Also bump pyproject requires-python >=3.8 -> >=3.10 to match the project's declared mypy python_version floor (pyproject.toml:221). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…ough the bridge The real RDNA WMMA matrix instruction now runs on the gfx1100/gfx1151 device and produces a numerically correct GEMM, routed through Tessera's C-ABI launch bridge — the first on-hardware execute-and-compare of a Tessera matmul on non-Apple silicon. - tests/unit/test_rocm_wmma_execute_compare.py: a hipcc harness whose launched kernel uses __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 (the same v_wmma_f32_16x16x16_f16 rocdl_emit.py emits) to compute a 16x16x16 f32<-f16 GEMM, routed through tsrLaunchKernel, compared to a host reference. The operand/accumulator fragment layout matches rocdl_emit.py's grounded mapping (col = lane&15, row = 2*e + lane>>4). maxerr ~3e-8 standalone, <1e-2 through the bridge (f16 rounding). f32<-f16 first (bf16 has documented gfx115x bugs). Honest status — NOT promoted to hardware_verified / backend_kernel complete. This clears the *numerical-proof* half of the backend_manifest hardware_verified contract (execute_compare_fixture), but that status also requires a *shipped* runtime_symbol (an auto-registered ROCm runtime launcher); today the kernel + launcher live in the test harness (like the Apple G7 proof), so flipping the status would be Decision #25 inflation. backend_kernel stays 474/0. The flip is gated on shipping the launcher (ROCM_AUDIT.md Next Work #4) and becomes mechanical once the symbol ships. Docs: STRIX_HALO_EXECUTION_PLAN.md + ROCM_AUDIT.md (Stage D proof + honest gate). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…er work The multi-op compiler-metadata P1 is closed: component-aware metadata (component_ops/program_executable/component_blockers + effects/ shape_envelope/layout_contracts/fusion_groups/outputs) is derived per-component and carried to fn.runtime_artifact().metadata (verified by inspection + 57 locking tests), and fusion dispatch is authoritative. Mark it done in the Current Truth Snapshot, Finished Work, and the P1 queue. Correct now-stale COMPILER_AUDIT claims: Graph-IR folders/canonicalizers are no longer "zero today" (7 ops carry folders/canonicalizers wired into the tessera_jit CPU canonicalize->cse pipeline, graph_ir_folders.mlir), and LayoutAssignmentPass exists/tested (not yet wired into named x86/GPU pipelines). Effect interfaces on the 8 non-pure ops remain the open Phase-1 remainder. Refresh the P2 batching/sharding long-tail counts against the live s_series_status dashboard (4+39=43 on 2026-06-17 -> 6+47=53 on 2026-06-22 as EDM/DiffusionBlocks primitives landed) and point prose at the dashboard as count authority per Decision #25. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Both halves of the backend_manifest hardware_verified contract now ship for ROCm matmul, so flip the row artifact_only -> hardware_verified: * runtime_symbol = tessera_rocm_wmma_gemm_f16 (libtessera_rocm_gemm.so, HIPRTC-compiled RDNA WMMA kernel) * execute_compare_fixture = tests/unit/test_rocm_wmma_runtime_symbol.py (dlopens the shipped symbol, f32<-f16 16x16x16 WMMA vs numpy, maxerr <1e-2) Implemented via a new _ROCM_HARDWARE_VERIFIED override table (the ROCm analog of _APPLE_GPU_KERNELS); when an op is in it the generic MFMA artifact row is replaced by the hardware_verified row. Honest dtype scope (Decision #25): the row claims fp16 + WMMA only (no CDNA MFMA shape/descriptor), shape_envelope documents the single-tile limit. No audit inflation: per-primitive backend_kernel stays 474 open / 0 complete (primitive_is_complete(matmul) is still False — x86/apple/nvidia/cpu rows are not hardware_verified). Only the rocm target row is hardware-verified. runtime_execution_matrix gets NO rocm row yet — that maps to runtime.launch() dispatch, and the shipped symbol is dlopened directly (not routed through launch()); ROCm stays in _UNIMPLEMENTED_TARGETS until an auto-registered executor lands. Regenerated rocm_target_map (matmul -> hardware_verified|fp16; artifact 32->31) and op_target_conformance. Doc drift gate green (17 in sync); manifest / numerical-check / s-series universal-gate / execution-matrix tests pass. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

The compiled path is now a real production-dispatch lane, not just a test chain. An artifact with compiler_path="rocm_compiled" routes through the execution matrix to _execute_rocm_compiled_gemm, which drives the Stage L3 in-process pipeline (tessera-opt -> hsaco, cached per (mt,nt) since the kernel is shape-generic) and launches it via HIP — the SAME runtime.launch() entry point the hand-written rocm_wmma lane uses; only which kernel runs differs. - runtime.py: _execute_rocm_compiled_gemm + helpers (tessera-opt path, hsaco cache, HIP module-launch ABI, size-adaptive (mt,nt) mirroring the oracle). f16 today; bf16 is a structured invalid_artifact, not a miscompute. - execution_matrix.py: ("rocm","rocm_compiled") row + KNOWN_EXECUTORS entry. - test_rocm_compiled_launch_execute.py: launch() executes vs numpy + bit-identical to the hand-written oracle across {16^3, 64x48x32, 256^3}; bf16 rejected. - Regenerated runtime_execution_matrix + test_coverage dashboards (drift in sync). Deliberately NOT flipped to default / promoted in the manifest (Decision #25): the hand-written rocm_wmma stays the default + reference oracle/fast fallback because the compiled path's masked ragged-edge tiles aren't yet perf-competitive (L2 caveat). Flipping the source of truth waits on the masked-edge optimization. Improvement (along the way): the ROCm lit suite is NOT broken — it passes 12/12 on lit 18 + llvm-22 (modern lit sorts substitutions longest-first, so the old %trop/%t collision no longer bites). Corrected the stale "currently broken" claim in ROCM_AUDIT.md and hardened lit.site.cfg.py.in to insert %trop ahead of the built-in %t so it's robust across lit versions regardless. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

@jit

…pe, L2 register-blocked perf, L3 in-process serialization) (#88) * feat(rocm): Stage L1 — problem-size-generic compiler-generated WMMA GEMM Generalize `generate-wmma-gemm-kernel` past the single 16x16x16 tile. The `tessera_rocm.wmma_gemm` directive's m/n/k now denote the WMMA *instruction* tile (16x16x16 — the only tile RDNA's V_WMMA exposes; other extents are a named error), and the emitted kernel is problem-size-generic: - signature (A,B,D : memref<?>, M,N,K : index) — one compiled kernel, any shape - 2-D grid: one wave per 16x16 output tile (blockIdx.y -> M, blockIdx.x -> N) - scf.for K-loop accumulating across 16-wide K panels (acc as iter_arg) - ragged-edge masking: clamp-and-select fragment loads + scf.if-guarded stores Mirrors the hand-written oracle's rung-0 (MT=NT=1) layout and accumulation order exactly, so it stays bit-identical. Executes on gfx1151 vs numpy (<5e-2) and **bit-identical to the hand-written oracle (0.0)** across square, rectangular, and ragged (non-multiple-of-16) shapes {16^3, 32^3, 48x64x32, 40x24x48, 17x15x31}. The 16^3 launch reduces to the Stage K single-tile case. Register-blocked macro-tiling (measured-best 3x4) is L2. - Pass: scf dialect dep + MLIRSCFDialect link; emit general body. - Tests: test_rocm_wmma_gemm_general.py (new, parametrized vs numpy + oracle); test_rocm_wmma_gemm_generated.py updated to the generic ABI (16^3 still oracle-bit-identical) + structural asserts (scf.for/scf.if, 32 inserts). - Docs: ROCM_AUDIT.md Stage L1 marked done. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): Stage L2 — register-blocked macro-tiling + interior fast path Each wave now computes an mt x nt grid of 16x16 output tiles (mt/nt on the tessera_rocm.wmma_gemm directive, default 1), reusing a loaded A fragment across the nt B-tiles and a B fragment across the mt A-tiles. To make register blocking actually pay off, the generated kernel splits: - interior fast path: whole macro-tile in-bounds AND K%16==0 -> each A fragment is a single contiguous vector.load (coalesced), B fragments are unmasked strided loads, stores are unguarded. No masking overhead, so the mt*nt accumulators stay in registers and blocking scales. - masked edge path: clamp-and-select loads + scf.if-guarded stores for ragged rows/cols or K not a multiple of 16 (correctness for any shape). Correctness: all (mt,nt) in {1x1,2x2,2x4,3x4} stay bit-identical to the hand-written oracle on ragged 100x96x64. Perf (gfx1151, kernel-only, vs hand-written _bench at the same mt/nt): - 1536^3 (aligned): compiled meets/exceeds hand-written at every tile, 1.06x-2.56x. - 2048^3: peak 4x4 = 18.7 vs 9.0 TF/s (2.07x). - The fast path is the win: the prior masked-everywhere kernel was 0.12-0.47x. - Honest caveat: a tile whose extent isn't divisible by 16*mt/16*nt has a ragged band on the slower masked path (e.g. 3x4 at 1024/2048, not /48). - Pass: mt/nt attrs (DefaultValuedAttr); fast/masked emit split. - New: benchmarks/rocm/benchmark_rocm_compiled_gemm.py (compiled-vs-handwritten sweep, honest gating, no fabricated numbers). - Tests: test_register_blocked_matches_oracle[1x1,2x2,2x4,3x4]; generated structural test updated to the fast/masked split. - Docs: ROCM_AUDIT.md Stage L2 marked done with measured numbers. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): Stage L3 — in-process MLIR→hsaco serialization (no mlir-opt) Stages I/K/L1/L2 close the compiler-generated-GEMM loop, but the serialization step (gpu-module-to-binary) rode the platform mlir-opt binary. A runtime launch lane can't shell out. L3 links the GPU/ROCDL serialization spine into tessera-opt itself, so the WHOLE chain runs in one invocation: tessera-opt - --pass-pipeline='builtin.module( generate-wmma-gemm-kernel, lower-tessera-target-to-rocdl, gpu.module(convert-scf-to-cf, convert-gpu-to-rocdl, reconcile-unrealized-casts), rocdl-attach-target{chip=gfx1151}, gpu-module-to-binary)' -> gpu.binary ELF Wiring (all gated behind a full ROCm build; the lean artifact driver stays lean): - register gpu-module-to-binary, rocdl-attach-target, convert-scf-to-cf, reconcile-unrealized-casts passes - register the LLVM-IR translations (builtin/llvm/gpu/rocdl) + the #rocdl.target serialization interface - register the cf/arith/func/memref/vector/index/ub ConvertToLLVM external models — what convert-gpu-to-rocdl needs to lower the full gpu.func body (cf block args included); this was the only missing piece vs mlir-opt - InitializeAll{TargetInfos,Targets,TargetMCs,AsmParsers,AsmPrinters} in main AMDGPU codegen comes from the shared libLLVM; ld.lld from the platform LLVM (the ROCDL serializer shells to it — no in-tree lld link needed). The in-process hsaco executes on gfx1151 bit-identical to the hand-written oracle. - tessera-opt.cpp: target init + pass/translation/interface registration (gated). - CMakeLists: link the ROCDL target + translation + ConvertToLLVM libs (gated on a real ROCm/CUDA toolchain build). - Test: test_rocm_wmma_gemm_in_process.py (one tessera-opt call -> hsaco -> execute vs numpy + oracle; mlir-opt never invoked). - Docs: ROCM_AUDIT.md Stage L3 marked done. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): Stage L4 — compiled GEMM as a runtime.launch() lane (opt-in) The compiled path is now a real production-dispatch lane, not just a test chain. An artifact with compiler_path="rocm_compiled" routes through the execution matrix to _execute_rocm_compiled_gemm, which drives the Stage L3 in-process pipeline (tessera-opt -> hsaco, cached per (mt,nt) since the kernel is shape-generic) and launches it via HIP — the SAME runtime.launch() entry point the hand-written rocm_wmma lane uses; only which kernel runs differs. - runtime.py: _execute_rocm_compiled_gemm + helpers (tessera-opt path, hsaco cache, HIP module-launch ABI, size-adaptive (mt,nt) mirroring the oracle). f16 today; bf16 is a structured invalid_artifact, not a miscompute. - execution_matrix.py: ("rocm","rocm_compiled") row + KNOWN_EXECUTORS entry. - test_rocm_compiled_launch_execute.py: launch() executes vs numpy + bit-identical to the hand-written oracle across {16^3, 64x48x32, 256^3}; bf16 rejected. - Regenerated runtime_execution_matrix + test_coverage dashboards (drift in sync). Deliberately NOT flipped to default / promoted in the manifest (Decision #25): the hand-written rocm_wmma stays the default + reference oracle/fast fallback because the compiled path's masked ragged-edge tiles aren't yet perf-competitive (L2 caveat). Flipping the source of truth waits on the masked-edge optimization. Improvement (along the way): the ROCm lit suite is NOT broken — it passes 12/12 on lit 18 + llvm-22 (modern lit sorts substitutions longest-first, so the old %trop/%t collision no longer bites). Corrected the stale "currently broken" claim in ROCM_AUDIT.md and hardened lit.site.cfg.py.in to insert %trop ahead of the built-in %t so it's robust across lit versions regardless. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * perf(rocm): masked ragged-edge parity — K-aligned vectorized edge path L2's first cut dropped any tile whose extent isn't divisible by 16*mt/16*nt to a per-element scalar masked path, so ragged-M/N bands were a perf cliff (3x4 at 1024/2048 -> 0.44x/0.69x of the hand-written kernel). Split the codegen three ways: kAligned ? (tileFull ? fast : edge) : masked. The new EDGE path handles the common ragged case (K%16==0 but the macro-tile straddles the M/N edge) WITHOUT giving up coalesced loads: - load A at a row clamped into range, then zero the whole fragment with ONE loop-invariant vector arith.select if the row is OOB; - load B at a clamped column likewise (one vector select per B tile); - only the once-per-kernel stores are masked (scf.if). The K-loop therefore stays vector-load speed. The per-element scalar path is now reserved for ragged K (K%16!=0) only. Measured on gfx1151 (vs hand-written _bench at the same mt/nt): - 3x4 @ 2048: 0.69x -> 1.82x (6.19 -> 16.25 TF/s) - 3x4 @ 1024: 0.44x -> 0.93x (essentially parity) - aligned unchanged/better: 1536^3 1.06x-3.82x, 2048^3 1.26x-3.93x Ragged-M/N tiles now reach parity-or-better. Bit-identical to the oracle preserved across square/rectangular/ragged shapes and all (mt,nt). This clears the L4 blocker: the compiled lane is perf-ready. Flipping the runtime default is held as an explicit decision pending bf16 in the generated kernel, the ragged-K fast path, and sign-off (ROCM_AUDIT L4). - GenerateWMMAGemmKernel.cpp: 3-path emit + shared masked/unmasked store helper; precompute clamped row/col offsets (arKsafe/colSafe). - test_rocm_wmma_gemm_generated.py: structural asserts updated to the 3-path IR. - ROCM_AUDIT.md: masked-edge parity recorded; L4 remaining updated. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): compiled GEMM — bf16 + ragged-K fast path Closes the two remaining technical gaps before the compiled lane can become the ROCm matmul default. bf16: - tessera_rocm.wmma_gemm gains a `dtype` attr (default "f16"); the generating pass parameterizes the fragment + memref element type, and Stage J emits the matching rocdl.wmma.f32.16x16x16.{f16,bf16} intrinsic. - runtime compiled lane (_execute_rocm_compiled_gemm) supports f16 AND bf16; hsaco cache keyed by (mt,nt,chip,dtype); f32 stays a structured invalid_artifact. Executes bit-identical to the hand-written bf16 oracle. ragged-K fast path: - the K-loop is split so masking never sits on the hot path: an aligned main loop over [0, kMain) (kMain = K rounded down to x16) using the fast/edge panel, then a SINGLE masked tail panel for [kMain, K) when K%16!=0. So ragged K costs one extra masked panel, not a masked K-loop. - collapses the old 3-path dispatch to: tileFull ? fast : edge, each = aligned main loop + ragged-K tail. Measured at 2040^3 (ragged K and M/N): every tile 1.42x-3.62x the hand-written -- no cliff (was the slow per-element path). Correctness: 30 rocm GEMM tests pass on gfx1151 — bf16 vs bf16-oracle and ragged-K (K=31,40) both bit-identical; f32 rejected. - GenerateWMMAGemmKernel.cpp: dtype param + fast/edge/maskedPanel + main+tail K-loop split + shared store helper. - TesseraROCMOps.td: dtype attr. - runtime.py: bf16 in the compiled lane + dtype-keyed cache. - tests: bf16 launch vs oracle, f32 rejection, ragged-K shapes (96x80x40), structural asserts updated to the 2-main-loop + tail-panel shape. - ROCM_AUDIT.md: ragged-K + bf16 recorded; L4 remaining is now just the flip-default sign-off. With these, the only thing between the compiled lane and being the ROCm matmul default is the explicit decision to flip it (held for sign-off). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): flip the default rocm matmul lane to the compiler-generated kernel Stage L4 is complete: @jit(target="rocm") matmul now EXECUTES through the compiler-generated WMMA GEMM by default, with the hand-written kernel as the reference oracle + availability fallback. Before: a jit'd rocm matmul was compiler_path="target_ir_artifact", executable=False — i.e. not executed at all (inspection-only). Both the compiled and hand-written lanes were opt-in. Now (host-gated): - jit.py stamps compiler_path="rocm_compiled" / execution_kind="native_gpu" / executable=True for a rocm single matmul/gemm WHEN the compiled lane can run on the host (tessera-opt built + a usable AMD GPU). _uses_rocm_compiled_default() is shared by the execution_kind property, so is_executable agrees with what launch() actually does (no metadata/property divergence). Off-device the artifact stays target_ir_artifact / not-executable exactly as before — no CI behavior change. - runtime: _execute_rocm_compiled_gemm degrades to the hand-written rocm_wmma oracle on _RocmCompiledUnavailable (no tessera-opt / no serialization spine / no GPU) so flipping the default never regresses availability; a genuine compiled-kernel failure (malloc/launch) is NOT masked — it surfaces. Manifest + matrix promoted: matmul-row notes + the rocm_compiled/rocm_wmma execution-matrix reasons record the compiled lane as the default and the hand-written symbol as the oracle/fallback (runtime_symbol stays the hand-written proof anchor; manifest stays hardware_verified). - jit.py: _rocm_compiled_lane_available() + _uses_rocm_compiled_default() + the rocm executable branch in runtime_artifact(). - runtime.py: _RocmCompiledUnavailable + impl/fallback split. - execution_matrix.py / backend_manifest.py: promoted wording. - tests: default-flip executes + monkeypatched fallback to oracle; test_target_ir_contract host-gated (executable on-device, artifact-only off). - ROCM_AUDIT.md: L4 flipped/promoted — Stage L complete. Pre-existing failures (test_apple_value_target_ir — Apple backend OFF in the ROCm build — and the halo pass_order tests) are unrelated: verified they fail identically with this change stashed. Drift in sync, ruff + mypy clean. 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>

…t4 deferred) Item #3. Measured the compiler-generated int paths rather than assuming. benchmark_rocm_compiled_gemm_dtype.py (new): kernel-only dtype sweep of the compiled WMMA GEMM (f16/bf16/int8/int4, best macro-tile), honest-gated, JSON schema. On gfx1151 at 2048^3: f16 ~23.2 TFLOP/s, bf16 ~23.1, int8 ~21.0 TOP/s, int4 ~23.8 TOP/s (within ~10%) Finding: RDNA 3.5 WMMA runs iu8/iu4 at the SAME matrix-op rate as f16 (no low-precision FLOP-rate multiplier), so the compiled int paths are already compute-competitive and the int4 in-kernel nibble-pack is amortized. Consequence (measured, not assumed — Decision #25): packed-memory int4 (2 int4/byte) would buy memory footprint (1/2) + bandwidth, NOT compute on this arch. Its large sub-byte-strided-B layout is therefore deliberately DEFERRED — unjustified by a compute speedup that doesn't exist on RDNA 3.5. Documented in ROCM_AUDIT with the numbers. drift in sync, ruff clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

@jit

…lue + int/FA perf ladders (#90) * feat(rocm): compiler-generated flash_attn (FA-2 forward) — second compiled op Brings flash_attn into the compiler-generated lane (was hand-written HIPRTC only). The Stage L machinery (directive -> generated WMMA kernel -> in-process hsaco) now covers a second op. - New tessera_rocm.flash_attn directive (head_dim, dtype) + the generate-wmma-flash-attn-kernel pass: a faithful MLIR re-emission of the hardware_verified hand-written FA-2 forward kernel — one wave per (16-query tile, b*h), LDS-staged Q (gpu.func workgroup attributions), S = scale*Q@K^T on WMMA over head-dim chunks, causal/ragged mask, online softmax (running max/sum, rescale), O += P@V on WMMA. Scores are staged in LDS so the QK^T accumulator layout is reread in the P@V A-fragment layout (the layout bridge). head_dim (mult of 16) is compile-time; Sq/Sk/scale/causal are runtime args. - tessera-opt: registered the math ConvertToLLVM external model (so convert-gpu-to-rocdl lowers the softmax math.exp -> llvm exp) + the math dialect; gpu.barrier + workgroup LDS already lower. The flash_attn pipeline is the same in-process chain as the GEMM lane (no mlir-opt). - Test: the compiler-generated FA-2 forward executes on gfx1151 matching a numpy attention reference (maxerr < 2e-2) across head_dim 16/64, causal/non-causal, and ragged Sq/Sk. test_rocm_flash_attn_compiled.py. Honest scope: forward only; the runtime.launch() executor-table lane (a flash_attn op-metadata contract + executor + matrix row) is the remaining glue, same additive step matmul took at L4 — not yet wired, so no execution-matrix row claimed. backward + perf ladder remain (audit item 10). drift in sync, ruff clean, rocm/wmma/flash_attn regression green. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): front-end glue — IR stack emits the wmma_gemm directive Item #2: close the Decision #19 gap. The Graph tessera.matmul -> Tile -> Target-IR lowering (_lower_rocm_op on tile.mma) now EMITS the executable tessera_rocm.wmma_gemm directive (m=n=k=16 WMMA tile + dtype) alongside the abstract tessera_rocm.mfma marker. So a @jit(target="rocm") matmul's target_ir contains the directive the generate-wmma-gemm-kernel pass consumes — the directive is produced by the IR stack, not only synthesized by the runtime. - target_ir.py: tile.mma -> [mfma (abstract marker, kept for the hardware-free contract + lit), wmma_gemm (concrete RDNA executable directive), async_copy, wait]. dtype threaded from the tile op (f16 default). - Test (GPU-free, CI-runnable): test_rocm_matmul_front_end_glue.py — the directive appears in target_ir with the right attrs AND the extracted directive feeds the generate pass into a gpu.func + WMMA op (directive consumed). - The abstract mfma marker stays (target_ir_contract / lit assertions unchanged). - The runtime lane still synthesizes a clean directive at launch for the per-shape mt/nt perf choice; the canonical lowering now owns directive production. Docs updated (op .td + ROCM_AUDIT). drift in sync, ruff + mypy clean, target_ir + rocm/wmma regression green (pre-existing test_apple_value_target_ir failures are Apple-backend-off, unrelated). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * perf(rocm): int8/int4 compiled-GEMM dtype sweep — measured (packed int4 deferred) Item #3. Measured the compiler-generated int paths rather than assuming. benchmark_rocm_compiled_gemm_dtype.py (new): kernel-only dtype sweep of the compiled WMMA GEMM (f16/bf16/int8/int4, best macro-tile), honest-gated, JSON schema. On gfx1151 at 2048^3: f16 ~23.2 TFLOP/s, bf16 ~23.1, int8 ~21.0 TOP/s, int4 ~23.8 TOP/s (within ~10%) Finding: RDNA 3.5 WMMA runs iu8/iu4 at the SAME matrix-op rate as f16 (no low-precision FLOP-rate multiplier), so the compiled int paths are already compute-competitive and the int4 in-kernel nibble-pack is amortized. Consequence (measured, not assumed — Decision #25): packed-memory int4 (2 int4/byte) would buy memory footprint (1/2) + bandwidth, NOT compute on this arch. Its large sub-byte-strided-B layout is therefore deliberately DEFERRED — unjustified by a compute speedup that doesn't exist on RDNA 3.5. Documented in ROCM_AUDIT with the numbers. drift in sync, ruff clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * perf(rocm): compiler-generated flash_attn forward perf ladder (measured) Item #4 (perf-ladder half). Moves the compiled FA-2 forward from "rung-0 correctness-only, no perf data" to a measured ladder. benchmark_rocm_flash_attn_compiled.py (new): kernel-only hipEvent ladder of the compiler-generated FA forward across (head_dim, seqlen), honest-gated, JSON schema. On gfx1151: ~4.0 TFLOP/s at head_dim 64, ~2.4 at 128 (FA-2 fwd FLOPs = 4*B*H*Sq*Sk*D). Modest by design — the kernel is correctness-first (one wave per query tile, LDS round-trips, online-softmax barriers, no KV pipelining / double buffering / multi-wave query tiles); the ladder quantifies the headroom. Audit item 10 updated: forward + forward-ladder done; flash_attn BACKWARD is the largest remaining attention piece (no hand-written oracle — a new kernel validated vs a numpy attention-backward reference; a focused standalone effort comparable to the forward), plus the runtime.launch() executor-table lane. drift in sync, ruff clean. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * feat(rocm): compiler-generated flash_attn BACKWARD (3 WMMA kernels) Expand a single tessera_rocm.flash_attn_bwd directive into the textbook FA-2 backward as three fragment-materialized RDNA WMMA kernels (no stored attention matrix; S/P recomputed per tile): _pre scalar logsumexp L + D=rowsum(O*dO) _dkdv per key-tile: dP=dO@V^T, dS=P*(dP-D), dV+=P^T@dO, dK+=scale*dS^T@Q _dq per query-tile: dQ+=scale*dS@K All use the same C[m,n]=sum_k A[m,k]B[n,k] WMMA primitive + Stage J->I lowering as the forward; P/dS staged in LDS and reread transposed (the layout bridge). Executes on gfx1151 vs a numpy attention-backward reference (itself checked against finite differences): rel-err ~2-4e-4 (f16 storage, f32 accumulate) across head_dim 16/64, causal+non-causal, ragged. flash_attn (fwd+bwd) is now the third compiler-generated op on ROCm after matmul. Backward perf ladder measured: ~1.1-1.3 TFLOP/s @ D=64 (correctness- first; scalar logsumexp pre-pass + 5 matmuls dominate — WMMA logsumexp / causal tile-skip / pipelining are the next rung). - ROCM_FlashAttnBwdOp ODS + GenerateWMMAFlashAttnBwdKernel pass (registered) - tests/unit/test_rocm_flash_attn_bwd_compiled.py (on-device, skip-clean) - benchmarks/rocm/benchmark_rocm_flash_attn_bwd_compiled.py (ladder) - ROCM_AUDIT item 10 updated 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>

…bring-up reality (#111) The Decision #26 entry point had drifted badly behind the bring-up: - MASTER_AUDIT claimed ROCm gfx1151 "executes two ops ... flash_attn forward-only / no perf ladder / no launch lane; the rest artifact_only". Reality on main: a compiler-generated matmul + flash-attention family executes via runtime.launch() — matmul (perf ladder + fused bias/relu/gelu/silu epilogue), flash_attn forward AND backward, GQA/MQA forward AND backward. (Sliding-window + Gemma-2 logit-softcap forward land via #109/#110.) - MASTER_AUDIT + NVIDIA_AUDIT claimed NVIDIA had "no executable runtime row" and that the sm_120 mma.sync emit path still needed building. #106 built it: a sm_120 mma.sync bf16 matmul is hardware-verified end-to-end on a real RTX 5070 Ti (emit -> assemble -> CUDA launch bridge -> execute-and-compare), CUDA 13.3. - Phase H is now framed as split: RDNA gfx1151 LIVE (this work), CDNA MI300X/ MI325X still hardware-gated (distinct MFMA table + FP4/FP6). Counts stay in the generated dashboards (Decision #25/#26) — this page references runtime_execution_matrix.md / rocm_target_map.md, never copies their numbers. Date bumped to 2026-06-24 with a reconciliation note. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…e/mla_decode_fused on gfx1151 (#130) The attention analog of the matmul-family lane — each op composes already-compiled kernels (the WMMA flash_attn kernel + the WMMA GEMM kernel for MLA's latent projections + an elementwise gate), no new MLIR pass: - gated_attention — flash_attn × elementwise sigmoid-gate - mla_decode — latent K/V projections (WMMA GEMM) + flash_attn - mla_decode_fused — down/up projections (c=x@w_dkv; K=c@w_uk; V=c@w_uv) + flash_attn Shared lane rocm_exotic_attn_compiled. f16/bf16, f32 softmax+accumulate. Validated on gfx1151 vs the numpy attention reference. The recurrent DeltaNet variants + block-sparse deepseek stay artifact_only with their technical blockers documented in ROCM_AUDIT (Decision #25 — not marked compiled). (Rebuilt clean onto post-#129 main; supersedes the original #127 branch which GitHub closed when its stacked base chain was deleted during the merge.) Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

… on gfx1151 + AVX-512 Adds the binary-arithmetic compiled lane — the 2-operand sibling of the elementwise unary-math lane (#135) — closing the S2 binary gaps sub/div/pow/ maximum/minimum on both devices we can run on. ROCm (gfx1151): - `tessera_rocm.binary` ODS op + `generate-rocm-binary-kernel` pass (flat 2-operand per-element kernel, mirrors the silu_mul 2-operand template); pow lowers via math→ROCDL; maximum/minimum are IEEE NaN-propagating. - runtime `rocm_binary_compiled` lane (_build/_execute + executor table + execution_matrix catalog/row); `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_binary_compiled.py: validated on gfx1151 vs numpy across kind × dtype × shape incl. rank-3 + NaN-propagation, + a GPU-free codegen gate. x86 (AVX-512): - avx512_binary_f32.cpp: direct-intrinsic subset sub/div/maximum/minimum (NaN-propagating via unordered-compare blend); pow stays numpy-reference (no fused x86 claim, as unary leaves the transcendentals). Validated standalone (test_binary.cpp, ALL PASSED on the AVX-512 host). No `x86:fused` manifest slot claimed — the Graph-IR→x86 lowering isn't wired yet (kernel-exists != end-to-end-wired, Decision #25). Regenerated drift dashboards (runtime_abi/runtime_execution_matrix/test_coverage) + ROCM_AUDIT.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

… on gfx1151 + AVX-512 (#136) * feat(s2): elementwise binary-arithmetic kernels (sub/div/pow/max/min) on gfx1151 + AVX-512 Adds the binary-arithmetic compiled lane — the 2-operand sibling of the elementwise unary-math lane (#135) — closing the S2 binary gaps sub/div/pow/ maximum/minimum on both devices we can run on. ROCm (gfx1151): - `tessera_rocm.binary` ODS op + `generate-rocm-binary-kernel` pass (flat 2-operand per-element kernel, mirrors the silu_mul 2-operand template); pow lowers via math→ROCDL; maximum/minimum are IEEE NaN-propagating. - runtime `rocm_binary_compiled` lane (_build/_execute + executor table + execution_matrix catalog/row); `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_binary_compiled.py: validated on gfx1151 vs numpy across kind × dtype × shape incl. rank-3 + NaN-propagation, + a GPU-free codegen gate. x86 (AVX-512): - avx512_binary_f32.cpp: direct-intrinsic subset sub/div/maximum/minimum (NaN-propagating via unordered-compare blend); pow stays numpy-reference (no fused x86 claim, as unary leaves the transcendentals). Validated standalone (test_binary.cpp, ALL PASSED on the AVX-512 host). No `x86:fused` manifest slot claimed — the Graph-IR→x86 lowering isn't wired yet (kernel-exists != end-to-end-wired, Decision #25). Regenerated drift dashboards (runtime_abi/runtime_execution_matrix/test_coverage) + ROCM_AUDIT.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(tests): refresh suite-by-suite test counts (fast ~11,686 / full ~12,464) The perf-doc drift gate (test_perf_doc_fast_and_full_counts_are_current) went red: the suite grew past the doc's stale ~9,776 fast / ~10,554 full figures (±15% gate). Measured 2026-06-26 via `pytest tests/unit --collect-only`: 11,686 fast / 778 deselected / 12,464 full. Pre-existing drift surfaced by this PR's added tests; updates the table + prose to current counts. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(audit): resync docs_freshness dashboard after rebase onto main consolidation The rebase onto main (5b06bb1 docs consolidation 34→22) shifted the freshness audit to 89 catalogued docs; regenerate so the drift gate matches the merged state. 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>

…51 + AVX-512 (#137) Adds the S2 comparison family — the first compiled lane with a non-float (boolean) result — on both devices this box runs. eq/ne/lt/le/gt/ge over float inputs, i8/bool output, NaN semantics matching numpy (ordered everywhere except ne, which is unordered-not-equal). ROCm (gfx1151): - `tessera_rocm.compare` ODS op + `generate-rocm-compare-kernel` pass: flat 2-operand kernel, arith.cmpf (OEQ/UNE/OLT/OLE/OGT/OGE) -> extui i1->i8, memref<?xi8> output. - runtime `rocm_compare_compiled` lane (_build/_execute with asymmetric in/out element sizes: float in, 1-byte bool out) + executor table + execution_matrix catalog/row; `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_compare_compiled.py: validated on gfx1151 vs numpy across kind × dtype × shape incl. rank-3 + NaN semantics, + a GPU-free codegen gate. x86 (AVX-512): - avx512_compare_f32.cpp: `_mm512_cmp_ps_mask` (matching _CMP_*_OQ / _CMP_NEQ_UQ predicates) + `_mm_maskz_set1_epi8` to expand the mask to 0/1 bytes. C's native float operators already match numpy's NaN rule, so the scalar reference uses them directly. Validated standalone (test_compare.cpp, ALL PASSED incl. NaN). No `x86:fused` manifest slot claimed — Graph-IR->x86 lowering still unwired (Decision #25). Regenerated drift dashboards + ROCM_AUDIT.md. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…loor/round) on gfx1151 + AVX-512 (#138) Extends the existing elementwise unary-math lane (#135) with 12 more ops, closing the S2 scalar-math trig/special + numeric-helper rounding gaps on both devices. ROCm (gfx1151) — generate-rocm-unary-kernel gains: - trig/special: cos, tan, sinh, cosh, asin, acos, atan, erfc (math.{Cos,Tan,Sinh,Cosh,Asin,Acos,Atan,Erfc} -> ocml) - rounding: floor, ceil, round, trunc; `round` = math.RoundEvenOp to match numpy's round-half-to-even. All 12 lower through the same ROCDL path; runtime _ROCM_UNARY_OPS + manifest + fixtures extended. test_rocm_unary_compiled.py validates all on gfx1151 vs numpy (tan domain bounded inside (-pi/2, pi/2); asin/acos to |x|<=1; erfc ref added) + the GPU-free codegen gate covers the new kinds. x86 (AVX-512): - avx512_unary_f32.cpp gains the rounding subset floor/ceil/trunc/round as direct `_mm512_roundscale_ps` intrinsics (round = ties-to-even / std::nearbyint). Trig/erfc are transcendental and stay numpy-reference on CPU (no fused x86 claim), as the lane already leaves exp/log/erf. Validated standalone (test_unary.cpp, ALL PASSED). No x86:fused slot claimed (Decision #25). Regenerated test_coverage dashboard + ROCM_AUDIT.md. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…VX-512 (#139) Adds the S2 logical family over i8 booleans — the mask-composition sibling of the comparison lane — on both devices. and/or/xor (binary) + not (unary). Inputs are normalized to bool via != 0 (numpy: any nonzero is true), so the kernel is correct for arbitrary i8 inputs. ROCm (gfx1151): - `tessera_rocm.logical` ODS op + `generate-rocm-logical-kernel` pass: normalize via arith.cmpi ne 0 -> i1, combine with arith.{andi,ori,xori}, extui i1->i8. Binary kinds emit a 4-arg signature; `not` emits a 3-arg (unary) signature. - runtime `rocm_logical_compiled` lane (arity-aware: 1 vs 2 operands, i8 bool in/out) + executor table + execution_matrix catalog/row; `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_logical_compiled.py: validated on gfx1151 vs numpy across kind × shape incl. rank-3 + nonzero-normalization, + a GPU-free codegen gate. x86 (AVX-512): - avx512_logical_i8.cpp: `_mm512_cmpneq_epi8_mask` + `_mm512_maskz_set1_epi8` to normalize inputs to 0/1, then `_mm512_{and,or,xor}_si512` (64 bytes/vector). Validated standalone (test_logical.cpp, ALL PASSED). No x86:fused slot claimed (Decision #25). Bitwise and/or/xor/not over integers is a separate follow-up (distinct int32 path). Regenerated drift dashboards + ROCM_AUDIT.md. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…VX-512 (#140) Adds the S2 bitwise family over i32 integers — the integer sibling of the logical lane (#139) — on both devices. and/or/xor (binary) + not (unary), acting on the full bit pattern (no normalization, unlike logical). ROCm (gfx1151): - `tessera_rocm.bitwise` ODS op + `generate-rocm-bitwise-kernel` pass: arith.{andi,ori,xori} on i32; not = a ^ -1. Binary kinds emit a 4-arg signature; not emits a 3-arg (unary) signature. - runtime `rocm_bitwise_compiled` lane (arity-aware, i32 in/out) + executor table + execution_matrix catalog/row; `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_bitwise_compiled.py: validated on gfx1151 vs numpy across kind × shape incl. rank-3 + signed inputs, + a GPU-free codegen gate. x86 (AVX-512): - avx512_bitwise_i32.cpp: `_mm512_{and,or,xor}_si512` (16 i32/vector); not = xor all-ones. Validated standalone (test_bitwise.cpp, ALL PASSED). No x86:fused slot claimed (Decision #25). Completes the S2 logical-family registry gaps (logical_* in #139 + bitwise_* here). Regenerated drift dashboards + ROCM_AUDIT.md. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

…86:fused on AVX-512 hardware Wires the first REAL x86 execution path for the elementwise families beyond GEMM. Until now the AVX-512 kernels (reduce/unary/binary/compare/logical/bitwise) were validated standalone only and never reachable through runtime.launch(); the x86 backend isn't even linked into the runtime. Infra (the CPU analog of the ROCm compiled lanes): - New CMake target `tessera_x86_elementwise` — a SHARED lib of just the AVX-512 elementwise kernels (no AMX, so it builds on gcc), exporting the C-ABI symbols. - Runtime `_load_x86_elementwise()` ctypes-loads libtessera_x86_elementwise.so (env override TESSERA_X86_ELEMENTWISE_LIB → CMake build dir); host-gated, skips clean where absent (e.g. CI's pytest lane that doesn't build the C++ backend). Reduce lane: - `_execute_x86_compiled_reduce` (compiler_path x86_reduce_compiled): folds an arbitrary axis to [outer,inner] (like the ROCm reduce lane), calls tessera_x86_avx512_reduce_f32, supports keepdims. f32 only. - Extended avx512_reduce_f32.cpp with `min` (kMin=3, NaN-propagating via _mm512_reduce_min_ps) so x86 covers sum/mean/max/min like ROCm; test_reduce.cpp updated. - execution_matrix `x86_reduce_compiled` row/catalog (new `x86` target); `_X86_KERNELS` sum/mean/max/min/amax/amin = `fused` with an execute_compare_fixture; `_NUMERICAL_FIXTURES` (op,"x86") entries. - test_x86_reduce_compiled.py: 41 on-device tests vs numpy (axis/keepdims/NaN), skip-clean without the .so. This makes the `x86:fused` reduce claim HONEST (real end-to-end exec + on-device numerical proof, Decision #25). unary/binary/compare/logical/bitwise ship in the same .so and earn their slots in follow-up PRs. Regenerated dashboards + BACKEND_AUDIT.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…86:fused on AVX-512 hardware (#141) Wires the first REAL x86 execution path for the elementwise families beyond GEMM. Until now the AVX-512 kernels (reduce/unary/binary/compare/logical/bitwise) were validated standalone only and never reachable through runtime.launch(); the x86 backend isn't even linked into the runtime. Infra (the CPU analog of the ROCm compiled lanes): - New CMake target `tessera_x86_elementwise` — a SHARED lib of just the AVX-512 elementwise kernels (no AMX, so it builds on gcc), exporting the C-ABI symbols. - Runtime `_load_x86_elementwise()` ctypes-loads libtessera_x86_elementwise.so (env override TESSERA_X86_ELEMENTWISE_LIB → CMake build dir); host-gated, skips clean where absent (e.g. CI's pytest lane that doesn't build the C++ backend). Reduce lane: - `_execute_x86_compiled_reduce` (compiler_path x86_reduce_compiled): folds an arbitrary axis to [outer,inner] (like the ROCm reduce lane), calls tessera_x86_avx512_reduce_f32, supports keepdims. f32 only. - Extended avx512_reduce_f32.cpp with `min` (kMin=3, NaN-propagating via _mm512_reduce_min_ps) so x86 covers sum/mean/max/min like ROCm; test_reduce.cpp updated. - execution_matrix `x86_reduce_compiled` row/catalog (new `x86` target); `_X86_KERNELS` sum/mean/max/min/amax/amin = `fused` with an execute_compare_fixture; `_NUMERICAL_FIXTURES` (op,"x86") entries. - test_x86_reduce_compiled.py: 41 on-device tests vs numpy (axis/keepdims/NaN), skip-clean without the .so. This makes the `x86:fused` reduce claim HONEST (real end-to-end exec + on-device numerical proof, Decision #25). unary/binary/compare/logical/bitwise ship in the same .so and earn their slots in follow-up PRs. Regenerated dashboards + BACKEND_AUDIT.md. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

… + vision (#222) The genuine remaining contract gap: three categories were ABSENT from the batching / sharding / transpose rule maps, so their ops resolved to an UNSPECIFIED axis (a `?` in the coverage scan, not a `partial`). P14 formalizes them honestly: - diffusion (edm_loss_weight, edm_precondition) — elementwise functions of σ → complete on all three axes (the elementwise rule: mapped axis broadcast- preserved, partition spec preserved, linear in the per-element cotangent). - diffusion_schedule (karras_sigma_schedule, equiprob_band_partition) — deterministic schedule generators from scalar config with no batchable / shardable data axis and no linear primal → not_applicable on all three. - vision (resize / crop / normalize / interpolate) — the previously-missing SHARDING rule is complete (batch-parallel per-image; batching + transpose were already complete). Every primitive's category now resolves in all three rule maps (no `?`). The structural/layout batching gap (the plan's original headline) was already closed (layout_transform/indexing/transform = complete); the remaining sharding `partial`s (attention/moe/spectral/linalg/sparse/state_space/recurrent/ga/ebm) are genuine mesh-gated distributed gaps and stay partial (Decision #25 — no dishonest flips). test_p14_diffusion_vision_contracts locks the rules + verifies the batch-parallel / vmap semantics behaviourally (edm_loss_weight elementwise batch; image_normalize batch-shard reconstruction). Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

github-advanced-security Bot found potential problems May 30, 2026

View reviewed changes

Comment thread src/compiler/codegen/Tessera_Apple_Backend/runtime/apple_gpu_runtime_stub.cpp

double acc = 0;

for (int32_t j = 0; j < Skv; ++j) {

double cv = 0;

for (int32_t l = 0; l < Dl; ++l) cv += ckvb[j * Dl + l] * wuvb[l * dv + d];

chatgpt-codex-connector Bot reviewed May 30, 2026

View reviewed changes

gstoner merged commit 0af38ca into main May 30, 2026
7 of 28 checks passed

gstoner deleted the apple-gpu-mla-weight-absorb branch May 30, 2026 14:27

gstoner mentioned this pull request May 30, 2026

Apple GPU MLA: paged-cache decoder (production-serving wiring) #26

Merged

gstoner mentioned this pull request Jun 15, 2026

Track L — genuine gated delta rule + Mamba2 SSM (L0–L4.1) #77

Merged

gstoner mentioned this pull request Jun 22, 2026

ROCm Stages A–D: RDNA WMMA path → emit/assemble → launch bridge → on-hardware execute-and-compare (gfx1151) #83

Merged

gstoner mentioned this pull request Jun 23, 2026

feat(rocm): Stage L — converge the compiled GEMM path (L1 general-shape, L2 register-blocked perf, L3 in-process serialization) #88

Merged

gstoner mentioned this pull request Jun 24, 2026

feat(rocm): compiler-generated flash_attn fwd+bwd + IR-stack matmul glue + int/FA perf ladders #90

Merged

gstoner mentioned this pull request Jun 25, 2026

docs(audit): refresh Phase G/H/I status to merged hardware-bring-up reality #111

Merged

This was referenced Jun 26, 2026

feat(s2): elementwise unary-math kernels (exp/log/sqrt/erf/…) on gfx1151 + AVX-512 #135

Merged

feat(s2): elementwise binary-arithmetic kernels (sub/div/pow/max/min) on gfx1151 + AVX-512 #136

Merged

gstoner mentioned this pull request Jun 26, 2026

feat(s2): elementwise comparison kernels (eq/ne/lt/le/gt/ge) on gfx1151 + AVX-512 #137

Merged

4 tasks

gstoner mentioned this pull request Jun 26, 2026

feat(s2): unary-math tail — trig/special/rounding (cos/tan/sinh/…/floor/round) on gfx1151 + AVX-512 #138

Merged

4 tasks

gstoner mentioned this pull request Jun 26, 2026

feat(s2): elementwise logical kernels (and/or/xor/not) on gfx1151 + AVX-512 #139

Merged

4 tasks

gstoner mentioned this pull request Jun 26, 2026

feat(s2): elementwise bitwise kernels (and/or/xor/not) on gfx1151 + AVX-512 #140

Merged

4 tasks

gstoner mentioned this pull request Jun 26, 2026

feat(x86): runtime-loadable AVX-512 elementwise lane — reduce earns x86:fused on hardware #141

Merged

4 tasks

This was referenced Jun 28, 2026

docs(roadmap): single S-series closure plan + Tier-0 compiler foundation + native fused-kernel inventory #187

Merged

feat(s-series): P14 — formalize vmap/sharding/transpose for diffusion + vision #222

Merged

Navigation Menu

Apple GPU MLA: weight-absorption decode (the bandwidth win) + KV-cache shape - #25

Apple GPU MLA: weight-absorption decode (the bandwidth win) + KV-cache shape#25
gstoner merged 1 commit into
mainfrom
apple-gpu-mla-weight-absorb

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!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

gstoner commented May 30, 2026

The win

Correctness anchor

KV-cache integration

Surface

Tests + benchmark

Verification (local, Apple Silicon)

Follow-on

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!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants