Rigel: Reverse-Engineering the Metal 4.1 Tensor Compute Path on the Apple M4 Max GPU
Quick Answer
Rigel reveals that Apple's Metal 4.1 fp8 matmul2d is emulated rather than hardware-accelerated, achieving only 0.94x the throughput of fp16.
Quick Take
The findings indicate that this operation runs entirely on GPU shader cores without a dedicated matrix datapath, impacting performance expectations for developers using the M4 Max GPU.
Key Points
- Metal 4.1's fp8 matmul2d sustains 0.94x throughput of fp16 despite lower operand bytes.
- The operation executes entirely on GPU shader cores without dedicated matrix datapath.
- Accumulation occurs in at least fp32, contradicting expectations from the specification.
- A hand-fused GEMM kernel outperforms the decomposed path by 6.5-12.9% in cache-resident scenarios.
- All findings are reproducible from MIT-licensed code and CSVs.
Paper Resources
📖 Reader Mode
~2 min readAbstract:Apple's Metal 4.1 exposes a tensor compute path: the Metal Performance Primitives (MPP) matmul2d operation over cooperative_tensor fragments, whose interface is documented but whose hardware behavior is deliberately hidden. The specification states which data-type rows are supported, never whether they are hardware-accelerated, where the operation physically executes, what its accumulator width is, or how it partitions matrix fragments across threads. We present Rigel, an empirical characterization of this path on a single Apple M4 Max (a pre-neural-accelerator generation). Using a checksum-gated, provenance-tracked microbenchmark harness, Rigel recovers eleven facts the v4.1 specification hides or contradicts. The headline finding: the Metal 4.1 fp8 (E4M3) matmul2d is emulated, not accelerated: it sustains 0.94x the throughput of fp16 despite reading half the operand bytes, so on M4 it is a memory-footprint feature, not a performance feature. We further show, via a three-signal triangulation (throughput ceiling, comparison against simdgroup_matrix, and per-rail power attribution), that matmul2d executes entirely on the GPU shader cores with no dedicated matrix datapath and no evidence of Apple Neural Engine routing; that it accumulates in >=fp32; and we reconstruct the opaque 8x8 cooperative_tensor fragment layout Apple documents nowhere. Acting on the characterization, a hand-fused GEMM + bias + GELU kernel beats the decomposed path by +6.5-12.9% in the cache-resident regime. All findings are reproducible from committed MIT-licensed code and per-cell CSVs.
| Subjects: | Computation and Language (cs.CL); Distributed, Parallel, and Cluster Computing (cs.DC) |
| Cite as: | arXiv:2606.12765 [cs.CL] |
| (or arXiv:2606.12765v1 [cs.CL] for this version) | |
| https://doi.org/10.48550/arXiv.2606.12765 arXiv-issued DOI via DataCite |
Submission history
From: Ramchand Kumaresan [view email]
[v1]
Thu, 11 Jun 2026 00:10:23 UTC (721 KB)
— Originally published at arxiv.org
Want this in your inbox every morning?
Daily brief at your local 8am — bilingual EN/中文, free.
More from arXiv cs.CL
See more →TriAgent: Divergence-Aware Committees for Cost-Efficient Financial Sentiment Analysis
TriAgent introduces a cost-efficient multi-agent system for financial sentiment analysis, combining VADER, FinBERT, and Qwen2.5. It achieves an F1 score of ~0.87 with significant savings of $9.3M/year at a 10M-user scale compared to GPT-4o-mini, while also detecting hallucinations with an AUC of 0.90.