Track papers across labs, topics, and trends.
Search papers, labs, and topics across Lattice.
This paper introduces Rigel, a tool for empirically characterizing the Metal 4.1 tensor compute path on the Apple M4 Max GPU, revealing critical information that Apple's documentation obscures. The authors discovered that the fp8 matmul2d operation is emulated rather than hardware-accelerated, achieving only 0.94x the throughput of fp16 despite a reduced memory footprint. Additionally, they demonstrate that matmul2d executes solely on GPU shader cores without a dedicated matrix datapath, leading to the development of a more efficient GEMM kernel that outperforms the standard implementation by 6.5-12.9%.
The Metal 4.1 fp8 matmul2d operation is a memory feature, not a performance feature, as it is emulated and not hardware-accelerated.
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.