MoE โหนดหลายจุดคอขวดประสิทธิภาพการอนุมานถูกพลิกโฉม: ระบบ Perseus ขจัดการซ่อนลำดับแบบอนุกรม บรรลุความเร็วเพิ่มขึ้น 10 เท่า

7 hours ago • การอนุมานโมเดลขนาดใหญ่ • 17 views

Perseus พลิกโฉมคอขวดประสิทธิภาพการอนุมาน MoE แบบหลายโหนด: กำจัดการเรียงลำดับที่ซ่อนเร้น เพิ่มความเร็วสูงสุด 10 เท่า

การค้นพบหลักจากบทความ Perseus ของมหาวิทยาลัย Cornell ได้เขียนความเข้าใจดั้งเดิมของอุตสาหกรรมเกี่ยวกับคอขวดของการอนุมาน MoE แบบหลายโหนดใหม่ทั้งหมด: “งานวิจัยของเรายืนยันว่า ปัจจัยสำคัญที่จำกัดประสิทธิภาพของ megakernel แบบหลายโหนดคือปัญหาการเรียงลำดับ ไม่ใช่การเลือกระหว่างการส่งผ่านแบบพร็อกซีกับ GPU-direct”

เป็นเวลานานที่สถาปัตยกรรม Mixture of Experts (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 在不同模型、不同传输方式以及不同规模下的端到端性能表现。
⚠️ หมายเหตุ: เนื้อหาได้รับการแปลโดย AI และตรวจสอบโดยมนุษย์ หากมีข้อผิดพลาดโปรดแจ้ง

☕ สนับสนุนค่ากาแฟทีมงาน

หากคุณชอบบทความนี้ สามารถสนับสนุนเราได้ผ่าน PromptPay

SCAN TO PAY WITH ANY BANK

本文来自网络搜集，不代表คลื่นสร้างอนาคต立场，如有侵权，联系删除。转载请注明出处：https://www.itsolotime.com/th/archives/35006