# DeepInfer 项目概览报告


<!--more-->

# DeepInfer 项目概览报告

> 阅读对象：本地仓库 `DeepInfer-main/`（对应 [deepseek-ai/DeepInfer](https://github.com/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：仓库顶层布局

```
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 自己的算子仓库，且**按硬件成对出现**：

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

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

### 2.2 L1：分层视图

| 层 | 主要文件 | 职责 |
| :--- | :--- | :--- |
| **入口层** | `deepinfer/entrypoints/llm.py`、`deepinfer/engine/llm_engine.py` | 离线 `LLM` 类、请求生命周期、`step()` 主循环 |
| **调度层** | `csrc/engine/scheduler.hpp`、`sequence.hpp`、`block_manager.hpp`、`zos.cpp` | batch 组批、抢占、chunk 切分、block/state 分配、零开销调度 |
| **执行层** | `deepinfer/worker/{worker,model_runner,cache_engine,cuda_graph_runner}.py` | 单卡执行封装、输入装配、KV/state 显存管理、CUDA Graph 捕获 |
| **模型层** | `model_executor/models/deepseek.py` → `v4_modules/*` → `layers/*` | V4.1 参数容器、前向实现、采样器 |
| **算子层** | `model_executor/layers/tilelang_ops/{cuda,asc}/*` + `vendor/` + `npu_ops/` | TileLang kernel 的成套双后端实现 + 第三方算子 |
| **平台层** | `deep_cluster/`、`deepinfer/platform.py` | 设备发现、NUMA 绑定、平台能力抽象 |
| **协议层** | `structure_output/`、`hugging_face/`、`distributed/shm_broadcast.py` | 结构化输出（xgrammar）、权重与 tokenizer 加载、进程间元数据广播 |

模型层的内部再分层非常清晰，`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：一次请求的调用链

```
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：仓库内 / 仓库外的物理边界

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

| 组件 | 是否在仓库内 | 证据 |
| :--- | :--- | :--- |
| 引擎、调度器、模型、算子、EPLB | ✅ 完整实现 | `csrc/`、`model_executor/`、`model_executor/eplb/` |
| 结构化输出（编译 + 运行时） | ✅ 完整 | `structure_output/`（xgrammar 封装 + 异步 builder） |
| 离线批量 API | ✅ 完整 | `entrypoints/llm.py`、`examples/offline_inference.py` |
| benchmark / Pareto sweep / 评测 | ✅ 完整 | `examples/benchmark.py`、`examples/sweep/`、`eval/` |
| **DP 级三层负载均衡** | ❌ 仅文档描述 | 仓库内无 Prefill/Decode/VE LB 实现；`doc/...:107-127` 只有设计 |
| **`kvcache-store` 服务端** | ❌ 仅客户端钩子 | `llm_engine.py:800-823` 的 `init_pd_split`/`do_save_kvcache`/`pe_prepare_save_kvcache` 均为 `"虚函数"`，空实现 |
| **P/D 分离编排** | ⚠️ 部分 | `EngineType.PROMPT`/`COMPLETION` 已实现，可按阶段独立部署；跨阶段 KV 搬运需自己接 |
| **HTTP / OpenAI 兼容层** | ❌ 无 | 全仓库唯一的 fastapi/uvicorn 用法在 `engine/profile_server.py`，且是默认关闭的调试用 profiling 端点 |
| **API 协议编解码** | ❌ 外部仓库 | 由 `deepseek-recipe`（MIT）承担，README 明确"caller must provide HTTP transport" |

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

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

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

### 2.5 L4：构建与依赖图

```
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 把它拆成三层介质，各有明确的归属、复制策略与管理方：

| 内容 | 介质 | 复制 | 管理方 |
| :--- | :--- | :--- | :--- |
| 在飞 block KV + state KV | 推理 GPU 显存 | — | P/D 实例 |
| SWA KV | 本机主机内存 | 单副本，不落 3FS | P/D 实例；跨机命中走 400G VLAN 回拉 |
| 分页 block KV | 专用 3FS 集群 | 无应用层复制，靠 3FS 自身冗余 | `kvcache-store` |

（表来自 `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/` 里有大量在通用引擎中找不到对应实现的东西：

| 模块 | 是什么 | 在通用引擎里的处境 |
| :--- | :--- | :--- |
| `hc.py` / `mhc_*.py` | mHC：把 `hc_mult` 路残差主干在每个子层前后做混合，末端 `HCHead` 再折回单路隐状态 | 无对应实现 |
| `engram.py` | Engram：把 token id 哈希成 n-gram，跨 rank 聚合分片表，经融合 KV 投影后门控进隐状态 | 无对应实现 |
| `compress.py` + `index_attn.py` | 压缩 KV（门控低秩）+ 稀疏 indexer（`relu(q@k)*w` 打分后 top-k 选块） | 无对应实现 |
| `cross_layer.py` | share-KV / share-index / share-index-K 三种跨层共享的拓扑推导与校验 | 无对应实现 |
| `yoco.py` | YOCO 层分段：`[0, n_prefill_layer)` 是 encoder，其后是 consumer 层 | 无对应实现 |
| `mlp.py` | 路由专家 + 共享专家 + gate，运行时重排 expert 权重 | 有，但布局不同 |

**证据。** `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` 的两个变量正是这个机制的旋钮：

```
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适配了昇腾"，而是同一仓库里两套并行实现。

**证据（四层都成对）：**

| 层 | NVIDIA | 昇腾 |
| :--- | :--- | :--- |
| TileLang kernel | `tilelang_ops/cuda/`（30 文件） | `tilelang_ops/asc/`（30 文件） |
| vendor 算子 | `DeepEP` / `DeepGEMM` / `FlashMLA` / `DeepSelect` | `DeepEP-Ascend` / `DeepGEMM-Ascend` |
| C++ 扩展 | `cuda_utils.cpp` + `graph_timer_cuda.cpp` | `asc_utils.cpp` + `graph_timer_asc.cpp` |
| 通信后端 | NCCL（`max_ctas`/`min_ctas` 钉到 `DEEPINFER_NCCL_NUM_SMS`=32） | HCCL |

`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，卡内/卡间重排专家副本） | ✅ 完整 | `model_executor/eplb/{__init__,policy,shuffle}.py`（约 37 KB） |
| **副本级**（Prefill / Decode / VE Load Balancer） | ❌ 仅文档 | `doc/...:107-127` |

**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`）：**

| 能力 | 开关 | 约束 |
| :--- | :--- | :--- |
| Batch-invariant 推理 | `DEEPINFER_BATCH_INVARIANT=1` | 仅 NVIDIA；不同 batch 组成下模型结果完全一致 |
| 确定性采样 | `DEEPINFER_BATCH_INVARIANT=1 DEEPINFER_ENABLE_MTP_VARLEN=0` + 固定 `seq_seed` | 仅 NVIDIA |
| 结构化输出 | `so_args={"schema_type", "schema"}` | regex / JSON Schema / EBNF |

**回归方式。** `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%）

---

## 附录：关键事实速查

| 项 | 数值 / 事实 | 来源 |
| :--- | :--- | :--- |
| 开源许可 | Apache-2.0（仓库）；配套 `deepseek-recipe` 为 MIT | `LICENSE`、业界公告 |
| 建仓 / 发布 | 2026-10-16 / 2026-10-20 | GitHub API |
| 支持模型 | 仅 `DeepSeek-V4.1-Flash`（默认 checkpoint `-1017`） | `README.md` |
| 硬件 | NVIDIA SM100、Huawei Ascend 950 | `README.md` |
| 环境 | Linux、GCC ≥13、Python 3.12、CUDA 13 / CANN 9.2 | `README.md` |
| 单集群吞吐 | 155 万亿 token/日 | `doc/inference_deployment_qa.md:17` |
| 单位经济 | $38 / GPU-hour（谷电价） | 同上 |
| 并行策略 | Prefill: EP2–8/DP2–8；Decode: EP64–128/DP64–128（Wide-EP）；不用 TP/CP | `doc/...:47-53` |
| 最小高效部署单元 | 64 卡 P + 64 卡 D（1P+1D 可部署，实际 P:D ≈ 1:1） | `doc/...:67` |
| KV 存储 | ~100 PiB SSD；TTL 3 天（Agent）/ 7 天（chatbot） | `doc/...:153` |
| KV 容量比 | ~14.4 TB / GPU（逻辑） | `doc/...:81` |
| SWA KV / 序列 | 20 层 × 134 token × 528 B = 1,415,040 B | `doc/...:169` |
| 落盘粒度 | 128 token（`KVSTORE_GRAIN`） | `cache_engine.py:787` |
| EPLB 实测 `moe_lb` | 1.074（prefill EP4DP4）/ 1.090（decode） | `doc/...:99,103` |
| 三层 LB 效果 | 约达完美均衡离线 benchmark 的 80% | `doc/...:127` |
| 变长投机验证 | 验证数 5.00→3.2–4.4，接受长度 3.29 vs 3.20，+4–7%（60 并发） | `doc/...:247` |
| 评测参考成绩 | AIME 2026 99.9% / GPQA-D 93.4% / HLE-Full 42.6%（Pass@1） | `eval/README.md:215-217` |
| 代码规模 | 342 文件（不含 vendor）；最大文件 `scheduler.hpp` 1329 行 | 本地统计 |
| 外部依赖 | torch 2.13.0、transformers 4.50.0、xgrammar 0.2.2、tilelang 0.1.15、tile-kernels 2.1.0、flash-attn-4 4.0.0b34 | `requirements*.txt` |


---

> 作者: [破戒僧](https://onlyblacktea.cn)  
> URL: https://onlyblacktea.cn/posts/deepinfer/