揭秘苹果M4 GPU张量计算路径，发现fp8运算实为软件模拟
Rigel: Reverse-Engineering the Metal 4.1 Tensor Compute Path on the Apple M4 Max GPU
- 通过微基准测试逆向分析Metal 4.1张量路径的硬件行为
- fp8运算吞吐仅为fp16的94%，证实其为内存优化而非性能加速
- 揭示了8x8矩阵片段布局及全在GPU着色器执行的机制，适合底层优化开发者
Apple的Metal 4.1引入了张量计算路径：基于cooperative_tensor fragments的MPP matmul2d操作。其接口虽有文档，但硬件行为被刻意隐藏。规范仅说明支持的数据类型行，未提是否硬件加速、执行位置、累加器宽度或分块方式。本文提出Rigel，对单颗Apple M4 Max（非神经引擎版）进行实证表征。借助校验和门控、溯源追踪的微基准测试框架，Rigel还原出11条规格未披露或矛盾的事实。核心发现：Metal 4.1 fp8（E4M3）matmul2d为软件模拟，吞吐量仅为fp16的0.94倍，尽管操作数字节仅一半，因此在M4上是内存占用优化，非性能提升。通过三信号三角验证（吞吐上限、与simdgroup_matrix对比、每轨道功耗归因），进一步证明matmul2d完全在GPU着色器核心执行，无专用矩阵路径，亦无神经引擎路由证据；累加器宽度≥fp32；并重建了苹果未公开的8x8 cooperative_tensor片段布局。基于此，手融合GEMM+bias+GELU内核在缓存驻留场景下比分解路径快6.5%-12.9%。所有结果均可由开源MIT许可代码与逐单元CSV复现。
原文摘要 · Abstract (English)
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.
Thank you to arXiv for use of its open access interoperability. PaperDance 不是 arXiv 官方产品；中文卡片由大模型生成，请以原文为准。