NCCL EP: Towards a Unified Expert Parallel Communication API for NCCL

作者/机构: Amos Goldman, Nimrod Boker, Maayan Sheraizin, Nimrod Admoni, Artem Polyakov, Subhadeep Bhattacharya, Fan Yu, Kai Sun, Georgios Theodorakis, Hsin-Chun Yin, Peter-Jan Gootzen, Aamir Shafi, Assaf Ravid, Salvatore Di Girolamo, Manjunath Gorentla Venkata, Gil Bloch, NVIDIA Corporation

---

A1 主要贡献

本文介绍了 NCCL EP (Expert Parallelism)，一个完全基于 NCCL Device API 构建的全新 MoE（Mixture-of-Experts）通信库。

核心问题与研究目标:
随着大型语言模型（LLM）的发展，MoE 架构已成为扩展其计算能力的关键。然而，专家并行（EP）引入了与传统集合通信不同的通信模式：细粒度、动态的 all-to-all 交换。现有的设备端发起的 MoE 通信库（如 DeepEP、Hybrid-EP）虽然性能优越，但面临生态系统碎片化和 API 复杂性的挑战。这些库独立于 NCCL（事实上的 GPU 集合通信标准），并且为不同模式（如低延迟、高吞吐）提供不同的 API，增加了应用开发和部署的复杂性。本文的目标是设计和实现一个统一的 MoE 通信库，将其整合到 NCCL 生态系统中，以解决上述碎片化和复杂性问题，同时保持高性能。

创新点与主要贡献:
本文的主要贡献在于 NCCL EP 的设计与实现，通过以下关键设计元素同时解决了 API 易用性和内核级性能问题：

1. 统一的 API 与算法模式：NCCL EP 为训练和推理应用提供了统一的 ncclEpDispatch/ncclEpCombine 接口。算法模式（低延迟 LL 或高吞吐 HT）在创建通信组时指定，未来可自动检测，允许应用程序无需修改代码即可切换模式。

2. 两级资源管理：通过 group/handle 架构，NCCL EP 将长生命周期的资源（如通信缓冲区、网络连接）与单次操作的状态（如路由决策）分离。ncclEpGroup_t 管理跨操作共享的资源，而 ncclEpHandle_t 捕获单次前向传播的路由决策。这种设计分摊了设置成本，同时适应了 MoE 架构固有的动态路由模式。

3. 设备端发起的通信：低延迟（LL）和高吞吐（HT）模式都依赖于通过 NCCL Device API 实现的 GPU 发起的通信，将 CPU 从数据路径中排除，实现了异步 RDMA 通信。

4. 混合通信路径：NCCL EP 自动选择最优传输方式：对节点内通信使用 NVLink，对节点间通信使用 RDMA，并在 NVLink 不可用时平滑回退到 PCIe。

5. 优化的低延迟（LL）设计：本文提出了一种内存优化的 LL 内核版本，其内存占用减少了一个数量级。

此外，为了在生产推理服务系统中评估 NCCL EP，本文提供了其与 vLLM 和 Megatron-LM 框架的集成。在 H100 集群上进行的全面评估表明，NCCL EP 在多节点配置下展现了有竞争力的 LL 内核性能，并展示了与 vLLM 集成的端到端结果。

---

A3 背景知识

Transformer 架构
现代大语言模型基于 Transformer 架构【索引14，Attention is all you need，2017，NIPS】。Transformer 使用由自注意力（self-attention）和前馈网络（FFN）块组成的层堆栈来处理 token 序列。自注意力捕捉 token 之间的长距离依赖和上下文关系，而 FFN 对每个 token 表示进行独立的非线性变换，以增强模型的表达能力。在前向传播中，token 嵌入通过 Transformer 层产生输出 logits；在训练的后向传播中，通过反向传播计算梯度以更新模型参数。

MoE 通信模式
MoE 模型扩展了 Transformer 架构，将 FFN 实现为分布在可用 GPU 上的专家网络【索引1，Outrageously large neural networks: The sparsely-gated mixture-of-experts layer，2017，ICLR】【索引2，GShard: Scaling giant models with conditional computation and automatic sharding，2021，ICLR】。在前向传播中，token 嵌入被路由到指定的专家子集，然后收集专家的响应。在后向传播中，损失梯度被路由到提供答案的专家以更新其参数。MoE 通信主要包括两个操作：
* Dispatch（分发）：将 token 从源 GPU 路由到托管其指定专家的 GPU。门控网络为每个 token 选择 top-k 个专家，产生定义了动态、不规则通信模式的路由索引。
* Combine（合并）：将专家输出聚合回原始 token 的位置。当 top-k > 1 时，Combine 根据门控权重对专家输出进行加权求和。

这些操作与标准集合通信有根本不同，其特点包括：动态路由、不规则消息大小、负载不均衡以及需要融合特定的计算（如 dispatch 中的量化和 combine 中的自定义 reduce）。标准的 NCCL AllToAll 无法高效处理这些不规则模式，因此催生了专门的 MoE 通信库。

NCCL Device API
NCCL 2.28 引入了 Device API【索引12，Fusing Communication and Compute with New Device API and Copy Engine Collectives in NVIDIA NCCL 2.28，2025，NVIDIA Blog】【索引13，GPU-Initiated Networking for NCCL，2025，arXiv】，允许 GPU 内核直接发起通信而无需 CPU 介入。该 API 包括三个模块：
* Load/Store Accessible (LSA)：通过 NVLink 提供对远程 GPU 内存的直接访问。应用程序可通过 ncclCommGetLocalSymm 获取指向对等 GPU 内存的本地可访问指针，并用于标准的 CUDA 加载/存储操作，非常适合节点内 token 传输。
* GPU-Initiated Networking (GIN)：使 GPU 内核能够对远程节点执行 RDMA 操作。GIN 提供了设备端可调用的原语，如 ncclGin_Put、ncclGin_Get、ncclGin_SignalAdd 和 ncclGin_WaitOnSignal。GIN 操作完全在 GPU 上执行，消除了 CPU 代理线程的开销。
* Multimem：提供跨多个 GPU 的集合加载/存储操作，适用于 NVLink 域内的广播和归约模式。

通过基于 NCCL Device API 构建，NCCL EP 继承了 NCCL 的拓扑检测、网络插件架构和生态系统兼容性，同时获得了设备端发起通信的能力。

MoE 通信现状
高效 MoE 通信的需求催生了多样化的专用库生态系统，但没有单一解决方案能在一个框架内覆盖所有用例。这导致了生产部署中的几个问题：
* 生态系统隔离：大多数库在 NCCL 之外运行，需要独立的初始化、内存管理和调试设施。
* 硬件耦合：一些解决方案需要特定的传输后端（如 NVSHMEM, IBGDA），将部署与特定的 NIC 配置绑定。
* API 复杂性：库为低延迟和高吞吐模式提供独立的接口，使框架集成复杂化。

相比之下，NCCL EP 提供了一个完全基于标准 NCCL 生态系统构建的统一 MoE 通信 API。

表 I MOE 通信库：特性比较

---

A2 方法细节

NCCL EP: 设计与 API

设计理念
NCCL EP 的设计遵循了几个源自 NCCL 编程模型的关键原则：
* 原生 C API 与 Python 绑定：与现有 MoE 库不同，NCCL EP 提供了一流的 C API，与 NCCL 的语法和惯例保持一致，支持直接集成到 C++ 运行时（如 TensorRT-LLM），同时通过基于 ctypes 的包装器支持 Python 工作流。
* 用户控制的资源生命周期：用户明确控制资源的分配和销毁，从而实现可预测的内存使用，并能与自定义内存分配器集成。
* 基于流的异步执行：所有 NCCL EP 操作都在 CUDA 流上异步执行，与 NCCL 执行模型一致，避免了影响性能的隐式同步点。
* 硬件抽象：NCCL EP 抽象了通信和计算组件，能够无缝支持未来的硬件特性或替代的网络后端。
* GPU 加速通信：NCCL EP 利用 NVIDIA GPU 的大规模并行性，以融合 CUDA 内核的形式加速通信和计算组件，将 token 处理分布到可用的流式多处理器（SM）上。

核心操作
NCCL EP 提供了四个核心操作来实现 MoE 通信模式。图 2 展示了典型的执行流程。
* Dispatch (ncclEpDispatch)：将 token 发送到其指定的专家。它接受输入张量列表（要路由的 token 和元数据）、输出张量列表（预分配的接收缓冲区）和本地张量列表（可选的每 rank 元数据，如 token 数量）。send_only 标志支持分阶段执行，允许计算与通信重叠。
* Combine (ncclEpCombine)：收集专家输出并将其归约回原始的 token 顺序。它利用存储在 handle 中的路由状态来反转 dispatch 的排列。当提供 top-k 权重时，combine 在返回 token 前应用加权归约。与 dispatch 类似，它也支持通过 send_only 标志进行分阶段执行。
* Complete (ncclEpComplete)：完成分阶段操作，等待前一个 send_only=1 操作发起的所有传入数据传输完成。
* Query (ncclEpHandleGetNumRecvTokens)：在 HT 模式下返回实际将接收的 token 数量，当 max_tokens_per_rank 设置为 auto 时，能够精确分配输出缓冲区。

表 II NCCL EP C API：资源管理函数创建和销毁组（长生命周期）和句柄（每次前向传播）；通信函数执行分发和合并操作，并可选地进行分阶段执行以实现重叠。

资源管理
NCCL EP 采用两级资源层次结构，将长生命周期的配置与每次操作的状态分开，从而实现资源在训练迭代或推理请求中的高效重用。
1. MoE 组 (MoE Group)：ncclEpGroup_t 是顶层资源容器，通过 ncclEpCreateGroup() 从一个现有的 NCCL 通信器创建。组的配置由算法模式（LL 或 HT）、专家数量、每 rank 的最大 token 数和缓冲区大小参数决定。该组封装了通信算法、预分配的发送/接收缓冲区和网络资源（如与所有对等方的连接）。组的创建是集体操作，通常在模型生命周期内持续存在。可选的分配器回调函数（alloc_fn, free_fn）允许与自定义内存池集成。
2. MoE 句柄 (MoE Handle)：ncclEpHandle_t 维护每次传递的状态，通过 ncclEpCreateHandle() 使用路由信息（topk_idx）创建。在 HT 模式下，句柄创建会触发跨 rank 的元数据交换以计算 token 分布；在 LL 模式下，这种交换在 dispatch 期间隐式发生。句柄在匹配的 dispatch 和 combine 操作之间共享，通常在每次前向传播之前创建，并只要路由不变就保留。

算法模式
NCCL EP 支持两种为不同工作负载优化的算法模式，在创建组时确定。
* LL 模式：为低延迟优化。它将 token 数据作为输入，并产生一个三维、专家主序（expert-major）的输出张量，可直接用于后续的专家 GEMM 操作（图 3）。LL 执行可以分阶段进行（send_only=1），以实现专家计算与数据传输的重叠。此外，LL 支持通过双缓冲重叠连续的 dispatch 和 combine 阶段。
* HT 模式：为高吞吐量优化，并利用聚合和网络层次结构。为了精确调整缓冲区大小，元数据交换在句柄创建期间跨所有 rank 启动。其二维输出格式（图 4）和每专家的 token 计数确保了可复现训练的确定性排序。

表 III HT 和 LL 内核模式的比较

张量抽象
NCCL EP 引入了 ncclNDTensor_t，一个编码了布局、数据类型和语义角色的 N 维张量描述符。它支持非连续张量，并包含 NCCL 数据类型、语义标签和指向 GPU 内存的指针。张量标签（通过 ncclEpTensorCreate 创建）在传递给 dispatch 和 combine 操作时识别每个张量的角色（见表 IV），用于验证张量形状并区分数据和元数据。

输入和输出张量布局因算法模式而异：
* Dispatch 输入：始终为二维，token 形状为 [num_tokens × hidden]，路由元数据形状为 [num_tokens × top_k]。
* Dispatch 输出：
* LL 模式 (图 3)：三维，形状为 [num_local_experts × tokens_per_expert × hidden]，token 按接收专家分组。这种专家主序布局可直接输入到分组 GEMM 操作。
* HT 模式 (图 4)：二维，形状为 [total_recv_tokens × hidden]，来自所有专家的 token 被连接在一起。一个单独的每专家 token 计数张量指示边界。

• Combine：输入形状与相应的 dispatch 输出匹配；输出形状与 dispatch 输入匹配，将 token 恢复到原始顺序。

表 IV 标识 ncclNDTensor_t 在 ncclEpTensorCreate、ncclEpDispatch 和 ncclEpCombine 中角色的张量标签。

使用模式：训练和推理中的 LL 和 HT
在 HT 模式训练中，MoE 句柄在前向和后向传播之间共享以维持路由状态，如图 5 所示。句柄在每层的前向传播开始时创建，在后向传播结束时销毁。

// 每个模型一次
ncclEpCreateGroup(&group, comm, &config, stream);
// 前向传播：在 ep[] 中为每层创建一个句柄
for (int l = 0; l < L; ++l) {
    topk = route(tokens);
    ncclEpCreateHandle(&ep[l], group, topk);
    ncclEpDispatch(ep[l], tokens, exp_in, ...);
    expert_ffn_forward(exp_in, exp_out);
    ncclEpCombine(ep[l], exp_out, tok_upd, ...);
    tokens = normalize(tokens, tok_upd);
}
// 后向传播：重用 ep[] 中的句柄
loss = calc_loss(ideal, tokens);
for (int l = L - 1; l >= 0; --l) {
    ncclEpDispatch(ep[l], loss, exp_in, ...);
    expert_ffn_backward(exp_in, exp_out);
    ncclEpCombine(ep[l], exp_out, loss_upd, ...);
    loss = update_loss(loss, loss_upd);
    ncclEpHandleDestroy(ep[l]);
}

在 LL 模式推理（解码）中，分阶段执行可以重叠多个微批次（micro-batches）：

// 微批次 0: 分发 (仅发送)
ncclEpDispatch(ep[0], tokens[0], exp_in[0], ..., /*send_only=*/1, stream);
// 微批次 1: 在 0 传输时分发
ncclEpDispatch(ep[1], tokens[1], exp_in[1], ..., /*send_only=*/1, stream);
// 完成微批次 0 的接收并运行专家
ncclEpComplete(ep[0], NULL, stream);
expert_ffn_forward(exp_in[0], exp_out[0]);
ncclEpCombine(ep[0], exp_out[0]], tok_upd[0], ...);
// 继续流水线...

Python API
NCCL EP 通过一个基于 ctypes 的包装器提供 Python 绑定，无需编译，可快速进行原型设计和框架集成。该包装器镜像了 C API，并能自动将 PyTorch 张量转换为 ncclNDTensor_t 描述符，从而与 vLLM 等基于 PyTorch 的框架无缝集成。

低延迟内核 (Low-Latency Kernels)

MoE 模型
NCCL EP 假设 E 个专家按块状分布在 N 个应用 rank 上，每个 rank 有 $L = \lceil E/N \rceil$ 个专家。在每个 rank $r$ 上，注意力子层产生一批 B 个固定大小的 token。对于 rank $r$ 上的 token $t \in [0, B)$，MoE 路由器定义了 t 被路由到的 K 个专家索引向量 $R(r, t)$。通信在数据并行（DP）侧和专家（E）侧之间进行，形成有效的专家-rank 对。

DeepEP 通信模型
NCCL EP 的 LL 模式采用了 DeepEP 的设计，并进行了优化。
* 拓扑：LL 内核采用完全的 N-to-N 网状连接，每个 GPU 可以直接与任何其他 GPU 通信。
* 通信缓冲区：LL 模式分配两个内部缓冲区，用于双缓冲方案以重叠 dispatch 和 combine 操作。每个缓冲区分为三个区域：协调区（coordination）、发送区（send）和接收区（receive）。为避免同步开销，每个有效的专家-rank 对在协调区和接收区都有一个专用的子区域。
* 通信模式：dispatch 和 combine 内核结构相似，分为发送（send）和接收（recv）阶段，通过细粒度的计数器更新和刷新操作进行协调。
* 发送阶段：rank $r$ 在其发送区域为每个输入 token 准备有效载荷消息，并直接写入对等方 $r'$ 的接收子区域。在 dispatch 中，写入的 slot 索引 $\alpha$ 是本地计算的下一个可用索引；在 combine 中，$\alpha$ 是 token 的源索引 $t$。发送完一个专家-rank 对的所有 token 后，通过更新和刷新操作通知对等方。
* 接收阶段：rank $r$ 监控所有有效专家-rank 对的计数器。一旦计数器更新，就从相应的子区域提取数据到输出张量。在 dispatch 中，输出张量布局与接收区域相同，同时缓存 slot 到 token 索引的映射以供 combine 使用。Combine 对每个 token 的 K 个专家输出执行加权求和。

DeepEP 并行化模型
LL 内核在一组 S 个 SM 上执行，每个 SM 进一步划分为 W=32 个 warp。一种常见的并行化策略是基于专家-rank 对进行划分。每个 SM 分配 $E_{SM} = \lfloor E/S \rfloor$ 个对，每个对分配一个由 $G = \lfloor W/E_{SM} \rfloor$ 个 warp 组成的组 $\epsilon_e$：

$$\epsilon_{e}=\{\omega_{ij}|(i,j):e=(i\cdot E^{SM}+\lfloor j/G\rfloor\cdot G)\}$$

• Dispatch/send：此阶段包括三个步骤：(a) 统计 token 数，(b) 发送 token，(c) 更新计数器。统计和发送并行执行。统计 token 数由每个 SM 的一个 warp 协作完成。发送 token 时，批次中的 B 个 token 均匀分布到 S 个 SM，每个 token 的 K 个目标专家由不同的 warp 处理。最后，为每个专家-rank 对分配的 warp 组 $\epsilon_e$ 等待前两步完成后发出更新和刷新操作。
• Dispatch/recv 和 Combine/send：这些步骤是对称的，并采用相同的并行化策略，根据专家-rank 对分配工作。在 dispatch 中，warp 原子地在输出张量中为专家 $e$ 保留一块连续的 slot，并缓存该信息以供 combine 使用。
• Combine/recv：在此步骤中，可用的 warp 被分成多个归约组 $\rho$，每个组包含 R 个 warp；每个 SM 托管 $G'$ 个组：

$$\rho = \{\rho_m = \{w_{ij} : \lfloor j/R \rfloor < G', m = i \cdot G' + \lfloor j/R \rfloor\}\}$$

所有接收到的 token 均匀分布到 $\rho$ 组。在每个组内，一个 warp 使用 NVIDIA Tensor Memory Accelerator (TMA) 将 K 个专家响应的块加载到 SM 的共享内存中，其余 warp 协作执行加权归约，并将结果存回输出张量。

NCCL EP 实现
NCCL EP 的实现基于上述 DeepEP 设计，并针对 NCCL 基础设施进行了适配和优化。
* 设备端发起的通信：NCCL EP 通过 NCCL Device API 实现通信通道。当目标在同一 NVLink 域内时，使用 LSA 的加载/存储指令；否则，通过 NCCL GIN API 使用 RDMA 网络。更新和刷新操作对于 NVLink 域通过 store-release/load-acquire 实现，对于 RDMA 则使用带有增量修饰符的零字节 GIN put 操作（ncclGin_SignalAdd）。
* 通信缓冲区布局优化：DeepEP 的通信缓冲区大小随 $O(E \cdot B \cdot P)$ 缩放，导致资源浪费，因为 B 个 slot 不会被所有专家完全占用。为了解决这个问题，NCCL EP 在 dispatch 消息头中扩展了路由信息 $R(r, t)$，从而避免了基于布局的路由。在发送 token 步骤中，一个 token 每个目标 rank 只发送一次，而不是每个专家一次。这使得 dispatch 缓冲区大小减少为 $O(N \cdot B \cdot P)$。Combine 过程也受益于此，通过缓存路由信息，实现了高效的子区域索引，缓冲区大小缩减为 $O(B \cdot K \cdot P)$。对于 N=64 ranks, E=512 experts, K=8 的配置，内存占用减少了约 14 倍。

$$\frac{2 \cdot E \cdot B \cdot P}{N \cdot B \cdot P + B \cdot K \cdot P} = \frac{2 \cdot E}{(N + K)}$$

高吞吐内核 (High-Throughput Kernels)

内核架构
NCCL EP 的 HT 模式针对训练和推理预填充工作负载，其内核改编自 Hybrid-EP【索引10，Optimizing Communication for Mixture-of-Experts Training with Hybrid Expert Parallel，2026，NVIDIA Blog】。Hybrid-EP 使用 warp 专用的持久性内核，其中每个 CUDA 块独占一个 SM 并运行一个完整的数据流水线。块内的不同 warp 组处理不同的流水线阶段：从全局内存加载数据到共享内存、通过网络传输、以及将结果写回。这种设计在只使用少量 SM 的情况下就能达到接近峰值的带宽，从而使大部分 SM 可用于重叠计算和通信。

NCCL GIN 适配
NCCL EP 的关键适配是将 Hybrid-EP 的 IBGDA 调用替换为 NCCL GIN 原语。IBGDA 提供低级的 GPU 发起 RDMA 接口，需要显式的 NIC 拓扑配置和内存注册。而 NCCL GIN 通过 NCCL 的可移植抽象层提供类似功能，自动处理拓扑发现并在不同网络结构上提供一致的语义。IBGDA 的 RDMA 写操作变为 ncclGin.put() 调用，流控制则通过 NCCL GIN 信号操作实现。基于 TMA 的节点内路径保持不变。这种适配简化了部署，因为 NCCL GIN 可以自动处理拓扑检测，并支持 NCCL 支持的各种网络。

缓冲区管理
HT 内核需要注册的缓冲区用于跨 rank 访问。节点间通信的内存使用 NCCL GIN 窗口注册；节点内通信则使用 CUDA IPC 句柄。由于缓冲区注册开销大，这些缓冲区在初始化时一次性分配，大小按所有 token 可能路由到单个 rank 的最坏情况确定。一个额外的设备到设备拷贝操作用于在这些注册缓冲区和标准的 PyTorch 张量之间移动数据，这个拷贝可以与专家 MLP 前的排列操作融合以隐藏开销。

与低延迟模式的比较
HT 模式的输出是一个二维张量，形状为 [num_recv_tokens × hidden]，并附带一个指示边界的每专家 token 计数的张量。这种格式适合处理可变大小专家批次的分组 GEMM 操作。相比之下，LL 模式产生一个三维专家主序布局，为解码阶段典型的小批量大小进行了优化。

---

A7 补充细节

框架集成

通用后端架构
集成堆栈由三层组成：NCCL EP C API、nccl_ep Python 包（ctypes 绑定）和一个缓冲区抽象层。该缓冲区抽象层将框架张量适配到 NCCL EP，并管理句柄、内存和流同步。
* 缓冲区接口：该接口暴露了 dispatch 和 combine 作为主要的 MoE API，直接映射到 ncclEpDispatch 和 ncclEpCombine。NDTensorWrapper 将框架张量转换为 ncclNDTensor_t 描述符。接口还提供了获取每专家 token 数、通信流和事件同步的功能。
* 内存和流管理：为避免频繁的 CUDA 分配，缓冲区将 NCCL EP 的分配器回调连接到 PyTorch 的 CUDA 缓存分配器。一个专用的通信流将通信与计算隔离开来，为未来的重叠执行奠定了基础。
* 专家计数：在 HT 模式下，缓冲区分配固定的、GPU 映射的主机内存，以便 NCCL EP 可以写入专家 token 计数，供后续的 GEMM 操作使用。

表 V 缓冲区操作：框架层面的 API 映射到 NCCL EP 原语。

Megatron-LM 集成
NCCL EP 作为一种新的设备端后端集成到 Megatron-LM【索引18，Efficient large-scale language model training on GPU clusters using Megatron-LM，2021，SC ’21】 的 Flex 插件架构中，支持 HT 模式。
* Dispatcher Manager：_NCCLEPManager 适配 Megatron 的 MoE 层，将 Flex 的 multi-hot 路由格式转换为 NCCL EP 的紧凑 top-k 格式。
* Autograd：NCCLEPDispatch 和 NCCLEPCombine 包装了缓冲区操作以支持自动微分。在前向传播中，dispatch 创建句柄并重新分配 token；在后向传播中，缓存的 dispatch 重新分配梯度，然后 combine 聚合梯度，最后销毁句柄。
* Buffer：NCCLEPBuffer 实现了与 DeepEP 缓冲区 API 兼容的通用接口。

vLLM 集成
NCCL EP 作为 vLLM【索引19，Efficient memory management for large language model serving with PagedAttention，2023，SOSP ’23】中一个额外的可插拔 All2All 后端（nccl_high_throughput 或 nccl_low_latency）进行集成，支持两种算法模式。
* All2All Manager：NCCLAll2AllManager 充当缓冲区实例的工厂和缓存。
* FusedMoE Layer：该层通过 vLLM 的 prepare/finalize 抽象来集成 NCCL EP。prepare 阶段调用 dispatch 并获取句柄；finalize 阶段调用 combine。
* 缓冲区和激活格式：缓冲区提供每专家的 token 计数以指导 GEMM 和 combine 操作。LL 模式下的批处理专家布局支持缓冲区预分配以降低延迟，而 HT 模式下的二维连接布局支持确定性排序。

相关工作

设备端发起的 MoE 库
* DeepEP【索引8，DeepEP: Efficient Mixture-of-Experts Communication Library，2025，GitHub】: 使用 NVSHMEM 和 IBGDA 的开创性工作，为 LL 和 HT 模式提供独立的 Python API。NCCL EP 的 LL 内核改编自其设计。
* Perplexity pplx-kernels【索引9，pplx-kernels: Perplexity GPU Kernels for MoE Communication，2025，GitHub】: 使用 NVSHMEM 和 IBGDA，专注于低延迟推理。
* NIXL EP【索引21，NIXL: NVIDIA Inference Xfer Library，2025，GitHub】: NVIDIA 推理传输库中的专家并行模块，源自 DeepEP。
* AMD MORI【索引17，MORI: Modular RDMA interface，2026，GitHub】: 为 AMD MI 系列 GPU 提供 ROCm RDMA 原语。

分层和可移植方法
* HybridEP【索引10，Optimizing Communication for Mixture-of-Experts Training with Hybrid Expert Parallel，2026，NVIDIA Blog】: NVIDIA 的 MoE 通信库，使用 warp 专用流水线、TMA 和 IBGDA。NCCL EP 的 HT 内核改编自其设计。
* UCCL-EP【索引15，UCCL-EP: Portable Expert-Parallel Communication，2025，arXiv】: 通过 GPU-CPU 控制通道实现可移植性，CPU 代理发出 RDMA，以延迟换取可移植性。
* Meta NCCLX AllToAllvDynamic【索引16，Collective communication for 100k+ gpus，2025，arXiv】: 支持 GPU 端的路由元数据，是 NCCL 的完全替代品，用于 Meta 的大规模服务。

推理和服务平台
* Mooncake【索引22，Mooncake: Trading more storage for less computation — a KVCache-Centric architecture for serving LLM chatbot，2025，FAST ’25】: Moonshot AI 的 KVCache 中心平台，其传输引擎已扩展用于 MoE。
* vLLM 和 SGLang: 支持多种 AllToAll 后端的推理框架，推动了可插拔 MoE 后端的发展。

NCCL EP 的定位
NCCL EP 与现有解决方案的不同之处在于三个关键方面：
1. NCCL 集成：与独立库或专有替代品不同，NCCL EP 通过 Device API 将 MoE 原语本地扩展到 NCCL 中，保持了生态系统兼容性和统一的资源管理。
2. 统一 API：与需要为 LL 和 HT 模式提供独立接口的库不同，NCCL EP 提供了统一的 ncclEpDispatch/ncclEpCombine 原语。
3. C 和 Python 接口：NCCL EP 同时提供 C API 和 Python 绑定，比仅提供 Python 或仅提供 C++ 的替代方案更具灵活性。

---

A4 实验环境

所有实验均在 NVIDIA EOS H100 集群上进行。

• 硬件配置：

  • 集群规模: 576 个节点。
  • GPU: 每个节点 8 x H100 80GB HBM3 GPU。
  • GPU 内存带宽: 3.35 TB/s。
  • 互联: 第 4 代 NVLink，900 GB/s，每个 GPU 18 个链路。
  • CPU: Xeon 8480CL。
  • 网络: 每个节点 8 x 400 Gbit/s InfiniBand。

• 软件配置：

  • LL 内核基准测试: DeepEP 配置为使用 NCCL v2.29 作为其网络后端，以进行公平比较。
  • vLLM 集成测试: vLLM v0.10。NCCL EP 与 DeepEP v1.2.1（使用其默认的 NVSHMEM 3.2.5/IBGDA 后端）进行比较。

• 模型与数据集：

  • vLLM 测试模型: Qwen3-30B-A3B，一个 30B 参数的 MoE LLM。
  • 请求: 1000 个请求，最大并发数为 32。
  • LL 内核测试参数: 256 个专家，隐藏维度 7168，128 个 token，top-k=8，BF16 精度。

表 VI EOS 评估集群的硬件规格。

---

A4 实验结果

低延迟（LL）内核性能

• 实验内容：在 1 到 8 个节点（8 到 64 个 GPU）上，比较 NCCL EP 和 DeepEP（两者都使用 NCCL 作为网络后端）的 LL dispatch 和 combine 操作的吞吐量。

• 实验结果：

  • Dispatch 吞吐量 (图 7)：在多节点规模（2-8 个节点）上，NCCL EP 的性能与 DeepEP相当或更高，在 8 节点（64 GPU）时有显著优势。在单节点上落后约 4%。
  • Combine 吞吐量 (图 8)：在 1-4 个节点上，NCCL EP 比 DeepEP 落后 5-12%，在 8 节点时性能趋于一致。

• 分析结论：NCCL EP 的计时包含了内核启动开销（主机端 C API 计时），而 DeepEP 报告的是纯 GPU 内核执行时间。这意味着 NCCL EP 的实际内核级吞吐量比报告的要高，因此 dispatch 的优势被低估，而 combine 的差距比显示的要小。

vLLM 集成性能
* 实验内容：在 vLLM 中比较 NCCL EP 和 DeepEP（使用其默认的 NVSHMEM/IBGDA 后端）的端到端服务指标。
* 实验结果 (表 VII)：在所有规模（1、2、4 节点）下，NCCL EP 在吞吐量方面比 DeepEP 落后 7-10%，在 token 间延迟（ITL/TPOT）方面落后 7-9%。
* 分析结论：这种差距不能完全由传输后端的差异来解释。在 LL 内核评估中观察到的 combine 开销是造成这种差距的一个因素。调查和优化这一差距是未来 NCCL EP 版本的重点工作。

表 VII vLLM 0.10 服务指标：NCCL EP 与 DeepEP v1.2.1 与 NVSHMEM 3.2.5/IBGDA 的比较

---

A5 结论

本文介绍了 NCCL EP，一个完全基于 NCCL 基础设施构建的全新 MoE 通信库。它通过统一的 ncclEpDispatch 和 ncclEpCombine 原语，以及 C 和 Python 接口，解决了现有 MoE 库生态系统碎片化的问题。NCCL EP 提供了一个单一接口，支持低延迟模式（用于推理解码）和高吞吐量模式（用于训练和预填充）。

在实现上，NCCL EP 采用了行业最佳实践，为低延迟场景提出了一种 DeepEP 设计的优化版本，性能相当但内存占用减少了一个数量级；为高吞吐量场景集成了 NVIDIA HybridEP 解决方案。GPU 发起的通信利用了 NCCL Device API，为 NVLink 通信提供 LSA 模式，为节点间 RDMA 提供 GIN。

评估结果显示，NCCL EP 在多节点配置下实现了与 DeepEP 相当的低延迟内核性能。端到端的 vLLM 集成展示了与 DeepEP 相当的服务性能。这些结果表明，在 NCCL 内部原生构建 MoE 通信，可以在保持性能的同时，带来生态系统的好处：拓扑感知、容错以及与 NCCL 生产基础设施的集成。

未来工作：
NCCL EP 目前是一个实验性 API。Megatron-LM 的集成已经完成，但训练性能受限于正在进行的多节点高吞吐量内核优化。社区的参与和反馈对于完善 API、发现性能优化机会，以及确保 NCCL EP 满足 MoE 工作负载在训练和推理部署中的多样化需求至关重要。