🤖 AI Summary
This work presents the first empirical investigation into the execution mechanism of fp8 matmul2d operations in Apple Metal 4.1 on the M4 Max GPU, revealing that they are implemented via software emulation on shader cores rather than dedicated hardware acceleration, utilizing an undocumented 8×8 tensor tiling layout. Through a meticulously designed microbenchmarking framework incorporating checksum-based validation, provenance tracking, throughput ceiling analysis, comparison against simdgroup_matrix primitives, and power attribution, we demonstrate that fp8 achieves only 0.94× the performance of fp16, confirming its primary benefit lies in memory footprint reduction. Leveraging these insights, we develop a hand-optimized fused kernel that attains performance gains of 6.5–12.9% in cache-resident scenarios.
📝 Abstract
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.