A few people have asked me what the Rigel paper is actually about, without the tables and without the appendix full of CSV checksums. This is my honest attempt to explain it.
The one-sentence version: Apple shipped a tensor matmul in Metal 4.1, documented almost nothing about how it behaves on real silicon, and I spent a few weeks measuring what the documentation refuses to say. The answer to the question everyone wants answered first — is the fp8 path fast? — is no. It is emulated. And the reason that matters is not that Apple did something wrong. It is that knowing it stops you from building a kernel around a speedup that does not exist.
The thing Apple gives you, and the thing it doesn't
Metal 4.1 exposes a new primitive: matmul2d, part of the Metal Performance Primitives. You hand it two matrices, it gives you the product, and the documentation tells you the API surface — the types it accepts, the shapes, the toolchain you need. What it does not tell you is the part that decides whether your kernel is good or garbage:
- Is the low-precision (fp8) path running on dedicated matrix hardware, or is it unpacked and run on the same shader ALUs as everything else?
- What precision does it accumulate in? fp16? fp32? Something in between?
- What is the fragment layout — the per-lane ownership pattern of a
cooperative_tensor— that you would need to know to fuse anything around the matmul?
Apple calls the cooperative-tensor partition "device specific" and stops there. That phrase is doing a lot of load-bearing work. "Device specific" means we are not going to tell you, and it may change. For anyone trying to write a fused kernel — the thing that actually wins performance on a memory-bound device — that is exactly the information you need and exactly the information you don't get.
So Rigel is a measurement project. No privileged access, no leaked headers. Just a microbenchmark harness pointed at the M4 Max, asking the hardware questions it will answer even when the docs won't.
The headline: the fp8 path is a memory feature wearing a performance costume
Here is the finding I keep getting asked about.
fp8 operands are one byte. fp16 operands are two. If the matmul were memory-bound — if it spent its time waiting on operand bytes to arrive — then halving the operand size should make fp8 faster. That is the whole pitch of low precision.
It does not. fp8 sustains 0.94× the throughput of fp16 at best, despite reading half the bytes. Measured across problem sizes:
The logic is a one-line proof once you have the number. If fp8 were memory-bound, it would be faster than fp16, because it moves half the bytes. It is slower. Therefore the matmul is compute-bound, and the fp8 variant is paying pure unpack overhead — converting one-byte operands up into the datapath the hardware actually computes in — with no low-precision compute path underneath to reward it.
I pre-registered the bar before I ran anything: fp8 had to clear 1.5× fp16 to count as a genuine acceleration path rather than emulation. This is the same cheap-baseline-falsification discipline I apply to every project — decide what would falsify the exciting interpretation before you have the data, so you cannot rationalize after the fact. fp8 peaked at 0.94×. It did not come close.
The practical translation, which is in the paper almost verbatim: on the M4 Max, the Metal 4.1 fp8 matmul2d is a memory-footprint and dynamic-range feature, not a throughput feature. The speedup waits for M5. The M4 Max predates the "Neural Accelerators" Apple added to the GPU with the A19 and M5 generations, so this is, in retrospect, exactly the prior you should have held. fp8 lets you fit a bigger problem in cache and gives you a wider dynamic range than fp16's exponent allows. It does not make the multiply faster. If your mental model was "switch to fp8, go 2× faster," that model is wrong on this chip, and you would have found out the hard way after building around it.
How do you prove there's no matrix unit?
A throughput number alone is suggestive, not conclusive. Maybe fp8 is slow for some unrelated reason. So the load-bearing methodological idea in Rigel is three-signal triangulation — three independent ways to ask "where does this operation actually execute?", which only converge if the answer is real.
- Throughput ceiling. fp16 matmul2d tops out at 14.8 TFLOP/s — about 46% of the ~32 TFLOP/s scalar-ALU fp16 roofline. It never exceeds the ALU ceiling. If there were a hidden matrix datapath, you would expect to break through that scalar roofline somewhere. It never does.
- Baseline comparison. A four-way GEMM shootout pits Apple's
matmul2dagainst a naive tiled GEMM, againstsimdgroup_matrix, and against stock fp16.matmul2dbeats the naive kernel by 2.9–5.5× — but beatssimdgroup_matrixby only 1.05–1.21×. That proximity is the tell: the new primitive is lowering onto the samesimdgroup_matrix+ FP32-ALU path that has been there all along, not onto a separate unit. - Power attribution. Under sustained matmul load the GPU rail draws ~42 W at full hardware-active residency. There is no second rail lighting up that would betray an accelerator doing the work elsewhere. (Honest caveat, and it is in the paper: the macOS 27.0 beta does not meter the ANE rail directly, so ANE exclusion is indirect — but ANE offload would add throughput, and it doesn't, so the indirect argument holds.)
Three signals, one conclusion: matmul2d runs on the shader cores. There is no dedicated matrix unit on this chip. Any one signal could be explained away. Three independent ones converging cannot.
The two findings I actually care about for writing kernels
The fp8 result is the headline because it is counterintuitive. But the two findings I will use are quieter.
The accumulator is at least fp32 — and I can prove it bit-exactly. I built a deliberately hostile test: a contraction of depth 128 with one large term (448) and 127 tiny ones (0.0625 each). The exact sum is 455.9375. An fp16 accumulator physically cannot represent that — near 448, fp16's spacing is 0.25, larger than the 0.0625 increments, so every small term vanishes and the sum pins at 448 (a 1.74% error). The fp8 matmul2d returned 455.9375, bit-exactly. Only an accumulator of fp32 width or greater preserves that sum. Apple never states this anywhere; it matches the at-least-fp32 convention documented for NVIDIA tensor cores, and now it is measured for Metal too. This is the difference between trusting your reductions and not.
The cooperative-tensor fragment is a unified 8×8 tile. This is the "device specific" layout Apple won't document. I recovered it by querying the partition rather than executing a GEMM — and the structure is clean: a repeating 8×8 tile across the output, each lane owning two vertically adjacent rows of one column per quadrant, and — the part that matters — the A, B, and C fragments share the identical 8×8 swizzle. The partition is unified across operands and output.
Why do I care? Because a unified layout is the lever for fusion. If the output fragment lives in the same per-lane arrangement as the inputs, you can do your bias add and your activation in-register, on the fragment, before you ever write to memory — no reshuffle, no round trip.
Putting the layout to work — and the fusion that didn't pan out
I wrote a fused GEMM + bias + GELU kernel using the recovered layout: matmul, add the bias, apply GELU, store once, all on the fragment in registers. Against the decomposed path — GEMM, then a separate bias pass, then a separate GELU pass, each making its own O(n²) trip through memory — the fused kernel wins:
+12.9% at 1024³, +6.5% at 2048³, in the cache-resident regime. The win is exactly where the theory says it should be: when the problem fits in cache, the epilogue memory traffic is a meaningful fraction of total time, so removing it matters. As problems grow and O(n³) compute swamps the O(n²) epilogue, the relative win shrinks. The mechanism is honest and the scaling story is honest.
And then — because I try to report the things that don't work with the same energy as the things that do — the same fusion principle failed for FlashAttention. Fusing the softmax into the attention matmuls forces those matmuls off the matrix path and onto scalar code, and the result is a 3.6–5× slowdown versus the decomposed path. At the M4's sequence lengths the S×S round trip is cheap on unified memory, and my flash kernels ended up overhead-bound. The lesson generalizes: fusion wins when it removes memory traffic that actually dominates, and loses when it drags compute off a fast path to do it. The fragment layout is a lever, not a magic wand.
The part nobody reads, which is the part that makes it true
Every number above could be a measurement artifact. GPUs ramp clocks, caches warm up, a single lucky run looks like a trend. The reason I trust these numbers is the harness, and the harness is four disciplines stacked on top of each other:
- A checksum gate. Every cell's reference is computed in float64 from the exact quantized bytes the GPU consumed, which cancels input-quantization error and bounds accumulation error alone. A cell that drifts past its per-type tolerance is marked garbage and never reported as a result.
- Provenance on every run. Before each measurement the harness records a config hash, the code commit, the input-generation recipe hash, the seed, and a full environment snapshot — device, GPU core count, OS build, toolchain versions. Every number traces back to the exact state that produced it.
- A cheap-baseline gate. Four mundane explanations — the scalar-ALU roofline, a naive GEMM,
simdgroup_matrix, stock fp16 MPP — have to be beaten by at least 1.10× before any "win" is allowed to be called a win. - Honest reporting. At least 30 repetitions per cell (R=300 in practice), medians with bootstrapped 95% confidence intervals, and single-threadgroup timings never reported as throughput.
That last discipline caught something real. At cache-resident sizes, the per-rep throughput split — 5 vs 6.5 TFLOP/s at 512³ — and the bootstrapped CI blew wide open to [5.21, 5.96]. That is not noise to average away; it is the GPU clock ramping mid-measurement. It is why every claim in the paper carries a confidence interval instead of a bare median. A separate Python re-verifier re-derives every verdict from the committed CSVs, and against the real harness output it agreed on all 50 cells with zero mismatches. Every figure in the paper redraws from a committed CSV through a single make reproduce — no GPU required to regenerate a single plot.
None of that is glamorous. All of it is the difference between "I measured the M4's tensor unit" and "I have a story about the M4's tensor unit." The whole thing is MIT-licensed: the kernels, the per-cell CSVs, the figure scripts, the re-verifier.
What this taught me
Three things I want to write down before they fade.
The most valuable measurement is the one that kills the exciting hypothesis. "fp8 is 2× faster" is the kind of claim that, left untested, sends you down a week of kernel work before you discover the speedup was never there. A one-afternoon throughput sweep with a pre-registered falsification bar saved that week. The cheap baseline is not a formality at the end of a project. It is the thing you run first.
"Device specific" is an invitation, not a wall. Apple declined to document the fragment layout. The layout is still there, and the hardware will tell you what it is if you ask with the right query. A surprising amount of what looks like a black box is a box nobody bothered to point a measurement at.
Negative results are load-bearing. The fp8-is-emulated finding and the FlashAttention-fusion-fails finding are both "no" answers. They are also the two findings most likely to save someone else real time. I would rather publish the no than let the next person rediscover it at their own expense.
Rigel is a small paper about a small primitive on one chip. But the tensor path is going to be where on-device inference lives for the next several years, and right now it is mostly undocumented. This is one corner of it, measured carefully, with the receipts attached.
Resources
- Paper: arxiv.org/abs/2606.12765
- Project page: murailabs.com
- Hardware: Apple M4 Max (40 GPU cores, 64 GB unified), macOS 27.0. MIT licensed.
If you build kernels on Apple silicon and any of these numbers contradict what you're seeing on your own hardware, I genuinely want to hear about it — the CSVs are public precisely so they can be argued with.