Characterizing and optimizing the Metal 4.1 tensor compute path on Apple Silicon GPUs — open the black box Apple ships, then beat its defaults with free, hand-tuned kernels on hardware people already own.
Codename RIGEL: the bright blue supergiant within the constellation Orion. Lineage callback — the prior paper was Orion (the ANE); this is a star inside it (the GPU tensor path). Methodology lineage: Orion (arXiv:2603.06728).
G0 skeleton complete; G1 (Phase 0 capability gate) is the next work — confirm what actually compiles and runs on this M4 Max before designing around it. No experiments have run.
- Act I — Characterize. Metal 4.1 documents the interface of its tensor/MPP path but
deliberately hides the hardware behavior: the
cooperative_tensorlayout is "opaque" and "device specific" (§2.22.3), the SIMD-group-matrix element-to-thread mapping is "unspecified" (§6.8), and acceleration is punted to the external Feature Set Tables. RIGEL recovers what the PDF refuses to state on Apple M4 Max — dispatch target, accumulator widths, the opaque cooperative-tensor fragment layout, the acceleration-vs-emulation dtype truth-table, and the microarchitectural cliffs. - Act II — Beat it. Use those discoveries to ship open-source hand-tuned Metal kernels that outperform the stock MPP path on M4 Max in the regimes characterization exposes.
- The discovery bar (§0.1): if a finding could have been written by reading the PDF, it is not a result; if a kernel isn't faster than the baseline it tuned against, it isn't a contribution.
- Null hypothesis the project must defeat: the MPP tensor path /
cooperative_tensorlayout / 4.1 low-precision rows / hand-tuned kernels outperform the cheap baselines — scalar-ALU roofline, naive threadgroup-tiled GEMM,simdgroup_matrix, and the stock fp16 path — on metric M (median GFLOP/s, J/FLOP). If the baselines match within ε, the performance claim is falsified (still publishable as characterization; no kernel ships against it).
- Paper — Act I characterization + Act II optimization (MLSys-class / arXiv-first, Orion cadence).
rigel-bench— MIT-licensed open microbenchmark suite; emits the CSVs and regenerates every figure.rigel-kernels— MIT-licensed hand-tuned Metal kernel library (the headline contribution): packed low-precision GEMM, layout-aware tiling, cliff-avoiding schedules, fused epilogues, ANE/GPU hybrid router.- Constraint catalog — Orion-format, numbered; headline metric = count of HIDDEN+CONTRADICTED entries.
- Decision guide — which Metal matmul path for which dtype/shape/fusion/energy budget, built on measured facts.
| Path | Purpose |
|---|---|
AGENTS.md / CLAUDE.md |
Operating contract (byte-identical) |
CODEX.md |
Codex pointer to AGENTS.md |
STATUS.md |
Current state |
CHECKPOINT.md |
Exact resume point |
tasks/atomic-task-list.md |
Canonical atomic task list |
spec/RIGEL-spec.md |
The science (design doc) |
docs/ |
Decision log, experiment log, reproducibility, risks, runbook, gate reviews |
notes/ |
Append-only audit trail |
configs/ |
Locked, versioned sweep configs (per phase) |
runs/ |
Run artifacts with full provenance (bodies gitignored) |
harness/ |
Swift/Obj-C++ host orchestration (rigel-bench host; language TBD — spec open-Q #1) |
kernels/ |
MSL kernels: simdgroup_matrix, MPP host-bound, MPP cooperative_tensor, rigel-kernels |
analysis/ |
Python: YAML config → CSV → stats / roofline / plots (the testable layer) |
- Device under test (FIXED): Apple M4 Max, 64 GB unified memory, Mac Studio. Pre-neural- accelerator generation — no GPU "Neural Accelerators" (those arrive on A19/M5). Every finding is reported "on M4 Max"; cross-generation claims are out of scope.
- Toolchain floor: macOS 26.x (latest installed, 26.4+ preferred), Xcode 26.1+ (26.0 had the MPP header bug), Metal Toolchain component installed, Metal 4 project template.
- Data posture: repo is text-free; the archived Metal v4.1 spec PDF, raw measurement bodies, and checkpoints live on the external drive, gitignored. Per-cell CSVs + run manifests are committed (they are the citable results).
- Source of truth: the shipped MPP headers, not the PDF, where they disagree — and every disagreement is logged as a constraint-catalog entry.
RIGEL is released under the MIT License — see LICENSE. This covers the rigel-bench
microbenchmark harness, the Metal kernels, the analysis layer, and the committed per-cell CSVs; the
forthcoming rigel-kernels library ships under the same license.