同一块 M 系 GPU，虚拟机里跑 LLM 比裸机慢 7–16 倍，只因自报能力老两代

如果你在 Apple Silicon 的 Mac 上开了一个 macOS 虚拟机（VM），然后在里面用 llama.cpp 跑本地模型，很可能会撞到一个困惑：同一个模型，在宿主上飞快，进到 VM 里却慢得不可理喻。慢到什么程度？2026 年 8 月，Cua 的工程团队发布了一组可复现的对照实验，把这个问题钉在了台面上：同一块 M1 Ultra GPU，虚拟机里的吞吐可以比裸机慢 7 到 23 倍。

这篇文章要讲的不是"虚拟化很慢所以认命"，而是一条精确的因果链：虚拟 GPU 向应用自报的能力太保守，让 llama.cpp 在每个决策点上都做出了合法却糟糕的选择。而解锁的方法，也不是把能力往高了乱报。

同一块 M 系 GPU，虚拟机里跑 LLM 凭什么慢 7–16 倍？

这一节只回答一个问题：虚拟 GPU 的能力自报，是如何让 llama.cpp 选上慢内核的。能力自报，就是虚拟设备向 guest 里的应用回答"你支持什么功能"的那套约定；llama.cpp 的 Metal 后端正是拿着这份回答，去决定矩阵乘法跑哪个内核。Cua 工程团队 2026 年 8 月发布的对照实验，把这条因果链的每一环都钉在了台面上，下面就用它当证据，一环一环拆。

先说这条链的第一环证据：慢，是真实存在的。实验条件很简单，一台 M1 Ultra 宿主跑着 macOS 26.6.1，上面开一个 Tahoe（macOS 26.5.2）虚拟机，用同一个 llama.cpp 版本跑同一个模型。TinyLlama 1.1B 的预填充阶段，stock 虚拟机是 431.86 tok/s，裸机是 4,871.99 tok/s，差了 11 倍多；逐 token 生成（decode）是 12.63 对 286.71 tok/s，差了 20 多倍（约 23 倍）。Gemma 4 12B 温和一点，预填充 71.66 对 517.88 tok/s，7.2 倍。Cua/Lume

先排除最容易背锅的解释：虚拟化开销。多数人的第一反应是 VM 多一层翻译，GPU 算力被吃掉一块，所以慢。这个直觉在这个场景里站不住，因为 macOS 虚拟机里根本没有"算力被吃掉一块"的机制。

Apple Silicon 上没有 x86 世界那种把物理 GPU 整块直通（passthrough）给虚拟机的路径，取而代之的是一条半虚拟化（paravirtualization）栈：guest 里跑着 Apple 自带的 Metal 驱动，host 这边的 ParavirtualizedGraphics 框架把 guest 发来的 Metal 命令翻译到物理 GPU 上执行。换句话说，workload 从头到尾跑在宿主那块 GPU 上，虚拟机只是多了一层协议翻译和一个调度时间片。所谓"没直通所以慢"，在 macOS 里本来就不成立，因为直通这条路压根不存在。ParavirtualizedGraphics

半虚拟化排掉了"算力被吃掉一块"这个解释，变量出在哪？看下面这张图。

图注：图上三层全部标着 UNCHANGED：Host Apple GPU、Virtualization.framework 桥、Guest 半虚拟化 GPU，唯一分叉发生在最下方的 capability query（单进程内）。左路是 stock 答案（Apple 5 · 32 KB，SIMD-group matrix 不可用，走慢内核），右路是 shim 的 tested profile（Apple 9 · 64 KB，走快内核）。图中没有任何性能数字，它说明的是一件事：性能变量不在桥上，在"答了什么能力"里。Cua/Lume

同一块物理 GPU、同一条桥，吞吐能差出 7–16 倍，这个量级不可能是翻译、调度开销能解释的。剩下的变量只有一个：虚拟设备向应用自报的能力。guest 里的应用会问 Metal 设备"你支持什么？"，而回答这个问题的，是宿主按自己的策略暴露出来的能力视图。这套"设备声明自己会做什么"的约定，就叫能力契约（capability contract）。

那么问题来了：它到底报了张什么样的牌？

保守的自报从哪来：虚拟 GPU 把能力藏到了哪一档？

先把"谁在报数"说清楚。Apple 的 Virtualization.framework 提供一个配置类 VZMacGraphicsDeviceConfiguration，挂在虚拟机的图形设备上；真正干活的是 host 端的 ParavirtualizedGraphics 框架，它创建 PGDevice 对象，由 PGDevice 决定 guest 看到一个什么样的 Metal 设备。WWDC22 的官方表述是"exposes the GPU capabilities to the virtual Mac"，把话说得非常明白：guest 看到的能力，是宿主主动暴露的视图，不是硬件事实。默认情况下，这个视图是刻意保守的。

保守到什么程度？需要引入 Metal 的世代概念。Metal 把 Apple 的 GPU 按功能分成家族（family），从 Apple 2 一路排到 Apple 10，每一代对应一批硬件，也对应一批能力。应用通过 MTLDevice 的 supportsFamily(_:) 方法查询"你是否支持某个家族的功能集"。Apple 官方文档给出的硬件对应关系是这样的：Metal Feature Sets

这张表是第一层：官方的硬件 family 映射，回答"某颗物理芯片属于哪一代"。照它看，M1 Ultra 的物理 GPU 属于 Apple 7。第二层是应用实际拿到的答案，Cua 实验里有两个版本，别混：stock 虚拟机中 llama.cpp 查到 Apple 5，同一块硬件只被允许看到低两代的牌面，镜头停在 A12 那个年代，这正是性能悬崖的源头；解锁时 shim 报给应用的 1009（Apple 9），是测试过的 profile，不代表 M1 的物理世代，M1 物理上始终是 Apple 7。

family 之外还有第二个被压低的数字：线程组内存（threadgroup memory）上限，一个内核工作组可以使用的片上共享内存。第三层事实是它的可运行性：Apple 官方文档里 Apple 5 到 Apple 10 的文档上限统一是 32 KB，这是规格表里的安全数字；shim 实测则证明，在本实验组合（M1 Ultra 宿主 + Tahoe guest）的这个 workload 下，上报 64 KB、调用 64 KB 线程组内存的内核能正常跑通。注意这句话的边界：64 KB 是"该 profile 在此工作负载下已验证可运行"，不是"这片硬件给得起 64 KB"的普遍事实，换 workload、换机器都得重新验证。默认虚拟设备报的是那个保守的文档值 32 KB。

下面这张图把"能力盖帽"画出来了。

图注：Apple 1→9 能力阶梯（文章自标 illustrative，示意用，不是 Apple 官方规格表）。左侧 Apple 5 标着 STOCK VM REPORTS HERE，右侧 Apple 9 标着 TESTED PROFILE；右侧大括号框住 Apple 5→9 之间的路径，标着 PATHS BEHIND THE CAP：Metal 3/4、SIMD-group matrix multiply、bfloat16、64 KB 线程组内存（相对 32 KB）。Cua/Lume

所以到现在我们知道了两件事：虚拟 GPU 的自报不是硬件事实，而是宿主策略；默认策略把能力压到了 Apple 5 加 32 KB。但这里有个关键的悬念：这张保守的牌，具体是怎么一步一步变成"慢内核"的？

llama.cpp 拿到保守答案，为什么"合理地"选错了内核？

答案在 llama.cpp 的 Metal 后端里。llama.cpp 不直接比较推理速度，而是先查能力，再把能力翻译成内核选择。注意这里的机制：同一份矩阵乘法代码，在不同能力声明下会落到不同的 GPU 内核（kernel）实现上，这个选择过程叫内核分发（kernel dispatch），llama.cpp 里的翻译码表落在 ggml-metal-device.m 的 ggml_metal_device_init 函数中，核心逻辑就几行：GGML Metal

dev->props.has_simdgroup_reduction = [dev->mtl_device supportsFamily:MTLGPUFamilyApple7];
dev->props.has_simdgroup_reduction |= [dev->mtl_device supportsFamily:MTLGPUFamilyMetal3_GGML];

dev->props.has_simdgroup_mm = [dev->mtl_device supportsFamily:MTLGPUFamilyApple7];

dev->props.has_bfloat  = [dev->mtl_device supportsFamily:MTLGPUFamilyMetal3_GGML];
dev->props.has_bfloat |= [dev->mtl_device supportsFamily:MTLGPUFamilyApple6];

三个关键开关的判定门槛并不一样：SIMD 矩阵乘法（has_simdgroup_mm）严格绑定 Apple 7；SIMD 归约（has_simdgroup_reduction）是 Apple 7 或 Metal 3；bfloat16 更宽松，Apple 6 或 Metal 3 即可。判定链走到这里，stock 虚拟机里那个 Apple 5 的答案让三个开关全部落空。把这条链画出来，就是下面这样：

这些开关落空意味着什么？意味着矩阵乘法永远走上一条绕开 SIMD 矩阵指令的慢路径。所谓 SIMD 矩阵指令（SIMD-group matrix），是一条指令直接算一个矩阵块的加速指令集，Apple 7 才引入。没有它，内核只能退回去一条条算，吞吐掉一个数量级不奇怪。

但这里还有一层，决定了"谁替这个错误买单"。llama.cpp 的 mat-vec 与 mat-mat 两条内核路径之间有个运行时闸门 ggml_metal_op_mul_mat_use_mm：只有当输入宽度 ne00 ≥ 64、且 batch 里 token 数 ne11 > 8 时，才走矩阵乘矩阵（mat-mat）的快路径。预填充（prefill）阶段，模型读入整段 prompt、产生第一个 token，token 数远超 8，走 mat-mat，恰好撞上被关掉的 SIMD 矩阵路径；而逐 token 生成时 ne11 永远等于 1，任何情况下都走 mat-vec，不经过这道闸门。所以这条 SIMD mat-mat 门槛能解释的，是它为何直接影响 prefill。反过来也要承认：表中的 decode 提升（TinyLlama 逐 token 12.63 → 206.60 tok/s，Gemma 4 12B 3.41 → 49.67 tok/s）不能由这条机制推出"无影响"的结论，它恰恰说明 decode 的收益来自别处，须另按线程组内存从 32 KB 抬到 64 KB、或其他已改声明逐项分析。这一点必须分清，否则很容易把倍数误套到整个推理过程。

还有一个值得知道的历史注脚：这个 Apple 7 门槛不是数学必然，而是人拍板的。2023 年 10 月 llama.cpp 的 PR #3524 把 mat-mat 硬绑定到 Apple 7 以上，原因为当时的实测显示矩阵乘矩阵在批次较小时反而更慢，而"分界点依赖不可查询的硬件细节"。换句话说，维护者选了保守的门槛，而虚拟设备又给了保守的答案，两头的保守撞在一起。

绕了一圈你会发现，llama.cpp 问的每一个问题都是 Apple 文档推荐的正规查询方式，答案保守，它就"合理地"落到了慢内核上。这不是 bug，而是能力契约的必然产物。那么预言就来了：如果答案改回高的，吞吐应该能回来。真的吗？

只改两个能力声明，预填充真的能追平裸机吗？

真的。这就是 Cua 实验的核心结果。他们做的不是改内核、不是换驱动、不是改虚拟化层，只改了"答给应用的两个数字"：把 family 报到对应用报为 1009 的测试 profile，把线程组内存上限报到 64 KB。这里要说明白：报 1009 不等于把 M1 认定为 Apple 9 硬件，M1 的物理 GPU 世代始终是 Apple 7；这只是证明在（宿主 × guest × 应用）这个组合里，这组声明能被虚拟设备兑现、确实可用。验证原语是一个 capability probe，同一个二进制跑两次：不注入时报 supportsFamily:1009 为 false、线程组内存 32,768 bytes；注入后 1009 变 true、线程组内存变 65,536 bytes。就这么两个声明的翻转。Cua/Lume

效果如下：

Gemma 4 12B 的预填充从 71.66 冲到 515.76 tok/s，达到裸机的 99.59%，几乎追平。TinyLlama 的预填充也回到裸机的 98.25%。全程仍然跑在 Virtualization.framework 那条半虚拟化桥上，没有物理直通，没有内核改动。因果链到这里被钉死了：变的只是两个声明，吞吐差就消失了。

必须诚实标注范围：这是单组合验证，只覆盖 M1 Ultra 宿主加 Tahoe 26.5.2 guest 加 Lume 0.5.1 这一套；TinyLlama 和 Gemma 4 的数字是 10 次采样的中位数，Muse Glimmer 是 3 次，且 Muse Glimmer 没有裸机对照行。另外 TinyLlama、Gemma 4 用的是 llama.cpp b10167，Muse Glimmer 用的是 b10359，跨表对比要谨慎。M2、M3、M4、M5 和其他 guest 版本都得自己再测。Cua/Lume

再把镜头转到裸机外侧，补一个交叉验证。BaseRT 是 2026 年 7 月挂到 arXiv 的原生 Metal 推理运行时，在 M4 Pro 裸机上对比 framework-based runtime：MoE 模型的预填充里，BaseRT 领先 llama.cpp 最高 1.81 倍（Qwen3-30B-A3B），Gemma 4 26B-A4B 上也是 1.59 倍。BaseRT

图注：M4 Pro 裸机对照。左图是 decode 吞吐，三家 runtime 的对比柱状图，小模型上 BaseRT 领先 llama.cpp 最多；右图是 Qwen3-30B-A3B（MoE）的 prefill 吞吐随 prompt 长度变化的曲线，pp128 处标着领先标注。右图数据专门针对 MoE 模型，不要外推到 dense 模型。BaseRT

这意味着 llama.cpp 的 GEMM 内核即使在裸机上也没到平台上限，还有一大截余量，而虚拟机的差是能力声明造成的。硅片本身更不缺能力。Apple Silicon 是统一内存架构（UMA）：CPU、GPU、神经引擎共享同一片内存池，没有 GPU 内部独立的显存。M2 Ultra 的 UMA 有约 800 GB/s 带宽、192 GB 容量。看这张示意图就明白了。

图注：Apple M 系列 SoC 示意：CPU、GPU、神经引擎（ANE）三个计算块共享同一个 Unified Memory 内存池，没有独立的 GPU VRAM 块。这是 UMA 的物理底座，也是"硬件不缺能力"这句话的根据。AppleSi Profile

于是"硬件本可更快，瓶颈在契约与内核选择"得到了双侧支撑。但这引出一个危险的想法：既然抬高声明这么爽，那把声明无限抬高、报出宿主全部能力，是不是就能解锁一切？AppleSi Profile

报得越高越好？MLX 的平局与 residency set 悬崖

第一个反直觉信号来自另一个主流运行时：MLX。在同一个 shim 下，MLX-LM 0.31.3 跑 Llama-3.2-3B 的结果几乎纹丝不动：Cua/Lume

为什么同一个 shim 对 llama.cpp 是性能开关、对 MLX 是中性？因为 MLX 的矩阵内核根本没有"Apple 7 守门人"这一环。llama.cpp 严格走 supportsFamily(Apple7) → has_simdgroup_mm 这条链，MLX 的 matmul 和 steel GEMM 却完全不查这个开关，它用的是自己的 shared-memory tile 路径。stock 虚拟设备报的 family 对它的内核选择本来就没有影响，抬高声明自然也换不来速度。MLX Residency

平局还只是"没赚到"。真正的悬崖在反方向：把声明往高了报，会踩到 residency set。residency set 是 Metal 3 引入的资源常驻登记机制，MLX 用它把常用缓冲区钉在显存里。它的启用判定在 mlx/backend/metal/resident.cpp 里，连写两道闸门：

if (!d->supportsFamily(MTL::GPUFamilyMetal3)) {
    return; // enabled_ stays false，以下全部 no-op
}
if (__builtin_available(macOS 15, iOS 18, *)) {
    enabled_ = true;
    // 尝试构造 residency set，失败直接抛 runtime_error
}

把这两道闸门画出来，能看到"冒报"是怎么引爆的：

你的 shim 一旦把 Metal 3 family 也报到高位，在像 MLX 这样会据此请求 residency set 的路径上，就可能触发兑现失败并崩溃：虚拟设备兑现不了这个请求，"[metal::Device] Unable to construct residency set" 被直接抛出，性能修复在一瞬间变成正确性事故。这是这条路径上已被实证的具体风险，而非所有冒报的必然结果。这段历史在 MLX 仓库里留有完整记录：PR #1537（2024 年 10 月）修的就是虚拟机里这第一次崩溃，加上了 supportsFamily(GPUFamilyMetal3) 这条短路；issue #1855 指出 Apple 官方规格里 residency set 的真门槛其实是 Apple family 6，Metal 3 是个更宽的集合，但 main 分支至今没有采纳这个修正；issue #2197 又暴露了第二条路径：PyPI wheel 默认 deployment target 15.0 加 Metal 框架静态链接，会让 __builtin_available 的运行时检查在编译期消失。三条记录指向同一个教训：声明宽于兑现能力，就一定会崩在某处。

这也是为什么 Cua 的 shim 把可改范围钉死在 Apple family 枚举（1001–1999），Metal family 保持 stock 不动，宁可少解锁一点，也不去碰那条会崩的边界。MLX Residency

所以"报得越高越好"不成立：同一个能力契约，报低了锁死性能，报高了触发兑现事故。那到底该怎么解锁，才既快又不崩？

收口：既别怪虚拟化，也别伪装高端 GPU，按应用验证可兑现的能力子集

先把视角拉远。这个"虚拟设备向 guest 报什么能力"的问题，在 GPU 虚拟化里已经争论了十年。AvA 是 2020 年 ASPLOS 上的一篇经典论文，它把加速器虚拟化的设计空间拆成三个互相拉扯的属性：interposition（中间层能拦截多少操作）、compatibility（接口能否扛住演进）、isolation（租户之间隔不隔离）。guest 看到的能力从来是一张被宿主筛选过的牌，关键就在这个筛子怎么筛。AvA

工业界最有名的实例是英特尔的 GVT-g（mediated pass-through）：把性能关键的帧缓冲和命令缓冲直通给 guest，把安全关键的 I/O 寄存器和页表操作 trap 回来由宿主模拟。同一个设备里，哪些能力"直接给你"，哪些能力"必须经过我"，是宿主逐资源定义的策略。下面的架构图就是它的样子。

图注：GVT-g 高层架构（Project ACRN HLD Figure 244）：Device Model 通过 Create vGPU 把 vGPU 注入 guest，guest 经 ACRN 触发 IO Trap 回到 Device Model，两者之间的 Emulation services 与直通共同构成分级能力报告。ACRN HLD

Apple 的 ParavirtualizedGraphics 与 GVT-g 实现层完全不同，思路却同构：host 决定暴露哪些能力，guest 在这张被筛选的牌面上做决策。而这套框架下有两个对称的失败模式，本专题的悬崖恰好在两个方向各踩过一次：隐报，默认报 Apple 5 加 32 KB，锁死了 SIMD 矩阵快内核；冒报，把 Metal 3 family 报到高位，让 MLX 请求兑现不了的 residency set。

正解既不是"认命，怪虚拟化"，也不是"伪装高端 GPU"。正解是 capability profile：以（应用 × 芯片 × 客体系统 × API）为单位，逐个应用实测哪些能力能被虚拟设备兑现，把验证过的能力子集写进声明。这才是把"Cua 的 shim 为什么只动这两个声明"讲清楚的方式：它不替用户做硬件发现，只报应用验证过的 profile。Lume Guide

落实到操作，Lume shim 的开关是三个环境变量：

注入方式是 per-process 的 DYLD_INSERT_LIBRARIES，只影响你启动的那个 workload；硬化或系统保护的二进制会拒绝注入，所以这要求"你掌控的进程"。host 侧还有一条路径：defaults write com.apple.gpusw.ParavirtualizedGraphics ForceUnrestrictedDeviceFeatureLevel -bool true，把虚拟设备的 feature level 限制抬起来，注意这是一个 Apple 官方文档里没有记载的偏好，版本敏感，而且必须在 VM 停止时写、重启后才生效。每一步都有独立回滚：进程级下次启动去掉注入即可；用 LaunchAgent 跑的用 launchctl bootout 外加 plist 移入废纸篓；host 偏好用 defaults delete 删掉。三条路互不依赖。

回到开头的问题：为什么说"没直通所以慢"和"伪装成高端 GPU 就能解锁"都不成立？因为慢的根本不是缺算力，而是声明让应用选了慢内核；而快的前提也不是报得高，而是报得"真实可兑现"。同一个悬崖，两个方向各有一个错误归因，正确的判断只有一种：能力声明既足以解锁优化，又必须由后端真实兑现。下一次再有人问"VM 里跑模型为什么慢"，先别急着怪虚拟化，去查一下虚拟 GPU 答给应用的那张牌，再对着你的工作负载验证一遍，而不是照着硬件型号想当然。

参考文献

本专题引用的核心资料如下，建议深入学习时阅读原文：

1. Apple Silicon and macOS VMs: 11–16× Faster LLM Inference with llama.cpp（Cua/Lume）： https://cua.ai/blog/gpu-passthrough-macos-vms
2. Metal Feature Set Tables + MTLGPUFamily + supportsFamily(_:)（Apple Developer）： https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf
3. llama.cpp GGML Metal Backend Source（ggml-metal-device.m + PR #3524）： https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-metal/ggml-metal-device.m
4. MLX Metal Backend Residency Set Code（resident.cpp）+ Issues #1537 / #1855 / #2197： https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/resident.cpp
5. Apple ParavirtualizedGraphics + VZMacGraphicsDeviceConfiguration Documentation： https://developer.apple.com/documentation/paravirtualizedgraphics
6. Lume GPU Passthrough: How-to Guide（+ ForceUnrestrictedDeviceFeatureLevel）： https://cua.ai/docs/how-to-guides/lume/gpu-passthrough
7. BaseRT: Best-in-Class LLM Inference on Apple Silicon via Native Metal： https://arxiv.org/abs/2607.00501
8. Profiling Large Language Model Inference on Apple Silicon: A Quantization Perspective： https://arxiv.org/abs/2508.08531
9. AvA: Accelerated Virtualization of Accelerators（ASPLOS 2020）+ GVT-g Mediated Pass-Through（Project ACRN HLD v2.5 / gVirt USENIX ATC 2014）： https://aakshintala.com/papers/asplos20-ava.pdf；GVT-g 架构图：https://projectacrn.github.io/2.5/developer-guides/hld/hld-APL_GVT-g.html