Apple SDK27 FP8/FP4 matmul2d lanes, measured toolchain pins, fleet-portable native tests, CLAUDE.md refresh - #749
Conversation
…nd tiled-scratch fixes macOS 27 / Metal 4.1 on the M1 Max. FP8 E4M3/E5M2 and FP4 E2M1 execute as MetalPerformancePrimitives matmul2d operands (both-lowp and half x lowp) bound as strided MTLTensor views with padded row strides and nonzero tile origins; fp32 accumulation proven bit-exactly; matched device timing shows FP8 at 0.77-0.93x fp16 (emulated), host packing 4-10x the matmul, fused epilogues 0.5-5% slower than decomposed on device time. No routing promotion. Measured Apple facts the lane enforces: 128-byte row strides and data-plane offsets for 8/4-bit tensors; 4-bit buffer offsets are applied by Apple at 2x the byte value (halved, 256 B granularity, canaried by the origin tests). Fixes found on the way: fast-math tanh made gelu NaN from pre-activation 10.25 in every fused epilogue (clamped, regression test); the tiled reduction scratch length was passed to Metal unrounded and the Python contract refused every ragged width since 2026-09-05 (both round to 16 B now); the packed-numeric probe and the new lanes go through the buffer pool and the dispatch breaker; off-Darwin stubs for the four new symbols; test files use the shared device markers; AIR fixtures follow the installed toolchain's air.version. runtime.py also carries the fleet-portable LLVM companion lookup and the ROCm toolkit detector fix that the next commit documents. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…ools, ROCm toolkit detection Pins bumped to the toolchains measured on the fleet 2026-09-15 (nvcc 13.4.59, PTX ISA 9.4, driver 610.88 on The-Super-Bear; ROCm 10.0.0 / HIP 7.15.26333 on Princess-Luna and Tajasarus) in gpu_target.py, rocm_target.py, the CMake pin file, AdapterVersionPin.h, the pipeline descriptions and the compile validators. The feature matrices were not re-evaluated; their version-named capability markers record the evaluation baseline and say so. The generated mfma_table.inc header follows the pin (gitignored build artifact). The runtime ABI audit read the CMake ROCm pin from a usage comment whose argument is the HIP version, and scanned to the first '=' after a constant name; both regexes now match the real assignments. Native lanes hard-coded /usr/lib/llvm-23/bin for mlir-opt/mlir-translate/llc, so every native package test was a false failure off Ubuntu. llvm_tools.py resolves the matched LLVM 23 companion per host, and the shared subprocess choke point maps a caller-supplied missing companion to it. The ROCm toolkit detector took Homebrew's generic ld.lld as a ROCm root; it now requires a ROCm marker, and backend='rocm' package tests gate on it. ROCm-lowering tests gate on the registered passes rather than on driver presence; the VNNI sanitizer test carries the AVX-512 marker. Sibling outcomes recorded in the rocm, nvidia and x86 queues (bare-pytest verification owed on each box). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…ecision amendments CLAUDE.md is refreshed against the 413 commits since its 2026-08-30 sweep: the execution-reality paragraph, the four-box fleet (Princess-Luna gfx1151, Tajasarus gfx1201 with the assertions-ON LLVM, The-Super-Bear RTX 5070 sm_120, the Mac on macOS 27 / Xcode 27), native NVIDIA scheduling, the foundation program (IR-NATIVE-FOUNDATION-1) as the status umbrella, dated amendments to Decisions #12, #19, #26, #26a and #31, a feature-flag GPU-only tier, the new source modules, scripts and dashboards, the AGENTS.md relationship, and the post-Windows-reboot WSL signature. GETTING_STARTED, PROJECT_STRUCTURE and the three compiler plans' fleet tables carry the same versions and hosts. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Recorded on the M1 Max (macOS 27.0, Xcode 27.0) with the dylib built from this branch's runtime source, after the recorder's clean-tree check; the e2e_fleet dashboards are regenerated from it. Measurement evidence for this revision only; no performance promotion. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 3f276719f8
ℹ️ 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".
| constant int4 &p [[buffer(2)]], | ||
| uint2 gid [[thread_position_in_grid]]) { | ||
| int n = int(gid.x), m = int(gid.y); | ||
| if (n >= p.z) return; |
There was a problem hiding this comment.
When M is not divisible by 8, the dispatch at tessera_apple_gpu_mtl4_bias_act_f32 rounds the Y grid up, but this kernel checks only n, so threads with m >= M access C[m * ldc + n]. Pool bucket rounding masks this for some shapes, but accepted shapes near a bucket boundary (for example M=127, N=ldc=513) exceed the actual Metal buffer and can produce a command-buffer fault or memory corruption. Pass M to the kernel and reject both out-of-range coordinates.
Useful? React with 👍 / 👎.
| ROOT = Path(__file__).resolve().parents[2] | ||
|
|
||
|
|
||
| @pytest.mark.hardware_avx512 # compiles the AVX-512 VNNI kernel: x86 hosts only |
There was a problem hiding this comment.
This test invokes only tessera_x86_reference_gemm_u8s8_s32, and the translation unit is compiled without AVX-512 flags, so it does not execute an AVX-512 instruction. Marking it hardware_avx512 makes PR_MARKER_EXPRESSION exclude the UBSan regression entirely and also skips it on ordinary x86 hosts, removing the portable coverage that detects signed-overflow regressions; gate only on the x86 compile environment rather than AVX-512 hardware.
Useful? React with 👍 / 👎.
PR #749 review fixes: guard bias/act rows past M; keep VNNI wraparound test in the PR lane
1 participant
Owning items
APPLE-METAL41-20260914(Apple SDK27 low-precision lanes, owner IR-NATIVE-FOUNDATION-1) andHOST-SWEEP-2026-09-15(fleet-portable native tests, owner COMPILER-DEVEX-1). Toolchain pins under Decision #11 / #19.What landed (four commits, one concern each)
matmul2doperands on the M1 Max (macOS 27 / Metal 4.1), bound as strided MTLTensor views with padded row strides and nonzero tile origins; fp32 accumulation proven bit-exactly; 99 owning-Mac tests. Measured: FP8 at 0.77–0.93× fp16 device time (emulated), host packing 4–10× the matmul, fused epilogue 0.5–5 % slower than decomposed → no routing promotion. Fixes found on the way: fast-mathtanhmade gelu NaN from pre-activation 10.25 in every fused epilogue (clamped, regression test); tiled reduction scratch was passed to Metal unrounded and the Python contract refused every ragged width since 2026-09-05 (both round to 16 B); pool/breaker/stub/marker governance gates satisfied for the new lanes and the packed-numeric probe; AIR fixtures follow the installed toolchain'sair.version.llvm_tools.pyresolves the matched LLVM 23 companion per host (the hard-coded/usr/lib/llvm-23/binmade every native package test a false failure off Ubuntu). The ROCm toolkit detector no longer takes a genericld.lldas a ROCm root. ROCm-lowering andbackend='rocm'package tests gate on registered passes / an actual toolkit.Shared contracts changed
Toolchain pins (three sources + ABI audit regexes);
runtime._rocm_toolkit_rootsemantics (requires a ROCm marker);apple_threadgroup.ThreadgroupSlot.lengthrounds instead of refusing; four new Apple runtime symbols with off-Darwin stubs; new capability gates intests/_support/rocm_build.py.Backend-plan updates
apple/todo.md (SDK27 lanes, measured envelope, Apple 4-bit offset defect, macOS 27 f16 regressions, ledger re-seal), rocm/nvidia/x86 todo.md (
HOST-SWEEP-2026-09-15sync entries with sibling outcomes), BACKEND_AUDIT.md (gfx1201 no longer artifact-only; Apple FP8 executes but is not an acceleration claim), NATIVE_RDNA4_COMMISSIONING.md and BLACKWELL_SM120_EXECUTION_PLAN.md stale banners, GETTING_STARTED fleet matrix.Validation performed (Mac M1 Max, macOS 27.0, unsandboxed for device lanes)
tessera-optrebuilt against HEAD: the 121 pre-existing failures of the stale build reduced to 4, all evidence/OS items recorded with reproductions (below). 672 passed / 164 skipped over the affected 41 files plus touched suites at the end.benchmarks/baselines/apple_macos27_20260914/lowp_matmul2d.json(Metal 4 counter-heap,route+timing_sourceper row).Still red on the owning Mac, by design or awaiting measurement
test_apple_legacy_retune_benchmark::test_strict_retune_ledger_admits_on_its_exact_live_apple_host: the strict route ledger was sealed on macOS 26.6.2 and is rightly inadmissible on 27.0; needs two independentbenchmark_legacy_retune.pyreports on this host +seal_strict_route_ledger.py.test_apple_gpu_slice::test_slice_f16_matches_numpyandtest_apple_gpu_control_flow_f16::test_f16_cond_native_selects_branch: macOS 27 f16 behaviour (MPSGraphsliceTensoron Float16 returns bf16-rounded values; f16condbranch off by up to 10 %). Not caused by any runtime edit; recorded in apple/todo.md with reproductions.Missing exact-device evidence
/usr/lib/llvm-23/binand no NVIDIA package test flipped from run to skip.🤖 Generated with Claude Code