ExpertPlex:共享专家、分离注意力,MoE 高吞吐推理服务的新架构

MoE 推理:权重有多大,计算有多稀疏#

先建立基本盘。MoE 模型把传统 Transformer 的 FFN 层替换成专家集合:所有 token 都会经过少数 shared experts,同时路由器为每个 token 在每个层动态挑选 top-k 个 routed experts 参与计算。稀疏激活让模型容量爆炸式增长,但计算量没有同比例膨胀。

MoE 层由路由器（router）为每个 token 挑选 top-k 个专家参与计算，再对专家输出加权合并；每个 token 只激活少数专家，稀疏激活让模型容量远大于单次计算量。（来源：HuggingFace Blog「Mixture of Experts (MoEs) in Transformers」）

代价是权重占据的显存。三款 2026 年的旗舰 MoE 模型中,专家权重占模型总参数的比例分别是：DeepSeek-V4-Pro 约 95%,GLM-5.1-FP8 约 96%,MiniMax-M2.7 约 98%。以 ExpertPlex 实验用的两个模型为例：MiniMax-M2.7 的 FP8 权重占用 230 GB 显存,每个 token 激活约 7.0B 路由专家参数;GLM-5.1-FP8 的 FP8 权重占用 756 GB,路由专家参数总量 724.8B,每个 token 激活约 22.6B。两个模型都是每层 256 个路由专家、每 token 激活 8 个。

显存只是第一重约束。推理的两个阶段有着截然不同的资源画像:

• Prefill 阶段并行处理全部输入 token,构建 KV cache 并产出第一个输出 token。它计算密集、对吞吐敏感,单次 kernel 可以跑几十毫秒。
• Decode 阶段每次迭代只为一个请求生成一个新 token,依赖前缀 KV cache。它对延迟敏感,因为逐 token 的时延直接暴露给用户,单次迭代的耗时比 prefill 短几个数量级。

两者混跑会互相干扰,所以现代推理系统普遍把权重分片到多张 GPU 上,用专家并行(Expert Parallelism,EP)承载 MoE 层。由于每个 token 的路由结果不同,激活张量需要经过 all-to-all 通信：dispatch 把激活发往持有对应专家的 rank,combine 把各专家的输出取回。通信可以用 two-batch overlap(TBO,一个 micro-batch 的通信与另一个 micro-batch 的计算重叠)或 single-batch overlap(SBO,同一 micro-batch 内共享专家计算与路由专家的通信/计算重叠)来隐藏。注意力层则因为权重小得多,通常用数据并行(DP,吞吐优先、无通信)或张量并行(TP,降单请求延迟、有通信)单独配置。

方案一:实例级 prefill-decode 分离(PDD)#

PDD 为两个阶段各部署一份完整模型副本,彻底消除跨阶段干扰。问题出在”完整副本”这四个字上。

MoE 权重动辄几百 GB,一份实例就要占几十上百张 GPU。DeepSeek-V3 的公开部署报告中,一个服务单元是 32 张 prefill GPU + 320 张 decode GPU;另一个不同负载的部署用 176 张 GPU 作为一个单元;Kimi-K2 则部署在 128 张 H200 上。PDD 的隔离性是用四重代价换来的:

1. 显存效率:两个阶段的专家权重完全重复,挤占了本可以存放 KV cache 的显存。
2. 弹性:资源分配粒度随模型增大而变粗。小集群根本凑不出 32:320 这样的 prefill：decode 比例,凑不齐就意味着一个阶段过载、另一个阶段闲置。
3. 扩容粒度:大集群凑得出比例,但每次扩缩容都以几百张 GPU 为单位,跟不上中等规模的流量波动。
4. 故障爆炸半径:分层通信把大量 rank 耦合进一个组,一个 rank 故障就能拖垮整个服务单元,且难以恢复。

方案二：prefill-decode 共置#

共置让两个阶段共享同一份模型,避免权重重复。早期方案是 chunked prefill(Sarathi 系列):把长输入切成小块,与 decode 迭代交错执行。但每个 chunk 都要重读前缀 KV cache 和模型权重,内存访问开销随输入变长而放大。

更近期的方案用空间划分：Green Contexts 机制(基于 SM 预留的 GPU 分区,MuxWise 等系统采用)把 GPU 的 SM 集合切给 prefill 和 decode 各一部分,在共享模型的同时提供计算隔离。MuxWise 在 ASPLOS 2026 上展示了 2.20 倍(最高 3.06 倍)的 SLO 保障吞吐提升,但它解决不了 MoE 特有的时变负载。ExpertPlex 论文把共置方案的局限归纳为三个维度的时变需求:

1. 层间变化：EP 下,每层激活的专家数量和各 rank 上的 token 数都不同,MoE 计算负载随层和 rank 波动。
2. 模块间变化:同一层内,attention 与 MoE 的资源需求不同。attention 要跨 token 混合信息、材料化 KV cache,计算更密集;MoE 的 FFN 是逐 token 变换。
3. 阶段内变化:每个 MoE 模块在 dispatch、专家计算、combine 之间轮换,瓶颈在通信与计算之间摇摆。

固定的 SM 划分无法跟随这些变化,于是出现两种失败模式。队头阻塞(Head-of-Line Blocking)：prefill 持有了太多 SM 而 decode 就绪时,延迟敏感的 decode 只能排在不可抢占的 prefill kernel 后面——prefill kernel 运行几十到几百毫秒,decode kernel 只要几百微秒,差了三个数量级。资源气泡:为 decode 预留的 SM 在 decode 无任务时闲置,prefill 也不能借用。重新配置需要 CPU 介入和 kernel 完成等待,实际系统只能在 prefill 层边界重划分,粒度太粗。

图源：ExpertPlex 论文 Figure 2（arXiv:2607.18002）

通信层面同样没有隔离。两个阶段的 dispatch/combine 共享同一批网卡和链路,一个大的 prefill 传输会挤占延迟敏感的 decode 传输。更隐蔽的问题是跨阶段死锁:传统双边通信要求收发双方同时推进,独立调度下可能出现”某些 MoE rank 在跑 prefill、另一些在跑 decode”的状态,每个阶段都在等对方 rank 上的接收 kernel,谁都无法清空 ring buffer 归还 credit,发送方也就无法完成——整个系统卡死。此外,把每张 GPU 都切成两半,每个阶段拿到的局部资源更少,要达到相同延迟目标就需要更宽的并行度,通信量随之上升,反过来加剧网络干扰。

两类方案各有一个无法解决的矛盾：PDD 的隔离以显存和弹性为代价,共置的显存效率以粗粒度隔离为代价。ExpertPlex 的观察是:矛盾出在”划分对象”的选择上。

GPU 执行模型铺垫#

理解 APK 之前,需要先铺一层 GPU 执行模型的底。一张 GPU 由多个流式多处理器(SM)组成,每个 SM 内有 Tensor Core、CUDA 核心、张量内存加速器(TMA)单元、寄存器、L1 cache 和软件管理的共享内存,所有 SM 共享 L2 cache 与全局内存。工作以 kernel 为单位通过 CUDA 流下发:同一流内的 kernel 顺序执行,不同流的 kernel 在资源允许时并行。kernel 内部,线程组成 warp,warp 组成协作线程阵列(CTA,一个 CTA 在一个 SM 上跑完),多个 CTA 可以组成 cluster 跨 SM 协作。Hopper 架构引入的两项特性对本文至关重要:分布式共享内存(DSMEM) 让 cluster 内的 CTA 协调,而 TMA 多播能把一块全局内存一次性搬给多个 CTA。高性能 GEMM kernel 把 warp 特化为数据搬运和 Tensor Core 计算两类,再用流水线隐藏访存延迟——这就是 warp specialization。

图源：ExpertPlex 论文 Figure 1（arXiv:2607.18002）

推理系统的另一项基础设施是 CUDA Graph:把一批 kernel 及依赖关系捕获成图,重放时省掉 CPU 的逐 kernel 启动与同步开销。decode 的 kernel 路径极短,CPU 开销占比不可忽视,所以现代推理引擎高度依赖 CUDA Graph。持久 kernel 是更进一步的形式：kernel 常驻 GPU,由 GPU 自己决定接下来做什么,完全不回 CPU。

为什么调度粒度必须是 tile#

APK 的起点是一个量化观察:同一模型里,decode 与 prefill 的 MoE GEMM 耗时差距有多大。在 MiniMax-M2.7 上以 EP4 配置测量,8 个激活专家的 decode grouped GEMM 耗时 17.7-34.7 微秒,而 16K token 的 prefill GEMM 耗时 1.8-2.9 毫秒——相差 84-101 倍。这个比例随模型而变,趋势是普适的。关键在于:每一层都会调用一次 MoE 操作,如果 decode 每层都被 prefill 阻塞 2 毫秒,阻塞会逐层累积成端到端的 decode 延迟灾难。反过来,静态划分避免了阻塞,却要在每个阶段通信或无任务时让资源空转。

图源：ExpertPlex 论文 Figure 6（arXiv:2607.18002）

因此 APK 选择了 tile 作为调度单元：tile 是这些 GPU 操作中最小可独立完成的单元。它的尺寸由寄存器、共享内存、TMA buffer 和 Tensor Core shape 决定——输入变长只是增加 tile 数量,不改变单个 tile 的尺寸。论文用的 SM90 DeepGEMM 配置中,一个 CTA tile 至多是 128×192 的输出块,配 128 宽的归约 tile。实测下来,MoE 侧操作的 tile 边界每 2.2-25.3 微秒出现一次(GEMM 本身不到 10.7 微秒),与操作总长度无关。

tile 边界还有一层更关键的性质:它是干净的状态边界。tile 提交时,没有存活的累加器、TMA 事务、共享内存 buffer 或通信状态,切换阶段不需要 checkpoint、restore 或重算。抢得更细(tile 内部)就要保存流水线状态;抢得更粗(kernel 或层级别)就会保留长阻塞区间。tile 恰好是工程上的最优点。

模板实例化与 cluster 级调度#

APK 的持久执行依赖一个前置工作:启动前为每一种 MoE 操作构建操作模板,包含预分配的 buffer 和 TMA 描述符。运行时,attention server 发出一个就绪操作信号,CTA 0 把对应模板实例化到 decode 或 prefill 队列里,并更新层号、token 数、SM 预算等运行时字段。

调度以 CTA cluster 为粒度,而不是单个 CTA。原因在于 TMA 多播把多个 CTA 绑在一起,它们必须执行同一个 tile。cluster 领导者挑选下一个 tile,把元数据广播给对等 CTA。不同 cluster 执行不同阶段,实现空间复用;每个 cluster 在 tile 边界切换阶段,实现时间复用和快速重分配。整个过程都在 GPU 上完成,不需要 CPU 介入、不需要重启 kernel,也就全程保持 CUDA Graph 兼容。

有界 tile 级抢占:为什么不能各自检查、如何用内存层级广播决策#

最直接的 tile 级调度实现是”每个 CTA 在自己的 tile 边界独立检查是否有更紧急的任务”,但论文指出这在流水线 kernel 里不安全。高性能 GEMM 同时流水化 warp 和 CTA:一个 CTA 内的 TMA warp 可能正在加载 tile k+1,而数学 warp 还在消费 tile k。如果有 warp 先切换了阶段,其他 warp 可能永远等不到旧阶段的数据或 mbarrier 事件。跨 CTA 同理:一个 CTA 擅自切换,会把对等 CTA 卡死在 TMA 多播或 cluster barrier 上。

在每个 tile 后停顿整个 cluster 可以避免死锁,但这会把微秒级的流水线串行化,毁掉细粒度调度本身。APK 的解法是把一个协作决策沿内存层级传播,各层只做一次检查:

1. attention server 通过系统级字 P 宣告紧急 decode 操作;
2. CTA 0 在当前 tile 边界检查 P,为每个 cluster i 写一个设备级字 p_i;
3. 每个 cluster 只读自己的 p_i,避免反复访问系统级内存,也不需要网格级 barrier;
4. cluster 领导者在下一个 tile 边界通过 DSMEM 广播决策,保证参与 TMA 多播的所有 CTA 留在同一操作上;
5. CTA 内部,第一个流水线 warp 读出决策、在 warp 内广播,并用 mbarrier 交接通知后续 warp 在认领下一个 tile 之前先看决策。

决策在一个检查周期内保持固定,观察到决策后任何 CTA 都不能再认领旧阶段的 tile,所以最慢的 CTA 至多跑完当前 tile,整个 cluster 就收敛到新阶段。抢占上界是一个 tile 的执行时间加一个本地 cluster 检查周期——与被打断操作的总长度无关。一个几千 token 的 prefill 请求可以被一个短 decode 操作在 25 微秒内打断,这正是队头阻塞问题的终点。

APK 沿内存层级（系统级字 → 设备级字 → DSMEM → warp 内广播）传播一次协作决策，实现有界快速抢占。（来源：arXiv:2607.18002）

时空多路复用#

tile 级调度同时给出空间与时间两个维度的复用:当只有一个阶段就绪时,APK 把所有 CTA cluster 都给这个阶段;两个阶段争用时,按预先搜索出的 SM 预算比例分配(第六节的在线重分配公式会细化这一点)。decode GEMM 短而间歇,APK 的做法是平时把所有空闲 cluster 交给 prefill,decode 一到就抢占,decode 一完成就把空闲 SM 还给 prefill——资源气泡在 tile 粒度上被抹平。

双边协议的死锁与浪费#

解决了 GPU 内的共享,还剩 GPU 之间的问题。两个阶段的 dispatch/combine 流量打到同一批 MoE servers,传统 MoE 通信是双边协议：token 路由在运行时才揭晓目的地和传输量,且通信算子与消费其输出的计算算子解耦,所以发送方把数据流式写入接收方的一个有界 ring buffer,接收方轮询、排空、把到达的数据散布到最终张量,再归还 credit 防止发送方覆盖未读槽位。ring buffer 控制了内存占用,代价是进展依赖双方。

共置场景下这套协议会死锁:独立调度让部分 MoE rank 跑 prefill kernel、部分跑 decode kernel,每个阶段都在等对方 rank 上的接收 kernel 来排空自己的 ring buffer,而阻塞的发送方无法 retire,也就无法释放 GPU 资源给对方的接收 kernel。这是一个环状等待,任何单边都解不开。预留接收侧轮询 SM 可以避免死锁,但 MoE 流量是突发的,这些 SM 在两次 dispatch 之间空转,不能执行就绪的专家 tile;让计算借用它们又不安全,因为接收进展会重新依赖于”每个 rank 是否有足够 SM 空闲”。

网络争用是独立于死锁的第二重问题。优化后的 MoE 传输可以打满互联带宽,而扩展网络带宽稀缺：DeepSeek-V3 部署报告里,H800 节点内 NVLink 带宽 160 GB/s,跨节点 InfiniBand 只有 50 GB/s,差距 3.2 倍。一个 prefill 突发就能让延迟敏感的 decode 传输排队,即使两个阶段用的是不同 GPU。

单边 push dispatch、单边 pull combine#

ExpertPlex 的解法是把 MoE 侧的进展责任彻底移除:由 attention server 发起单边传输。前提条件是 APK 已经为每个操作按最大路由 token 量预分配了 buffer——持久执行要求模板、张量地址和通信描述符在运行前固定,这个要求恰好与单边传输的语义吻合。

于是 ExpertPlex 把最终的 dispatch/combine buffer 直接暴露给 attention server,删掉了中间 ring buffer:

• Dispatch 用 push:路由确定每个专家的目的地和最终偏移后,attention server 通过 NVLink peer store 或单边 RDMA write 把激活直接写进目标 MoE buffer,payload 全部可见后才发布 ready 信号。APK 完全不参与协调,只在 tile 边界观察到信号后调度对应的 MoE 任务。
• Combine 用 pull:专家计算写入最终 combine buffer 并发布 done 信号后,attention GPU 上一个单线程的 WaitDone kernel 观察到完成,attention server 再用 NVLink load 或单边 RDMA read 把结果拉回来。

没有了接收侧的排空、散布和 credit 反馈,也就没有轮询 SM 和双边进展要求。WaitDone 是单线程 kernel,跑在通信流上,只阻塞依赖它的 pull,计算流继续执行当前 overlap 策略暴露的工作,与 TBO/SBO 天然兼容。

ExpertPlex 用单边 push dispatch / 单边 pull combine 取代双边 ring buffer 协议，消除死锁与轮询 SM 开销。（来源：arXiv:2607.18002）

删掉 MoE 侧通信 kernel 还带来一个结构性红利:跨阶段的计算-通信重叠。APK 可以在 decode 等待 combine 时执行 prefill tile,或在 prefill 传输在途时执行紧急 decode tile。两个阶段各自有独立的依赖链,这种重叠不受 TBO/SBO 的相内依赖限制,只要有活干,SMs 就不会闲下来。

分层 prefill 路径与流量隔离#

单边传输消除了协调开销,但 prefill 和 decode 仍会争抢稀缺的 RDMA 带宽。ExpertPlex 的隔离策略是让两类流量尽量走不同的物理路径:

• Decode:延迟敏感,每个 MoE server 都通过本地 NVLink 或远端 RDMA 直达。
• Prefill:扩展流量尽量在 prefill attention server 之间走。发往远端节点的激活,先经 RDMA 发一次给该节点上同 local rank 的 prefill server,再由接收方通过 NVLink 逐 chunk 多播给所有被激活的专家;combine 反向汇聚,走 prefill 专用的 RDMA 路径返回。这条路径保留了分层通信的去重收益——发往同一远端节点的多个专家的激活只跨扩展链路一次。放置优化器会尽量把 prefill server 散布到各节点,让这条路径成为常态。
• 兜底:当目标节点没有 prefill server 时,prefill 用直连单边 RDMA,但分配比 decode 更低的 InfiniBand 虚拟通道优先级,保证共享链路时 decode 优先。

拓扑在可能时分离两个阶段,虚拟通道优先级在必须共享时保护 decode。实现上,论文在 InfiniBand GPUDirect Async(IBGDA)路径上定制了 nvshmemi_ibgda_get_nbi_warp,支持从远端 combine buffer 做细粒度拉取。

Goodput 的统一定义#

搜索需要一个目标函数。一个配置记为一个布局 $\ell$(MoE server 数、attention server 数、prefill：decode 比例、各服务器摆放)加一个争用时的 decode SM 预算 $q$。对每个 $(\ell, q)$,设 $B_p$、$B_d$ 分别是满足各自 SLO 的最大 prefill、decode batch,对应的建模迭代延迟为 $T_p$、$T_d$,平均输出长度为 $\bar{O}$,则请求级 goodput 定义为:

这个 $\min$ 结构捕捉的是流水线平衡:每个请求恰好消费一次 prefill 迭代和 $\bar{O}$ 次 decode 迭代,两个阶段任何一个成为瓶颈都会拖累整体。优化器只在两个阶段都满足 SLO 的可行配置里最大化 $G$,而不是孤立地为某一阶段过度配置。

Tile 感知的延迟模型#

在全集群上评估每个候选配置昂贵到不可行。ExpertPlex 只在少数代表性 GPU 上 profile attention 和 MoE kernel(含 APK 开销),插值出各 rank 的延迟;网络时间不需要 profile GPU,直接从路由字节数、实测 RDMA/NVLink 带宽和分层路径计算。对每个 profile 过的组件 $c$,用少量输入样本拟合:

其中 $x$ 是局部 batch 大小、$s$ 是序列长度,四个系数 $\alpha_c$、$\beta_c$、$\gamma_c$、$\delta_c$ 分别捕捉固定开销、随 batch 的线性增长、batch 与序列长度的交互项和二次项。对 attention 直接代入这两个量即可;对 MoE 计算,论文强调了一个反直觉的观察:原始路由 token 数不够用,MoE 延迟取决于执行的 tile 数。如果专家 $e$ 收到 $m_e$ 个 token 行,kernel tile 高度为 $M_t$,MoE 的计算足迹定义为:

代入拟合式时取 $x = x_{\mathrm{moe}}$、$s = 1$。为什么?Tensor Core 和 TMA 流水线按对齐的 tile 操作,每个被激活的专家即使只收到几个 token,也要触发至少一个 tile 的元数据处理、权重搬运和计算。所以相同 token 数的两个 batch,激活的专家数不同,延迟可以显著不同。这不是某个模型或某代 GPU 的特例,而是 tiled MoE kernel 的固有行为。预测每个专家的 token 向量、再按 $x_{\mathrm{moe}}$ 插值,就能用小规模 GPU profile 集捕捉这个效应,而不必 profile 每一种全集群布局。

离线搜索与在线重分配#

离线搜索遍历布局空间:生成器枚举所有布局 $\ell$,$\mathrm{FitsMemory}$ 先过滤掉装不下权重和足够 KV cache 的候选;对每个 $\ell$ 再枚举 decode SM 预算 $q$ 的取值;$\mathrm{MaxFeasibleBatch}$ 对每个阶段在配置和 KV cache 界限内二分搜索最大可行 batch,每次探测都用延迟模型估算 attention 延迟、$q$ 个 SM 下的 tile 感知 MoE 计算延迟和通信延迟。关键路径的合成取决于 overlap 策略：TBO 下,一个 micro-batch 的 attention 与另一个 micro-batch 的通信加 MoE 计算重叠,迭代延迟取两条路径的最大值;SBO 下,共享专家计算与路由专家通信/计算重叠,迭代延迟取最大值再加 attention 延迟。最后保留两阶段均可行的配置中 $G$ 最大者,固定并行度、放置、overlap 策略和期望的 decode SM 预算。

搜索流程用伪代码表示如下:

1
best ← ⊥
2
for ℓ in ℒ do                              // 枚举全部布局
3
    if not FitsMemory(ℓ) then continue     // 装不下权重/KV cache 直接剪枝
4
    for q in Q_ℓ do                        // 枚举 decode SM 预算
5
        (B_p, ok_p) ← MaxFeasibleBatch(ℓ, q, prefill)
6
        (B_d, ok_d) ← MaxFeasibleBatch(ℓ, q, decode)
7
        if ok_p and ok_d then              // 两阶段都要满足 SLO
8
            g ← G(ℓ, q)
9
            if best = ⊥ or g > best.g then
10
                best ← (ℓ, q, B_p, B_d, g)
11
return best

三个内层动作对应前文的三个组件:$\mathrm{FitsMemory}$ 检查 attention/专家权重和 KV cache 容量约束;$\mathrm{MaxFeasibleBatch}$ 在配置与 KV cache 界限内做二分,每次探测调用 tile 感知延迟模型;目标函数 $G$ 只在双阶段可行的配置上比较。整个搜索是离线的一次性过程,产出的配置固定下来后,再由 APK 的在线重分配消化运行时波动。

离线解针对平均负载,而运行时路由会让 $x_{\mathrm{moe}}$ 逐层、逐服务器漂移,某个阶段还可能暂时没有就绪的 MoE 计算。APK 因此把 $q$ 当作争用策略而非静态划分:只有一个阶段就绪时,所有 CTA cluster 都归它;两个阶段争用时,按当前 decode 足迹与离线期望之比缩放:

其中 $\lceil \cdot \rceil_c$ 向上取整到 CTA cluster 的倍数,$Q_{\max}$ 是为 prefill 保留进度所需的上限,剩余 cluster 全部给 prefill。优先保护 decode 份额的理由是：decode 对延迟更敏感,且已经开始的请求的输出应当优先于新请求的输入。prefill 拿到剩余的全部 SM,吞吐保持在最高可能水平。

实验设置#

ExpertPlex 的实现建立在三个开源组件之上：APK 基于 DeepGEMM 和 DeepEP v1 改造,服务框架用 SGLang。具体工程动作:把两个 grouped GEMM、激活函数、MoE 数据前后处理和 CTA cluster 调度器融合成每个 MoE GPU 上的一个常驻持久 kernel;主机在启动前确定 tile 尺寸并 JIT 编译为常量以降寄存器压力;用 attention 发起的 M-to-N 路径替换 DeepEP 的接收方驱动 all-to-all;在 SGLang 里新增分离 attention 加共享 MoE 的布局支持、该布局的 MoE token 分发器和通信重叠策略。

模型与硬件:单节点实验用 MiniMax-M2.7(FP8 权重 230 GB,每 token 激活约 7.0B 路由专家参数,full attention),多节点实验用 GLM-5.1-FP8(FP8 权重 756 GB,每 token 激活约 22.6B 路由专家参数,使用 DSA 注意力)。每个节点 8 张 H800、NVLink 互联、8 个 200 Gbps InfiniBand 网卡,多节点实验最多用 3 个节点(24 张 GPU)。

基线全部基于 SGLang 实现,以隔离系统级差异:

• SGLang-ChunkedPrefill:共置加 chunked prefill,chunk 大小取先前调优的最优值;
• SGLang-Colocated:原始 SGLang 共置模式,attention 用 DP、MoE 用 EP,显存允许时开 TBO;
• SGLang-PDD:实例级 prefill-decode 分离。MiniMax-M2.7 用 1P1D 部署;GLM-5.1-FP8 在 24 GPU 上因权重重复导致显存不足,直接不报告;
• SGLang-PDMux:基于 MuxWise 的开源实现改造以支持 MoE,用 Green Contexts 划分 GPU 资源,attention 受限只能用 TP。GLM 上部分基线无法表达 ExpertPlex 的 24 GPU 细粒度布局,改在最大的 16 GPU 兼容布局上运行,按每节点请求数对比。

负载用 ShareGPT(短请求)和 LooGLE(长请求)采样输入输出长度,到达服从泊松过程。主指标是 P90 goodput:至少 90% 请求同时满足 TTFT 与 TPOT 两个 SLO 的最高到达率。四组 SLO 分别为：MiniMax-M2.7 上 ShareGPT 取 TTFT 1 s、TPOT 50 ms,LooGLE 取 TTFT 10 s、TPOT 100 ms;GLM-5.1-FP8 上 ShareGPT 取 TTFT 2 s、TPOT 100 ms,LooGLE 取 TTFT 20 s、TPOT 100 ms。

端到端结果#

四组实验的 P90 goodput 汇总如下:

绝对数上,MiniMax-M2.7 的 ShareGPT 负载下 ExpertPlex 达到每节点 11.3 请求/秒。

图源：ExpertPlex 论文 Figure 7（arXiv:2607.18002）

图源：ExpertPlex 论文 Figure 8（arXiv:2607.18002）

每个 baseline 的失利原因都对应前文分析的机制,值得逐一对照:

• ChunkedPrefill 垫底:每个 chunk 重读模型权重和 KV cache、重复 MoE 通信。chunk 大则干扰 decode,chunk 小则 prefill 效率崩塌,LooGLE 长输入放大了内存访问开销,整个评估负载区间内撑不住 SLO。
• Colocated 次之：SGLang 为了吞吐激进地堆积 prefill batch,让长 prefill 持续推迟 decode 迭代。LooGLE 上差距拉到 4.12 倍,长请求正是队头阻塞的重灾区。
• PDD 中游:消除了跨阶段干扰,但重复的 MoE 权重挤占了 KV cache 显存(实验里因容量不足不得不截断采样序列长度),且实例级分离导致资源比例错配——decode 不是瓶颈时 prefill 也只用得上分配给它的那部分资源。
• PDMux 最接近：Green Contexts 提供了 SM 隔离,但固定划分跟不住每个 MoE rank 上实际激活的专家数,还引入了跨阶段通信干扰,TP attention 又要求更宽的并行度。它在 GLM 短请求上与 ExpertPlex 打平(约每节点 1.5 请求/秒)——TP attention 给短请求更多并行、TTFT 占优——但长请求场景(1.66× 差距)暴露了网络干扰和固定划分的代价。

APK 的微观有效性#

论文用一个双阶段 GEMM 微基准直接验证 APK 的 GPU 共享机制：GLM-5.1-FP8 的单 GPU 上,decode 形态的 grouped GEMM(128 token、8 专家)在 prefill 形态 GEMM(8192 token、8 专家)启动 10 微秒后到达,对比独占执行、CUDA 流优先级、MPS、Green Contexts 四种方案。

CUDA 流优先级只在启动时排序,kernel 一旦运行就没有隔离,decode 排在运行的 prefill GEMM 后面,队头阻塞显著——decode 延迟比独占执行恶化 13.79 倍。MPS 和 Green Contexts 把 decode 延迟保持在独占执行附近,但固定划分让 prefill 分别变慢 3.33 倍和 4.07 倍。APK 是唯一同时进入低 decode 延迟区并保住 prefill 性能的机制：decode 仅增加 8% 开销,prefill 只变慢 1.12 倍。原因就是前面说的策略：decode GEMM 短而间歇,APK 平时把所有空闲 cluster 给 prefill,decode 到达才抢占,decode 完成立刻归还空闲 SM。

图源：ExpertPlex 论文 Figure 11（arXiv:2607.18002）

各机制的固有开销#

tile 级调度在 grouped GEMM 上增加的队列检查、任务选择和 CTA 预算执行,开销按布局分：prefill 用的连续布局里,调度每 tile 组执行一次、不动 TMA/Tensor Core 流水线,开销低于 12%;decode 用的掩码布局里,调度总开销低于 20 微秒,8 专家激活时相对开销低于 10%。用不到 20 微秒换掉排在 prefill kernel 后面的长队,是笔划算的买卖。

图源：ExpertPlex 论文 Figure 13（arXiv:2607.18002）

attention 发起的单边通信在 16 GPU 上与 DeepEP v1 对比:普通模式下 dispatch/combine 与 DeepEP v1 差距在 5% 以内,保留了分层通信的带宽收益;低延迟模式下差异在测量噪声量级的 45 微秒以内。也就是说,去掉 MoE 侧协调和轮询没有牺牲通信效率,却把 MoE server 的 SM 解放给了另一阶段的计算。

图源：ExpertPlex 论文 Figure 14（arXiv:2607.18002）

抢占间隔的测量给出上限的实证：MiniMax-M2.7 全部 MoE 操作的 tile 间隔都在 25.3 微秒以下,GEMM 低于 10.7 微秒。论文与 LithOS、Bless、GPreempt、PipeSwitch、REEF 报告的抢占延迟做了参考对比——这些系统不面向 MoE 负载,也不支持 TMA 多播、CTA cluster、warp specialization 和 CUDA Graph,其中最好的 REEF 也要 35 微秒且需要重算被抢占的 kernel。APK 在自然 tile 边界调度、每个间隔都在做有用功,无需 checkpoint、restore 或重算,对频繁被 decode 打断的毫秒级 prefill GEMM 而言优势明显。

图源：ExpertPlex 论文 Figure 15（arXiv:2607.18002）