OSDI’26 | 双路 CPU + 两张 5090，本地跑满血 671B DeepSeek：云端 SLO 是怎么在消费级硬件上达成的

原文：Achieving Cloud-Grade SLOs for Local Mixture-of-Experts Inference through CPU–GPU Hybrid Design

---

1. 前言

先交代下背景。想在自己机器上跑 DeepSeek-R1 这种 671B 的满血 MoE 模型，最出圈的方案是清华系的 KTransformers：把 routed expert 放 CPU 内存、CPU 算，attention 等密集部分放 GPU 算。它确实能跑起来，但体验和云端 API 差一截——INT4 量化模型 decode 约 16 tokens/s，8K prompt 的首 token 要等 20 秒，12K 以上直接奔着半分钟去了。

这篇 OSDI’26 的工作（清华 + 芯云集成电路）把目标钉得非常明确：在”双路服务器 CPU + 1-2 张消费级 GPU”这个价位的机器上，把本地推理拉到云端服务的 SLO 水平。它先把”云端水平”量化成四条硬指标，也就是论文识别的四个 gap：

1. 模型完整性：本地方案普遍靠量化/蒸馏/改路由（rerouting）换速度，模型能力打折——目标是跑原始精度的 FP8 满血模型
2. Prefill（TTFT）：云服务的参考线是 30 秒内出首 token；现有本地方案 12K token 以上就超标
3. Decode（TPOT）：云端基线 ≥20 tokens/s；KTransformers 量化后才 16
4. 并发：本地系统 prefill 和 decode 互相踩踏，多一个 decode 请求 TPOT 就明显劣化

结果先放这：32K prompt 30 秒内完成 prefill（双卡 45K）、原始 FP8 精度 21.5 tokens/s decode、并发干扰 <15%。四条全部达标，而且是在满血模型上。下面拆它是怎么做到的。

2. 为什么 CPU 值得被认真对待

先看一张定调的图，理解”本地跑大 MoE”这件事的硬件账：

这张图横轴是内存容量、纵轴是内存带宽（都是对数刻度），散点是各代模型，折线是三类硬件（绿色 GeForce、红色 Tesla 数据中心卡、蓝色 EPYC CPU）。模型点的位置这样解读：横坐标是全部参数的存储需求（决定你要多大内存才装得下），纵坐标是每秒激活权重访问量——按 20 tokens/s 的 SLO，每秒要串行读 20 次激活参数，这就是带宽需求。

看模型的演化轨迹（颜色越亮越新）：总参数一路涨到 1TB 量级（Kimi-K2 已到 1T 参数），但稀疏激活把带宽需求摁在约 1 TB/s——DeepSeek-V3 每 token 只激活 37B/671B 的参数。这个”容量疯涨、带宽稳住”的剪刀差正好撞进 CPU 的舒适区：

• 容量：消费级 GPU 32 GB 封顶，模型点全在它右边够不着；双路 EPYC 插满 DDR5 轻松上 TB（本文平台 1.15 TB）
• 带宽：24 通道 DDR5-6400 理论聚合带宽 1228 GB/s——和 A100 的 HBM 一个量级，正好压住 1 TB/s 的需求线
• 成本：DDR5 约 $3/GB，HBM 约 $30/GB，差 10 倍

所以结论不是”CPU 能替代 GPU”，而是大 MoE 的稀疏激活模式恰好把带宽需求降到了 CPU 内存也够得着的水平，低并发场景下这笔账成立。剩下的问题全是系统问题：CPU 的带宽怎么用满、GPU 的算力怎么喂饱、两边怎么不互相等。

3. Prefill：把 GPU 从”装不下”里解放出来

3.1 先理解为什么 prefill 难

prefill 和 decode 的资源属性完全相反。decode 一次算 1 个 token，是 memory-bound（瓶颈在读权重）；prefill 一次算几千上万个 token，同一份权重被所有 token 复用，是 compute-bound（瓶颈在算力）。KTransformers 让 CPU 算 MoE 的策略在 decode 端合理，在 prefill 端就露馅了：CPU 的矩阵算力比 GPU 差一两个数量级，8K/16K/32K prompt 的 TTFT 分别约 20/40/100 秒——长 prefill 必须上 GPU。

但 GPU 装不下模型（671B 的 FP8 也要 ~700 GB，显存 32 GB）。传统 weight-offloading 方案（权重现用现搬）都是为”边缘小模型 + 低延迟”设计的，搬运和计算的重叠做得粗，在大 MoE 的高强度 prefill 上吞吐上不去。

3.2 Stream-Loading Prefill：三线程流水线

SLP 的思路：prefill 全程在 GPU 上算，权重以子层粒度从 DRAM 流式搬进显存，用完立刻扔——显存里永远只驻留”正在算的那一小片”模型。

对着图走一遍。三个线程（各配一条 CUDA stream）：

• Loader（左，绿）：不停歇地把下一个模块（attention 权重、expert 0、expert 1……）从 DRAM 拷进显存，拷完一个就设置该模块的 ON 事件
• Model（中，蓝）：执行前等 ON 事件（权重到位才算），算完设置 OFF 事件（宣告”这份权重我用完了”）
• Unloader（右，灰）：等到 OFF 事件就释放对应显存

三条流水线错开推进：model 在算第 0 层 attention 时，loader 已经在搬第 0 层的 expert；model 算 expert 0 时，loader 在搬 expert 1……权重传输完全藏在计算后面。每个事件是”CUDA event + 主机线程 event”的成对设计，保证设备侧和主机侧的依赖都严格有序。

两个关键工程细节：

1. Experts ring buffer。DeepSeek-V3 有 58 个 MoE 层 × 256 个 expert × 3 个线性层 = 44.5K 个权重张量，每个都走 cudaMalloc/cudaFree 会被 API 开销淹死。解法是一块跨层复用的手工管理环形缓冲（图中深红框）：expert 权重进环、用完槽位直接复用，不与分配器打交道。环的长度还会自适应——短 prompt 时传输是瓶颈，环开到一层的 expert 数（256）拉满预载深度；prompt 超过约 50K 后计算成瓶颈、中间激活开始抢显存，环缩到 2（ping-pong）保住重叠即可
2. 为什么不用 CUDA UVM/cudaMemPrefetchAsync？统一虚拟内存的换页是隐式的，何时搬、搬多少不可控，长尾停顿会打穿流水线。显式三线程换来的是可预测的显存占用和精确的重叠控制

效果：单卡 RTX 5090 prefill 冲到 1200 tokens/s，32K prompt 的 TTFT 控制在 30 秒内——而且跑的是 FP8 原始精度，不需要量化。

3.3 双卡怎么分：SmallEP 砍一半通信

两张卡加速 prefill，自然想到 context parallelism（CP，把序列切两半）+ expert parallelism（EP，把 expert 切两半）。但本地机器没有 NVLink，卡间通信走 PCIe，这就有讲究了。

先看标准 EP 的问题。看下图（a）：

标准 EP 流程：每张卡对自己那份 token 做 gate 和排序 → All-to-All Dispatch 把 token 发给”拥有对应 expert 的卡” → 各卡算本地 expert → All-to-All Combine 把结果发回原卡 → 加权求和。通信量算笔账：N 个 token、hidden 维度 D、每 token 激活 A 个 expert、EP 大小 S，dispatch 的每链路流量是 N·A·D/S。云端 S=32、64 时这个数被摊薄；本地 S=2，每链路要扛 N·A·D/2——DeepSeek-V3 的 A=8，相当于 4·N·D，PCIe 直接被打爆。这就是”云端方案不能直接搬到本地”的具体形态。

SmallEP 的改法（图 b）：先做一次 All-Gather 让两张卡都拿到全部 token（每链路 N·D），gate 和排序各卡本地重复做，各自只算分给自己的 expert，算完先在本卡做部分加权求和得到一个 [N, D] 的部分和，最后 All-to-All 交换部分和再本地归约。两段通信都是 N·D 量级——当 S ≤ A（本地必然如此，2 ≤ 8），SmallEP 的每链路峰值流量严格更低。代价是 gate/排序算两遍，但 profiling 显示这部分 <10ms/层（<5% 端到端延迟），比省掉的通信便宜得多。

分解图印证这一点：

从上到下：单卡 Baseline 约 280ms（Attn + MoE 串行）；CP2+EP2 降到约 225ms，但橙色（Dispatch）和深紫（Combine）两段通信肉眼可见地厚；CP2+SmallEP2 把通信压掉约一半，整层约 185ms。最终 DSLP 双卡 1800 tokens/s，45K prompt 30 秒内搞定。

TTFT 全景图（注意纵轴对数）：

几个读点：256-2K 的短 prompt 区间 CPU 方案反而最快（SLP 有权重搬运的固定成本，所以系统在 ≤4K 时自动走 CPU AVX 路径）；4K 之后所有纯 CPU 线（灰、橙）快速穿破 30s 红线，SLP/DSLP 长得慢得多——32K 时 SLP 29.7s 贴线过关，DSLP 18.1s 余量充足；64K 时 DSLP 58.9s，KT AMX 的乐观估计（橙色虚线）是 158s。

4. Decode：把 CPU 的带宽榨到 947 GB/s

4.1 CPU 上的 FP8：先补背景

DeepSeek-V3 原生以 FP8（E4M3FN：1 符号 + 4 指数 + 3 尾数，无 NaN 编码）存权重，每 128 元素的 block 共享一个 FP32 scale 还原动态范围。GPU 有原生 FP8 tensor core，直接算；CPU 没有 FP8 运算单元，已有系统只有两条路：要么换量化格式（Q4_K_M 等，掉精度），要么把 FP8 升格成 BF16 存起来（内存占用直接翻倍，700 GB 变 1.4 TB，带宽需求同步翻倍）。这就是”本地只能跑阉割版”的技术根源。

本文的答案是第三条路：FP8 原样存，算的时候用 AVX-512 整数位操作现场解码。decode 是 memory-bound，多花点计算指令换一半内存流量，稳赚。

4.2 从 36 cycle 到 27 cycle：post-scaling 的巧思

Baseline 做法（下图）符合直觉但很贵：FP8 → 扩展成 FP32 → 乘 scale、修指数偏置 → 截断成 BF16 → 和激活做 BF16 点积。

问题出在 FP32 中间格式：一条 512-bit 寄存器只装得下 16 个 FP32，而且整条链 15+ 条向量指令，估算延迟 ~36 cycle、吞吐 12 cycle/16 元素。

优化版（下图）把顺序换了——先点积，后缩放（post-scaling）：

1. FP8 直接位操作扩展成 BF16（E4M3 的指数加上偏置差 120 就是合法的 BF16 指数，尾数补零——纯整数移位拼接，不经过 FP32）
2. 激活本来就是 BF16，装进 __m512bh——一条寄存器 32 个元素，槽位利用率翻倍
3. 用 AVX-512 BF16 扩展的 vdpbf16ps 指令做融合点积，直接累加进 FP32 累加器
4. 整块点积完了，scale 只在最后乘一次

为什么数学上等价？因为 scale 是整个 block 共享的常数，scale·Σ(w·x) = Σ(scale·w·x)，乘法可以提出求和号。效果：延迟 27 cycle、吞吐 4.96 cycle/迭代，实测带宽 815 GB/s。数值精度实测 L1 误差 95 分位只有 0.0017（对比 GPU 的 BF16 参考实现），端到端 benchmark 也验证无损（见第 6 节）。

再往上叠一层循环重构：M 轴 tiling 对齐 scale block 边界、K 轴按 128 元素 block 迭代（每 block 只乘一次 scale_inv，比逐次乘再省 4 倍乘法）、block 内 SIMD 展开流水。最终 947 GB/s——理论带宽 1228 GB/s 的 77%。对比现成库（论文 Table 1）：

库	格式	延迟 (µs)	带宽 (GB/s)
OpenBLAS	FP32	69.5	844
AOCL-BLAS	FP32	59.2	993
AOCL-BLAS	BF16	69.5	422
Ours (Optim.)	FP8	15.5	947

延迟 4-5 倍优势，且是唯一原生吃 FP8 的——别家先要把 FP8 解出来，内存流量就已经输了。

4.3 细粒度并行：同步和格式转换的抠法

每个 expert 是 gate/up/down 三个投影，gate 和 up 无依赖可并行，down 要等两者的结果合并（SiLU(gate)⊙up）。朴素做法是全局 barrier——所有核等所有 expert 的 gate/up 都算完再进 down。本文改成逐 expert 的细粒度 barrier：每个 expert 的 gate/up 沿输出维切成 m 段、摊给 NUMA 节点内的核，该 expert 自己的段全完成就立刻做合并并放行自己的 down，不等别的 expert；激活函数和逐元素乘还直接融进 stage 1。量化推理路径（Q4_K_M）里的一串格式转换（hidden→Q8_K→F32→Q8_K→F32）也全部融进相邻 kernel，每个 MoE 层只剩两次全局同步。

decode 端到端结果：

671B DeepSeek-R1 满血 FP8：21.5 tokens/s，到 32K context 还稳在 ~20；1T 参数的 Kimi-K2 也有 22.4。对照第 1 节的 SLO：云端 20 tokens/s 基线，达标，而且是原始精度。量化模型上同样领先——INT4 28 tokens/s（KTransformers 22.4，llama.cpp 只有 8）；128K 长 context 下别家掉到 15 以下，本文还有 19：

注意图里一个扎心的对比：本文引擎跑满血 FP8 的速度 ≈ KTransformers 跑 INT4 的速度。同样的硬件，别人跑阉割版的速度你跑原版。

5. 并发：两个正交的干扰源，两个解法

5.1 Prefill 挤 decode：节点内 PD 分离

云端早就流行 prefill-decode disaggregation（两类请求分属两组机器）。本地只有一个节点怎么分离？SLP 恰好创造了条件：prefill 被它做成了”纯 GPU 消费、CPU 只出带宽”的负载，于是一张卡专职 prefill、另一张卡加 CPU 跑 decode，两边资源画像几乎不相交。权重在 DRAM 里只存一份，两个进程通过 experts ring buffer 零拷贝共享——prefill 走 PCIe 拉权重（十几 GB/s），decode 吃 DRAM 带宽（>1 TB/s），互不挤兑。

调度策略也直白：<2K 的短请求走 chunked prefill；长请求看窗口内有没有 decode 在跑——没有就双卡 DSLP 拉满吞吐，有就单卡 SLP 保 decode 延迟。

5.2 Batch decode 挤自己：dual-batch 重叠

batch decode 在 MoE 上有个特有的坑：batch 内每个请求激活的 expert 几乎不重叠（256 选 8，两个请求撞 expert 的概率很低），所以 batch=2 时要读的 expert 权重接近 2 倍——带宽需求随 batch 线性涨，cloud 那套”batch 越大越赚”在带宽已经打满的 CPU 上失灵。

另一个浪费：CPU-GPU 混合推理是串行交替的——GPU 算 attention（~350µs/层）时 CPU 闲着，CPU 算 MoE（~450µs/层）时 GPU 闲着。

dual-batch 方案一石二鸟：

两个 microbatch A、B 错开半拍：GPU 算 A 的 attention 时，CPU 在算 B 的 MoE；A 的 attention 完了进 CPU 算 MoE，B 则进 GPU 算下一层 attention——两个硬件都不闲。看图中箭头，A（黄）和 B（紫）在 GPU 列和 CPU 列上正好互补填空。

效果（下图）：batch 1→2 系统吞吐 21.5 → 33.6 tokens/s（1.56 倍），每请求还有 ~16 tokens/s；batch 6 达 34.7 tokens/s。作为对照，KTransformers 的 batch 扩展全靠带宽硬扛，每请求 TPOT 掉得更狠：

并发微基准（1P+1D、2D、1P+2D 等混合负载）的结论：干扰控制在 <15% 延迟增加，而 KTransformers 在 1P+1D 下 decode 的 TPOT 直接劣化数倍——prefill 一来 decode 就卡住，这正是没有 PD 分离的代价。

6. 精度：满血就是满血

全文最省心的一张表：DeepSeek-V3.1 在 MMLU-Redux 上 91.00（官方 91.8）、MMLU-Pro 83.41（官方 83.7）；Kimi-K2 也在 1 个百分点级别的误差内。这点差距作者归因于生成配置的细微不一致，而不是系统引入的能力损失。对比所有靠 INT4/改路由换速度的方案，这是定性的差别——你跑的就是官方那个模型。

7. 一点个人 take

作为一直做 MoE 推理的牛马，这篇我读得很过瘾，说几点：

1. 最大的观念贡献是把”本地推理”从’能跑就行’拉到了’SLO 工程’。四个 gap 的定义（完整性/TTFT/TPOT/并发）本身就值得抄——以后评价任何本地部署方案，先拿这四条量一遍。尤其”模型完整性”这条：过去大家默认本地=量化，这篇证明了满血 FP8 在消费级平台是可行的，而且速度不输别人的 INT4。
2. SLP 本质上是把”权重驻留”这个假设翻掉了。传统推理引擎的出发点都是”模型在显存里”，offloading 是打补丁；SLP 反过来，显存只是权重的流水车间，DRAM 才是家。这个视角对 prefill 特别成立——compute-bound 负载下，只要搬运藏得住，”模型装不装得下显存”根本不重要。和我之前解读的那篇 Shannon 无损压缩正好互补：一个把权重变小，一个让权重流动，都是在打”显存容量不够”这同一场仗。
3. SmallEP 的通信账值得所有做小规模并行的人算一遍。N·A·D/S vs N·D 这个对比说明：并行方案的优劣是参数区间相关的，S ≤ A 时最优解和 S » A 时完全不同。云端方案搬到本地失效，多数时候不是工程没做好，是渐进条件变了。
4. 冷静的部分：平台是双路 EPYC 9355 + 24 条 DDR5-6400 + 双 5090，攒下来七八万人民币量级——比 8×H100 便宜一个数量级，但离”消费级”的日常语义还有距离，叫”工作站级”更诚实。另外 1.15 TB 内存对 Kimi-K2（1T 参数）已经贴脸，下一代模型再涨就得 24 条 64 GB 或者双节点了，容量墙只是被推远没有消失。

欢迎评论区交流，尤其是自己攒机跑大模型的朋友——这套配置的性价比确实到了值得认真考虑的临界点。

---

顺带扯一句题外话：这篇论文的功夫全在”帕累托前沿上找配置”——ring buffer 长度随 prompt 长度自适应、≤4K 走 CPU / >4K 走 SLP 的路由策略、batch 上限 Y=6 的取舍，每一处都是在延迟/吞吐/资源的多目标空间里做决策。这种系统调优问题和 AutoML 的超参搜索在方法论上是同一件事。我们把 AutoML 方向的积累整理成了《动手学 AutoML：从 NAS 到大语言模型优化实战》，进阶篇讲的贝叶斯优化、进化搜索正是自动化这类决策的工具，感兴趣的读者可以翻翻。