MoE 多节点推理性能瓶颈被颠覆：Perseus 系统消除隐藏序列化，实现 10 倍加速

康奈尔大学 Perseus 论文中的一项核心发现，彻底改写了业界对多节点 MoE 推理瓶颈的传统认知：“我们的研究证实，限制多节点 megakernel 性能的关键因素是序列化问题，而非在代理式与 GPU-direct 传输之间做出何种选择。”

长期以来，混合专家（MoE）架构凭借其能将模型参数规模与计算成本解耦的独特优势，成为大语言模型规模化发展的主流方案。从 Qwen3、DeepSeek-V3 到最新的 Llama4，几乎所有前沿大模型都采用了 MoE 设计。 然而，随着模型参数量的持续膨胀，单节点 GPU 的显存容量已无法容纳全部专家参数，多节点部署因此成为不可避免的趋势。

令人意外的是，在单节点环境中表现出色的 megakernel 设计，一旦扩展至多节点，便遭遇了严重的性能滑坡。论文实验数据显示， 通信 bound 类型的 Qwen3-30B 模型在 8 节点部署时，实际性能比单节点趋势预测值慢了近 10 倍，且随着节点数量增加，性能退化愈发严重 。这一瓶颈极大地制约了 MoE 模型的大规模高效部署。

通信 bound 指系统的整体执行时间主要由通信（数据传输、网络交互、同步等）环节决定，而非计算（ALU 运算）或内存访问（带宽/延迟）环节。

Perseus 系统的出现，为这一棘手难题提供了根本性的解决方案。 研究团队通过深入剖析多节点 RDMA 传输路径，发现性能崩溃的根源并非网络带宽不足，而是代理式 RDMA 传输中存在的“隐藏序列化”问题 。针对这一发现，Perseus 提出了两项互补的核心技术：解耦信令技术将 fence 数量减少了 8 倍，而 NIC 端排序则彻底消除了代理线程的阻塞。

图 3 | Megakernel 执行模型。每个 GPU 运行由主机一次性启动的持久内核。协作线程阵列（CTA）专门作为处理器或操作系统（订阅者+调度器）。GPU 上对称内存构成分区全局地址空间（PGAS，绿色虚线），支持单边通信。工作流：❶ 一个 GPU 上的处理器向对等端对称内存发出带信号的 PUT；❷ 订阅者轮询标志并通知调度器；❸–❹ 调度器入队并分配任务；❺–❻ 处理器出队、读取到达的分块并计算。左侧 GPU 的操作系统内部、任务队列和主机 CPU 为简洁省略；两个 GPU 运行相同的 megakernel 代码。该图清晰展示 megakernel 区别于传统 CPU 驱动模式的核心执行逻辑，其以持久 GPU 内核运行，CTA 分工明确，处理器负责计算，操作系统管理依赖与调度。PGAS 内存模型支撑单边通信，分块级调度实现数据到达即启动计算，无需主机参与，大幅提升 GPU SM 利用率，从架构层面解决传统 MoE 通信与计算串行的痛点，为后续多节点优化奠定基础。

实验数据显示，在代理式传输场景下，Perseus 实现了最高 10.3 倍的端到端加速；更令人惊叹的是，基于 IBRC 代理传输的 Perseus，其性能甚至超越了 IBGDA GPU-direct 传输高达 1.2 倍。

这一成果不仅显著提升了多节点 MoE 推理的性能表现，更证明了代理式传输架构本身并非性能瓶颈，为大规模 AI 基础设施的设计开辟了全新的思路。

本文目录

一、背景：MoE 推理的性能革命与多节点困境
1.1 混合专家模型与专家并行
1.2 传统 CPU 驱动执行模型的致命缺陷
1.3 Megakernel：单节点 MoE 推理的性能革命
二、问题根源：代理式 RDMA 传输中的隐藏序列化
2.1 设备发起 RDMA 的两种提交路径
2.2 Put-with-signal 与 fence 的必要性
2.3 隐藏序列化的本质与量化分析
2.4 为什么简单的解决方案行不通
三、Perseus 设计：双重机制消除隐藏序列化
3.1 核心设计理念
3.2 解耦信令：减少 fence 频率
3.3 NIC 端排序：消除 fence 阻塞成本
3.4 两种机制的互补性
四、Perseus 实现与实验评估
4.1 实现细节
4.2 实验设置
4.3 端到端性能
4.4 消融实验
4.5 鲁棒性与通用性
4.6 可扩展性恢复
五、相关工作
5.1 分布式 MoE 系统
5.2 Megakernel 研究
5.3 GPU 发起通信
六、结论与展望
6.1 结论总结
6.2 进阶分析
6.3 未来工作

一、背景：MoE 推理的性能革命与多节点困境

1.1 混合专家模型与专家并行

混合专家模型的核心设计理念是“按需激活”，即每个 token 仅激活模型中一小部分专家参数，从而在保持模型容量的同时，大幅降低推理计算成本。

这一设计使 MoE 模型能够以远低于稠密模型的计算开销，实现更高的模型质量。例如：

DeepSeek-V3 模型拥有 671B 总参数，但每个 token 仅激活 37B 参数；
Qwen3-235B 和 Llama4 Maverick 则分别使用了 128 个专家。

然而，MoE 架构的效率优势也带来了独特的通信挑战。由于专家参数被分布在多个 GPU 上，每个 token 必须被路由到其分配的专家所在的 GPU，计算完成后再返回原始 GPU 进行聚合。 这一过程被称为专家并行（Expert Parallelism, EP），其核心是两次全局的 all-to-all 通信：分发（dispatch）和合并（combine）。

专家并行的四个核心步骤如下：

本地门控计算：为每个 token 分配 k 个专家

分发 all-to-all：将 token 发送到对应专家所在的 GPU

专家计算：在接收的 token 上执行两次 GEMM 和激活函数

合并 all-to-all：将专家输出返回原始 GPU 进行加权聚合

在这四个步骤中，分发和合并通信位于每一层的关键路径上，其延迟直接决定了 MoE 模型的整体推理性能。随着专家数量和 GPU 数量的增加，通信开销会呈线性增长，成为制约 MoE 模型扩展性的主要因素。

1.2 传统 CPU 驱动执行模型的致命缺陷

传统的 ML 框架（如 PyTorch）采用 CPU 驱动的执行模型，即由主机 CPU 逐个向 GPU 流提交内核任务，每个内核运行完成后才会启动下一个。

这种执行模型在处理稠密模型时表现尚可，但在 MoE 推理中却暴露出了严重的效率问题。首先是内核启动开销：一个简单的 MoE 层可能需要启动数百个短生命周期的内核，每个内核启动都有固定的提交成本。当内核数量足够多时，主机 CPU 无法及时向 GPU 提交任务，导致 GPU SM 长时间处于空闲状态。

更严重的问题是集体通信的全局同步屏障。传统的 all-to-all 集体通信是一个批量同步操作，要求所有 GPU 在开始计算之前完成全部 token 的接收，在开始合并通信之前完成全部专家计算。这种层粒度的序列化完全排除了更细粒度的计算与通信重叠的可能性，而这正是 MoE 推理中最具潜力的性能提升空间。

论文引用的测量数据显示，在两个 A100 GPU 上进行 MoE 前向传播时，GPU 有 60-90%的时间处于空闲状态，其中内核启动间隙和非重叠的 all-to-all 通信各占了很大一部分。如此低的资源利用率，显然无法满足大规模 MoE 模型的部署需求。

1.3 Megakernel：单节点 MoE 推理的性能革命

为了解决传统执行模型的效率问题，一种名为“megakernel”的新设计应运而生。它将整个 MoE 层甚至整个模型融合成一个持久化的 GPU 内核，在 GPU 内部实现任务调度和通信发起。

Megakernel 设计的核心优势在于实现了 tile 粒度的计算与通信重叠。它将输入张量划分为固定大小的 tile（例如 128×64），每个 tile 独立地经历分发、计算、合并的完整流程。一旦某个 tile 的输入准备就绪，它就可以立即进入下一个计算阶段，而无需等待其他 tile 完成。

为了实现这一细粒度的流水线，megakernel 将 GPU 的协作线程阵列（CTA）划分为不同的角色：

处理器 CTA：负责执行 tile 的计算任务（如 GEMM）
操作系统 CTA：包括订阅者和调度器，负责管理 tile 的依赖关系和任务分配

订阅者 CTA 持续轮询远程 GPU 的信号标志，当检测到某个 tile 的数据到达时，通知调度器将该 tile 的计算任务加入就绪队列。调度器则将任务分配给空闲的处理器 CTA 执行。这种完全在 GPU 内部的任务调度，避免了主机往返的开销，能够实现高达 93%的 SM 利用率，远高于传统 CPU 驱动模型的 9-60%。

[[IMAGE_X]]

图 1 展示了不同 MoE 模型在单节点到 8 节点的弱缩放性能。可以看到，在单节点上，megakernel 设计表现出色，所有模型都能实现接近理想的性能。然而，当扩展到多节点时，性能出现了明显的分化。计算 Bound 的 Llama4 模型仍然能够保持接近线性的缩放，而通信 Bound 的 Qwen3-30B 模型则出现了严重的性能崩溃，在 8 节点时比单节点趋势慢了 10 倍。

这一现象引发了研究人员的深入思考：为什么在单节点上表现优异的 megakernel 设计，在多节点环境中会遭遇如此严重的性能退化？问题究竟出在哪里？

二、问题根源：代理式 RDMA 传输中的隐藏序列化

2.1 设备发起 RDMA 的两种提交路径

为了实现 GPU 内部的通信发起，megakernel 依赖于设备发起的 RDMA（Remote Direct Memory Access）技术。RDMA 允许一个 GPU 直接写入另一个 GPU 的内存，而无需涉及任何一方的 CPU。

设备发起的 RDMA 有两种不同的提交路径：

GPU-direct 提交路径：GPU 线程通过 NIC 暴露给 GPU 的硬件接口直接提交请求，无需 CPU 线程参与。典型代表是 NVIDIA 的 IBGDA。
代理式提交路径：GPU 线程无法直接访问 NIC，而是将请求写入主机驻留的队列。一个专用的 CPU 代理线程负责排空队列，并代表 GPU 将请求转发给 NIC。

虽然 GPU-direct 路径看起来更高效，但它需要特殊的 NIC 硬件支持，而许多生产环境的网络架构（如 AWS EFA、HPE Slingshot-11 和基于 Broadcom 的云部署）都不支持 GPU-direct。因此，大多数生产环境中的多节点 megakernel 部署都依赖于代理式提交路径。

[[IMAGE_X]]

图 4 展示了代理式 RDMA 的提交路径。NVSHMEM（megakernel 最常用的通信库）为每个处理单元（PE，即每个 GPU）分配一个代理通道。GPU 上的所有 CTA 都将它们的网络请求写入这个通道对应的主机队列，代理线程按照 FIFO 顺序排空队列，并逐个将请求提交给 NIC。

在 MoE megakernel 中，每个 PE 需要向每个远程 PE 发送多个传输请求。具体来说，如果总共有 P 个 PE，每个 PE 托管 E/P 个专家，那么每个 PE 在每次分发阶段需要向每个远程 PE 发送 E/P 个传输请求。在 Qwen3-30B 的 4 节点 16GPU 配置中，每个 PE 需要向 12 个远程 PE 各发送 8 个传输请求，总计 96 个并发传输通过单个代理通道提交。

2.2 Put-with-signal 与 fence 的必要性

Megakernel 实现 tile 粒度重叠的关键原语是 put-with-signal。它将数据传输和完成信号组合成一个原子操作，确保接收方只有在数据完全到达后才会开始计算。

put-with-signal 原语执行两个操作：

PUT：将数据非阻塞地写入远程对称缓冲区
SIGNAL：在数据保证可见后，写入远程 PE 上的一个小标志变量

为了保证 SIGNAL 在 PUT 操作完成后对接收方可见，必须在 PUT 与 SIGNAL 之间插入一个内存屏障（fence）指令。该屏障会等待通道上所有先前的 PUT 操作执行完毕，之后才会提交 SIGNAL。

在单节点环境下，数据传输通过 NVLink 进行，代理不在数据路径上，因此这个屏障操作不会产生可观测的性能开销。然而，在多节点环境中，代理式传输路径上的屏障操作却成为了性能瓶颈。

2.3 隐藏序列化的本质与量化分析

在代理式传输中，屏障操作会强制代理停止处理新请求，直到通道上所有正在进行的 PUT 操作都从 NIC 返回完成状态。这导致了严重的序列化，使得 NIC 无法充分利用其流水线能力。

当多个 CTA 并发提交 put-with-signal 请求时，代理队列中会出现“PUT→FENCE→SIGNAL”三元组的交错排列。每个屏障都需要等待它前面所有的 PUT 操作完成，包括其他 CTA 提交的 PUT。随着并发传输数量的增加，每个屏障遇到的前置 PUT 数量也会增加，从而导致总的排空成本急剧上升。

更糟糕的是，单个屏障的排空成本还会随着节点数量的增加而增长。因为随着部署规模的扩大，PUT 操作会指向更多不同的远程目的地，屏障需要等待所有这些目的地中最慢的一个完成，这是一种随着目的地数量增长的尾部延迟。

图 5 | 跨消息大小、节点数和并发级别的信号开销。(a) 信号效率：耦合 put+信号吞吐量相对于流水线仅 put 吞吐量归一化，随并发数和节点数增加急剧下降。(b) 96 次并发传输时，栅栏成本随节点数和消息大小增长。(c) 栅栏开销占总通信时间比例高达 98%，小消息场景下占比更高，且随节点数增长。信号成本极低（<0.2%总时间），计入传输部分。该图通过微基准测试量化信号序列化开销，并发传输越多、节点数越多，耦合信号传输效率越低，8 节点 96 并发时仅为无信号传输的 2%。栅栏成本随节点与消息大小递增，小消息场景下栅栏开销占通信时间 98%，证明代理侧栅栏是性能损耗的核心。信号本身成本可忽略，优化核心应聚焦消除栅栏引发的序列化延迟，为 Perseus 技术设计提供关键数据支撑。

图 5 展示了研究人员对屏障开销的量化分析结果。从图 5a 可以看到，在 8 节点 96 个并发传输的情况下，耦合 put-with-signal 的吞吐量仅为纯 put 上限的 2%。这意味着，原本设计用于实现细粒度重叠的机制，反而成为了阻碍重叠的最大瓶颈。

图 5c 进一步揭示了问题的严重性：在小消息大小下，屏障开销占总通信时间的比例高达 98%； 即使在 4MB 的大消息下，这一比例仍然高达 19%。而 megakernel 恰恰工作在小消息 regime，因为它将数据划分为小的 tile（例如 128×64 BF16 tile，大小为 16KB）以实现细粒度重叠。

多节点 megakernel 性能崩溃的根本原因，不是网络带宽不足，也不是代理线程的提交速率不够，而是代理式 RDMA 传输中 put-with-signal 原语所必需的屏障操作导致的隐藏序列化。这种序列化在小消息、高并发的 megakernel 工作负载下尤为严重，使得 NIC 的流水线能力几乎完全无法发挥。

图 6 | 96 个远程专家跨 12 个远程 PE 时，CTA 与代理交互。(a) 原始：每个 CTA 发出 PUT+栅栏+信号，96 次传输全串行。(b) Perseus：阶段 1 连续流水线化 96 个 PUT，NIC 自由流水线处理；阶段 2 发出 96 个信号，仅 12 个携带 NIC 栅栏标志（每组 1 个）。领导者 CTA（★）发出带标志信号；其余继续计算。该图对比原始与 Perseus 的 CTA-代理交互流程，原始模式下 96 笔传输因逐个栅栏完全串行，代理反复阻塞。Perseus 分两阶段执行，先批量提交 PUT 让 NIC 流水线并行处理，再由领导者统一提交带栅栏标志的信号，将栅栏次数从 96 次降至 12 次。非领导者 CTA 无需等待信号提交，可提前回归计算，兼顾通信并行与计算重叠，从流程层面大幅降低序列化开销。

2.4 为什么简单的解决方案行不通

面对这一问题，人们很自然地会想到一些简单的解决方案，但这些方案都存在致命的缺陷：

批量合并传输：将多个小传输合并成少数几个大传输，可以减少屏障的数量。但这会导致信号延迟，接收方必须等到整个批量传输完成后才能开始计算，从而失去了 megakernel 设计所追求的 tile 粒度重叠优势，退回到了传统的粗粒度同步模型。
增加代理并行度：使用多个代理线程或多个代理通道可以消除单 FIFO 瓶颈。但为了保持 put-with-signal 的语义，需要跨通道的排序机制，这会重新引入同步开销，并且增加了系统的复杂性。

这些方案的失败表明，问题的核心在于细粒度信令与屏障序列化之间的根本矛盾。要解决这个问题，必须同时重新设计信令协议和代理传输层，在保留细粒度重叠的同时，消除序列化开销。

unsetunset三、Perseus 设计：双重机制消除隐藏序列化unsetunset

3.1 核心设计理念

Perseus 的核心设计理念是：通过在协议层减少屏障的频率，在传输层消除每个屏障的阻塞成本，从而彻底解决隐藏序列化问题。

针对隐藏序列化的两个来源——屏障数量过多和每个屏障的排空成本过高，Perseus 提出了两项互补的技术：

解耦信令（Decoupled signaling）：将数据传输与信号分离，按目标 GPU 分组信号，将屏障数量从每个专家一个减少到每个远程 GPU 一个。
NIC 端排序（NIC-side ordering）：将排序执行从软件代理转移到 NIC 硬件，使用硬件屏障标志代替代理阻塞，使代理永远不会停止排空请求。

图 2 | 3 个带信号传输的代理-NIC 交互（=PUT，=信号）。(a) 原始：每个耦合的 put+信号插入代理栅栏（3 次代理暂停）。(b) 解耦信号：连续提交 PUT；所有信号前仅 1 个代理栅栏（1 次代理暂停）。(c) NIC 侧排序：代理通过每个信号的栅栏标志（）将排序委托给 NIC；代理永不阻塞，但每个栅栏标志序列化 NIC 流水线（0 次代理暂停，3 次 NIC 停滞）。(d) Perseus（两者结合）：仅首个信号携带栅栏标志（0 次代理暂停，1 次 NIC 停滞）。整体而言，对比四种信号传输模式的代理与 NIC 交互逻辑，原始模式每笔传输均触发代理栅栏，多次阻塞导致性能暴跌；解耦信号合并栅栏，大幅减少代理阻塞次数；NIC 侧排序将排序压力转移至硬件，消除代理阻塞但增加 NIC 停滞；Perseus 融合两种技术，仅产生 1 次 NIC 停滞，最大程度降低序列化开销，为解决多节点通信瓶颈提供了核心设计思路。

图 2 直观地展示了这四种不同设计的代理-NIC 交互过程。可以看到，原始方案需要 3 次代理停止；解耦信令将其减少到 1 次；NIC 端排序消除了所有代理停止，但引入了 3 次 NIC 停顿；而 Perseus 结合两者，实现了 0 次代理停止和仅 1 次 NIC 停顿，达到了最优的性能。

3.2 解耦信令：减少屏障频率

解耦信令的关键洞察是：屏障只需要在信号之前，而不是在每个 PUT 之后。如果所有 CTA 先提交它们的 PUT，然后由一个指定的领导者 CTA 提交所有信号，那么只需要一个屏障就可以保证所有信号都在对应的 PUT 之后。

基于这一洞察，Perseus 设计了一个两阶段的协议：

算法 1 | 解耦信令机制详解。该算法完整定义了“解耦信号”的两阶段执行流程。第一阶段：所有CTA并行完成数据传输，通过原子计数器同步各分组进度。非领导CTA无需等待，可直接返回执行计算任务，从而最大化计算与通信的重叠。第二阶段：仅由领导CTA执行一次栅栏操作，并为整个组批量发送信号。这使得栅栏操作次数从每个专家一次降低至每个组一次。本算法无需改动megakernel的核心调度逻辑，仅需调整通信提交流程，易于集成且高效。它为Perseus的核心通信优化提供了可落地的执行规范。

第一阶段，每个CTA暂存其token，发起一个不带信号的非阻塞PUT，随后递增一个基于分组的原子计数器，通知领导CTA该PUT已完成提交。非领导CTA在完成上述操作后，立即返回至megakernel调度器，成为可用的计算资源，用于处理本地路由或通过NVLink传输的token所对应的专家计算任务。这种通信与计算的重叠机制，进一步提升了CTA的利用率。
第二阶段，领导CTA等待原子计数器达到组大小，这确保了组内所有PUT都已进入代理FIFO。之后，领导CTA发出一个单一的fence，紧接着为组内的每个专家发出信号。

该设计的正确性由以下机制保障：

原子计数器确保所有PUT在fence之前已进入代理FIFO。
单一代理线程维护提交顺序，因此fence会跟随组内所有PUT。
fence保证了在所有PUT被NIC处理完成后，才会处理任何信号。
根据传递性，组内所有信号都在组内所有PUT之后执行。

分组粒度的选择

解耦信令的分组大小存在一个权衡：较小的组会保留较多的fence，而较大的组则会迫使每个fence等待更多未完成的PUT。 为找到最优分组大小，研究人员将分组大小作为可调参数，对远程专家数量的所有约数进行了扫描。

图7 | 解耦信号效果（S=1K，8节点）。左图：栅栏数量与分组大小的关系；右图：延迟与分组大小的关系，虚线为耦合基线。加速效果可分解为：❶ 提交重排（从耦合到解耦，分组大小为1），❷ 栅栏数量减少（从112降至4，分组大小从1增至28）。该图量化了分组大小对解耦信号优化效果的影响。栅栏数量随分组增大快速下降，边际效益逐步递减。当分组大小为1时，仅提交重排就降低了12%的延迟；当分组大小增至28时，栅栏数从112降至4，延迟再降低38%。Perseus默认按PE分组，处于效益曲线的拐点，平衡了栅栏减少与等待延迟，适配主流MoE模型，兼顾了通用性与性能最优。

图7展示了不同分组大小下的性能结果。 可以看到：

即使在组大小为1（fence数量与耦合基线相同）的情况下，解耦本身也能将延迟降低12%（从22.7ms降至19.9ms），因为NIC可以连续流水线处理PUT，而不会被交错的fence打断。
随着组大小增加到28，fence数量从112减少到4，带来了额外38%的延迟降低（从19.9ms降至12.3ms）。进一步增加组大小的收益则逐渐递减。

Perseus默认采用按PE分组，即一个组由发往同一个远程PE的所有专家组成。 这种分组方式位于解耦收益曲线的拐点附近，能够捕获大部分fence减少的收益，同时限制每个fence必须等待的传输数量。实验证实，在从1K到64K的所有序列长度下，这个默认值都在最优分组大小的2%以内。

3.3 NIC端排序：消除fence阻塞成本

虽然解耦信令减少了fence的频率，但每个剩余的fence仍然需要支付完整的代理排空成本。NIC端排序技术通过将排序执行从软件代理转移到NIC硬件，彻底消除了这种阻塞成本。

现代RDMA NIC都暴露了一个按请求的fence标志（例如Libfabric中的FI_FENCE和InfiniBand verbs中的IBV_SEND_FENCE）。当NIC遇到一个带有fence标志的请求时，它会推迟该请求的处理，直到同一连接上的所有先前请求都完成。这种排序是通过NIC内部的硬件寄存器跟踪的，而不是软件计数器。

Perseus利用这一硬件特性，将代理的阻塞排空替换为NIC fence标志。当请求一个fence时，代理设置一个挂起标志，而不是轮询完成，并且继续提交工作。下一个发送到NIC的信号会携带这个标志，指示NIC推迟该信号，直到其连接上的所有先前请求都完成。

这种设计的正确性由以下几点保证：

NIC fence标志提供了与代理端排空相同的排序保证。
代理按提交顺序发送PUT和带标志的SIGNAL。
在多队列对（QP）传输上，对等哈希路由将一个对等体的PUT和SIGNAL固定到同一个连接，因此每个连接的排序足以保证每个对等体的正确性。

3.4 两种机制的互补性

解耦信令和NIC端排序是高度互补的，它们分别针对隐藏序列化的两个不同方面，共同覆盖了从开销主导到传输主导的所有工作负载场景。

如果单独将NIC端排序应用于原始的耦合信令，它会消除所有代理停止，但每个信号仍然需要一个NIC fence标志，导致112次NIC停顿（在8节点Qwen3配置中）。而结合解耦信令后，只有每组的第一个信号需要携带fence标志，因此只需要28次NIC停顿。

图8 | 解耦信号与NIC侧排序的综合效果（4节点，Qwen3-30B）。基线：原始（红）、耦合+NIC排序（绿）为水平线；解耦曲线遍历分组大小并启用NIC侧排序。(a) S=1K：NIC侧排序主导，分组大小影响极小；(b) S=64K：三种效果均有贡献。该图验证了Perseus两项核心技术的互补性。小序列长度（S=1K）时，NIC侧排序消除代理阻塞，成为加速核心，分组大小影响微弱；大序列长度（S=64K）时，传输耗时增加，栅栏数量减少的效益凸显，两种技术协同发力。这说明Perseus适配不同负载场景，无论是开销主导的小批量推理，还是传输主导的大批量推理，均能显著降低延迟。

图8展示了两种机制在不同序列长度下的相对贡献。

在小消息（S=1K）情况下，NIC端排序将每个fence的成本降低到接近零，因此性能对组大小不敏感，加速主要来自NIC排序和提交重排序。
在大消息（S=64K）情况下，NIC必须推迟每个带标志的SIGNAL直到前面的PUT完成，而较大的传输需要更长的时间才能完成，因此每个fence的成本不再可以忽略，fence数量的减少重新获得了价值。

解耦信令减少了fence的频率，在fence数量占主导的场景下最为有效；NIC端排序消除了每个fence的阻塞成本，在每个fence开销占主导的场景下最为有效。两者结合，能够在所有工作负载和规模下实现最优的性能。

四、Perseus 实现与实验评估

4.1 实现细节

Perseus的实现具有极高的非侵入性，仅对现有系统进行了少量修改，确保了易于采用和移植。

解耦信令：该机制基于最先进的开源 MoE megakernel FlashMoE 实现，其核心修改仅限 RDMA 请求逻辑，而关键 megakernel 组件（如对称张量布局、订阅者、调度器、工作模块）均保持不变。在节点内部，分发逻辑保持不变；对于跨节点通信，则划分为两个阶段：
- 第一阶段：每个 CTA 先将 token 暂存至本地堆，随后调用 nvshmem_putmem_nbi() 发出不带信号的指令，并通过递增原子计数器通知对应 PE 组的领导者。
- 第二阶段：每个组领导者 CTA 依次调用 nvshmem_fence() 和 nvshmemx_signal_op()。
NIC 端排序：此优化在 NVSHMEM v3.5.21 的传输模块中仅修改了不到 100 行代码。它将原有的基于代理的安静排空机制（例如 Libfabric 的 fi*cntr_wait，IBRC 的 check_poll_avail），替换为基于请求的 NIC fence 标志（例如 Libfabric 的 FI_FENCE，IBRC 的 IBV_SEND_FENCE）。所有修改均限定在传输共享库（nvshmem_transport*.so）内，用户只需替换该库文件，无需修改或重新编译任何应用程序代码。任何以 NVSHMEM 作为通信后端的系统，均可通过此库替换来继承该优化。
多 QP 传输适配：在使用多个队列对（QP）的 IBRC 环境中，默认的轮询策略会将发往同一对等体的操作分散到不同 QP 上。然而，IBV_SEND_FENCE 仅在单个 QP 内部强制排序，这可能导致相互依赖的 PUT 操作与信号操作落到不同 QP，从而破坏预期的执行顺序。Perseus 通过将所有发往同一对等体的操作固定到一个确定性 QP 来解决此问题。这样一来，依赖操作共享同一个 QP，继承了其 FIFO 排序特性；同时，不同对等体仍可分散到不同 QP，以维持聚合带宽。

4.2 实验设置

研究团队在两个不同的硬件平台上对 Perseus 进行了评估：

NERSC Perlmutter 超级计算机：每个节点配备 4 块 A100 GPU，采用 Slingshot-11 网络，最多可扩展至 16 个节点（共 64 块 GPU），并使用 Libfabric 代理传输。
商业 GPU 云：每个节点配备 8 块 H100 GPU，搭载 ConnectX-7 NIC，使用 InfiniBand NDR（400 Gb/s）网络，最多可扩展至 4 个节点（共 32 块 GPU）。在此平台上，评估了 IBRC 代理传输和 IBGDA GPU-direct 传输两种方式。

评估选取了三个覆盖计算-通信光谱的 MoE 模型：

Qwen3-30B：通信密集型
GPT-OSS-120B：计算与通信均衡型
DeepSeek-V3：计算密集型

表 1 | MoE 模型配置。H：隐藏维度，I：中间维度，E：专家数量，k：top-k 路由数量。该表明确了实验所使用的三款 MoE 模型的核心参数，覆盖了不同的计算通信比：Qwen3-30B 因隐藏维度小、中间维度低，属于通信密集型；DeepSeek-V3 维度大、专家数多，计算与通信较为均衡；GPT-OSS-120B 的隐藏与中间维度一致，计算占比更高。这三款模型覆盖了主流大语言模型的主流配置，确保实验结论能适配不同场景，提升了研究的普适性与参考价值。

序列长度设定在 256 到 64K tokens 之间，覆盖了开销主导（小 S）和传输主导（大 S）两种典型场景。

4.3 端到端性能

图 9 展示了 Perseus 在不同模型、不同传输方式以及不同规模下的端到端性能表现。

图 9 | 端到端前向延迟对比：原始 FlashMoE vs Perseus。(a)–(c) 在 Perlmutter 集群（A100，Slingshot-11，Libfabric）上，2–16 节点（8–64GPU）：Perseus 最高实现了 10.3 倍加速。(d)–(e) 在商用 GPU 云（H100，ConnectX-7，InfiniBand）上，2–4 节点（16–32GPU）：在 IBRC 上，Perseus 最高实现了 2.47 倍加速，性能甚至匹配或超越了 IBGDA GPU-direct。图中标注了 Perseus 相对于原始系统的提速倍数。该图清晰地展示了 Perseus 在不同硬件和传输层下的端到端性能：在 Libfabric 代理传输下，通信密集型 Qwen3 的加速比最高可达 10.3 倍，且节点数越多，加速效果越明显，有效解决了弱扩展衰退问题。在 IBRC 上，Perseus 超越了原始 GPU-direct，证明了序列化（而非传输方式）才是多节点 megakernel 的性能瓶颈。在不同模型中，通信占比越高，加速效果越显著，验证了 Perseus 对通信密集型 MoE 模型的优化价值。

核心结果如下：

Libfabric 传输上的巨大加速：在 Perlmutter 的 Libfabric 代理传输上，Perseus 实现了最高 10.3 倍的端到端加速（针对 Qwen3-30B，S=256，16 节点）。对于通信密集型的 Qwen3 模型，加速比随节点数增加而增加，这与原始系统开销随节点数增长的趋势一致。对于均衡型 GPT-OSS 和计算密集型 DeepSeek-V3，加速比也分别达到了 2.8 倍和 2.2 倍。
IBRC 传输上的显著提升：在商业云的 IBRC 代理传输上，Perseus 实现了最高 2.47 倍的加速（针对 Qwen3-30B，S=64K，4 节点）。
代理传输反超 GPU-direct：最引人注目的结果是，基于 IBRC 代理传输的 Perseus，其性能不仅匹配，甚至超越了 IBGDA GPU-direct 传输，最高可达 1.2 倍。这一结果彻底颠覆了业界关于 GPU-direct 传输必然优于代理传输的传统认知。研究人员解释，IBGDA 需要消耗 SM 周期来进行 NIC 提交，这在 megakernel 中会与并发计算产生竞争；而 CPU 代理虽然能避免这种干扰，但此前一直被 fence 诱导的序列化所限制。Perseus 消除了这个瓶颈，使得代理传输的优势得以充分发挥。
接近线性的弱缩放：Perseus 使所有模型都实现了近乎线性的弱缩放。在原始系统中，Qwen3 模型在 16 节点时出现了 19 倍的性能退化，而 Perseus 将其降低到了仅 3.5 倍。

4.4 消融实验

为了量化解耦信令和 NIC 端排序各自的贡献，研究团队进行了消融实验，比较了三种配置：仅解耦信令、仅 NIC 端排序以及完整的 Perseus 方案。

图 10 | 左图：在 Perlmutter 集群 Qwen3-30B 上，不同优化方案相对于原始系统的提速消融实验。节点数越多，单栅栏成本越高，栅栏减少的占比越小，NIC 侧排序的贡献越大；节点数越少则相反。右图：8 节点下的延迟分解，包含栅栏数量与单栅栏成本（↑=全流水线排空，↓=轻量 NIC 侧标志）。该图通过消融实验验证了两项技术的互补性：少节点时，栅栏数量多但单栅栏成本低，解耦信号主导提速；多节点时，单栅栏成本飙升，NIC 侧排序成为核心优化。8 节点的分解数据显示，原始模式下栅栏数量多且成本高，而 Perseus 将栅栏数与单栅栏成本均降至最低，两项技术叠加实现了最高提速，明确了各自适用场景，为后续硬件适配优化提供了依据。

图 10 展示了消融实验的具体结果。

在 2 节点场景下，解耦信令（1.2-1.5 倍）的效果优于 NIC 端排序（1.1-1.4 倍）。这是因为此时每个 fence 的成本相对较小，减少 fence 数量带来的收益更为显著。
在 8 节点场景下，情况发生了逆转：NIC 端排序（1.3-2.6 倍）的效果超过了解耦信令（1.2-1.6 倍）。这是因为剩余的 28 个 fence 每个都代价高昂。而两者结合的完整 Perseus 方案实现了 1.5-3.5 倍的加速，证实了这两项技术具有互补性。

图 15 | Libfabric 上延迟分解为开销（α）与传输成本（β）。左侧展示了原始方案中 α 随节点数增长的趋势，而 Perseus 则能保持其稳定；右侧显示 Perseus 还将 β 降低了 25%–38%。该图借助 α-β 模型量化了延迟的构成。在原始模式下，固定开销 α 随节点数急剧上升，这是导致性能衰退的核心原因；Perseus 则将 α 稳定在较低水平，在 16 节点时削减幅度超过 90%。同时，Perseus 还使每字节传输成本 β 降低了 25%–38%，这是因为消除了代理阻塞，使 NIC 流水线能够高效并行。该分解结果精准解释了提速机制，为不同传输层的优化提供了量化参考，有力支撑了 Perseus 设计的合理性。

4.5 鲁棒性与通用性

专家路由倾斜的鲁棒性

在实际的 MoE 推理过程中，token 到专家的路由通常并非均匀分布，而是存在倾斜现象。为评估 Perseus 在此场景下的表现，研究人员引入了遵循 Zipf 分布的路由模式，其倾斜指数从 0（均匀）到 1.5（前 10 个专家接收了 82%的 token）。

图 12 | 该图展示了专家路由倾斜的鲁棒性（Qwen3-30B，Perlmutter 2–8 节点）。图中描绘了 token 到专家的路由遵循齐普夫分布，倾斜指数在 0–1.5 范围内时，Perseus 相对于原始方案的提速倍数。该图验证了 Perseus 对专家路由倾斜的鲁棒性：当路由均匀（倾斜指数为 0）时，提速表现稳定；当倾斜加剧（指数达到 1.5）时，仍能保持 2 倍以上的加速。在小序列长度下，均匀路由的提速效果更佳；而在大序列长度下，倾斜路由的提速反而更高，这是因为倾斜使得传输更加集中，Perseus 的批量优化效果也因此更加显著。这证明了 Perseus 能够适配真实 MoE 部署中路由不均的场景，具备很强的实用性。

图 12 显示，Perseus 在所有倾斜级别下都保持了显著的加速比。

在 S=1K 时，8 节点的加速比范围从 2.7 倍（均匀）到 2.0 倍（Zipf 1.5）；
在 S=8K 时，加速比随着倾斜程度的增加而上升（从均匀的 1.5 倍到 Zipf 1.5 的 2.1 倍），这与 Perseus 的每字节成本优势在更大传输下被放大的特性相一致。

通用性：Triton-distributed 案例研究

为证明 Perseus 的通用性，研究人员将 NIC 端排序技术应用于 Triton-distributed 的 ALLTOALL 内核，这是一个不含计算重叠的纯通信工作负载。

图 11 | Triton-distributed 全对全微基准测试（H=2048，Perlmutter）。(a) 延迟：原始方案的延迟与消息大小无关，基本保持不变；(b) 开销占比（α/T）：原始方案始终高于 65%；(c) α 对比：Perseus 将序列化开销减少了约 99%。该图验证了 Perseus 的通用性。在无计算重叠的纯通信 Triton-distributed 基准测试中，原始方案的延迟不随消息大小变化，序列化开销占比超过 65%，成为性能的核心瓶颈。Perseus 将固定序列化开销（α）削减了 99%，使延迟回归到由传输主导的正常趋势，且无需修改应用代码即可生效。这证明了该优化能适配各类 GPU 发起的 RDMA 通信，其应用范围不限于 MoE megakernel 场景。

图 11 展示了令人震惊的结果：原始 Triton-distributed 的延迟几乎与消息大小无关，序列化开销占总延迟的 65-100%。Perseus 将序列化开销减少了约 99%（4 节点时从 68.8ms 降至 0.7ms），带来了最高 79 倍的加速（4 节点时平均为 59.6 倍）。

更重要的是，在没有 Perseus 的情况下，GPU 发起的 ALLTOALL 平均比 NCCL 集体 ALLTOALL 慢 18.7 倍。而 Perseus 使其比 NCCL 快了最高 11 倍，恢复了细粒度通信在小消息大小下的延迟优势。

图 13 | Triton-distributed GPU 发起全对全通信 vs NCCL 集体全对全通信（H=2048，Perlmutter）。Triton 原始模式受序列化开销影响，当 S≤512 时性能弱于 NCCL；Perseus 消除开销后，GPU 发起通信比 NCCL 最高快 11 倍。4 节点结果因内存限制仅测试了 S=2048。该图对比了 GPU 发起通信与传统 NCCL 集体通信。原始 GPU 发起模式因序列化开销，在小批量时不如 NCCL；经 Perseus 优化后，GPU 发起通信凭借细粒度调度优势，最高比 NCCL 快 11 倍。NCCL 集体通信的开销随数据量递增，而 Perseus 的固定开销极低，凸显了 GPU 发起通信在多节点 MoE 场景下的优势，为分布式 AI 通信提供了新的选择。

4.6 可扩展性恢复

最后，研究人员重新审视了最初的微基准测试，评估 Perseus 在高并发下的吞吐量恢复情况。

图 14 | Perseus 带来的可扩展性恢复。上图：8 节点微基准测试吞吐量，复现了图 5a；下图：端到端弱扩展，归一化至 1 节点基线（S=1024），复现了图 1。该图直观展示了 Perseus 对多节点可扩展性的修复效果。在微基准测试中，Perseus 将低并发时仅 2%的吞吐量恢复至 74%，在大消息场景下匹配了无信号传输的性能。在端到端弱扩展中，Qwen3 的 16 节点性能衰退从 19 倍降低至 3.5 倍，GPT-OSS 实现了近线性扩展。同时，TensorCore 利用率恢复至单节点水平的 95%以上，全面解决了多节点 megakernel 性能崩塌的问题。

从图 14 上图可以看到，在 96 个并发传输和 4KB 消息的情况下，Perseus 将吞吐量从原始的 2%恢复到了 PUT-only 基线的 74%；在更大的消息下，Perseus 完全匹配了基线。

图 14 下图展示了端到端的弱缩放恢复情况。Perseus 将 Qwen3 的 16 节点弱缩放退化从 19 倍降低到了 3.5 倍，并在 GPT-OSS 上实现了接近平坦的缩放。此外，Perseus 还将 TensorCore 利用率恢复到了单节点水平的 95-98%， 而原始系统中 Qwen3 的 TensorCore 利用率仅为 31%。

表 2 | 4 节点下 TensorCore 利用率（S=1K），归一化至各负载单节点执行（100%=单节点利用率）。该表量化了 Perseus 对 GPU 计算资源利用率的提升。在 4 节点小序列场景下，原始模式中通信密集型 Qwen3-30B 的 TensorCore 利用率仅 31%，大量计算资源因等待信号而闲置；Perseus 优化后利用率升至 95%，提升幅度达 64 个百分点。即使是计算密集型 GPT-OSS，其利用率也从 75%增至 98%，证明该优化消除了通信阻塞，让 GPU 计算单元得到充分利用，实现了通信与计算的高效重叠。

unsetunset五、相关工作unsetunset

5.1 分布式 MoE 系统

分布式 MoE 系统的研究主要分为两个流派：基于集体的系统和基于 GPU 发起通信的系统。

基于集体的系统，例如 DeepSpeed-MoE、Megatron-LM、Tutel 和 FasterMoE，均基于 all-to-all 集体通信原语来优化专家并行。 这些系统依赖于 CPU 驱动的执行模型，因此受限于层粒度的同步，无法实现细粒度的计算与通信重叠。
基于 GPU 发起通信的系统，如 DeepEP、Mooncake EP 和 NCCL EP，采用 GPU 发起的点对点通信替代集体通信，以实现细粒度重叠。 然而，这些系统都继承了底层传输中的隐藏序列化问题，在多节点部署时会遭遇严重的性能退化。

Perseus 与这些系统是互补的：它解决了当 GPU 发起通信遍历代理式 RDMA 时出现的传输级序列化瓶颈，这是现有系统均未解决的问题。

5.2 Megakernel 研究

Megakernel 设计的核心思想是将多个算子融合成一个持久化的 GPU 内核，以消除内核启动开销和粗粒度同步。

FlashMoE、Spector 等人的工作、MPK 和 Triton-distributed 都展示了 megakernel 在单节点 MoE 推理中的巨大性能优势，实现了高达 6 倍的延迟降低和 9 倍的 SM 利用率提升。

5.3 GPU 发起通信

NVSHMEM、NCCL 设备 API、NCCLX 以及 MSCCL++ 为 GPU 主动发起的通信奠定了底层基础。

TransferEngine 与 UCCL-EP 致力于解决跨异构 GPU 和 NIC 的可移植节点间通信问题，它们利用基于 WriteImm 的通知机制，成功规避了发送方的 fence 开销。但需要注意的是，这些系统所需的接口与 megakernel 所依赖的基于内存的 put-with-signal 模型存在本质差异。

相比之下，Perseus 选择保留现有的 put-with-signal 接口，并在该接口内部彻底消除了由 fence 引发的序列化问题，从而实现了无需修改任何应用程序代码的即插即用部署。这一特性使得 Perseus 能够立即惠及所有基于现有接口构建的 megakernel 系统。

六、结论与展望

6.1 结论总结

Perseus 系统通过两项互补的核心技术，成功瓦解了多节点 megakernel 通信中隐藏的序列化瓶颈：

解耦信令：将数据传输与信号分离，并按照目标 GPU 对信号进行分组，从而将 fence 数量减少了 8 倍。
NIC 端排序：将排序执行逻辑从软件代理线程转移到 NIC 硬件上，从根本上消除了代理线程的阻塞问题。

实验数据有力地表明，在代理式传输场景下，Perseus 实现了最高达 10.3 倍的端到端性能加速；更令人瞩目的是，基于 IBRC 代理传输的 Perseus，其性能甚至超越了 IBGDA GPU-direct 传输，最高领先 1.2 倍。这一结果清晰地证明：制约多节点 megakernel 性能的根源在于序列化本身，而非代理式传输与 GPU-direct 传输之间的选择。

Perseus 的实现具备极高的非侵入性，用户仅需替换 NVSHMEM 传输库即可获得性能提升，无需改动任何应用程序代码。在 Triton-distributed 上，它实现了最高 79 倍的加速，充分验证了其通用性。

Perseus 的核心贡献在于，它首次揭示了多节点 megakernel 性能崩溃的真正根源，并提供了一个优雅、高效且易于落地的解决方案。 它不仅大幅提升了现有多节点 MoE 推理系统的性能，更为未来大规模 AI 基础设施的设计提供了全新的思路与理论基础。

6.2 进阶分析

尽管 Perseus 的成果令人瞩目，但我们仍需客观审视其方法存在的局限性与边界条件：

首先，Perseus 的 NIC 端排序技术依赖于现代 RDMA NIC 对按请求 fence 标志的支持。尽管这一特性在大多数当代数据中心 NIC 中都已具备，但在一些较旧的硬件上可能无法使用。在此类情况下，Perseus 只能依赖解耦信令技术，无法获得 NIC 端排序带来的全部收益。
其次，解耦信令的性能高度依赖于合理的分组大小选择。虽然 Perseus 默认的按 PE 分组策略在大多数场景下都能接近最优性能，但在某些特殊的工作负载或硬件配置下，可能需要手动调优分组大小才能获得最佳效果。
第三，在极端专家路由倾斜的情况下，Perseus 的性能会有所下降。尽管实验表明，在 Zipf 1.5 的倾斜度下它仍能保持显著的加速比，但在更极端的倾斜场景中，分组大小的不平衡可能导致某些 fence 需要等待过长时间，从而影响整体性能。
最后，Perseus 目前主要针对 MoE 推理工作负载进行了专项优化。虽然其传输层优化具有通用性，但解耦信令协议是专门为 megakernel 的 tile 级执行模型设计的。将其应用于其他类型的分布式工作负载，可能还需要进一步的适配与修改。

6.3 未来工作

论文作者明确规划了以下几个未来研究方向：

与服务引擎的集成：将 Perseus 集成到 vLLM、SGLang 等主流 LLM 服务引擎中，让更多生产部署能够受益于这一优化。
与训练框架的集成：探索将 Perseus 的技术应用于 MoE 训练场景，优化分布式训练中的通信性能。
GPU-direct 传输的进一步优化：初步评估表明，Perseus 的解耦信令技术在 GPU-direct 传输上也能带来最高 1.25 倍的加速。未来将深入研究 GPU-direct 传输的优化空间，以进一步提升其性能。
更大规模的评估：在更大规模的集群上（如 128 节点以上）评估 Perseus 的性能，验证其在超大规模部署下的可扩展性。

从更宏观的领域发展趋势来看，Perseus 的成果可能会催生以下几个重要的技术方向：

首先，Perseus 证明了代理式传输架构在消除序列化瓶颈后，其性能可以与 GPU-direct 传输相媲美甚至超越。这将重新点燃业界对代理式传输架构的兴趣，推动更多针对代理式传输的优化研究。未来可能会出现统一的 GPU 发起通信传输层，能够在代理式和 GPU-direct 传输之间自动选择最优路径，或结合两者的优势。
其次，Perseus 的成功为 megakernel 设计在更大规模集群中的应用铺平了道路。随着多节点性能问题的解决，megakernel 有望从单节点优化技术，发展成为大规模分布式 AI 系统的标准执行模型。未来可能会出现全栈的 megakernel 框架，支持自动算子融合、分布式任务调度和通信优化，从而大幅简化大规模 AI 系统的开发与部署。
第三，Perseus 的传输层优化技术具有广泛的适用性，不仅可以应用于 MoE 推理，还可以推广到其他类型的分布式 AI 工作负载，如大模型训练、分布式推理、图神经网络计算等。未来的研究可以探索如何将 Perseus 的核心思想应用于这些领域，解决它们面临的类似通信瓶颈问题。
最后，Perseus 的成果也对 AI 硬件设计提出了新的启示。随着 GPU 发起通信成为主流，NIC 与 GPU 之间的接口设计将变得越来越重要。未来的硬件设计可能会提供更丰富的硬件原语，支持更高效的细粒度通信和排序，从而进一步消除软件层的开销，实现真正的零延迟通信。

总之，Perseus 是多节点 MoE 推理领域一个里程碑式的工作。它不仅解决了一个长期困扰业界的关键性能问题，更为未来大规模 AI 基础设施的发展指明了新的方向。我们有充分理由相信，在 Perseus 及其后续工作的推动下，分布式 AI 系统的性能和效率将迎来新一轮的飞跃。

如需交流或加群，请在 NeuralTalk 公众号后台回复：加群

关注“鲸栖”小程序，掌握最新AI资讯

本文来自网络搜集，不代表鲸林向海立场，如有侵权，联系删除。转载请注明出处：https://www.itsolotime.com/archives/35005

MoE多节点推理性能瓶颈被颠覆：Perseus系统消除隐藏序列化，实现10倍加速