BaseRT：通过原生Metal在Apple Silicon上实现最佳LLM推理

# 基于原生 Metal 在 Apple Silicon 上实现最佳 LLM 推理
来源：https://arxiv.org/html/2607.00501
Prabod Rathnayaka、Fabian Waschkowski、Lukas Wesemann，Base Compute，澳大利亚墨尔本

###### 摘要

我们提出了 BaseRT，一个面向 Apple Silicon 上大型语言模型（LLM）的原生 Metal 推理运行时，并报告了迄今为止该硬件上的最高推理吞吐量。现有的运行时，包括 llama.cpp 和基于 MLX 的框架，由于使用的抽象层并非针对 Metal 执行模型或 Apple Silicon 的统一内存拓扑设计，因此产生了开销。通过基于 Metal 原生构建，结合芯片特定内核融合、统一内存感知优化和自定义调度逻辑，BaseRT 恢复了基于框架的方法所浪费的性能。BaseRT 支持多种模型系列，涵盖八种量化格式（Q2 至 FP16），适用于所有 Apple M 系列设备。在本文中，我们使用 M3 和 M4 Pro 设备，在 Q4 和 Q8 量化级别对 Qwen3、Llama 3.2 和 Gemma 4 系列进行了评估。BaseRT 的解码吞吐量相比 llama.cpp 最高提升 1.56×，相比 MLX 最高提升 1.35×，在混合专家模型的前缀填充方面优势更大，为从小于 1B 到 30B 参数的模型提供了一致的最高吞吐量。这些结果表明 Apple Silicon 是一个比先前报告更强大的推理平台，对新兴的边缘推理范式具有直接影响：随着隐私要求、延迟限制和云成本压力推动推理向设备端部署转移，性能优化的本地运行时是实现这一转变的关键使能层。BaseRT 已在 https://github.com/basecompute/baseRT 公开提供。

## 1 引言

Qwen3-0.6B Llama-3.2-1B Llama-3.2-3B
0 200 200 400 400 +35% +35%
模型 解码吞吐量 (tok/s)
128 256 512 1024 2048 400 400 600 600 800 800 +78% +78%
提示长度 (tokens) 前缀填充吞吐量 (tok/s)
BaseRT MLX llama.cpp

图 1：BaseRT 在 Apple Silicon 上提供了最佳吞吐量。在 Apple M4 Pro 上测试的 Q4 模型中，解码吞吐量（左）相比 llama.cpp 提升 1.15–1.56 倍，相比 MLX 提升高达 1.35 倍。对于前缀填充吞吐量（右），在 4 比特量化的 Qwen3-30B-A3B 模型上，BaseRT 在所有提示长度下均领先，相比 MLX 提升高达 1.78 倍。

AI 推理需求的增长速度对集中式云基础设施的可扩展性提出了挑战。主要提供商的 Token 使用量每年大约增长一个数量级，这从 Google Cloud API 的 Token 使用量以及 Open Router 上更广泛的提供商组合可以证明[1, 20]。然而，大多数企业采用仍处于早期阶段，这一趋势很可能会持续[25]。全球推理算力预计将在 2027 年首次超过训练算力，到 2030 年 AI 推理预计将占全球数据中心算力的 40% [4]。这一趋势给纯云推理的经济和架构假设带来了越来越大的压力，并为向边缘推理的结构性转变提供了有力论据：在消费和专业硬件上本地运行 LLM。支持这一转变的论据基于四个独立的驱动力。

**隐私与数据驻留**：在云中处理的提示词面临多租户、第三方基础设施风险以及复杂的法律管辖权问题[17]，Gartner 预计到 2027 年，由于对跨境提示路由监督不足，超过 40% 的 AI 相关数据泄露将由这些风险导致[7]。本地推理将所有数据保留在受控环境中，从而满足隔离和零信任要求。

**延迟**：云推理增加了请求处理、排队和 GPU 调度开销，尤其在负载下，可能将首个 Token 生成时间（TTFT）推至数百毫秒[12]。设备端推理避免了网络和排队延迟，通常足以将延迟降低到支持亚百毫秒的交互式响应，并使实时机器人[10]和响应式智能体工作流[14, 21]成为可能。

**连接性**：对外部网络可用性的硬性依赖意味着，网络、认证或共享云基础设施的故障可能使 AI 系统完全不可用，无论底层模型是否正常运行。本地推理消除了这种依赖，在低连接性或隔离环境中实现确定性运行，并确保现场延迟敏感应用程序的连续性。最近的云服务中断事件突显了这种潜在影响[2, 19]。

**成本**：云推理的成本与使用量线性相关，而本地推理则将成本转化为固定的硬件成本，在规模上每次推理的边际成本趋近于零[15]。

实现这一转变的硬件前提条件已基本满足。在消费和专业设备类别中，专为 AI 设计的加速器现在支持有用的设备端 LLM 推理：高通的 Snapdragon X Elite 包含 Hexagon NPU，能够在设备上运行超过 130 亿参数的 LLM；而英特尔的 Core Ultra 系列 2 和 AMD 的 Ryzen AI 平台也集成了专用 NPU，OpenVINO 和 ONNX Runtime 等运行时提供了在这些设备上执行量化模型的路径[22, 11, 18]。

在这些平台中，Apple Silicon 因其高内存带宽、大统一内存容量和成熟的 GPU 计算栈而脱颖而出[6]。内存历来是本地 LLM 推理的主要限制因素：模型权重必须为每个生成的 Token 从内存流式传输到计算单元，这对哪些模型能够在本地运行设置了硬性容量和带宽限制。Apple 的统一内存架构消除了 CPU 和 GPU 内存池之间的传统隔离，使 GPU 能够以高带宽直接访问整个系统内存，这一结构性变化使得在本地运行大型模型变得可行。现代 M 系列设备能够以交互速度运行量化的 7B–70B 参数模型，而最近的开源权重模型已将最先进的能力压缩到适合本地部署的参数数量中[13]。最新分析还表明，本地运行的模型在处理常见查询时，其准确率与云模型相当，同时能耗显著降低[24]。

本文的工作重点是专门在 Apple Silicon 上最大化推理吞吐量，其统一内存架构和 Metal GPU API 提供了其他平台上不存在的独特优化机会。尽管硬件潜力巨大，但现有的 Apple Silicon 推理运行时各有局限，限制了可实现的吞吐量。我们提出了 BaseRT，一个直接构建在 Apple Metal GPU API 之上的 LLM 推理运行时，无需任何中间框架。通过实现芯片特定的内核融合、统一内存感知数据布局以及自定义 GPU 调度逻辑，我们恢复了基于框架的方法所浪费的性能。在一系列开源 LLM 上，BaseRT 实现了比 Apple Silicon 上任何先前报告运行时都更高的吞吐量，为在该硬件上进行设备端推理建立了新的性能基线。

### 1.1 贡献

我们的贡献如下：

- 一个专为 Apple Silicon 统一内存架构设计的原生 Metal LLM 推理运行时，无需依赖 MLX 或任何中间框架。
- 芯片特定的内核融合策略，消除了基于框架的运行时中存在的调度开销。
- 一种统一内存优化方案，利用了 Apple Silicon 共享的 CPU-GPU 内存拓扑。
- 在多个开源 LLM 上进行的全面基准评估，在 Apple Silicon 上建立了新的最先进吞吐量数据，并表明先前的运行时在系统性地未充分利用该硬件。

## 2 相关工作

Apple Silicon 上主流的开源运行时各自存在架构约束，限制了它们对该拓扑的充分利用。llama.cpp[8] 框架最初是作为跨平台、以 CPU 为先的代码库开发的，后来扩展了 Metal 后端。其架构反映了这些起源。MLX[9] 虽然是专门为 Apple Silicon 构建的，但它是一个通用数组框架，其惰性求值和框架级调度引入了专用推理运行时中不存在的开销[9, 3]。基于 MLX 构建的方法，包括 vllm-mlx，已在这些约束下取得了显著成果，报告称相比 llama.cpp 吞吐量提高了 21-87%[5]。

最近，出现了绕过这些框架的原生 Metal 方法。MetalRT[23]（RunAnywhere）直接针对 Metal，无中间抽象层。开源运行时 uzu（Mirai）采用了混合方法，将工作负载分配到 GPU 内核和 MPSGraph 上，以访问 GPU 和 Apple 神经引擎[16]。两者都表明，框架抽象层会带来可测量的性能代价。在 BaseRT 中，我们沿着这一方向进一步推进：通过芯片特定的内核融合、统一内存感知优化和自定义调度逻辑，我们实现了超过这些先前原生方法的吞吐量。

## 3 BaseRT 运行时设计

BaseRT 是一个直接针对 Apple Metal GPU API 的 C++ 推理运行时，不依赖 MLX、PyTorch、CoreML 或任何中间数组框架。该运行时暴露了 C API 以实现跨语言互操作性，并支持基于 Transformer 的解码器模型，使用 BaseRT 的原生 *base* 权重格式。设计围绕四个原则组织，如下所述，每个原则都基于在现有运行时中观察到的特定开销来源。此外，BaseRT 包含两项直接针对解码热路径上 CPU 端开销的创新：一种低开销的命令调度机制，消除了基于框架的运行时中每个操作符的调度、分配和图评估工作；以及一种解码调度方案，将 CPU-GPU 同步开销分摊到多个生成的 Token 上。

### 3.1 数据驱动的架构描述符

在基于框架的运行时中，支持新的模型架构通常需要修改推理循环、添加条件分支并连接新的算子序列。这种耦合使核心引擎变得脆弱且难以统一优化。BaseRT 将所有特定于架构的变化捕获为数据：架构不同的点，例如激活和归一化变体、混合专家细节以及注意力和位置编码约定，都通过紧凑的架构描述符表示，而不是编码为控制流。核心引擎使用此描述符，并且永远不会根据架构身份进行分支，因此所有模型的热路径都是相同的。BaseRT 目前支持 LLaMA、Qwen3、Gemma、Whisper 和 BERT 系列，扩展到新架构只需进行一次独立的声明性更改，而无需修改推理循环。

### 3.2 零分配解码循环

模型加载完成后，解码循环分配零字节。前向传递使用的所有中间缓冲区，即残差、注意力和前馈暂存缓冲区，以及 logits 和 Token 缓冲区，都在加载时预分配并跨每个 Token 重用。KV 缓存也以最大化注意力期间 GPU 合并内存访问的布局，在加载时预分配到最大上下文长度。即使是错误处理也使用静态的线程局部缓冲区。结果是热路径完全由 GPU 命令分发和轻量级 CPU 端修补组成，没有分配器争用或内存管理开销。

### 3.3 内核融合与专业化

BaseRT 实现了一个大型的手写 Metal 着色器库，涵盖矩阵乘法、注意力、归一化、RoPE、嵌入、激活和采样类别。矩阵乘法内核占大多数，为每种支持的量化格式（2 位到 16 位宽度）提供了专用的 GEMV（解码，M=1）和 GEMM（前缀填充，M>1）变体。每个内核将反量化直接集成到内循环中，而不是将反量化后的权重物化到全局内存中，从而按压缩比例减少内存流量。

除了每个内核的优化之外，BaseRT 还融合了在其他运行时中通常作为单独内核分发的算子序列，包括注意力路径和前馈块内。每个融合减少了一次内核启动和一次全局内存往返，这在解码期间尤其重要，因为单个内核运行时间短，启动开销占比较大。

内核选择是硬件自适应的：基于 GPU 核心数量和系列（M1 到 M5）得出的每芯片配置决定了线程组大小、内核选择阈值和融合决策。每种量化格式都映射到其专用的内核和启动几何形状。

### 3.4 前缀填充

提示词处理（前缀填充）使用 GEMM 而非 GEMV，在形状为 [seq_len × dim] 的激活张量上运行。算术强度随序列长度缩放，使前缀填充成为计算密集型而非内存密集型。BaseRT 使用 Metal 的 simdgroup_matrix 内联函数实现分块 GEMM 内核，其分块几何形状针对序列长度范围进行了调整，并应用分块前缀填充，使用有界最大分块大小来限制暂存缓冲区内存。FlashAttention 内核使用在线 softmax 和运行统计信息以分块方式计算注意力，实现 O(n) 内存复杂度，而不是物化完整的 O(n²) 注意力矩阵。

## 4 评估

### 4.1 实验设置

BaseRT 支持广泛的模型系列、八种量化格式（Q2、Q3、Q4、Q5、Q6、Q8、BF16 和 FP16）以及所有 Apple M 系列（M1 到 M5）。为保持评估的可控性，我们研究了具有代表性的子集：在 M3 和 M4 Pro 设备上，以 Q4 和 Q8 量化级别测试了 Qwen3、Llama 3.2 和 Gemma 4 系列。

我们将 BaseRT 与三个基线进行比较：llama.cpp[8]，这是 Apple Silicon 上最广泛部署的开源推理运行时，使用 GGUF 格式权重；MLX[9]，Apple 专门