Flint 在 Apple Silicon 上的推理表现：rec Megakernel 与 Tempo9 对比

2026-10-27

English · 中文

Recursive 开源的 rec 是一个针对 Flint 优化的 Apple Silicon 推理引擎。它采用持久运行的 Metal Megakernel，在一次 dispatch 内完成模型的解码计算。对单流解码而言，这是一种非常有效的设计。

我们在同一台 M5 Pro MacBook Pro 上测试了 rec 和通用推理引擎 Tempo9，比较单流 decode、并发、冷 prefill、跨会话 prefix cache 和 CPU 使用情况，并通过逐层 profiling 与 Metal 实验分析性能来源。

这次对比不以评判哪个引擎“更好”为目标，而是想弄清两个问题：rec 为什么能实现如此高的单流解码速度？当多个会话交替运行时，缓存和请求调度会如何影响整体性能？ 测试配置与测量方法见文末。

1. 性能测试：不同场景下各有优势

主要测试均在一台配备 M5 Pro、24 GB 统一内存的 MacBook Pro 上完成。

为了让对比更公平，我们选择了每次生成 token 时权重读取量接近的模型版本：Tempo9 约为 1.385 GB，rec 约为 1.39 GB。

这样可以尽量减少量化格式带来的干扰，让测试更接近对引擎执行效率的比较。

单流解码：rec 快约 26%

在单请求、贪心解码 256 token 的测试中，rec 达到了 151–155 tok/s，而 Tempo9 为 123–127 tok/s。

也就是说，在单流解码场景下，rec 的吞吐量约为 Tempo9 的 1.26 倍。

rec 在单流解码中的优势十分明显，体现了专用 Megakernel 的执行效率。不过，多请求服务还需要看总吞吐和响应时间。

并发推理：Tempo9 的优势逐渐显现

rec 的执行模型面向单个活动会话，不支持将多个独立请求进行连续批处理。当多个客户端同时发起请求时，后续请求需要等待当前请求完成。

Tempo9 则采用 continuous batching，在每一步解码时动态组织正在运行的请求，让多个会话共享 GPU 计算资源。

我们测试了 1、2、4、8、16 个并发客户端，每个客户端请求生成 256 token。

只有一个客户端时，rec 更快。从两个客户端开始，Tempo9 的总吞吐量超过 rec。随着并发数增加，两者的差距进一步扩大。

在 16 路并发下，Tempo9 达到 378 tok/s，而 rec 为 141 tok/s。Tempo9 的合计吞吐约为 rec 的 2.68 倍。

最慢请求的完成时间也明显缩短：Tempo9 为 12.2 秒，rec 为 29.0 秒。

这些结果反映了两种架构在请求调度上的差异。只运行一个会话时，单流生成速度是关键指标；多个会话同时运行时，总吞吐和完成时间也变得重要。

例如，IDE 助手、Agent 和聊天窗口可能同时使用同一块 GPU。此时，引擎能否有效安排多个请求，会直接影响整体体验。

冷 Prefill：不同输入长度下的表现

我们比较了四种输入长度下的冷 prefill 延迟——没有前缀命中。

在 562 token 时，Tempo9 用时 0.37 秒，rec 为 0.76 秒，Tempo9 约快 2.1 倍。在 1,971 token 时，rec 略有优势：0.80 秒对 0.84 秒。

在 8,022 token 时 rec 领先：三轮交替测试中 rec 2.8–3.0 秒，Tempo9 3.7–4.1 秒。此前单独测得的 4.12 秒受内存压力影响，因此本文采用了重新交替测试的数据。在 9,289 token 时顺序又反过来：Tempo9 4.5–4.9 秒，rec 5.3–5.6 秒。rec 按 4,096 token 一块做 prefill，一旦 prompt 需要第三块，耗时陡增；Tempo9 的耗时随长度大致线性增长。

冷 Prefill 只是推理延迟的一部分。多轮对话中，能否复用已有上下文同样重要，因此我们进一步测试了前缀缓存。

前缀缓存：连续对话与多会话交替

前缀缓存可以避免重复处理已经出现过的上下文。不过，同一个对话连续进行和多个独立对话交替运行是两种不同的使用场景，需要分开评估。

rec 的 PromptCache 会保存当前对话的 token 和 KV 状态，并逐 token 比较新旧前缀；只要新请求沿着原来的对话继续追加，就能充分利用缓存。Tempo9 则采用 128 token 分段的共享前缀缓存。下面先看单会话连续请求，再看切换会话后的情况。

单会话连续请求：rec 的缓存命中更完整

我们先向对话 A 发送约 8,000 token 的文档，再在同一个对话中追加一个问题。这相当于常见的多轮聊天或 Agent 持续追加工具结果的场景。

rec 在单会话下的缓存效果非常好。第二轮请求复用了全部 8,022 个已有 token，只需处理新追加的 40 个 token。相比之下，Tempo9 以 128 token 为单位缓存，末尾不足一段的部分需要重新计算，因此复用了 7,936 个 token。这组测试中，rec 的命中更完整，Prefill 也更快。

rec 还会在请求结束后通过 prepare_prefix() 将生成的 token 纳入缓存，适合上下文不断增长的连续对话。对于修改历史内容的分叉请求，它会尝试回退到已有 checkpoint；这与下面的跨会话切换问题不同。

多会话交替：切换后能否复用原来的缓存？

再考虑同一台服务器交替处理两段独立对话的情况。请求按 A → A → B → A 的顺序发送：先开启 A、继续 A，再开启使用另一份文档的 B，最后回到 A。每组测试重复三轮，图中为引擎报告的 Prefill 耗时。

第二张图重点展示从 B 返回 A 时的 Prefill 耗时。

返回 A 时，rec 的单份 PromptCache 已被 B 替换，因此需要重新处理 A 的上下文，耗时 2.86–2.88 秒，与冷 Prefill 接近。Tempo9 仍能复用 A 的 7,936 个 token，耗时 0.28–0.31 秒。

rec 的单会话前缀缓存表现很好，限制主要在跨会话复用。 在测试版本中，切换对话会替换当前缓存，因此无法同时保留多个独立会话的前缀。主 Agent 与子 Agent 交替执行、或者用户同时打开多个聊天窗口，都会遇到这种缓存管理需求。

另需说明：rec 在第三轮实验的最后一个请求中出现了 Metal 内存不足错误；相关 9.3k token Prefill 在这台 24 GB 机器上的峰值内存为 23.2 GB。

CPU 占用：两种等待与调度方式

CPU 使用情况也存在明显差异。

在解码期间，Tempo9 大约占用 0.15–0.17 个 CPU 核，而 rec 占用 0.87–1.00 个 CPU 核。

通过采样，我们发现 rec 的主要 CPU 开销来自 MLX stream 线程在 Fence::wait 中轮询等待 GPU 完成。这是 rec 默认的 --fast-sync（MLX_METAL_FAST_SYNCH=1）行为，用 CPU 时间换更低的 CPU–GPU 同步延迟；我们没有改动它。

Tempo9 采用了不同的实现方式。它的推理引擎由 C++ 实现，生成的 token 通过无锁队列传递给服务端，通过 command buffer 的完成机制协调 GPU 执行，而不是依靠持续自旋轮询。

CPU 编码下一步解码计算所需的时间约为每 token 1.3–1.5 ms，并与 GPU 当前步骤的执行重叠。

对于本地 AI Agent，这一点尤为重要。CPU 不只需要运行推理引擎，还承担工具执行、文件处理、浏览器操作以及其他应用任务。推理引擎占用更少的 CPU 资源，就意味着可以为这些任务留下更多执行空间。

2. rec 的单流解码为什么更快？

为了定位性能差距，我们结合 rec 的开源实现和逐层性能分析，观察 Metal 上的计算和同步方式。

传统 GPU 推理引擎通常会把一个 Transformer layer 拆分成多个操作，例如 RMSNorm、QKV projection、RoPE、attention、output projection、MoE routing 和 expert computation。这些操作通常由一个或多个 GPU kernel 实现，再按照计算图的顺序执行。

对于熟悉 CUDA 的读者，Metal 的 SIMD group 大致对应 CUDA 的 warp，threadgroup 对应 thread block，threadgroup memory 对应 shared memory；threadgroup_barrier() 类似 CUDA 的 __syncthreads()，只能同步同一线程组内的线程。Metal 的一次 compute dispatch 可类比一次 CUDA kernel launch。

rec 在 M5 Pro 上启动 36 个 threadgroup，每组 256 个线程；按 CUDA 术语理解，就是 36 个 block、每个 block 256 个线程。由于普通组内 barrier 无法同步不同 threadgroup，rec 使用 device-scope atomic 和自旋等待协调跨组执行。这要求参与同步的工作组能同时驻留，不能假设任意 grid 都适用。Apple GPU 核心与 CUDA SM、Metal command buffer 与 CUDA stream 都不是严格的一一对应关系。

下图展示了 M5 Pro 上的执行层级：从统一内存到 GPU，再到 SIMD group 和 threadgroup。不同层级的同步范围，是理解 rec 设计的关键。

理解 rec 的实现，需要先了解 Metal 的线程组织与同步方式。一次 compute dispatch 会启动多个 threadgroup，每个 threadgroup 由若干 SIMD group 组成。同组线程可以共享 threadgroup memory，并通过 threadgroup_barrier() 同步；但这个接口无法让不同 threadgroup 彼此等待。

rec 的 Megakernel 需要在执行过程中协调多个 threadgroup，因此使用 device-scope 原子计数器和自旋等待来实现跨组同步。这样做要求参与同步的 threadgroup 能同时驻留在 GPU 上，而 Metal 并未提供适用于任意 dispatch 的全局同步保证。rec 也因此针对不同 Apple 芯片配置相应的工作组数量。

相关机制可参考 Apple 的 Threads and Threadgroups、Metal Shading Language Specification 和 WWDC 的 Apple GPU 计算优化讲解。

具体到 rec，它在 M5 Pro 上启动 36 个常驻 threadgroup，每个包含 256 个线程。这些 threadgroup 在一次 Metal dispatch 中共同完成 Flint 的全部 49 层 Transformer 计算。

不同 threadgroup 之间通过 device-scope atomic counter 和自旋 barrier 进行同步。

下图概括了 rec 如何在单次 dispatch 内执行一个解码步骤。多个常驻 threadgroup 通过同步机制推进计算，并在内部完成大量融合操作。

这种设计的主要收益是让一个 threadgroup 连续处理多个相关操作。例如，一个负责特定 attention head 的 threadgroup，可以在同一个执行上下文中完成 projection、normalization 和 KV 写入，同时尽可能保留寄存器和 threadgroup memory 中的中间结果。

这样就能减少不同 kernel 之间反复读写和重新组织数据的开销。由于 rec 只针对 Flint 进行优化，因此可以围绕 Flint 固定的模型结构设计执行路径，实现通用计算图不容易表达的深度融合。

这次 dispatch 覆盖整层计算，而不只是 attention：输入 RMSNorm、attention 之后的 RMSNorm、MoE router、top-6 选择、专家的 gate/up 与 down 投影（第 0、1 层是稠密 MLP，kernel 里为它留了一个分支），以及最后一层之后的 final norm、lm_head 和 argmax。其中两处 RMSNorm 和 top-6 不需要额外的跨组同步：每个 threadgroup 的寄存器里都放着完整的 2048 维残差（每线程 8 个值），自己把 norm 重算一遍——每个 threadgroup 多做几千次乘加，换来不用等另外 35 个。专家按小组执行，gate/up 到 down 的交接是组内的，不是全网格栅栏。剩下的栅栏都在数据真正跨 threadgroup 的地方：qkv 之后、attention 部分结果之后、合并之后、o_proj 之后、router 之后、专家之后。

逐层分析：差距主要来自 Attention

为了找出 rec 单流解码更快的原因，我们将每层 Transformer 的耗时分为 Attention 和 MoE 两部分，并分别测量。

在 Attention 阶段，rec 每层约需 44 µs，Tempo9 约为 63 µs；到了 MoE 阶段，Tempo9 反而略快（94–100 µs，对比 rec 的约 109 µs）。

这表明，rec 的性能优势主要来自 Attention 路径，而不是专家计算。其 Megakernel 将归一化、RoPE、KV 写入等操作与相邻计算紧密融合，在 threadgroup 内复用寄存器和中间数据，减少了独立 kernel 之间的数据搬运与同步。

对于 49 层的 Flint，这些细小的逐层差异会累积成明显的单 token 延迟差距。需要注意，图中的 rec 逐阶段耗时来自可进行 profiling 的版本；文章开头的 151–155 tok/s 则来自进一步优化过的执行路径，两组数据不能直接逐项相加。下面我们再通过独立实验，考察减少 dispatch 本身究竟能带来多少收益。

3. Megakernel 的优势：计算融合比减少 Dispatch 更重要

直觉上，rec 更快的原因似乎很简单：一次 GPU dispatch 总应该比很多次 dispatch 更高效。但 dispatch overhead 究竟占了多少？

为了回答这个问题，我们构建了独立的 Metal 实验，使用相同的 device functions 实现两条 attention 执行路径：一条在同一 command buffer 中使用七次 kernel dispatch，另一条在 persistent kernel 内通过 atomic barrier 同步。

两种主要执行方式的耗时十分接近。

在这颗 M5 Pro 上，一次跨 threadgroup barrier 的开销约为 1.0–1.5 µs，而一次 kernel dispatch boundary 约为 1.4 µs。

把多个 kernel 合并成一次 dispatch，并不会自动让计算更快。

关键在于融合后减少了多少中间计算、数据读写和同步。rec 的优势来自计算阶段内部更高效的数据复用、更少的中间内存操作，以及多个原本独立的计算步骤之间更紧密的融合。

这些收益也可以通过局部算子融合获得，并不一定需要把整套模型放入一个 Megakernel。

4. Tempo9 为什么选择通用执行架构？

当模型架构、执行形状和目标硬件都提前确定时，Megakernel 是一种自然的设计选择。但 Tempo9 的目标不同：我们希望构建支持多种模型架构、量化格式、硬件后端和动态并发 workload 的通用推理引擎。

Tempo9 不只面向 Apple Silicon，也支持 NVIDIA CUDA GPU 和纯 ARM CPU 等不同计算设备，并通过统一的模型执行框架支持 Qwen、Gemma、Llama、Mistral、Flint 等多种架构。引擎核心采用原生 C++ 实现；在 CUDA 后端，通过 CUDA Graph 尽可能将推理执行路径组织为可复用的图，减少主机调度和重复提交开销。这里的 CUDA Graph 是 CUDA 后端的执行机制，并不意味着 Metal 或 ARM CPU 使用相同的 GPU 图接口。

这样的选择是在性能与软件复杂度之间寻找平衡：高度专用化的 Megakernel 可以针对某个模型和设备做更深的融合，而共享的 C++ 执行框架、可复用算子及 CUDA Graph 路径，则更便于在不同模型和硬件之间复用实现。它并非没有性能代价，但有助于避免为每一种模型、设备和请求形状维护一套独立的 Megakernel。

下图展示了 Tempo9 的执行方式：调度器每一步确定活跃请求，再将动态 batch 交给一系列算子执行。与 rec 相比，核心区别是动态请求调度与固定工作分配。

动态批处理

真实的 Agent workload 往往动态变化。新的请求随时加入，已有请求完成生成，也可能等待工具执行而暂停；有的请求正在 prefill，有的正在 decode。

Tempo9 在每个解码步骤重新组织 batch，并根据执行形状选择合适的 GPU kernel。较小的 batch 可以采用 GEMV，较大的 batch 可以切换到 tiled GEMM 或 grouped GEMM。

rec 的 Megakernel 则主要针对固定执行结构优化；不支持 continuous batching 是这种设计取向的体现。

不同模型共享底层 kernel

Tempo9 通过可复用的计算 kernel 支持 Qwen、Gemma、Llama、Mistral 和 Flint 等模型家族。模型架构定义 operator graph，底层 kernel 实现其中的计算操作。

因此，支持新架构时，通常无需从头编写完整的 GPU 推理 kernel。Megakernel 能实现更深的模型专用融合，但需要更多针对具体模型和硬件的实现工作。

支持会改变执行流程的推理功能

Prefix caching、speculative decoding、KV cache 管理和动态调度都可能改变推理过程的组织方式。保留 operator boundary，可以灵活组合这些能力，而无需为每种执行变化构建新 Megakernel。

这种灵活性也有代价：在 batch size 为 1 时，Tempo9 的通用执行路径存在额外开销。但它也带来了更好的并发支持、模型覆盖范围和服务能力。

两种路线并不互斥。Megakernel 中有效的局部融合技术，同样可以应用于通用执行架构。

5. 测试环境与方法

硬件和软件

• 硬件：MacBook Pro，M5 Pro，24 GB 统一内存
• 操作系统：macOS 26
• Tempo9：1.2.0，引擎提交 754721a8f
• rec：upstream main 的 0dcf43c（2026 年 10 月 22 日）加两处小的本地修改，已公开在 thinkspread/rec 的 local/flint-e3-24gb 分支。4eacf0a 去掉启动时的 48 GiB 内存检查、接受 INT3 专家 checkpoint、并用 mx.set_wired_limit 锁定权重——没有它 rec 在 24 GB Mac 上无法启动；325d244 把服务端的准入信号量从 9 提到 32，让并发测试的 16 个客户端都能连上。两处都不涉及 kernel、调度器或 prefix cache。
• 测试时间：2026 年 10 月 23–26 日

量化配置与权重读取量

主要速度对比使用每 token 权重读取量接近的模型版本，而不是简单比较量化名称。

Tempo9 发布版与等字节性能测试版使用不同量化配置，不能直接混为一谈。模型质量另以 Q8_0 参考实现评估。

测试负载

单流解码使用短 prompt、贪心生成 256 token，最终上下文长度约 300 token。并发测试使用不同 prompt，每个客户端生成 256 token。

测量方法

所有主要对比在同一机器上进行。测试期间停止本地 CI runner，并在同一会话内交替测试配置，至少重复三轮。

单流吞吐使用引擎报告的数据，并发吞吐在客户端测量。我们也监控 macOS 内存状态，因为压缩内存压力可能明显影响绝对性能。

这些控制不能消除所有误差，但可以降低单次测试、环境变化和量化差异对结果的影响。

6. 在 Tempo9 上运行 Flint

通过 Homebrew 安装：

brew install thinkspread/tap/tempo9

或者运行安装脚本：

curl -fsSL https://raw.githubusercontent.com/thinkspread/tempo9/main/install.sh | bash

启动 Flint：

tempo9 --hf thinkspread/Flint-Q6K-IQ3XXS-GGUF

服务默认监听 localhost:11435，提供兼容 OpenAI 和 Anthropic 的 API。也可以加载本地 GGUF 文件：

tempo9 --gguf <file>

7. 总结：单 Megakernel 是一个很好的开始

rec 展示了在 Apple Silicon 上针对单一模型深度优化的潜力。通过持久运行的 Metal Megakernel，它在我们测试的 M5 Pro 上达到 151–155 tok/s 的单流解码速度，在接近相同的每 token 权重读取量下约为 Tempo9 的 1.26 倍。逐层 profiling 也说明，针对 attention 路径进行细致的计算融合与数据复用，确实能够带来明显收益。

在 Metal 上将一个模型的解码过程组织成单个 kernel，是一个很好的技术起点。 它提供了一种值得探索的执行方式，也让我们更清楚地看到模型专用优化可以做到什么程度。不过，单次 dispatch 本身并不是全部：在我们的实验中，跨 threadgroup barrier 与 kernel dispatch boundary 的成本接近，真正值得关注的是融合后减少了哪些计算和数据搬运。

与此同时，多请求并发和 prefix cache 也应当是推理系统设计中同样重要的考虑因素。当前测试版本的 rec 不支持 continuous batching；多个客户端的总吞吐无法随着并发数像动态批处理引擎那样提升。rec 对单条线性增长的会话做了出色的缓存优化：A 的第二轮复用了全部 8,022 个 token，只需 prefill 新增的 40 个 token，耗时 0.13–0.14 秒，比 Tempo9 的 0.22 秒还快。它的限制在于跨会话保留：两个独立对话按照 A→B→A 交替时，当前 rec 实现会用 B 覆盖 A 的 PromptCache，返回 A 时需要重新 prefill。这是当前 PromptCache 实现的限制，而不是单 Megakernel 在原理上无法支持这些能力。

对于主要运行一个 Flint 会话、强调持续 decode 速度的场景，rec 的设计很有吸引力。对于多个 Agent、交互式会话以及工具调用频繁交替的工作负载，调度能力和跨会话 KV cache 复用则会显著影响端到端体验。

因此，单 Megakernel、连续批处理和多会话 prefix cache 并不是互相排斥的方向。rec 已经在单流 GPU 执行效率上迈出了很有价值的一步；如何在此基础上兼顾并发与缓存复用，是同样值得继续探索的问题。

感谢 Recursive 以 Apache-2.0 许可证开源 rec。完整的实现让这类架构讨论有了可验证的技术依据，也让不同推理引擎之间的比较能够超越简单的吞吐数字。