把专家并行从启动参数变成运行期资源：FastEEP 如何在不重建实例的前提下把 EP 从 4 改到 16 - 技术文章

除另行标注外，本文引用的外部事实与页面核验均以 2026-10-21（北京时间）为时点。

先看两组公开 trace 的数字，它们说明负载本身有多不均匀。

论文引用 BurstGPT trace（Wang 等，2025，两个月 140 万条请求）：其中 59% 的十秒窗口里一条新请求都没有；忙时到达率的 99 分位相对中位数有 25 倍摆幅。

论文另引用微软 Azure 的公开对话 trace（Azure/AzurePublicDataset，采集于 2024 年 5 月 10–19 日，CC-BY-4.0，一周 2700 万条请求）：按论文的表述，波动是 1.7 倍。

这两个数是两种不同统计口径（前者是「十秒窗口的空闲比例 + 到达率分位比」，后者是论文给出的一周尺度波动倍数），所以只能一起说明「流量极不均匀」，不能直接换算成「尖峰持续多久」。

真正支撑「扩容动作必须快」的是论文另一句话：现有的实例内 resize 要 36.4–92.0 秒，明显长于它所应对的抖动窗口的中位持续时间。这是论文的判断，本文沿用它的口径。

这篇文章要讲的 FastEEP，就是冲这个缝来的。作者是上海交通大学 IPADS（Institute of Parallel and Distributed Systems，并行与分布式系统研究所）的 Yunfei Feng、Jingkai He（共同一作）、Shiqi Liu、Cheng Zhang、Dong Du、Erhu Feng，2026 年 10 月 13 日提交到 arXiv，编号 2610.18451，分类 cs.DC，通讯作者邮箱前缀 fengerhu1。

本文按「先给直觉 → 符号约定 → 形式化 → 拆机制 → 读实验 → 划边界」的顺序讲，并把能自己算的数算一遍、能核的来源都点出来。

结论先行：它同时改变了变更的粒度与开销

一句话概括：FastEEP 把「改一次 EP 规模」从销毁并重建整个服务实例，降级成搬动几个 worker 的运行期状态。

这里先解释 EP：专家并行（Expert Parallel，EP）是把 MoE 的专家权重切到多张卡上，改 EP 就是改变「同一个实例里，承载专家的 rank 集合有多大」，而不是新起一个实例。

具体到数字，论文给出三组对照，口径必须先分清，否则很容易把三个倍数混成一个：

一是相对 vLLM 现有的 Elastic EP 实现：FastEEP 单次 EP 变更 4.2–13.4 秒，基线是 36.4–92.0 秒；在所有「基线可运行」的模型-方向组合上，逐项加速比的区间是 4.5–11.7 倍。

二是把上述逐项倍数取几何平均，得到 7.0 倍；它的全异步变体（FullAsync）是 9.3 倍。

三是相对「杀掉实例再拉起」：整实例重启要 126–293 秒，几何平均加速 22.0 倍（FullAsync 29.3 倍）。

这三组数不是同一个东西：4.5–11.7 是区间，7.0 与 22.0 都是几何平均，而 22.0 换的是另一个基线——不能读成「某一次变更快 22 倍」。

顺带说明 FullAsync：按论文的表述，它把权重与专家的搬运、缓冲的预取/卸载都挪到那次极短的全局停顿之外，与推理并发执行，代价是短暂与推理争用带宽；所以它的加速比更高，但并不是「完全不停顿」。

比加速比更能说明问题的是停服时长：在真实负载下做 8→4 与 4→8，FastEEP 的零输出窗口是 4.2 秒与 7.3 秒；vLLM Elastic EP 是 31 秒与 44 秒；整实例重启是 141 秒与 149 秒。

代价写在论文的限制里：变更只能在启动时定下的最大拓扑内部进行；缩容后的 worker 虽然不占 GPU 算力，但进程、线程与元数据还活着，要占 CPU 侧资源与一小块恢复状态。

读之前先对齐符号与术语

这篇文章的叙述依赖几个记号，先把它们钉住，后面读公式会顺很多。表里的「后文出现处」用小节标题，方便回查。

记号 / 术语	含义	后文出现处
$\mathit{EP}$	专家并行度数，专家权重被切成的份数	全文
$\mathit{DP}$	数据并行度数，attention 侧按请求切分	「两条公式打架」一节
$\mathit{TP}$	张量并行度数	「两条公式打架」一节
$S$	当前参与计算的 rank 集合（活跃集）	「父通信域」与「影子 worker」两节
shadow worker（影子 worker）	进程、线程、虚拟地址与恢复所需设备状态都保留，但前向计算跳过、可回收缓冲的物理显存已释放的 worker	「影子 worker 与一次原子提交」一节
parent domain（父通信域）	启动时按最大 worker 集合建立的 NCCL 通信域，全程保留 rank 空间与 bootstrap（成员连接的初始化关系）上下文	「父通信域」一节
NIXL 掩码	设备端读的可达性掩码，控制 EP dispatch/combine 的 P2P 可见性	「父通信域」一节
CuMem 池	用 CUDA 虚拟内存管理（VMM）分配、让虚拟地址与物理页解耦的显存池	「地址稳定」一节
commit point（提交点）	一次性发布专家映射、数据面通信视图与请求路由归属的那个安全点	「影子 worker 与一次原子提交」一节
excursion（抖动窗口）	一次流量偏离持续的时间段，是扩容动作要匹配的时间尺度	「两条公式打架」一节
SPMD	Single Program Multiple Data：同一段程序在所有 rank 上并行执行，集合通信要求成员齐步进入	「三条刚性绑定」一节
TTFT / TPOT	首 token 延迟 / 除首 token 外的每 token 输出间隔，是服务级目标（SLO）常用的两个指标	「实验结果怎么读」一节
MaaS	Model-as-a-Service，按模型对外提供推理服务的形态，特点是多模型共卡、需求独立波动	「多模型共存」一节

还有一个约定要在开头说清：本文里「重建」指销毁实例并重新初始化运行期环境；「重配」指保留实例、只改变参与计算的资源。

FastEEP 主张的是后者，但它不是「零成本」，而是把成本从「不可掩盖的停服」挪到「可异步的准备与回收」。

为什么大 EP 又省钱又难扩：两条公式打架

先讲清楚 EP 到底在买什么。

MoE 模型的参数绝大多数是专家参数。论文引用 DeepSeek-V3 的说法是专家参数占总参数的 97%——这是一条转述（FastEEP 论文 §2.1 正文），本文另外直接核对了 DeepSeek-V3 技术报告摘要里的模型规模：671B 总参数、每 token 激活 37B。

专家权重按 EP 切到多张卡上，单卡权重占用随 EP 增大而下降，省下来的显存拿去做 KV cache。

要注意「按 EP 切分的只是专家参数」，非专家部分是每个 DP 副本都要重复的，所以单卡权重占用应写成两项之和：

M_{w}(\mathit{EP}) = \frac{P_{\text{exp}} \cdot b}{\mathit{EP}} + M_{\text{repl}},\qquad M_{\text{repl}} = (P_{\text{total}} - P_{\text{exp}}) \cdot b

其中 $P_{\text{total}}$ 是总参数量， $P_{\text{exp}}$ 是专家参数量， $b$ 是每参数字节数（FP8 为 1，BF16 为 2）， $M_{\text{repl}}$ 是不随 EP 切分的那部分（attention 权重、embedding、共享专家）。

举个数量级的例子（本文自行推导，假设全部参数都按 EP 切分、即忽略 $M_{\text{repl}}$ ，只为说明趋势）：671B 参数按 FP8 存储是 671 GB， $\mathit{EP}=8$ 时每卡摊到约 84 GB， $\mathit{EP}=16$ 约 42 GB， $\mathit{EP}=64$ 约 10.5 GB。

真实部署里 $M_{\text{repl}}$ 会把这条曲线整体抬高，所以上面的数字只是乐观值（真实占用更大、是下界），不能用来算显存预算。

KV cache 容量与并发的关系是：

M_{kv}(\mathit{EP}) = C_{\text{HBM}} - M_{w}(\mathit{EP}) - M_{\text{ctx}}

N_{\text{conc}} = \frac{\mathit{EP} \cdot M_{kv}(\mathit{EP})}{s \cdot L}

这里 $C_{\text{HBM}}$ 是单卡显存， $M_{\text{ctx}}$ 是运行期上下文占用（下一节会看到它有多大）， $s$ 是每 token 跨全部层的 KV 字节数， $L$ 是序列长度。

$s$ 可以再展开成模型结构： $s = 2 \cdot N_{\text{layer}} \cdot n_{kv} \cdot d_{\text{head}} \cdot b_{kv}$ ，其中 $N_{\text{layer}}$ 是层数、 $n_{kv}$ 是 KV 头数、 $d_{\text{head}}$ 是每头维度、 $b_{kv}$ 是 KV 每元素字节数，系数 2 对应 K 与 V。

还要提醒这条并发估算的隐含假设：它把 $\mathit{EP}$ 个单位都当成能提供 KV 容量的单元，也就是默认 attention 侧按 DP 复制、没有额外的 TP 切分。如果你的部署是 $\mathit{TP} > 1$ ，每个 rank 的 KV 只是其中一片，这条式子要重写成按 DP 计。

$\mathit{EP}$ 出现在分子上、又通过 $M_{w}$ 反向影响 $M_{kv}$ ，所以「加卡」在 KV 容量上是双重收益——这也是宽 EP 成为生产默认配置的原因。

但宽 EP 会把弹性掐死：一个实例横跨几十张卡，任何按实例粒度的扩缩容，搬动的资源量都远大于需求变化量。

论文用一张示意图讲这件事：宽 EP 实例动辄占用几十张卡（论文原文为 tens of GPUs），实例级伸缩搬动的资源单位远粗于需求变化——需求只差一档，你却只能整块搬。

于是生产上形成一个二选一：按峰值配，低谷大量闲置；按中位数配，尖峰打穿服务级目标（SLO）。

论文据此提出弹性化的三条要求：粒度要细（到单卡）、要快（快过它要对付的抖动窗口）、要可逆（尖峰过去后能立刻还卡）。

第三条常被忽略，但在多租户集群里它是决定性的：加进来的卡如果不能快速还回去，那就等于把成本永久固化。

一次冷启动的账单：单 rank 剖析里，只有 10 秒在搬权重

论文的事先分析很扎实，但测量对象要说准：它把 DeepSeek-V2-Lite 的一个 rank 从进程启动到稳态服务拆成六个阶段，量出各自耗时。

这六项是单 rank 冷启动的剖析，不是「一次在线 resize 的现场计时」；论文另外还给了全实例启动的区间 156–267 秒（跨不同配置），以及 §5.2 里「整实例重启」基线的 126–293 秒。三个数不是一回事，不能互换着引用。

阶段	耗时	说明
进程启动（worker、executor、kernel 加载）	24 秒	与 EP 无关
通信域构造（含 RDMA 缓冲注册）	13 秒	只有成员集与 EP 有关
权重加载	10 秒	唯一搬模型数据的阶段
算子 profiling 与编译	32 秒	与 EP 无关
KV 容量协商	6 秒	与 EP 无关
CUDA Graph 捕获	36 秒	与 EP 无关（这一项是本次启动样本的值；论文另给的「变更时重捕」区间是 25–82 秒）

一个细节必须点出来：这六项相加是 121 秒，而论文正文给的总时长是 126 秒，中间有 5 秒没有在这些条目里单列。

本文无法确定这 5 秒的构成（论文列项与总时长未完全对齐），所以两个数都保留、不做凑整——差不影响结论，但把它抹平会让复算的人对不上账。

另一个更能说明问题的事实：在 t=47 秒、权重还没落地之前，这张卡上已经吃掉了 6.32 GiB 显存；到稳态时，运行期上下文占 10.90 GiB，比 9.62 GiB 的模型权重还大。

换句话说，这次单 rank 启动里真正在搬模型数据的时间只占 $10/126 \approx 7.9\%$ ，而光 CUDA Graph 捕获一项就是 $36/126 \approx 28.6\%$ 。

把这件事写成公式：

T_{\text{rebuild}} = T_{\text{proc}} + T_{\text{comm}} + T_{w} + T_{\text{prof}} + T_{kv} + T_{\text{graph}}

\frac{T_{w}}{T_{\text{rebuild}}} \approx 7.9\%,\qquad \frac{T_{\text{graph}}}{T_{\text{rebuild}}} \approx 28.6\%

这两个比值都是用 126 秒做分母算的；若改用六项相加的 121 秒，则分别是 8.3% 与 29.8%。

这就是全文的支点：在这个样本上，重建的成本主要来自「不需要随 EP 变更而重建」的那部分状态。

按这个观察，工程上该问的问题就变成：既然这些东西与 EP 无关，为什么改变 EP 要把它们全部销毁？

论文用一张表把状态逐项对照：只有最后一行（专家放置与请求路由）在每次变更时都必须重写，其余各行都能以某种形式复用。

状态	重建成本	是否随 EP 变化	今天的做法	FastEEP 的做法
worker 进程、executor、已编译 kernel	24 秒	否	销毁	保留
通信域、RDMA 缓冲	13 秒	仅成员集	重建	派生子域
CUDA context	0.32 GiB	否	销毁	保留
CUDA Graph	36 秒、1.47 GiB	否	重新捕获	直接回放
权重、KV cache、workspace	9.62 / 71.04 GiB	仅尺寸	重新加载	重新填页
expert 放置与请求路由	—	是	重建	就地改写

表里「是否随 EP 变化」这一列要读细一点：它问的是「为了改变 EP，这份状态的结构是否必须重建」，而不是「尺寸是否随 EP 变」。

所以通信域那一行写「仅成员集」——成员要换，但通信域本身、rank 身份、bootstrap 上下文可以沿用；权重那一行写「仅尺寸」——字节数变，但分配池与地址可以沿用。

顺着这张表，也就解释了为什么 vLLM 的 Elastic EP 仍然要 47 秒缩、75 秒扩：它避免了整个服务的进程重启，但仍会新建/移除 Ray engine、重建进程组与通信域、重捕 CUDA Graph——正是表里那些不需要重建的行。

三条刚性绑定：为什么今天的引擎只能重建、不能重配

论文把「必须重建」的原因归成三条结构性绑定，这三条都不是硬件限制，而是软件实现约定的结果。

第一条是通信成员的不可变绑定。NCCL 集合通信是 SPMD 语义：同一次集合调用，所有成员 rank 必须都进入，否则先到的 rank 就在原地等。

把「等」写实一点：设一次调用在一次统一的发起时刻开始， $t_{r,i}$ 是 rank $r$ 完成第 $i$ 次集合调用的时刻，则这次调用从发起到全员完成的时间是：

T_{\text{coll}, i} = \max_{r \in S} t_{r,i}

只要 $S$ 里有任何一个 rank 永远不进入，上式在理想模型下没有有限解（ $T_{\text{coll},i} = \infty$ ）；真实系统里则会先挂住、再由超时或看门狗报错。

所以对集合通信来说，把不活跃 rank「屏蔽掉」是无效的：屏蔽只对点对点有效，对集合通信无效。这是后面第一个设计的分水岭。

第二条是地址的静态绑定。CUDA Graph 的执行体里把 kernel 参数与设备指针固化了；RDMA 缓冲注册也绑定虚拟地址。

只要缓冲被释放后重新分配、地址变了，即使 shape 与内容完全不变，图也不能回放，必须重新捕获——而重新捕获是 25–82 秒（论文给的区间，与上表里那 36 秒是不同场景：一个是本次启动的捕获，一个是变更时的重捕）。

第三条是 worker 执行的紧耦合。一个 worker 同时承担三件事：跑本地请求的 attention、维护本地 KV cache、替远端 rank 执行被路由过来的专家计算。

所以「把这个 worker 的本地请求排空」并不等于「它不再参与通信」——远端还在往它的专家发 token。

结论很硬：缩容时必须同时关掉两条路径（请求路径与专家路径），而且不能靠「等它自己闲下来」。

设计一：父通信域 + 子组切换，把通信成员变成可变量

第一个设计解决「通信成员」这一层。

FastEEP 在启动时按最大 worker 集合建立一个父通信域，并且全程保留它的 rank 空间、bootstrap 关系与控制面状态——论文给的量级是控制面不到 300 MB。

每次 resize 时，从父域派生一个只含活跃 worker 的工作子组；缩容装小一点的子组，扩容装大一点的子组。

NCCL 本身提供了这个能力：ncclCommSplit(comm, color, key, *newcomm, config) 是一个集合调用，把同色（同 color）的 rank 组成新通信域。

这里有个关键细节，也是很多人第一次读会卡住的地方：split 是「在父域上」的集合调用，所以父域的所有成员（包括当前休眠的影子 worker）都必须进入。

影子 worker 之所以必须保活，正是为了这件事：它要能参加控制面的成员变更集合操作，只是不参加活跃子组的数据面集合通信。

换句话说，「不参与」要分清是哪一层：数据面（前向计算依赖的 DP 集合通信与 EP dispatch/combine）它不参与；控制面的成员变更它仍要参与。

对 EP 的 dispatch/combine 这类点对点语义，FastEEP 换另一种办法：保留最大规模的逻辑 rank 空间，用一个设备端掩码控制可见性。

NIXL（NVIDIA Inference Xfer Library，仓库在 ai-dynamo/nixl，跨节点走 UCX）把远端元数据交换与端点创建分开，端点是懒创建的，所以增删一个 rank 只影响它自己的端点，其他连接不动。

于是同一份「活跃集」有两种表达：NCCL 子组（让集合通信只等活跃者）+ NIXL 掩码（让 P2P 只发给活跃者）。

两种表达必须在同一个安全点一起发布，否则会出现「集合通信认为 A 在、P2P 认为 A 不在」的分裂视图。

flowchart TB
  P["父通信域：最大 rank 空间 + bootstrap 上下文（控制面 <300 MB）"]
  subgraph ACT["活跃集 S：数据面成员"]
    A0["rank0：attention + 本地专家槽"]
    A1["rank1：attention + 本地专家槽"]
  end
  subgraph SH["影子 worker"]
    S2["rank2：进程/线程/虚拟地址保留，可回收缓冲的物理显存已释放"]
    S3["rank3：同上"]
  end
  P -->|"ncclCommSplit(color, key) 派生子组"| ACT
  P -.->|"保留身份，仍参加父域上的成员变更"| SH
  M["NIXL 掩码：dispatch/combine 的可见集"]
  ACT --- M
  SH -.->|"被掩码隐藏，收不到 token"| M

这张图说明了什么：一份最大规模的通信骨架在启动时建立一次，之后「谁在算」变成了可切换的视图；缩容的 rank 不是被销毁，而是从数据面两个视图里同时被摘出去，身份与地址原地保留。

flowchart LR
  subgraph L0["EP=8（最大规模）"]
    U["rank0..rank7 均在父域"]
  end
  subgraph L1["收缩到 EP=4"]
    V["子组 A：rank0..rank3（active）"]
    W["子组外：rank4..rank7（masked）"]
  end
  subgraph L2["再扩张到 EP=6"]
    X["子组 B：rank0..rank5"]
  end
  L0 -->|"split + 掩码"| L1
  L1 -->|"再次 split，复用父域身份"| L2

这张图说明了什么：层次化派生让「缩」和「扩」都不需要重新 bootstrap——每次都是从父域再切一刀，而不是把整个通信域拆了重建。

关于 NCCL 版本这里补一句本文核过的边界：NCCL 官方文档当前版本（2.32.3，2026-10-21 打开）的 API 索引里已经出现 ncclCommShrink、ncclCommGrow、ncclCommRevoke 等函数名，但本文只核对了这份索引与 ncclCommSplit 的语义，没有实测 shrink/grow 的行为，也没有确认它们在论文实现的 2.28.9 上是否已存在；FastEEP 论文自述走的是 ncclCommSplit 这条路径。

设计二：地址稳定，物理页可回收

第二个设计解决「显存」，目标很精确：保留逻辑身份，回收物理页。

手段是 CUDA 虚拟内存管理。FastEEP 按用途建立几个独立的 CuMem 池：共享权重、专家权重、KV cache、CUDA Graph 缓冲。

缩容时不「释放对象」，而是「解映射」：物理页还给设备，虚拟地址、张量元数据与对象引用一律保留；扩容时在同一批地址上映射新页。

VMM 的动作要分清楚，不能压成一个「释放」：先预留一段虚拟地址区间，再把物理分配句柄映射上去；回收时先解映射（物理页归还驱动），再决定是否释放句柄，而预留的虚拟地址区间可以继续留着。

本文的推测是论文的实现就落在这条路径上：解映射 + 释放句柄，但保留虚拟地址，于是「同一个指针」在唤醒后仍然有效。

CUDA 驱动 API 里能核到的是其中两项：cuMemMap 把分配句柄映射到一段预留的虚拟地址范围；cuMemRelease 释放的是句柄，而不是地址（本轮在官方文档核对的就是这两条语义，完整的生命周期调用本文没有逐项实测）。

地址稳定带来的好处是连锁的：kernel 参数不用改、张量指针不用改、分配器对象不用重建，于是「唤醒」不需要重建运行期结构。

哪些内容必须保住、哪些可以丢，论文按语义分类处理：

非专家权重（attention 等，本来就在 DP 副本上重复）直接丢掉，唤醒时从活跃 rank 走 NVLink 拷回来；
专家权重同理走设备到设备拷贝，控制面先算出目标放置与源-目的 rank 对，保证每个逻辑专家在活跃集里有副本；
请求相关的 KV cache 按块三选一：丢弃后用 prompt + 已生成 token 重算、异步迁移到活跃 rank、或者换出到 pinned host 内存（页锁定内存，可在 CPU 与 GPU 间异步拷贝）；
CUDA Graph 缓冲单独走专用池：捕获期分配就落在这个池里，休眠时备份内容、解映射物理页，唤醒时在同地址映射新页并恢复内容，于是图可直接回放、无需重捕；
通信缓冲的回收必须等远端访问停下来，否则「本地释放、远端还在写」会把写打到重新分配的页上。

KV 三种策略的选择可以写成一个比较：

T_{\text{recompute}} = n_{\text{tok}} \cdot \tau_{\text{prefill}}, \qquad T_{\text{migrate}} = \frac{S_{\text{kv}}}{B_{\text{nvlink}}}

$n_{\text{tok}}$ 是块内 token 数、 $\tau_{\text{prefill}}$ 是重算的每 token 前向成本、 $S_{\text{kv}}$ 是该块的字节数、 $B_{\text{nvlink}}$ 是链路有效带宽。

工程上还有一个隐含的舒服之处：权重回填走的是活跃 rank 之间的 NVLink，而不是主机内存或网络存储，因此不受「重启时抢同一块存储」的带宽坍缩影响。

不过这里可以算一笔账，看看阶段耗时大概落在什么量级（本文自行推导的假设，不是论文数据）。

扩容阶段论文给「回填权重」0.85 秒；若假设回填量接近 9.62 GiB 的单卡权重足迹，则等效速率约 12 GB/s 量级，而测试床的节点内 NVLink 标称是 900 GB/s。

两者差近两个数量级，据此最多能提一个待验证的猜测：这个阶段很可能不是被裸带宽限住，而是被分批拷贝、注册与元数据开销限住。

但仅凭阶段计时无法证实这一点——要确认得测真实字节数与有效路径。所以这条只作为「要往哪里查」的线索，不作为结论。

设计三：影子 worker 与一次原子提交

前两个设计把「通信视图」和「物理页」变成可变量，第三个设计负责「让 worker 在不变身份的前提下退出/回归」。

FastEEP 引入影子 worker 这个状态：进程、CPU 线程、executor、虚拟地址、控制状态全部保留；权重、KV cache、Graph 缓冲、通信缓冲的物理显存被解映射或换出；只留下恢复所需的那部分设备状态。

所以它的显存不是归零，而是「可回收部分基本归还、恢复所需的一小部分保留」——把影子 worker 理解成「显存清零」会算错容量预算。

影子 worker 被排除在活跃子组之外、被 NIXL 掩码隐藏，因此收不到数据面的集合通信与 dispatch/combine 流量。

它的 CPU 线程只处理控制命令与屏障，不启动任何 attention/专家/通信 kernel——所以它不占 GPU 算力，但确实还活着（这也是论文承认的资源代价）。

缩容时要在两条路径上同时收口，而且要分清三个阶段：

第一阶段是准入控制：控制面把目标标成 draining，调度器不再给它派新请求；这一步只改「谁来接新请求」，不改变其他 rank 的数据面视图。

第二阶段是存量排空：近完成的请求本地跑完，其余请求要么迁移、要么按 KV 三策略处理，直到没有活跃请求与 KV 状态依赖这个 rank。

第三阶段才是数据面成员切换：重算逻辑专家到物理槽位的映射，优先复用活跃 worker 上的冗余副本，缺的专家走 NVLink 拷进空闲槽位；等每个逻辑专家都有副本了，再把新视图一次性发布。

这里出现一个约束：可用专家槽位决定了最小可行活跃规模（下式用 $E$ 表示逻辑专家数，避免与序列长度 $L$ 混用）：

N_{\min} = \left\lceil \frac{\rho \cdot E}{h} \right\rceil

$\rho$ 是冗余系数（为了热专家有副本，生产上常大于 1）， $h$ 是每个活跃 rank 能提供的专家槽位数。

这条式子是个容量约束：如果某个模型有 256 个路由专家、每 rank 只放 16 个槽、又想要 2 倍冗余，那么活跃 rank 少于 32 个就装不下全部必需专家。

注意这是本文用公式举的假设例子，不是论文的实验选点原因；也不意味着可以丢弃逻辑专家——同一模型要能用，所有必需专家都必须有位置，缩容撞到这条线就只能停、或者先改放置方案。

两条路径都收口之后，目标才能休眠并释放算力、通信与 HBM。

扩容严格按相反顺序：先恢复内存、再恢复专家放置、再恢复通信、最后才把这个 rank 纳入活跃 DP 集。

sequenceDiagram
  participant C as 控制面
  participant A as 活跃 worker
  participant T as 目标 worker
  participant N as 父通信域 / NIXL
  C->>T: 第一阶段：标记 draining，停止新请求准入
  T->>A: 第二阶段：迁移或重算 KV，迁走未完成请求
  C->>A: 第二阶段：拷贝缺失 expert 到空闲槽位（NVLink）
  Note over C,N: 第三阶段提交点：expert 映射 + 数据面通信视图 + 请求路由归属 一起发布
  C->>T: 允许休眠（跳过前向，保留进程与地址）
  C->>N: 先撤销可达性，再回收通信缓冲
  C->>T: 卸载权重/KV/Graph 缓冲，释放可回收物理页

这张图说明了什么：一次 resize 被拆成「可异步的准备（排空、迁 KV、拷专家）」与「一个极短的提交」，只有提交那一刻需要全局停顿；这三阶段的划分也解释了为什么「停止新请求」和「原子发布」并不矛盾——前者只动准入，后者的原子性只覆盖数据面三份视图。

整个事务可以写成三段：

T_{\text{scale}} = T_{\text{prepare}} + T_{\text{commit}} + T_{\text{reclaim}}

这里刻意不写「零输出窗口 = commit」：在线测量与分段实验的计时边界不同（前者是端到端从发出变更到恢复出 token，后者是按阶段打点），没有证据说两者等价。

机制层面的目标只是「把必须同步的那一段压到最短」，实际零输出窗口要看在线实验。

论文在 Qwen3-30B 上给的分段实测正好对得上这个结构：

方向	FastEEP 总耗时	主要分段
8 → 4（缩容）	4.21 秒	resize 2.81 秒（装新 expert 布局与通信视图）、卸载影子 worker 缓冲 0.92 秒、销毁 NIXL 缓冲 0.24 秒
4 → 8（扩容）	7.15 秒	唤醒 1.09 秒、回填权重 0.85 秒、恢复 NIXL 连通 2.50 秒、resize 2.55 秒

同一模型上 vLLM Elastic EP 的分段是：缩容 47 秒（standby 建组 10、专家重切分 9、编译与预热 11、图捕获 15），扩容 75 秒（standby 17、权重与 KV 同步 7、专家重切分 11、编译预热 14、图捕获 18）。

两份分段对照着看，差异的来源一清二楚：FastEEP 的账上根本没有「编译与图捕获」这两行。

顺带说明：vLLM 这两组分段相加是 45 秒与 67 秒，而论文给的总时长是 47 秒与 75 秒（原文以「其余控制步骤很小」概括差额），与前面 121 vs 126 的差额同类，本文同样只保留、不解释其构成。

至于上表 4.21 秒与在线实验 4.2 秒、7.15 秒与 7.3 秒为何如此接近：前者是分段打点的总耗时，后者是端到端零输出窗口，论文没有声明两者是同一场景的同一计时，所以本文只把它们并排列出，不据此下「零输出窗口就等于某一段」的结论。

正确性从哪里来：三条不变量

这类「成员变更 + 内存回收」的系统，正确性几乎总是归结为几条顺序不变量。论文把它们隐含在机制里，我把它显式写出来，便于评审自己的实现。

第一条：可达性不得长于内存生命周期。撤销可达性必须发生在释放内存之前；恢复内存必须发生在授予可达性之前。论文用一句话总结了这条规则，两个方向（集合通信、点对点）都遵守它。

第二条：提交点上，一个 rank 对三份数据面视图（专家映射、通信视图、请求路由归属）要么同时可见、要么同时不可见。这里的「请求路由归属」指「请求该由哪个 rank 接」，而不是「当前是否接受新请求」——后者是第一阶段就改掉的准入控制。

第三条：请求的生命周期与 worker 的资源生命周期解耦。调度器只在请求完成、或拿到了另一个执行位置之后才解绑请求，因此可以安全回收原 worker 的 KV cache。

这三条听起来像分布式系统的老生常谈（配置变更、成员视图、租约），但 GPU 场景更硬：没有日志可以回放，也没有第二次机会——KV cache 一旦被覆盖就是真的丢了，CUDA Graph 一旦地址失效就必须重捕几十秒。

这也是为什么论文反复强调「一次提交」而不是「逐步生效」：逐步生效在人看来更平滑，在正确性上却是把窗口开大了。

实验结果怎么读

测试床是两个节点，每节点 2 颗 Intel Xeon Platinum 8575C、8 张 NVIDIA H20-3e GPU（每张 140 GiB HBM）；型号与数量按论文记载，本文没有独立核验硬件配置。

节点内用 NVLink（标称 900 GB/s），跨节点用 8 张 BlueField-3 DPU 集成的 ConnectX-7 网卡（每张 400 Gb/s）——DPU 与网卡是集成关系，不是两套独立设备。

实现基于 vLLM v0.27.0、Python 3.12、PyTorch 2.11.0+cu130、NCCL 2.28.9、NIXL 1.2.0，新增约 5K 行生产代码，覆盖弹性扩缩、请求迁移、NIXL rank 掩码、专家重映射、权重回填。

被测模型六个：DeepSeek-V2-Lite、Qwen3-30B-A3B-Instruct-2507、DeepSeek-V4-Flash、Qwen2-57B-A14B、Qwen3-235B-A22B、Mistral-Small-4-119B-2603，EP 取值覆盖 {2, 4, 8, 12, 16}，既有节点内也有跨节点。

第一个结论是「没有最优 EP，只有匹配负载的 EP」。在并发 4 时，把 EP 从 4 拉到 16，三类请求形状的吞吐分别只有 0.88×、0.99×、0.95×——宽 EP 在低负载下是浪费的；在并发 256 时，同样动作带来 2.30×、2.08×、1.85× 的提升，其中 prefill 型受益最大、decode 型最平。

这条结论的价值不在数字，而在它给出一个可执行的策略：FastEEP 可以选「满足当前 TTFT/TPOT 目标的最小 EP」，压力上来再扩。

第二个结论是 EP 同时是 KV 容量的旋钮：论文用 H20、nixl_ep 后端、128 个冗余专家的配置画了六个模型的内存预算曲线，直观展示「EP 越宽、百万 token 级请求的可实现并发越高」。

第三个结论是迁移延迟。全模型、六个迁移方向上，FastEEP 4.2–13.4 秒，FullAsync 3.2–9.8 秒，Elastic EP 36.4–92.0 秒，整实例重启 126–293 秒。

这里必须点明基线缺口：vLLM 的 Elastic EP 在 DeepSeek-V4-Flash 与 Mistral-Small-4 上跑不起来（论文图 9 图注对这两个模型统一标注 NaN；正文脚注另外指出 BatchedMarlinExperts 实现不支持 Elastic EP，论文并未按模型分别归因），所以这两个模型没有 Elastic EP 的数字。

也就是说，4.5–11.7 倍这组区间只统计「基线可运行」的模型-方向组合；这两个模型的加速比是相对重启算的，不能和那组区间混看。

第四个结论是「总耗时」与「停服时长」是两个指标，都不等于「服务损失」。

总耗时 4.21 秒的缩容，在线测到的零输出窗口是 4.2 秒；扩容总耗时 7.15 秒，零输出窗口 7.3 秒；对照之下 Elastic EP 的总耗时 47 秒对应 31 秒零输出，重启是 141 秒。

要判断真实服务损失，还得看窗口内吞吐、延迟与 SLO 违约——只看总耗时会漏掉「期间吞吐掉了多少」，只看零输出会漏掉准备阶段的资源占用。

多模型共存：卡在进程之间怎么交接

论文最后做了一个 MaaS 场景的实验：把 Qwen3-30B-A3B 与 DeepSeek-V2-Lite 放在同一批卡上，需求独立波动。

过程是这样的：服务 1 从 EP 8 缩到 EP 4，把 4 张卡让给服务 2；等它需求回来，服务 2 先把自己从 4 缩到 2，服务 1 再从 EP 4 长到 6——两个实例都没有被重建。

驱动侧看到的空闲显存从 34 191 MiB 涨到 64 613 MiB，第一次 8→4 的 drain 用 3 秒、重配用 4.4 秒，服务 1 的 4→6 扩容用 6.9 秒，服务 2 的 4→2 缩容用 4.2 秒。

这里暴露一个论文只提了一句、但生产上一定会踩的坑：跨服务的扩容有顺序依赖。

释放的页返回到节点的分配器，而不是直接在两个进程之间转让，所以「服务 1 想扩」必须排在「服务 2 先缩」之后。

单机两服务时这只是个排序问题；一旦集群里有几十个服务、每个都在等别人先让卡，这就变成一个分布式调度问题，需要防死锁（比如按优先级与超时回退）。

论文把这块留给调度层，没有展开。

边界与未解

把限制摆清楚，才能判断这篇论文能不能直接搬到自己机房。

第一，变更只能在启动时定下的最大拓扑内进行。想超过这个上界，仍然要重建；这是一个明确的工程取舍：用「启动时按峰值预建」换「在拓扑与容量约束内伸缩」。

第二，实验规模是 2 节点 16 卡。论文自己说设计可以扩展到更大规模（把集中式的连接重建换成分布式的状态卸载），但这是可扩展性论证，不是实测。

第三，评测是 prefill/decode 共置部署，含跨节点实验，但没有评测 PD 分离（prefill/decode disaggregation）。论文说机制上兼容两侧，可是没有给 disaggregated 部署的数据。

第四，硬件与通信栈是单一口径：H20-3e + NCCL + NIXL。跨厂商加速器（比如昇腾）或别的传输库上的行为没有验证，而这一层恰恰是同类系统最容易翻车的地方（ElasticMoE 那类工作是直接在昇腾 NPU 上做的实例级方案）。

第五，作者称实现已开源并向 vLLM 提交了 PR，但本文本轮没有在 vLLM 仓库检索到可公开核对的那个 PR 链接（检索时点为 2026-10-21），因此这一条按「作者自述、未核实」处理。

第六，也是最容易被忽略的一条：影子 worker 的 CPU 线程、进程与元数据都在，还保留少量设备状态。省的是 GPU 算力，不是所有资源——按卡时计费时收益成立，按节点租用时缩容并不能把节点还回去。

第七，控制策略缺失。论文解决的是「怎么快速改」，没有解决「什么时候改、改到多少」。它把这块交给自动扩缩容研究——而现实里，误判一次扩缩容的代价（抖动 × 频繁变更）可能比不做更大。

第八，论文的对比集中在 vLLM 生态，SGLang 侧的相关工作是容错导向（rank 失效后用专家重映射把故障 peer 掩掉），目标与 FastEEP 不同：它保图、但明确不归还物理内存，所以不能直接当作同类基线比较。

如果你在做推理平台，这篇能带走什么

即使不打算复刻 FastEEP，这篇论文也给出了几个可以直接用的判断。

一是先量「与配置无关」的那部分成本。任何「换个并行度就要重启」的系统，都该先把自己拆一遍：多少时间在编译、多少在图捕获、多少在通信域初始化。如果八成以上的耗时与配置无关，那重建就是设计缺陷，不是必然。

二是把「成员视图」和「资源占用」解耦。成员视图可以靠层次化通信域（父域 + 子组）与掩码来表达；资源占用可以靠 VMM 的地址/物理页分离来回收。这两件事一旦解耦，很多「必须重启」都能变成「必须提交」。

三是把停服窗口压到「提交点」，其余尽量异步。以 Qwen3-30B 的分段为例，真正需要全局一致的只有 2.81 秒的 resize，其余是卸载与恢复连通——这些可以放在请求继续服务的窗口里做。

四是扩容与缩容不是对称的。缩容只需单向卸载，扩容要回填权重、恢复连通、重新装布局，所以扩容天然更慢（4.21 秒 vs 7.15 秒）。排期与容量规划应该按扩容代价留余量。

五是给自己留一个「最小可行规模」的公式。如果没有 $N_{\min}$ 这个概念，缩容到某个点会突然失败（专家放不下），而失败往往发生在压力最大的时候。

小结

FastEEP 的贡献不在于发明了某个新原语，而在于把一个「分布式系统的成员变更」问题正确地搬到了 GPU 推理栈上，并且证明三条绑定都可以拆。

拆法是三件事：通信视图靠父域派生子组 + P2P 掩码；物理内存靠 CUDA VMM 解映射、保留地址；worker 生命周期靠影子状态 + 一次原子提交。

收益是 4.2–13.4 秒完成一次 EP 变更，相对现有 Elastic EP 几何平均快 7.0 倍，在线零输出窗口从 31/44 秒降到 4.2/7.3 秒，稳态吞吐基本不变。

代价是最大拓扑定死、元数据常驻、只在 16 卡两节点上验证过。

如果要用一句话概括它的工程含义：让「EP 是多少」从启动参数变成运行期资源——而这一步一旦成立，MoE 服务的成本曲线就不再只由「峰值配置」决定，还会由「多快能跟上需求」决定。

再补一句容易被跳过的判断：这篇论文的价值有一半在「把账拆开」上。

它把一个 rank 的冷启动拆开，说明在这个样本里模型权重只占 7.9%——这个占比是样本值，方法才是可复用的：你自己的硬件、模型与加载路径必须重测。

它也顺带标出了 vLLM Elastic EP 在两类模型上跑不起来（图注统一标注 NaN，脚注另提 BatchedMarlinExperts 实现不支持），提醒读基线数字时先确认基线真的跑过。

它还暴露了多服务共卡时的顺序依赖（页返回节点分配器、不能进程间直转），这条路绕不开调度器。

所以更稳妥的读法是：把它当成一份「MoE 服务的状态清单 + 生命周期顺序约定」，能复用的地方比加速比更多。

📋 来源与核验记录

以下来源均为本文本轮实际打开并核对的页面（核验时点：2026-10-21，北京时间）。

FastEEP 论文（arXiv:2610.18451v1，2026-10-13 提交，cs.DC）摘要页：https://arxiv.org/abs/2610.18451
FastEEP 论文全文（HTML 版；本文的实验数据与机制描述主要据此页核验）：https://arxiv.org/html/2610.18451v1
vLLM 官方博客《Elastic Expert Parallelism in vLLM》（2026-05-14，作者含 Itay Alroy、Yongji Wu、Rui Qiao 等）：https://blog.vllm.ai/blog/2026-05-14-elastic-expert-parallelism
SGLang / LMSYS 博客《Elastic EP in SGLang: Achieving Partial Failure Tolerance》（2026-03-25）：https://www.lmsys.org/blog/2026-03-25-eep-partial-failure-tolerance/
NVIDIA NCCL 官方文档 API 索引与 ncclCommSplit 语义（页面当前版本 2.32.3）：https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/comms.html
CUDA 驱动 API 虚拟内存管理（cuMemMap / cuMemRelease 语义）：https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__VA.html
DeepSeek-V3 技术报告（671B 总参数 / 37B 激活，用于专家参数占比的背景口径）：https://arxiv.org/abs/2412.19437
NIXL 仓库（NVIDIA Inference Xfer Library，跨节点传输层，UCX 之上的点对点语义；本文只核对了仓库首页的平台与依赖说明，未逐文件阅读）：https://github.com/ai-dynamo/nixl

关键数字的一手出处逐条对应如下，便于复核：

BurstGPT trace（140 万条请求 / 两个月 / 59% 的十秒窗口无请求 / 25 倍摆幅）：FastEEP 论文 §2.1 引用 Wang 等 2025 的工作，本文为二手引用（经论文转述），未直接读取原始 trace。
Azure 对话 trace（2700 万条请求 / 一周 / 1.7 倍波动）：FastEEP 论文 §2.1 引用 Azure/AzurePublicDataset，采集期 2024-05-10 至 2024-05-19，CC-BY-4.0；本文同为二手引用。
「现有 resize 36.4–92.0 秒，长于其应对的抖动窗口中位持续时间」：FastEEP 论文 §2.1，属论文判断，本文沿用其口径。
专家参数约占总参数 97%：FastEEP 论文 §2.1 的转述；本文另直接核对了 DeepSeek-V3 技术报告摘要的 671B 总参数 / 37B 激活。
24 / 13 / 10 / 32 / 6 / 36 秒六个阶段、126 秒单 rank 总时长、6.32 GiB @ t=47 秒、10.90 GiB 运行期状态、9.62 GiB 权重、156–267 秒全启动区间：FastEEP 论文 §2.2 与图 2；六项相加 121 秒与正文 126 秒的差异为本文核对时发现并如实保留。
47 秒 / 75 秒 / 31 秒 / 44 秒 / 141 秒 / 149 秒：FastEEP 论文 §2.2 与 §5.3（vLLM Elastic EP 与整实例重启的对照）。
4.2–13.4 秒、3.2–9.8 秒、36.4–92.0 秒、126–293 秒、4.5–11.7 倍、几何平均 7.0 / 9.3 / 22.0 / 29.3 倍：FastEEP 论文 §5.2 与摘要。
4.21 秒 / 2.81 秒 / 0.92 秒 / 0.24 秒 / 7.15 秒 / 1.09 秒 / 0.85 秒 / 2.50 秒 / 2.55 秒：FastEEP 论文 §5.2 图 10 的分段（计时为分段打点，与 §5.3 的端到端零输出窗口不是同一测量）。
0.88× / 0.99× / 0.95×（并发 4）与 2.30× / 2.08× / 1.85×（并发 256）：FastEEP 论文 §5.1。
34 191 → 64 613 MiB 空闲显存、3 秒 drain、4.4 秒重配、6.9 秒扩容、4.2 秒缩容：FastEEP 论文 §5.4 与图 11。
控制面 <300 MB、通信缓冲 2–4 GB、图重捕 25–82 秒：FastEEP 论文 §4.1、§6 与 §3.2。
测试床（2 节点 × 8 张 H20-3e、140 GiB、NVLink 900 GB/s、8× BlueField-3 集成的 ConnectX-7 400 Gb/s、双路 Xeon Platinum 8575C）与实现版本（vLLM 0.27.0、PyTorch 2.11.0+cu130、NCCL 2.28.9、NIXL 1.2.0、约 5K 行）：FastEEP 论文 §5 测试床与实现段；型号与配置按论文记载，本文未独立核验。
ncclCommSplit(comm, color, key, *newcomm, config) 的集合语义与同色成组：NCCL 官方文档（2026-10-21 打开，页面版本 2.32.3）。
cuMemMap 把分配句柄映射到预留虚拟地址范围、cuMemRelease 释放句柄：CUDA 驱动 API 文档（2026-10-21 打开）。
671B 总参数 / 37B 每 token 激活：DeepSeek-V3 技术报告摘要（2026-10-21 打开）。

本文自行推导（非论文或官方给出的数字）：

671B 参数按 FP8 存储时，EP=8 / 16 / 64 对应的单卡权重摊薄约 84 / 42 / 10.5 GB（只算按 EP 切分的部分、忽略复制项，是真实单卡占用的乐观值与下界）。
7.9%（权重加载占 10/126）与 28.6%（图捕获占 36/126）两个比值：以论文正文 126 秒为分母；若改用六项相加的 121 秒则分别为 8.3% 与 29.8%。
0.85 秒回填对应约 12 GB/s 的等效速率：假设回填量接近 9.62 GiB 的单卡权重足迹。该数只是提出「可能不是带宽瓶颈」的待验证猜测，不构成结论。
六阶段相加 121 秒与正文 126 秒的差值 5 秒：本文只做保留与提示，不解释其构成。
256 个路由专家 / 每 rank 16 槽 / 2 倍冗余 ⇒ 活跃 rank 至少 32 个：本文用 $N_{\min}$ 公式举的假设例子，非论文实验选点。
文中全部公式（ $M_w$ 、 $M_{kv}$ 、 $N_{conc}$ 、 $s$ 的展开式、 $T_{rebuild}$ 与两个占比、 $T_{coll}$ 、 $T_{recompute}/T_{migrate}$ 、 $T_{scale}$ 、 $N_{\min}$ ）都是本文为讲清机制而建的模型；论文原文只给出 $\mathit{EP}=\mathit{DP}\times\mathit{TP}$ 这一条关系，并未给出上述公式。
三条正确性不变量是本文对论文机制的归纳表述，论文本身以机制描述呈现，未以「不变量」形式编号列出。