解构 DeepSeek 的基座：DeepInfer

发表于 分类于 ，
拆开 DeepSeek 开源的推理框架 DeepInfer，看一台专用推理内核如何用 C++ 调度、步级流水和双后端镜像换取吞吐，又在哪里留了工程化的白。

DeepSeek 把生产环境用的推理框架开源了，名字叫 DeepInfer。README 第一句话写得很直白：这是支撑 DeepSeek 大规模线上服务的推理框架，开源版本提供离线推理能力。

这句话本身就值得琢磨。”支撑线上服务”和”只给离线推理”之间，隔着的正是 KV cache 系统、PD 分离、多副本负载均衡这几块——换句话说，能决定服务质量的那一层没有放出来，放出来的是计算内核。

那就顺着这个边界去拆：一个为了把 V4 系列模型跑满硬件而生的内核，究竟长什么样，哪些设计是真正的取舍，哪些地方又是明显的欠账。

一、先看体量与出身

  • Python 约 5.5 万行，C/C++ 约 5,900 行，均不含 vendor/ 下的子模块。
  • 派生自 vLLM，文件头保留着 Derived from vLLM 的声明，但只保留了推理链路。
  • Git 历史只有三个提交，属于整理后的一次性放出。
  • 目标硬件只有两种：NVIDIA SM100 和 Ascend 950。
  • 目标模型只有一个系列：DeepSeek-V4.1 / V4.1-Flash。

这几个数字放在一起，结论很清楚：它不是一个通用推理框架，而是围绕一个模型结构深度定制、围绕两种硬件深度优化的执行系统。通用框架追求的是”什么模型都能跑”，它追求的是”这个模型在这些卡上跑得最快”。

二、骨架：四层加一颗 C++ 内核

graph TB
  subgraph 接入
    L[LLM 门面]
  end
  subgraph 编排
    E[LLMEngine]
  end
  subgraph 调度内核
    S[Scheduler]
    B[BlockManager]
    Q[Sequence]
  end
  subgraph 执行
    W[Worker]
    R[ModelRunner]
    C[CacheEngine]
  end
  subgraph 模型与算子
    M[DeepseekForCausalLM]
    K[Kernel: FlashMLA / DeepGEMM / DeepEP]
  end
  L --> E
  E --> S
  S --> B
  S --> Q
  E --> W
  W --> R
  R --> C
  R --> M
  M --> K

分层本身不稀奇，稀奇的是切分的依据不是职责，而是延迟。

请求进来到出 token，要经过：tokenize、入队、调度、张量化的元数据准备、前向、后处理、采样、状态回写。这里面绝大多数步骤是”每请求一次”的循环，在 Python 里做意味着 GIL 和解释器开销随请求数线性上涨。所以设计上做了一个很硬的切分：

  • 调度决策、block 分配、元数据准备、输出对象构造，全部下沉到 C++；
  • Python 只保留请求注册与编排，以及模型结构这种变化快、需要灵活表达的部分。

Scheduler.hpp 1,329 行，llm_engine.py 1,029 行，两个千行文件正好卡在这条语言边界的两侧。跨过去的是 SchedulerOutputs、InputMetadata、SamplerOutput 这几份契约，靠 pybind11 手工对齐。

三、性能主轴：把一步拆成四段，再让相邻两步错位

DeepInfer 最核心的设计叫 ZOS，Zero Overhead Scheduling。名字有点夸张，做法其实很克制：一步推理拆成四段，相邻两步的同名段错位重叠。

sequenceDiagram
  participant S as Scheduler
  participant P as 元数据准备
  participant F as Forward
  participant O as 后处理 + 采样
  S->>P: 第 N 步 schedule
  P->>F: InputMetadata
  F->>O: hidden
  Note over P,F: 第 N+1 步的元数据准备与第 N 步的采样并行
  S->>P: 第 N+1 步 schedule
  P->>F: 下一批 InputMetadata
  O-->>S: SamplerOutput（第 N 步）

重叠的对象是”上一步的采样”和”当前步的前向”。听起来只是并行，实际上它给整个系统加了三条硬约束：

  1. 元数据准备不能依赖采样结果。因为准备第 N+1 步时，第 N 步的 token 还没出来。这条约束直接决定了 prepare_input_metadata 必须只依赖调度器已经确定的信息，而不能顺手把采样结果拼进去。
  2. 缓冲区必须双份。相邻两步同时持有输入输出，单缓冲一定会互相踩。
  3. 资源释放的时点要往后推。请求中止时不能立刻回收它的 KV block，因为当前步的 forward 可能还在读。

第三条在代码里落成了一个专门的对象 DeferredRelease：中止的请求先移进延迟集合，不再参与调度，每个 step 结束后统一 flush，flush 点必须晚于”最后一次引用该 block 的 forward”。如果序列化理解这段逻辑，它就是典型的时间耦合——一个看起来很自然的”中止即释放”，在流水线里会变成 use-after-free 或者 block 被重复分配。

这也是我在读这类引擎时最看重的地方：性能不是加出来的，是把原本串行的约束重新排布出来的。每拿到一份重叠收益，就要在别处还一份时序约束。ZOS 的价值不在于”快”，而在于它把这些约束显式地写成了几个字段和一个延迟释放队列，而不是散落在各处的 sleep 和注释里。

四、双后端：用镜像换清晰

同时支持 NVIDIA 和 Ascend，常见的做法是抽一层 HAL，把算子差异藏在接口后面。DeepInfer 走的是另一条路：Platform 只负责能力查询、dtype、通信后端这类轻量差异，真正的算子实现放在 tilelang_ops/cuda 和 tilelang_ops/asc 两个目录里，结构镜像，代码各写各的。

重复度检测的结果很能说明这是有意为之：Python 侧重复率 0.25%，四处克隆里有三处是 CUDA/Ascend 的镜像实现。换成抽象接口，这些重复会消失，但两个后端各自的优化空间也会被接口形状锁死。

这个选择是有代价的：镜像代码没有统一的一致性测试，两端的语义差异只能靠人记。对一个要长期维护双硬件的团队来说，这笔账是否划算，取决于后端数量会不会继续增加——两种硬件时镜像划算，五种硬件时就会变成灾难。

五、模型结构反向塑造了引擎

真正让 DeepInfer 区别于通用框架的，是 V4 系列那几个专有结构，以及它们逼出来的引擎改造：

模型能力 引擎侧的代价
mHC 多路残差 显存按 hc_mult 放大，KV 预算要重算
YoCo 跨层 hidden 复用 状态管理变复杂，prefill 可截断成有界回放
Engram n-gram 稀疏表 多一张表要存、要通信、要均衡
KV 压缩 C1/C2 压缩解压开销换显存，缓存层级变多
稀疏 indexer top-k 多一条 index cache，且要与主 KV 对齐
DFlash MTP 变长验证 投机解码的验证开销与接受率绑定
EPLB 冗余专家 专家迁移，负载均衡变成在线动作

最能说明问题的是缓存体系。通用框架基本只有一种 KV cache，DeepInfer 有四种：主 KV、滑动窗口 SWA、压缩层的 state、稀疏层的 index，分别由 block 池、窗口 block、独立 state 池和对齐的 index 池管理，再统一挂在 BlockManager 的页表下。

这不是为了炫技。每一种 cache 都对应模型里的一类层，如果硬塞进一种 cache，要么浪费显存，要么让 kernel 拿到不符合假设的形状。模型结构决定了显存布局，显存布局决定了调度器的 block 管理方式——这条传递链是专用引擎和通用引擎最本质的分野。

另外两个容易被当成”功能点”的东西，其实是服务一致性的底线：

  • 批不变性（batch invariant）：同一个请求，不管和谁拼在一个 batch 里，输出必须一致。否则线上做压测、做 A/B、做灰度时，结果无法复现。
  • 确定性采样：同样的输入和采样参数，必须得到同样的 token。这要求固定 seq_seed，并且关掉 MTP 变长验证这类会引入非确定性的路径。

这两条在论文里不显眼，在工程上是刚需：一个不能复现的推理服务，出了问题连回放都做不到。

六、三条值得抄的判断

抛开具体实现，DeepInfer 里有三个判断我认为是可以直接迁移到其他系统的。

第一，按延迟切语言边界，而不是按职责切。 调度、元数据、对象构造放 C++，因为它们是每请求一次的循环；模型结构留 Python，因为它变化快、需要表达力。很多系统把”性能敏感”当成一个整体标签，结果要么全下沉导致僵化，要么全留在上层导致慢。

第二，用重叠隐藏开销，而不是优化开销本身。 调度再快也有成本，与其把调度优化到零，不如让它在设备计算的时候发生。步级重叠的收益上限比微优化高得多，代价是要付出一整套时序约束。

第三，为模型定制，而不是为通用性买单。 一个只服务一个模型系列的引擎，可以把模型里的稀疏结构、压缩方式、跨层共享全部写进缓存布局和 kernel 选择里。这种定制换来的性能，通用框架很难追上。代价也很明确：换一个模型，就要重做一遍。

七、四处明显的欠账

分析文档里对工程化的评价很直接：内核成熟度明显高于外围。具体到代码，有四处值得记住。

校验靠 assert。 代码里有 1,045 处 assert 承担参数校验职责。这在开发环境没问题，但 python -O 会把它们整体优化掉，等于生产环境没有校验。对于一个加载几 GB 权重、跑在特殊硬件上的系统，”配置错了但没报错”是最难查的故障。

测试与体量严重不匹配。 16 个测试文件对应 5.5 万行 Python，覆盖集中在 sampler、EPLB、sequence、userbuffer 这些点状能力上，缺端到端回归和长稳测试。也就是说，改动调度或缓存这类核心路径时，没有安全网。

跨语言契约没有机器可校验的定义。 SchedulerOutputs、InputMetadata 这些结构由 C++ 生产、Python 消费，字段靠 pybind 手工同步。改一个字段语义，编译器不会提醒你，只有跑到线上才会发现形状或 dtype 对不上。这是双语言架构的固有成本，但可以用 schema 和双向校验把它压下去。

没有 CI、指标导出、健康检查。 运维层面基本是”日志加重启”。这跟开源版只放离线推理的定位有关，但也意味着拿到它的人，要自己补一整套可观测性。

还有三个千行级类（Scheduler、LLMEngine、CacheEngine）和固定容量 10,000 的 RequestOutputPool。后者用手工归还的裸内存池，省掉了每步的堆分配，也把”重复归还”变成了一类只能靠人守住的错误。

八、怎么读这个项目

如果把它当成”又一个 vLLM”，会失望：它不支持一堆模型，也不提供 OpenAI 兼容的服务端。

如果把它当成”一个模型在两种卡上怎么跑满”的参考答案，收获会很大。ZOS 的步级重叠、C++ 调度内核与 Python 模型层的切分、四种 cache 的组织方式、双后端的镜像策略，这些都是从真实负载里长出来的，不是设计文档里推演出来的。

而它留下的空白同样有价值：内核可以靠少数高手写到极致，服务化必须靠机制兜住。断言换成显式校验、测试补上端到端、契约加上 schema、指标接上导出，这四件事做完，它才从”一台锋利的发动机”变成”一台能交给别人开的车”。

用一句话总结这个项目：性能已经开刃，工程化还没开刃。