DeepInfer 项目概览报告

阅读对象：本地仓库 DeepInfer-main/（对应 deepseek-ai/DeepInfer main 分支） 代码规模：排除 vendor/ 后 342 个文件；最大单文件为 csrc/engine/scheduler.hpp（1329 行 / 66 KB），其次 deepinfer/engine/llm_engine.py（1029 行 / 45 KB） 许可：Apache-2.0；Python ≥ 3.12；构建需 CUDA 13（NVIDIA）或 CANN 9.2（昇腾）

一、Introduction

1.1 背景

大模型的能力边界由训练决定，但它能否以可接受的成本服务海量用户，取决于一个长期不出现在论文里的中间层：推理框架。当负载从"一问一答"演化为"长前缀 + 多轮 + 工具调用"的 Agent 形态后，系统优化的目标函数发生了结构性变化——单次请求的算力不再是瓶颈，前缀的复用率、KV cache 的容量与生命周期、以及解码阶段的每步延迟成为一等约束。在这一阶段，推理框架的吞吐不再主要由 kernel 决定，而由调度、缓存、路由、通信四者的耦合方式决定。

DeepSeek V4.1-Flash 的线上推理系统给出了一个极端样本：单集群日吞吐 155 万亿 token，按低谷电价折算单卡每小时收入 38 美元（doc/inference_deployment_qa.md:17）。DeepInfer 是该系统的开源版本，以 Apache-2.0 发布，2026-10-16 建仓、2026-10-20 正式公开。同期开源的 deepseek-recipe（MIT）承担 API 协议与 token 序列之间的编解码，二者构成"协议层 / 引擎层"的上下分工。

1.2 问题陈述

本文档所面对的工程问题可以形式化为：在由商用硬件组成的大规模集群上，如何在给定 TTFT（首 token 延迟）与 TBT（token 间隔）约束下最大化有效吞吐，同时使 KV cache 的复用窗口达到天级、且不被单一硬件厂商锁定。

这个问题的难点不在任何单一环节，而在于几个目标互相拉扯：

1. P/D 资源配比冲突。Prefill 是计算密集、Decode 是访存密集，混布会让二者互相拖累；分离部署则引入跨阶段 KV 传输与两套独立扩缩容逻辑。
2. 前缀复用与负载均衡冲突。前缀亲和路由能提高命中率，却会破坏 DP 副本之间的负载均衡；反之亦然。
3. 模型结构与通用引擎冲突。MoE 超稀疏 + 超长上下文 + 层间共享 KV 的结构，会击穿通用引擎"以 TP 为主、以进程内 prefix cache 为缓存"的两个隐含假设。
4. 硬件异构。同一套服务要同时跑在 NVIDIA SM100 与华为昇腾 950 上。

1.3 与现有工作的关系

DeepInfer 的 README 明确承认自己是"originally built primarily on vLLM, with additional inspiration from SGLang and TensorRT-LLM"；仓库里 22 个文件保留 Derived from vLLM; modified by DeepSeek. 的 SPDX 头，覆盖 config.py、sequence.py、sampling_params.py、engine/llm_engine.py、engine/arg_utils.py、entrypoints/llm.py、worker/{base,worker,model_runner,cache_engine}.py、model_loader.py、layers/{sampler,layernorm,rotary_embedding,vocab_parallel_embedding}.py、distributed/shm_broadcast.py、setup.py 等。

三条既有路线各自解决了问题的一部分：vLLM 用 PagedAttention 与 continuous batching 解决显存碎片与批处理；SGLang 用 RadixAttention 与前后端分离解决前缀复用与结构化生成；TensorRT-LLM 走编译期图优化路线。但三者在设计取向上有共同点：以通用性为前提（一套引擎跑很多模型）、以 TP 为默认并行策略、KV 复用主要发生在引擎进程自身可见的内存里；vLLM 与 SGLang 的调度器还留在 Python 侧。这三个默认值恰好是 DeepSeek 线上负载最不能接受的。DeepInfer 的贡献可以理解为：把这三个默认值逐一替换掉，并承担替换后的全部工程代价。

1.4 本报告要说明的贡献

对应第三部分的详细论证，DeepInfer 的差异化贡献可归纳为六项：

1. C++ 调度器 + “步型投票"的同引擎 P/D 协同：把 batch 调度从 Python 搬到 C++，并用 DP 组内 all-reduce 投票保证同组各 rank 的步型（prefill / decode）一致。
2. 分层 KV 存储：block KV 落 3FS 专用存储集群、SWA KV 留主机内存、state cache 按 128-token 粒度落盘，三层介质各司其职。
3. 模型-系统协同的推理机制：YoCo 前截断、SWA Bounded Replay、mHC / Engram / 稀疏 indexer 的原生 kernel——为特定结构写引擎，而非削足适履。
4. 置信度驱动的变长投机验证（DSpark），并把 per-request 长度决策下沉到 GPU kernel，绕开异步调度与置信度回传的结构性矛盾。
5. 同源双平台：不是移植分支，而是同仓库两套 TileLang 后端（tilelang_ops/{cuda,asc}/）+ 两套 vendor 算子（vendor/Deep*-Ascend/）+ 一层平台抽象。
6. 可复现性作为一等公民：batch-invariant 推理、确定性采样、逐样本精确比对的 PPL 黄金基线，以及一套可复跑的 AIME/GPQA/HLE 评测流水线。

1.5 需要预先声明的边界

必须明确指出：这份开源代码不是一套完整的服务平台，而是一个离线推理引擎库加上一份带操作手册的生产部署剖面。 以下三块在文档中被详细描述、但代码中只有扩展点：

• DP 级三层负载均衡（Prefill / Decode / Vision Encoder Load Balancer）
• kvcache-store 服务端
• HTTP/OpenAI 协议服务层

证据见 §2.4。官方给出的理由是这些组件"tightly coupled to our internal infrastructure”（doc/inference_deployment_qa.md:21）。

二、项目结构（自顶向下视角）

2.1 L0：仓库顶层布局

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

DeepInfer-main/
├── deepinfer/          # Python 引擎主体（218 个 .py），含 C++ 扩展的 Python 侧胶水
├── csrc/               # C++ 核心：调度器、序列管理、输入装配、pybind（33 文件）
├── deep_cluster/        # 硬件发现与拓扑绑定，刻意与加速器运行时解耦
├── vendor/             # 6 个 git submodule：算子库
├── examples/           # 离线推理、benchmark、Pareto sweep、YoCo 等价性验证
├── eval/               # AIME 2026 / HLE / GPQA-Diamond 评测流水线
├── tests/              # 16 个 test_*.py，按 engine/models/worker/ve/distributed 划分
├── doc/                # 部署 Q&A（312 行）+ 4 张架构与 Pareto 图
└── scripts/            # 构建、格式化、stub 生成

vendor/ 的 6 个子模块全部来自 deepseek-ai 自己的算子仓库，且按硬件成对出现：

1
2
3
4

vendor/DeepEP         + vendor/DeepEP-Ascend
vendor/DeepGEMM       + vendor/DeepGEMM-Ascend
vendor/FlashMLA       （仅 NVIDIA）
vendor/DeepSelect     （MoE 稀疏选路）

这个成对关系是整个仓库"双平台"设计的物理证据：跨平台不是条件编译分支，而是两份独立构建的算子仓库。

2.2 L1：分层视图

模型层的内部再分层非常清晰，v4_modules/__init__.py 用一段 docstring 把 24 个文件划成四种角色：

1. 参数容器：只分配参数，声明无 forward（containers.py、mla_attn.py、moe.py、hc_blocks.py、engram_parameters.py、heads.py）
2. 运行时适配与前向：包装容器、转成运行时布局、拥有"模型怎么算"（attn.py、compress.py、index_attn.py、mlp.py、hc.py、engram.py）
3. 张量算子：无参数无 forward，只操作张量（mhc_ops.py、mhc_modeling.py、yoco.py、hc_lazy.py）
4. 配置派生的拓扑与装配：从 config 推导并校验结构，再搭参数树（cross_layer.py、engram_metadata.py、builder.py）

这种"容器 / 前向 / 算子 / 装配"四分离，是为"参数按 HF 命名对齐、运行时按 kernel 布局重排"这一需求服务的：builder.py:93 的 build_model() 严格按 MultimodalEmbedding → N×MoeDecoderBlockHC → MainPredictHeadHC → M×DFlashNextPredictHeadHC 组装，参数名保持与 HF checkpoint 一一对应，转换只在运行时发生。

2.3 L2：一次请求的调用链

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

LLM.add_request()                              entrypoints/llm.py
  └─ 构造 Sequence（Python 对象，含 token_ids / sampling_params / so_args）
       ↓
LLMEngine.step()                               engine/llm_engine.py:860
  ├─ scheduler.schedule(seq_metadata_pool)     → C++，见下
  │    ├─ _schedule()                          csrc/engine/scheduler.hpp:275
  │    │    ├─ EngineType::PROMPT     → _schedule_prefill()
  │    │    ├─ EngineType::COMPLETION → _schedule_decode()
  │    │    └─ EngineType::NORMAL     → DP 投票决定本步全组做 prefill 还是 decode
  │    └─ 返回 SchedulerOutputs + seq_metadata_list（pinned buffer）
  ├─ clone_states / clone_blocks / zero_states 的 GPU 侧下发（如抢占发生）
  ├─ post_schedule()                          虚函数，子类钩子
  └─ _execute_model_async_step()              → run_workers("execute_model_step")
       ↓
ModelRunner.execute_model_step()               worker/model_runner.py
  ├─ _prepare_input()                          → prepare_tensor_input_v4（C++）
  ├─ execute_main_model()                      → 模型 forward（mHC/MLA/MoE/Engram）
  ├─ execute_sample()                          → fused_sample kernel（含 SO mask 应用）
  └─ post_process()                            → 写回 pinned buffer
       ↓
LLMEngine.post_process()                       engine/llm_engine.py:984
  ├─ pe_do_save_kvcache()                     虚函数：KV/state 落盘钩子
  ├─ _process_model_outputs()                 停止条件、输出组装
  ├─ scheduler.update_mtp_runtime()           用本步实测 forward/sampler 耗时校正投机预算
  └─ deferred_release.collect_*()             跨步延迟释放（生命周期管理）

链路中有一处值得单独注意：输出的所有权是跨步延迟的。step() 的 docstring 明确写着返回的 RequestOutput 由 RequestOutputPool 管理，在下一次 step() 时释放，调用方若要保留必须显式 clone/release。这是零开销调度（§3.6）的直接后果——第 N 步的后处理必须在第 N+1 步的 GPU 计算期间完成，因此不能在第 N 步就回收。

2.4 L3：仓库内 / 仓库外的物理边界

这是理解本项目最关键的一张表。

EngineType 的定义（engine/engine_type.py，注释为原文）把这件事讲得最清楚：

1
2
3

NORMAL     = 0            # process prompts & completions
PROMPT     = enum.auto()  # only process prompts
COMPLETION = enum.auto()  # only process completions

同一个引擎二进制，通过这一个枚举就能扮成 P、D 或 P/D 混布三种角色。分离部署的"引擎侧"确实是完备的，“编排侧"留给了使用者。

2.5 L4：构建与依赖图

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

build.sh  ($DEEPINFER_ACCELERATOR = nvidia | ascend，必填，无默认值则拒绝执行)
  ├─ pip install -r requirements-${ACCELERATOR}.txt
  │    ├─ 公共：torch 2.13.0 / transformers 4.50.0 / xgrammar 0.2.2
  │    │        tilelang 0.1.15 / tile-kernels 2.1.0 / tokenizers 0.21.4
  │    ├─ nvidia：torch==2.13.0+cu132、flash-attn-4==4.0.0b34
  │    │          并强制把 NCCL 钉到 2.32.3
  │    └─ ascend：torch==2.13.0+cpu、torch-npu==2.13.0rc1
  ├─ scripts/build_vendor.sh   → 按平台编译 vendor/ 下 5～6 个算子库
  └─ scripts/build.sh build    → setup.py 编译 deepinfer._C（NVIDIA 路径 13 个 C++ 源文件）
                                 产出软链到 deepinfer/_so/<platform>-<arch>/

setup.py 用 accelerator_kind() 一次性分流到 CUDAExtension 或 NpuExtension，并强制要求 ${CONDA_PREFIX} 或 ${VIRTUAL_ENV}（NVIDIA 路径）。平台差异的收口点只有一个环境变量，但代价是这个变量的取值会改变依赖集、算子库、编译产物、kernel 后端目录四件事。

三、项目重点：它和其他 Infra 到底不一样在哪

以下 11 点按"现象 → 代码证据 → 与同类系统的差别 → 代价"组织。前 6 点属于系统架构层面，后 5 点属于工程实践层面。

3.1 调度器在 C++ 里，而且引擎有三种人格，靠 DP 投票保证步型一致

现象。 组批逻辑不在 Python，而在 csrc/engine/scheduler.hpp——全仓库最大、最复杂的单个文件（1329 行）。它导出的 Scheduler 类被 pybind 包装，Python 侧每步只是一次调用。

证据。 scheduler.hpp:275 的 _schedule() 按 EngineType 分三路。最值得注意的是 NORMAL 分支（:295-341），它解决了一个通用引擎里不存在的问题：P/D 混布时，一个 DP 组内的所有 rank 必须在同一步做同一类型的计算，否则 all-reduce 会死锁或陪跑。做法是在 DP 组内做 all-reduce 投票，并区分两类尝试点：

• ALL 点（每 prefill_attempt_interval=70 个 decode step）：MIN(want) ⇒ 只有全员都想 prefill 才做，避免"一个 rank 干活、其余 rank 陪跑空 prompt step”。
• ANY 点（每 prefill_any_attempt_interval=100 个 decode step）：MIN(!want) ⇒ 任一 rank 有活就全组一起 prefill，防止孤立 rank 被饿死。
• 第二轮投票（:332）：即使投票决定做 prefill，_schedule_prefill 也可能空手（slot 满、KV 块不够）。若全组都空手必须转回 decode，否则"running 满 + waiting 非空"会死循环。

代码注释把这段死循环风险写得非常明白，这不是设计文档里的推导，是踩过坑之后的注释。

差别。 vLLM 与 SGLang 的 scheduler 都在 Python 侧，且默认"一个引擎只做一件事"，P/D 分离靠外部编排。DeepInfer 的做法是让同一个引擎在 DP 维度上自同步步型，把 P/D 混布变成一个引擎内的调度问题而非部署问题。这让"1P+1D 可部署，实际 P:D 约为 1:1"（doc/...:67）成为可能。

代价。 调度器一旦落到 C++，所有业务逻辑的调试成本上升一个数量级；而且步型投票引入了对 all-reduce 延迟的敏感性——每个投票点都是一次跨 rank 同步。

3.2 KV cache 被当成分布式存储系统，而不是引擎内的缓存

现象。 大多数框架里 KV cache 是"显存里的一块池 + 可选的 CPU 换出"。DeepInfer 把它拆成三层介质，各有明确的归属、复制策略与管理方：

（表来自 doc/inference_deployment_qa.md:131-135）

关键设计判断。 为什么 block KV 要放专用存储集群而不是本地盘分层？文档给出的理由很硬：让每一台 prefill 机器都能命中任意前缀，从而消除"前缀亲和路由"与"负载均衡路由"之间的冲突；同时把存储扩缩容与计算扩缩容解耦，并让 KV 不受计算节点故障影响——“compute nodes fail at a rate more than an order of magnitude higher than storage nodes”（:141）。

规模。 3FS 提供约 100 PiB SSD；Agent 流量 TTL 3 天、chatbot 7 天；实测每 GPU 逻辑容量约 14.4 TB（:81）。单个存储服务器推荐配置 20×30.73 TB SSD、≥40 GB/s 读、5 GB/s 写，容量足够保留 24 小时（:77）。数据按 10+2 纠删码存放，读请求同时打到 11 个节点、任 10 个返回即可解码（:87-89）。

SWA KV 的粒度设计。 这是全项目最精细的一段工程。SWA KV 不进分页块，而是每个序列占一个固定大小的 StateBlock，槽位内按"是否需落盘"切成 [stored region: SWA][dropped region]；只在 128 的倍数位置落盘；decode 期间另开一块 replica 区域，每跨一次 128 边界就整窗拷贝一次，请求结束时从最近一次完整拷贝取，拷贝位置最多超出边界 5 个 token，靠 ring buffer 多留的 6 个位置吸收。最终落盘的东西是精确的：20 layers × 134 tokens × 528 B = 1,415,040 B/序列（:163-169）。

代码里 KVSTORE_GRAIN = 128 的对齐断言出现在 cache_engine.py:787，与文档一致。

搬运动作的分工。 匹配（Entrypoint 对 kvcache-store 做 CHECK 角色的哈希链前缀匹配）→ 取 KV（独立进程异步，3FS → 主机共享内存）→ swap-in（引擎进程内线程，主机内存 → HBM，用 GDR 而非 CUDA copy engine，因为 copy engine 在细粒度块拷贝上表现差，而 RDMA 网卡可以打满 PCIe 且不占 SM）。三段各用不同机制，目标都是把开销藏进前向计算里（:157）。

差别。 vLLM 的 prefix cache 默认在进程内/主机内存，跨实例复用需要外部 KV connector；SGLang 的 RadixAttention 同样是本地树结构。DeepInfer 把 KV 提升为一个有独立 SLO、独立扩缩容、独立故障域的存储服务，并用"每台机器都能命中任意前缀"这个性质反过来简化路由。

代价。 需要 3FS 这样的分布式文件系统作为前提。这不是一个"装上就能跑"的设计，而是"如果你已经有 100 PiB 级分布式存储"的设计。

3.3 为特定模型结构写引擎：mHC / Engram / 稀疏 indexer / YoCo / 跨层共享 KV

现象。 v4_modules/ 里有大量在通用引擎中找不到对应实现的东西：

证据。 cross_layer.py 的 docstring 说得很直接：SharedLayerTopology 从 config 派生并校验 share-KV/share-index/share-index-K 如何对层分组，不兼容的 segment map 会 fail fast。builder.py:67-72 根据 layer_kv_map[runtime_layer_id] == runtime_layer_id 判断某层是"KV 源层"还是"消费者层"，并让消费者层直接复用源层的 k_proj/k_norm 模块对象（:89-91）。

这就是说：这不是"KV cache 共享"，是"权重与 KV 双重共享"。一个 60 层的模型，实际持有独立 KV 的层可能远少于 60 层——这直接改变了 KV 显存的估算模型，也解释了为什么 decode 能放到 64～128 卡的 Wide-EP 上。

差别。 vLLM/SGLang 的模型实现是"通用退化"的：任何结构差异都要表达成通用接口的组合，表达不了就加 hook。DeepInfer 反过来——先有 V4.1 的结构，再写最贴合的引擎。v4_modules/__init__.py 用四角色约定（参数容器 / 运行时前向 / 张量算子 / 拓扑装配）来管理这种紧耦合，避免它退化成意大利面。

代价。 换模型就要改引擎。README 里"Supported models: DeepSeek-V4.1-Flash"只有一项，不是谦虚，是架构决定的。

3.4 SWA Bounded Replay：把"缓存缺失"的代价从 O(上下文) 压成 O(1)

现象。 滑动窗口注意力（SWA）的 KV 一旦缺失，精确重建的代价与上下文长度成正比。DeepInfer 的做法是：只重算最近 128 个 token，并把注意力的可见域截断在被重放段内——从位置 s 重放时，位置 i 的 query 只与 [max(s, i-127), i] 的 key 计算。

证据。 文档（:183）明确：“The cost of recomputation is thereby reduced to a constant, independent of context length and layer count.” 产出的是近似 SWA KV，不是精确值。

它被用在两个地方，第二个更狠：

1. Encoder SWA bounded replay：前缀命中长度只由 block KV 决定。SWA KV 缺失时，重算命中前缀的最后 128 token，与未缓存后缀一起前向；重放段只重建 SWA KV，block KV 原样复用、不重算也不覆盖。
2. Decoder SWA bounded replay：CED 架构下 decoder 层的 block KV 是从最后一层 encoder 的隐状态投影出来的，所以只有 decoder 的 SWA KV 需要逐层算。精确重建需要把最后 2560 个 token 跑过 20 个 decoder 层；实际做法是每次 prefill 只重算 prompt 的最后 128 token，结果只供 decode 用、永不进前缀缓存——“nearly halving total prefill computation”。

命中策略也随之变形（:191）：命中长度由 block KV 决定，但"最后一块要不要重算"由 SWA KV 决定——SWA 缺失就退一个块重算。在 YOCO 截断要求下，长度 L 的输入最大命中长度不超过 L - 134（128 窗口 + 6 个投机回溯位）。

代码对应。 environ.py:188-194 的两个变量正是这个机制的旋钮：

1
2

DEEPINFER_YOCO_PREFILL_TRUNCATION_MODE = off | mask | compact
DEEPINFER_YOCO_PREFILL_KEEP_TOKENS

mask 保留全部 query 只收窄 SWA 可见域；compact 真正移除截断区 query 以减少 consumer 层计算。注释还写死了约束：KEEP_TOKENS >= window_size + DEEPINFER_MAX_BACKTRACKING_TOKENS。

差别。 通用引擎把"缓存未命中"当作必须精确处理的正确性问题；DeepInfer 把它当作可以换精度的性能问题——用近似的 SWA KV 换掉一半 prefill 计算。只有当你同时拥有模型和引擎（能证明近似误差可接受）时才敢这么做。

代价。 引入了正确性上的"软边界"：SWA KV 是近似的，且 compact 模式会改变可见 query 集合。examples/yoco/yoco_equiv.py（51 KB）加上 examples/yoco/run_yoco_equiv_l2.sh 就是为验证 mask/compact 与原始语义等价而存在的——用了近似，就配一份等价性测试。

3.5 DSpark：置信度驱动的变长验证，以及为什么决策必须下沉到 GPU

现象。 投机解码的草稿 token 验收概率随请求、随位置变化，固定验证 5 个会在低概率 token 上浪费验证算力。DSpark 用 Confidence Head 预测每个草稿 token 的验收率，按草稿 token 的预期边际收益排序，只验证收益最大的前缀子集。

收益。 每步验证的草稿数从 5.00 降到 3.2～4.4（约 −30%），而平均接受长度基本不变（3.29 vs 3.20）。在 EP64/DP64 的基准上折算，生产工作点（约 60 并发/rank）下 +4～7% 吞吐；并发 96 时到 +11%（:247-256 有完整表格）。

真正的架构难题在这里。 变长验证需要"下一步验证几个 token"这个决策，而决策需要置信度；但开启异步调度后，schedule(N+1) 执行时第 N 步的置信度还没回来。文档（:264）和代码给出了同一个答案：

当 schedule(N+1) 运行时，第 N 步的置信度尚未返回，因此 CPU 调度只固定总预算；每个请求的长度由 GPU 上的 TileLang kernel 重算（累计乘积、全局排序取 top-k、融合 remap scatter），结果写回 pinned buffer。

代码对应：worker/worker.py 引入 compute_seqlens_q_by_confidence_rates（来自 tilelang_ops/schedule_mtp_varlen.py），并且 scheduler.update_mtp_runtime(num_seqs, graph_size, forward_ms, sampler_ms)（llm_engine.py:997）会把本步实测的 forward/sampler 耗时回灌给调度器——控制器在闭环里用实测性能曲线校正自己的预算。

差别。 多数实现的投机解码是"固定 k + 批内统一"；DeepInfer 是"per-request 变长 + 置信度排序"，且为了不被 CPU 同步卡住，把这段决策做成了 GPU kernel。这是"模型层的置信度头"与"系统层的异步流水线"两个设计互相逼迫出来的产物。

代价。 决策逻辑从 Python 搬到 TileLang kernel，可调试性大幅下降；environ.py 里 DEEPINFER_DFLASH_CONFIDENCE_HEAD_TEMPERATURE 甚至支持按 nextn 逐位传不同温度（逗号分隔且长度必须等于 nextn），说明这套参数是靠实验扫出来的，不是推出来的。

3.6 ZOS（零开销调度）：把"不能有隐式同步"变成 CI 断言

现象。 Decode 一步是 schedule(N) → forward(N) → sample(N) → post-process(N)。其中 schedule 和 post-process 是纯 CPU 逻辑，但取 sample(N) 的结果必须做一次 CPU-GPU 同步，于是 GPU 在 post-process(N) 和 schedule(N+1) 期间空转。

做法。 把流水线错开一拍：GPU 跑 forward(N) → sample(N) 时，CPU 并发跑 schedule(N+1) → launch_forward(N+1) → post-process(N-1) → launch_sample(N+1)。文档承认这是"following existing implementations (SGLang, NanoFlow)"（:274），environ.py:179 给了开关 DEEPINFER_ZERO_OVERHEAD_SCHEDULING，而 EngineType.default_zero_overhead_scheduling 让它对 NORMAL 与 COMPLETION 引擎默认开启。

工程上真正硬的部分。 一旦开启异步，任何隐藏在模型代码里的隐式 CUDA 同步都会把流水线打回串行，而这种 bug 不会报错、只会变慢。所以仓库里有：

• deepinfer/utils.py 的 forbid_cuda_sync() / allow_cuda_sync() 作用域装饰器，基于 torch.cuda.set_sync_debug_mode("error")
• tests/test_cuda_sync_guards.py（213 行）专门测这个守卫：嵌套、内层覆盖、退出泄露、多线程交错
• environ.py:198 的 DEEPINFER_DISABLE_SYNC_GUARD 作为逃生口

差别。 绝大多数框架把"性能"当作 benchmark 数字；DeepInfer 把一条性能不变量（“forward/sample launch 期间不许同步”）变成测试里会失败的错误。这是生产环境的痕迹，不是论文实现的味道。

3.7 同源双平台：不是移植，是一个仓库两个 target

现象。 支持 NVIDIA SM100 与华为昇腾 950。注意这不是"社区PR适配了昇腾"，而是同一仓库里两套并行实现。

证据（四层都成对）：

platform.py 是收口点，它做的不只是"选后端"，而是语义对齐：

• 把 ASCEND_RT_VISIBLE_DEVICES 映射到 CUDA_VISIBLE_DEVICES，不一致直接 RuntimeError
• transfer_to_npu 把 torch.cuda.* 映射到 torch.npu.*，于是 torch.cuda.is_available() 在纯 NPU 机器上返回 True——这会骗过 TileLang 的目标自动探测，所以必须 setdefault("TILELANG_DEFAULT_TARGET", "ascend")
• 把 torch.npu.Stream.cuda_stream 做成 property，让读取 .cuda_stream 的 CUDA 风格代码能拿到 NPU stream
• pin_host_memory 在 GPU 用 cudaHostRegister，在 NPU 用 aclrtHostRegisterV2(ACL_HOST_REG_PINNED)，并把"已注册/未注册"两个良性错误码静默，其他错误每次都告警（“whole-node degradation canary”）
• FP8 scale factor 的类型都不同：GPU 是 int32，NPU 是 int16

差别。 一般的"多硬件支持"是一个抽象层加多个后端实现，功能对齐靠测试。DeepInfer 是把 CUDA 语义当成方言，在 NPU 上做尽可能完整的语义模拟，代价是 npu_ops/patch_imports.py 那种猴子补丁（把 deep_select.interface.topk 和 flash_attn.cute.interface._flash_attn_fwd 替换掉）。

代价。 大量 TODO(NPU): Waiting for torch_npu bug fix, hack it temporarily.——这些注释诚实得刺眼，也说明双平台的真实成本。

3.8 开源边界画在"专家级"和"副本级"之间

这是一个容易被忽略但极其重要的产品决策。负载均衡被切成两半：

EPLB 的实现细节（来自文档 + 代码）： 每张卡除自己负责的专家外多持几个"冗余专家"承载热门专家的副本；收集最近 N 步的专家负载，每 M 步重新分配；专家权重在组内通过 NVLink symm_mem 就地 shuffle（shuffle.py 的 expert_model_parallel_shuffle_inplace）。参数很具体：

• Prefill：4 个冗余专家/GPU 起步（向上取整使总数能被 ep_size 整除），100 步统计窗，每 1000 步重排；实测 moe_lb（最慢 rank/均值）= 1.074（EP4DP4）/ 1.045（EP2DP2）
• Decode：1 个冗余专家/GPU 起步（EP64 时取整成每卡 7 个，不含共享专家），1000 步统计窗，每 10000 步重排；实测 moe_lb = 1.090

副本级 LB 的设计（只有文档，但从描述能看出复杂度）：

• Prefill LB：按总上下文长度分 4 档（<16K / 16–64K / 64–256K / ≥256K），按上一轮各档算力需求份额动态重分配引擎数；跨 DP 组靠"算力配额"拉活；组内请求分给当前配额负载最低的 rank；单卡内 FIFO 出队，chunk 大小不是固定值，而是从预算反解——每步给一个时间预算 compute_quota_limit（纳秒），用拟合的二次代价模型 cost(q, c) 预测耗时，解二次方程取"不超过 32768 token 预算且刚好用完时间预算"的最大 q，且该预算还要在 DP 组内做 all-reduce MIN。
• Decode LB：同时平衡两个维度——请求数（决定每步 batch size）与总 KV 长度（决定 attention 代价）。启发式是"短请求发给并发低但 KV 长的卡，长请求发给 KV 短的卡"，细节指向 DualPath 的组内调度算法。
• VE LB：按机器粒度路由，被选中的机器上再分给当前处理图像 patch 数最少的 GPU（patch 数直接从宽高除以 patch_size 算出）。
• 整体效果：约达到"完美均衡的离线 benchmark"的 80%（:127）。

为什么这样切。 EPLB 是纯引擎内部的优化——只需要看专家负载、重排权重，不依赖外部基础设施。三层 LB 是跨实例的——需要知道集群里有哪些实例、每条 KV 在哪台机器上、哪个 Entrypoint 持有目录。后者与内部基础设施强耦合，所以官方选择不开放。

这也给出了清晰的社区分工：想补全这套系统，缺的恰好是"副本级路由 + KV 目录"这一层，而这一层的接口点已经在代码里留好了（init_pd_split、do_save_kvcache、pe_prepare_save_kvcache、select_state_slice_to_store、load_state_from_store_to_prefill_cache）。

3.9 可复现性：batch-invariant、确定性采样、PPL 黄金基线

现象。 大规模在线服务里，“同样输入不同输出"通常是可接受的；DeepInfer 把它当作需要开关和测试的一等特性。

机制（README + environ.py）：

回归方式。 tests/models/offline_decode_ppl.json 是按"平台/负载"为键的 PPL 黄金基线，逐样本精确比对；examples/refresh_offline_decode_ppl.py 负责刷新（--long_context 刷长上下文那一组）。多节点时每节点只更新/比对自己拥有的 rank，并靠跨节点文件锁 offline_decode_ppl.json.lock 避免互相覆盖。--output（保存候选基线）在多节点下被显式禁止，因为候选文件必然缺其他节点的 rank。

差别。 这是"数值回归测试"而非"性能回归测试”。多数推理框架的 CI 只能测 smoke，因为输出本身不确定；DeepInfer 通过 batch-invariant + 确定性采样，把输出变成了可以逐位断言的量。没有前者，后者无从谈起——两个开关必须成对设计。

3.10 结构化输出：编译与推理解耦，mask 在 sampler kernel 里落

现象。 用 xgrammar 0.2.2 把格式约束编译成语法自动机，每步 decode 推进自动机、求出合法 token 集合、生成 mask 施加到采样逻辑上。

两个非标准做法：

1. 编译在引擎之外。文档（:284）：“SO compilation runs asynchronously, and a request retrieves the compiled result before it enters the engine”，编译在专用 SO 机器上进行。代码对应 structure_output/so_matcher_builder.py：一个 ThreadPoolExecutor + Future，submit/pump 都在主线程，worker 线程只填 future、不碰入调度逻辑，编译完由 pump 挂到 seq 上再 on_ready(seq) 入调度。DEEPINFER_SO_BUILDER_CPUS 默认 4 核，DEEPINFER_XGRAMMAR_TIMEOUT 默认 20 秒。
2. mask 应用在融合 kernel 里。structure_output/mask.py 懒加载 xgrammar 的 allocate_vocab_mask，再走 tilelang_ops/structure_output 的 kernel；llm_engine.py:913-928 甚至专门处理"最后一个 prefill chunk 的 SO 状态异常"——只判死本请求（置 FINISHED_SO_FAILED），不让簿记 bug 扩大化。

还有一条安全细节值得记：compiler.py:17 的注释写着 “xgrammar <= 0.1.x had production incidents with cache enabled; keep this env as the kill switch”，并配了 DEEPINFER_XGRAMMAR_DISABLE_CACHE 与缓存字节上限（默认 10 GB）。把上游库的历史故障做成自己的开关，是生产团队的典型动作。

差别。 大多数实现把 grammar 编译放在请求路径上（首个 chunk 时编译），长 schema 会直接推高 TTFT。DeepInfer 把编译移到请求进入引擎之前并放到独立机器上——因为它的调度器设计前提是"每个进入引擎的请求都已经可调度"（_has_prefill_waiting() 只看队头，队头不可调度就 break）。两者是配套的。

3.11 前端契约：这个仓库故意不提供 HTTP

现象。 全仓库只有一个 LLM 类（离线批量接口）和一个默认关闭的 profiling HTTP server。没有任何 OpenAI 兼容端点、没有流式响应、没有鉴权与限流。

理由。 协议编解码被完整地切给了 deepseek-recipe：它把 Chat Completions / Responses / Messages 三种请求归一成 conversation，渲染成 DeepSeek 的 prompt，编码成 token ID；输出的反向解析也在那里。README 明确"the caller must provide model inference, tool execution, and HTTP transport"。

这个切分带来的实质好处（在 deepseek-recipe 的 README 里）：特殊 token 的注入被结构性阻断——构建 prompt 时记录"显式插入的特殊 token 及其位置"，普通文本段独立编码，因此用户输入里出现的 <｜User｜>、<｜image｜> 只会被编成普通 token 序列，无法伪造控制符；输出侧同理，模型用普通 token 拼出的 </think> 不会被当成真正的思维结束标记。

差别。 通用框架倾向于"一站式"：引擎自带 server、自带模板、自带 tokenizer 逻辑。DeepInfer 选择引擎只管算子与调度，协议层完全外置，并用"官方 API 逻辑逐行对齐"来保证第三方部署与官方服务行为一致。

代价。 想跑通一个"能对外服务的实例"，你至少还要自己写 HTTP 层、接 deepseek-recipe、实现 DP 级路由、准备 KV 存储。§2.4 的表就是这份工作量的清单。

四、总结及展望

4.1 总结

第一，它的类型不是"又一个推理引擎"，而是"一份带操作手册的生产部署剖面"。 155 万亿 token/日、$38/卡/小时的数字、3FS 的 100 PiB 容量、EPLB 的 1.074 负载均衡比、SWA 的 1,415,040 B/序列、变长验证从 5.00 降到 3.2–4.4——这些不是论文里的 best-case 曲线，是线上系统的运营参数。开源代码 + 312 行部署 Q&A 的组合，把"一个团队怎么用硬件"这件事写得比"一个算法怎么更快"更具体。

第二，它的技术核心是三个"默认值的替换"。 调度器从 Python 搬到 C++ 并引入 DP 步型投票；KV cache 从引擎内缓存升级为分布式存储服务；模型实现从"通用引擎 + 适配层"改为"为特定结构写引擎 + 四角色约定管理耦合"。三处替换各自都带来可观收益，但也都把系统从"可移植"推向"不可移植"——这是本项目的性格。

第三，它最值得学的不是具体数字，而是几条工程纪律：

• 近似要配等价性测试（YoCo 用了近似，于是有 51 KB 的 yoco_equiv.py）
• 性能不变量要变成会失败的测试（forbid_cuda_sync + 213 行专项测试）
• 数值回归要逐样本精确断言（PPL 黄金基线 + 跨节点文件锁）
• 上游库的历史故障要变成自己的开关（xgrammar 缓存 kill switch）
• 模型的假设要写进 fast-fail 校验（SharedLayerTopology 对不兼容 segment map 直接报错）

第四，它的适用面是窄的，但窄得有理由。 硬件要求 SM100 或昇腾 950；只支持 V4.1-Flash 一个模型；不用 TP/CP，Decode 推荐 EP64 起，最小高效部署单元是 64 卡 prefill + 64 卡 decode；缺 KV 存储服务端与三层 LB。这些约束共同指向一个结论：它服务于"已经有 64 卡以上集群 + 分布式存储"的团队，以及想理解该量级系统如何设计的人，而不是想在自己机器上跑一个模型的个人开发者。

4.2 展望：还差什么

以下按"缺口 → 为什么 → 需要做什么 → 难度"组织。前四项是补齐开源边界的必答题，后六项是向前走的选答题。

（1）补齐 kvcache-store 服务端 —— 最大、也最明确的缺口。 现在只有客户端钩子：select_state_slice_to_store（按 128 粒度切片）、load_state_from_store_to_prefill_cache（回填）、llm_engine 里的三个虚函数。缺的是服务端：哈希链前缀匹配的元数据索引、CHECK 角色的匹配接口、GC 与 LRU 驱逐。 文档已经把规格写死了：单机 20×30.73 TB SSD、≥40 GB/s 读、5 GB/s 写、10+2 EC、任一 10 个节点返回即可解码、每 GPU 约 14.4 TB 逻辑容量、TTL 3/7 天。 可落地的路径：先用本地 NVMe + 内存索引做一个单机参考实现，把 KVSTORE_GRAIN=128 的对齐语义和哈希链匹配跑通，再逐步替换成 3FS 或任何支持 EC 的分布式存储。难度：中高，但接口齐备——这是社区最值得先做的事。

（2）补齐 DP 级三层负载均衡。 缺 Prefill/Decode/VE LB，以及 Entrypoint 侧的 KV 目录（文档说它是"logically unified, physically partitioned by key across Entrypoint machines"）。 文档给了两个可直接实现的算法：Decode LB 的 DualPath 双维度启发式；Prefill LB 的"4 档分桶 + 二次代价模型反解 chunk 大小 + 组内 all-reduce MIN 预算"。后者尤其可复用——代价模型 cost(q, c) 的形式是公开的，csrc/engine/flops_estimator.hpp 里已有 FP8 等效 FLOPs 的估算器可以作为拟合起点。 难度：高，因为它天然需要集群视角（服务发现、实例健康、KV 位置）。但一个"单机多副本"的简化版就能验证算法。

（3）从"单模型引擎"走向"模型族引擎"。 v4_modules/__init__.py 的四角色划分（参数容器 / 运行时前向 / 张量算子 / 拓扑装配）已经是一个成熟度不低的内部契约，builder.py 也把"从 config 派生拓扑"独立出来了。把它提升为公开的模型接入接口（类似 vLLM 的 register_model 但保留"结构可表达"的能力），是让 V4.2、V5 以及第三方模型复用的前提。 官方已承诺"new models’ inference code and subsequent optimizations will continue to be open-sourced here"，所以这件事有明确的需求方。难度：中，主要是接口设计而非算法。

（4）并行策略补全：TP 与 CP。 文档原话是"TP and CP are not used for now"（:53）。Prefill 推荐 EP2–8/DP2–8，Decode 推荐 EP64–128/DP64–128。这个选择在 DeepSeek 的集群里是对的（Wide-EP 显著降低 MoE 专家的每步访存、配合 EPLB 提升 TPS），但它把门槛钉死在"64 卡起"。 context_parallel_size 在 ParallelConfig 里已经存在并被当作 model_parallel_size（config.py:171），tilelang_ops/prepare_cp_mirror.py 也有 CP 相关 kernel——CP 的骨架在，TP 几乎没有。 难度：高。补 TP 不是加一个 all-reduce，而是要重新审视"每层只存一份 KV"的跨层共享假设。

（5）硬件面：让非 SM100 的卡也有路。 当前只支持 SM100 与昇腾 950，而 FP4 量化假设写得很细（main KV per-16 e4m3，相当于去掉 global scale 的 NVFP4；indexer KV 与 MoE 路由专家都是 per-32 ue8m0 MXFP4）。这些格式在 H100/A100 上没有原生支持。 可行的中间态是增加一个 BF16/FP8 路径：主流 kernel 已有 FP8 实现（DeepGEMM），主要工作量在 KV 格式与 fused transform 的分支。tilelang_ops/ 的 cuda/ 与 asc/ 双目录结构已经证明了这套代码组织能容纳"第三套后端"。 难度：高，但收益是数量级的用户面扩大。

（6）服务层：HTTP、流式、多租户、限流。 目前完全外置。一个务实的做法是不自建，而是与 deepseek-recipe 的 server-rs / server-py 示例对齐，产出一个"DeepInfer + recipe + 简单路由"的参考部署（docker compose 级别），把 §2.4 那张缺口表填掉。这比在引擎里造一个 server 更符合本项目"分层清楚"的性格。难度：中。

（7）观测性：从日志到指标。 已有 engine/stats_logger.py（20 KB）、deepinfer/timer.py（cpu/cuda/cuda_graph 三套计时器）、engine/profile_server.py（可选 fastapi 端点）、examples/_trace.py（torch profiler）。 缺的是统一口径的指标导出：KV 命中率、SWA 命中率、moe_lb、每步 bubble 时间、ZOS 重叠率、chunk 实际大小 vs 预算、投机接受长度。这些量代码里都算出来了（python_timer.record("bubble")、update_mtp_runtime 等），只是没有 Prometheus/OTel 出口。难度：低，是性价比最高的社区贡献点之一。

（8）KV 生命周期的策略空间。 现在的 TTL 是人工设定的 3 天/7 天，驱逐完全交给 3FS 的 LRU GC（“There is no expiry logic on the kvcache-store side”）。文档自己给了反例：3 天 TTL 会让 chatbot 流量损失 0.5%–1% 的命中率，因为 chatbot 的命中长尾很重（少量 KV 被读非常多次）。 这里有一个明确的研究问题：按"每字节 KV 的边际命中收益"做分区 TTL/淘汰，而不是全局 LRU。命中统计已经在 Entrypoint 的目录里，做策略替换的接口也清楚。难度：中，学术价值高。

（9）投机解码的最优性。 DSpark 现在是"按预期验收率排序取前缀"，是一个有界启发式：DEEPINFER_USE_NEXTN_DECODE 默认 5，DEEPINFER_MAX_BACKTRACKING_TOKENS 取 max(设定值, nextn+1)。表里的数据也显示低并发时变长验证甚至略负（并发 16 时 −1.5%，并发 32 时 +1.6%）。 真正的最优问题应该是"在给定验证算力预算与硬件吞吐曲线下，联合决定每请求的草稿长度与批内路由"，而现在的实现是"先按批组好，再在 GPU 上按置信度分配"。把这两步联合优化是一个有发表价值的方向，且 decode_perf_cache.py（把实测吞吐曲线按 TTL 1 小时缓存）已经提供了硬件曲线这一半的输入。难度：高。

（10）工程门槛与可复现性。 构建链路要求：Linux + GCC ≥13 + Python 3.12 + CUDA 13/CANN 9.2 + 编译 6 个 vendor 子模块 + 强制 NCCL 2.32.3。这几乎排除了"试用"级别的用户。 最直接的改善是提供预编译 wheel 与容器镜像（按 nvidia/ascend × 架构分发，正好对应 deepinfer/_so/<platform>-<arch>/ 现有的目录约定）。其次是让 load_format="dummy"（已在 config.py:115-119 支持，用于 profiling 合成权重）成为"无 GPU 也能跑通流程"的正式路径。 难度：低到中，但对采用率影响最大。

4.3 一句话结论与阅读路线

结论： DeepInfer 的价值不在于"能跑多快"，而在于它把"一个团队如何用商用硬件把一个特定模型服务到 155 万亿 token/日"这件事的完整推理链摊开给你看——包括它刻意不开源的那两块，也在文档里说了为什么不开、接口留在哪。

推荐阅读路线（按性价比排序）：

1. doc/inference_deployment_qa.md（312 行）——先读这个，它比 README 信息密度高一个数量级
2. doc/images/architecture_overview.png——一张图看清 E/P/D + 三层 LB + 3FS 的位置关系
3. csrc/engine/scheduler.hpp:275-345（_schedule 的 NORMAL 分支）——理解 P/D 混布为什么要投票
4. deepinfer/model_executor/models/v4_modules/__init__.py（42 行）——理解模型层的四角色约定
5. deepinfer/worker/cache_engine.py 的 select_state_slice_to_store / GPUCacheInterface._carve_state_views——理解 SWA state 的 128 粒度布局（KVSTORE_GRAIN = 128 定义在第 32 行）
6. examples/benchmark.py + examples/README.md——理解三类 Pareto 曲线（prefill / decode / encoder）怎么测、口径是什么
7. eval/README.md——理解 Pass@1 的完整评测流程与参考成绩（AIME 2026 99.9% / GPQA-D 93.4% / HLE-Full 42.6%）

附录：关键事实速查