Rigel：逆向工程Apple M4 Max GPU上的Metal 4.1张量计算路径

# 逆向工程 Apple M4 Max GPU 上的 Metal 4.1 张量计算路径  
来源：https://arxiv.org/html/2606.12765（2026年6月）  

###### 摘要  

Apple 的 Metal 4.1 暴露了一条张量计算路径：Metal Performance Primitives (MPP) 的 `matmul2d` 操作，它通过 `cooperative_tensor` 片段处理两个张量操作数。其*接口*有文档说明，但*硬件行为被有意隐藏*。规范说明了支持哪些数据类型行，但从未说明它们是否经过硬件加速、操作在物理上于何处执行、其累加器宽度是多少，或者如何在线程间划分矩阵片段。我们提出了 Rigel，一种在单个 Apple M4 Max（预神经加速器一代）上对该路径的经验性特征描述。通过一个校验和门控、来源追踪的微基准测试框架，Rigel 恢复了 v4.1 规范隐藏或矛盾的 11 项事实。最重要发现：Metal 4.1 的 fp8 (E4M3) `matmul2d` 是*模拟*的，而非加速的：尽管读取的操作数字节数只有一半，其吞吐量仅为 fp16 的 0.94 倍，因此对于 M4 而言，这是一个内存占用功能，而非性能功能。我们进一步通过三信号三角测量（吞吐量上限、与 `simdgroup_matrix` 的比较、以及每轨功耗归因）证明，`matmul2d` 完全在 GPU 着色器核心上执行，没有专用的矩阵数据路径，也没有证据表明路由到 Apple Neural Engine；它累加在 ≥ fp32；我们重建了 Apple 在任何地方都未文档化的透明 8×8 `cooperative_tensor` 片段布局。基于该特征描述，一个手动融合的 GEMM + bias + GELU 内核在缓存驻留区域比分解路径提升了 +6.5–12.9%。所有发现均可从已提交的 MIT 许可代码和每个单元的 CSV 中重现。  

## 1 引言  

量化和注意力密集的机器学习工作负载越来越多地在消费级 Apple Silicon GPU 上运行。Metal 4.1（随 Xcode 27 / macOS 27.0 发布）增加了一条张量计算路径（Metal Performance Primitives (MPP) 的 `matmul2d` 操作，它通过设备分布的 `cooperative_tensor` 片段收缩两个张量操作数），以及低精度前沿的 fp8、fp4 和块缩放 MXFP4 格式。对于任何在 Mac 上部署 LLM 或视觉模型的人来说，三个问题决定了一切：张量操作*在哪里*执行，每种低精度格式*有多快*，以及它提供*什么*数值保证。Metal Shading Language 规范[1 (https://arxiv.org/html/2606.12765#bib.bib13)] 对这些一个都没有回答。问题：规范记录了*接口*但隐藏了*硬件行为*。其特性表说明某种数据类型行是*支持*的；它们从未说明是否*加速*。它声明 `cooperative_tensor` 布局是“不透明的”和“设备特定的”（§2.22.3.1）。它对 `matmul2d` 的分派目标、累加器宽度以及内核必须满足的微架构对齐约束保持沉默。仅凭规范，开发人员无法判断 fp8 是否会比 fp16 快两倍，还是根本不会更快。差距：没有公开资料描述特定 Apple GPU 上的 Metal 4.1 张量路径。工作负载级研究端到端地基准测试 Apple Silicon[4 (https://arxiv.org/html/2606.12765#bib.bib10),5 (https://arxiv.org/html/2606.12765#bib.bib11)]，但没有打开张量原语本身。社区在一个负载关键问题上存在分歧：Apple 的唯一“张量核心”是传统的 `simdgroup_matrix` 指令（一种 ALU 利用率优化），还是类似于 NVIDIA 张量核心[8 (https://arxiv.org/html/2606.12765#bib.bib9),12 (https://arxiv.org/html/2606.12765#bib.bib8)] 的专用矩阵单元。M4 Max 处于 Apple 为 A19/M5 系列报告的 GPU“神经加速器”*之前*。洞察：这些隐藏的事实是*可以通过经验恢复的*，采用剖析 NVIDIA GPU[7 (https://arxiv.org/html/2606.12765#bib.bib6),6 (https://arxiv.org/html/2606.12765#bib.bib7)] 的微基准测试传统，并且在预神经加速器芯片上，强烈的先验是 fp8/fp4 行在功能上受支持但*被模拟*。我们采用一个发现标准：只有当结果相对于规范是*隐藏*或*矛盾*的，才算数；一个细心读者可以从规范推导出的事实是对照，而非贡献。贡献：  

1. 一个校验和门控、来源追踪的框架（第3节 (https://arxiv.org/html/2606.12765#S3)），将每个断言转化为一个可重现的单元，带有 float64 参考和一个预先注册的廉价基线证伪门。  
2. **标题发现**：在 M4 Max 上，fp8 (E4M3) `matmul2d` 是*模拟的*（0.94× fp16 吞吐量），使其成为占用功能，而非性能功能（第5节 (https://arxiv.org/html/2606.12765#S5)）。  
3. 三信号分派目标结果（第4节 (https://arxiv.org/html/2606.12765#S4)）：`matmul2d` 通过传统的 `simdgroup_matrix` 路径在 GPU 着色器核心上运行，没有专用矩阵单元，也没有 ANE 路由的证据。  
4. 重建不透明的 8×8 `cooperative_tensor` 片段布局（第7节 (https://arxiv.org/html/2606.12765#S7)），以及累加器宽度（≥ fp32）和 fp8/fp4 数值语义。  
5. 一个**矛盾的**版本门：Metal 4.1 需要 Xcode 27 *和* macOS 27.0（而非规范的“Xcode 26.1+”），并且 4.1 二进制文件拒绝在 macOS 26.5 上加载（第2节 (https://arxiv.org/html/2606.12765#S2)）。  
6. 11 项隐藏/矛盾的发现作为一个清晰可读的目录（表1 (https://arxiv.org/html/2606.12765#S1.T1)），是低可见性约束（128 字节对齐、SFINAE 雷区、版本门）的自然归属。  
7. 一项优化研究（第8节 (https://arxiv.org/html/2606.12765#S8)）：一个手动融合的 GEMM+bias+GELU 内核在缓存驻留区域比分解路径提升了 +6.5–12.9%（当 O(n³) 计算占主导时，优势衰减），而融合注意力研究表明该优势不扩展到放弃矩阵单元的操作。  

表 1：满足我们发现标准的 11 项发现：每项要么被 Metal 4.1 规范隐藏，要么直接与其矛盾，并且无法仅从文档推导。八项是隐藏的 (H)，三项是矛盾的 (C)。对照项，如符合 OCP 的 fp8 和 fp4 网格，被排除在外。“规范依据”列说明了规范对每一点的说法或缺乏说法。  

## 2 背景  

本节为读者提供理解结果所需的三件事：Metal 4.1 张量路径究竟是什么、控制对其访问的版本门、以及我们测量的设备。  

**Metal 4.1 张量路径：** 本文中心原语是 Metal Performance Primitives 操作 `matmul2d`，它将两个矩阵相乘得到结果 C = A·B。其操作数是 `tensor` 对象，Metal 提供两种相关风格。`tensor_inline` 是对普通 `device` 缓冲区的一种轻量级、非拥有视图，在着色器内部创建并仅描述一个连续区域；而宿主绑定的 `tensor_handle` 则在 CPU 上构造并可携带显式跨度。目标还可以声明为 `cooperative_tensor`，这是一个寄存器片段，由执行范围内的线程共同持有，永不写入设备内存，其内部布局规范有意保持不透明。每次调用由描述符 `matmul2d_descriptor(m, n, k)` 塑造，该描述符固定单个线程组产生的输出瓦片。将 `dynamic_extent` 作为收缩长度 k 传递，告诉操作在内部循环该维度，而非期望调用者驱动它。执行范围写为 `execution_simdgroups`，然后绑定 N 个各含 32 个线程的 SIMD 组协作计算一个瓦片。这些少数类型，以及规范省略的硬件行为，是本文的唯一主题。  

**版本门：** 在测量任何张量特性之前，使用它的二进制文件必须编译*并*运行，而这里规范已经不准。它声明 Metal 4.1 从 Xcode 26.1 起可用，但我们发现访问受限于工具链和操作系统。Xcode 26.5 和 26.6 候选发布版仍然发出 Metal 4.0；只有 Xcode 27 产生语言版本标记 `__METAL_VERSION__ 410`。在 Xcode 27 下以 4.1 版本编译的库随后拒绝在 macOS 26.5 上加载，运行时失败并显示“语言版本 4.1 在此操作系统上不受支持”。因此，使用 fp8 和 fp4 路径需要 Xcode 27 和 macOS 27.0 一起，而非文档记录的 Xcode 26.1。这是表1 (https://arxiv.org/html/2606.12765#S1.T1) 中目录化的矛盾类发现的第一个。  

**被测设备：** 所有测量在单个 Apple M4 Max（40 个 GPU 核心，64 GB 统一内存，Mac Studio）上运行，macOS 27.0，构建 26A5353q。M4 代早于 Apple 随 A19 和 M5 部件引入的 GPU“神经加速器”，因此入职期望是一颗功能上支持低精度格式但未加速它们的芯片；结果证实了这一点。每个数字报告为“在 M4 Max 上”，我们不声称适用于其他 Apple GPU，并将跨代行为视为未来工作。  

## 3 方法  

特征描述的可信度取决于其防止自欺欺人的防御能力，因此在呈现任何结果之前，我们描述 Rigel 如何在一个测量中获得信心。框架特意跨两种语言分割。轻量级的 Objective-C++ 宿主只做分配缓冲区、分派 Metal 内核和计时的事情，而依赖轻的 Python 层拥有所有需要检查的内容：测量单元的架构、每数据类型的误差容忍度、正确性门和来源记录。一个数字在通过以下四个规范后才能进入本文。  

**校验和门：** 关于任何 GPU 内核的第一个问题是它是否计算了正确答案，而低精度使这个问题变得微妙。fp8 结果与双精度参考出现差异有两个原因（*输入*被量化，而*累加*损失了精度），但只有第二个能告诉我们关于硬件的任何信息。因此，我们在 float64 中从 GPU 消耗的完全相同的量化字节计算参考，因此输入量化误差抵消，残差单独界定了内核的累加误差。任何误差超过其每数据类型容忍度的单元被标记为「垃圾」并从所有平均值中丢弃，而不是悄然混合。作为独立的交叉检查，一个单独的 Python 重新验证器直接从已提交的 CSV 重新推导每个通过或失败判定；针对真实的框架输出运行，它在所有 50 个单元上一致同意，零差错。  

**来源：** 为了结果在几个月后仍然可引用，它必须是可重建的，因此在运行开始前，框架记录一个配置哈希、代码提交、输入生成配方的哈希、随机种子以及环境的完整捕获（设备、GPU 核心数、操作系统构建和工具链版本）。缺少其中任何一项的运行不被使用。论文中的每张图随后通过一个不需要 GPU 的 `make reproduce` 目标从其已提交的 CSV 重新生成，因此读者可以从我们的数据重新绘制我们的图表。  

**廉价基线门：** 任何“张量路径很快”故事中的核心风险是速度来自平凡的事情（更好的循环，或者仅仅是普通 ALU）而非特殊硬件。我们通过拒绝做出性能声明来防范这一点，直到最简单的非机制解释实际上已被运行并且未能解释效果。四个基线代表这些平凡解释：(A) 标量 ALU 屋顶线，(B) 一个不使用张量原语的朴素线程组瓦片化 GEMM，(C) 传统的 `simdgroup_matrix` 指令，以及 (D) 标准 fp16 MPP 路径。阈值预先固定：内核必须击败最强基线至少 1.10 倍才能算作胜利，而低精度格式必须达到其匹配 fp16 吞吐量的 1.5 倍才能算作*加速*而非仅仅*模拟*。  

**报告：** 每个吞吐量单元在丢弃预热运行后在至少三十次重复上测量（此处 R=300），我们报告中位数以及自举的 95% 置信区间（种子固定，10⁴ 次重采样）；单个线程组的“是否运行”计时从不报告为吞吐量。置信区间并非套话。如下一段所示，M4 GPU 时钟在小问题规模上是双峰的，而那里区间的宽度本身就是发现。  

**双峰性本身值得描述。** 每次重复的吞吐量在缓存驻留大小处分裂（图1 (https://arxiv.org/html/2606.12765#S3.F1)）。在 512³ 处，R=300 个样本分裂为两个几乎相等的簇，约 4.65 和约 6.5 TFLOP/s（分别为 145 vs. 155 次重复），产生一个宽的自举 CI [5.21, 5.96] TFLOP/s；1024³ 运行稳定在约 12.7 TFLOP/s，带有一个罕见的低时钟尾（8/300 次重复约 6.4 TFLOP/s）；2048³ 保持稳定最大时钟（14.80 TFLOP/s，CI [14.79, 14.80]）。机制是时钟状态：短内核在 GPU 时钟斜坡附近结束，并落入低或高 P 状态，而长内核维持高状态。这个时钟状态悬崖是为什么论文中的每个吞吐量声明都携带自举 CI 而非裸中位数：在 512³ 处的点估计会默默地报告 4.65 或 6.5 TFLOP/s，取决于运气。  

参阅图注  
图 1：时钟状态双峰性。(a) 中位数 fp16 吞吐量及自举 95% CI：区间在 512³ 处宽，一旦时钟稳定则极薄。(b) 每次重复的分布 (R=300)：512³ 是双峰的（两个时钟状态），较大尺寸是单峰的。  

## 4 `matmul2d` 在哪里执行？  

分派目标是负载关键问题：如果 `matmul2d` 降级到现有 GPU ALU/simdgroup_matrix 流水线，那么没有专用硅片就无法加速 fp8 行。我们通过三个独立信号进行三角测量，这个标准是规范自身发现问题所要求的。  

**信号 1，吞吐量上限：** 一个经校验和验证的瓦片化 fp16 `matmul2d` 在 2048³ 处维持 14.8 TFLOP/s（图2(a) (https://arxiv.org/html/2606.12765#S4.F2.sf1)），约为标量 ALU fp16 屋顶线 ~32 TFLOP/s 的 46%，从未超过它，因此没有隐藏的数据路径击败 ALU 上限。  

**信号 2，比较：** 一个四路 GEMM 对决（图4 (https://arxiv.org/html/2606.12765#S5.F4)）显示 MPP 比朴素线程组瓦片化 GEMM 快 2.9–5.5 倍（原语重要），但仅比 `simdgroup_matrix` 快 1.05–1.21 倍。这种接近度与 MPP 降级到 `simdgroup_matrix` 和 FP32-ALU 路径而非独立单元一致性。  

**信号 3，功耗归因：** 在持续 `matmul2d` 负载下，GPU 轨功耗约 42 W，而 GPU 硬件活动驻留率达到 100%（图2(b) (https://arxiv.org/html/2606.12765#S4.F2.sf2)）；通过一个无根 IOReport“能源模型”读取器读取，GPU 能源轨从 0.27 上升到 47.6 W。*警告*：在此 macOS 27.0 beta 上，`powermetrics` 和 IOReport 表都不测量 CPU/ANE 功耗（CPU 在 84% 驻留时仍读取 0mW），并且没有 CoreML 工作负载会在 ANE 上执行（视觉塔 ANE 编译失败），因此 ANE 排除一半的论证是间接的（它依赖于信号 2：ANE 卸载会*增加*超出 GPU 单独路径提供的吞吐量，而它没有）。  

综上所述，这三个信号设定