GPU-Initiated Networking for NCCL

作者/机构: Khaled Hamidouche, John Bachan, Pak Markthub, Peter-Jan Gootzen, Elena Agostini, Sylvain Jeaugey, Aamir Shafi, Georgios Theodorakis, Manjunath Gorentla Venkata {khamidouche, jbachan, pmarkthub, pgootzen, eagostini, sjeaugey, ashafi, gtheodorakis, manjunath}@nvidia.com NVIDIA Corporation

A1 主要贡献

现代AI工作负载，尤其是专家混合（Mixture-of-Experts, MoE）架构，对低延迟、细粒度的GPU间通信以及设备端控制提出了越来越高的要求。传统的GPU通信遵循主机启动模式，即由CPU协调所有通信操作，这带来了同步开销。为了解决这一问题，NCCL 2.28引入了Device API，旨在通过设备启动的通信来消除CPU的协调开销，从而实现计算与通信的紧密集成。

本文重点介绍了NCCL Device API中的GPU启动网络（GPU-Initiated Networking, GIN）功能。GIN是Device API的三种操作模式之一，专为通过InfiniBand和RoCE网络进行的节点间通信而设计。另外两种模式是用于节点内NVLink/PCIe通信的Load/Store Accessible（LSA）和用于NVLink SHARP硬件多播的Multimem。

图1. NCCL Device API架构，展示了三种操作模式及其底层互连技术。Load/Store Accessible (LSA) 使用PCIe和NVLink进行节点内内存操作，Multimem利用NVLink SHARP进行硬件多播，而GPU启动网络 (GIN) 为基于InfiniBand和RoCE的网络通信提供双后端实现 (GDAKI和Proxy)。

GIN的核心设计是一个三层架构：
1. NCCL核心层主机端API：用于设备通信器设置和集合内存窗口注册。
2. 设备端API：提供可从CUDA核函数中调用的远程内存操作。
3. 网络插件架构：具有双重语义（GPUDirect Async Kernel-Initiated和Proxy），以支持广泛的硬件。

这种双后端设计是GIN的一大特色。GDAKI后端利用DOCA GPUNetIO实现GPU到NIC的直接通信，而Proxy后端则通过标准的RDMA网络，利用无锁的GPU到CPU队列提供等效功能，从而在不同硬件环境中保持功能兼容性。

图2. 高层架构比较：NCCL Device API（左）具有三种操作模式（用于NVLink/PCIe的Load/Store Accessible，用于NVLink SHARP的Multimem，用于网络RDMA的GIN），通过集合对称内存实现单次（single-shot）集合算法；而传统NCCL（右）使用主机启动的算法，通过流水线原语在常规内存上操作。

本文的主要贡献如下：
i) GIN的设计与实现：在NCCL中设计并实现了GIN，包括统一的主机和设备API、用于异步完成的模块化同步原语（signals和counters），以及两种可互换的后端架构：使用DOCA GPUNetIO进行直接GPU到NIC通信的GDAKI，以及通过标准RDMA进行CPU辅助操作的Proxy。
ii) 与DeepEP的集成：将GIN与专门的MoE通信库DeepEP集成，展示了GIN在实际应用中的适用性及其与现有基于NVSHMEM的设备启动通信的兼容性。
iii) 全面的性能评估：通过微基准测试和在DeepEP核函数中的应用级实验，对GIN的性能进行了全面评估，并进行了比较分析以确定其性能特征。

A3 背景知识

传统单向通信库的局限性。像OpenSHMEM这样的传统通信库提供了单向原语（put, get, atomics），这些原语在对称内存区域上操作，实现了无需发送方-接收方协调的异步数据移动【12，OpenSHMEM Application Programming Interface, Version 1.0, 2012, OpenSHMEM】。然而，这些规范假定执行是以CPU为中心的，所有通信原语都从主机代码中调用。当GPU集成到HPC系统中时，这种模式的低效性就暴露出来了：细粒度的GPU到GPU通信会产生核函数启动开销、通过主机内存暂存数据的PCIe传输开销以及CPU调度延迟【13，GPU-Centric Communication on NVIDIA GPU Clusters with InfiniBand: A Case Study with OpenSHMEM, 2017, 24th IEEE International Conference on High Performance Computing】。这些瓶颈推动了支持GPU感知的扩展，使得通信可以直接从CUDA核函数中发起【13，GPU-Centric Communication on NVIDIA GPU Clusters with InfiniBand: A Case Study with OpenSHMEM, 2017, 24th IEEE International Conference on High Performance Computing】,【14，Exploring OpenSHMEM Model to Program GPU-Based Extreme-Scale Systems, 2015, OpenSHMEM and Related Technologies. Experiences, Implementations, and Tools. Springer】。

GPUDirect技术的发展。GPUDirect RDMA（2013年）【15，GPUDirect RDMA, 2013, NVIDIA Corporation】,【16，Efficient Inter-Node MPI Communication using GPUDirect RDMA for InfiniBand Clusters with NVIDIA GPUs, 2013, 2013 42nd International Conference on Parallel Processing】使得支持RDMA的网络接口卡（NIC）能够通过PCIe基地址寄存器（BAR）映射直接访问GPU内存，从而消除了节点间传输数据路径中的CPU和主机内存。NIC的DMA引擎执行到GPU BAR的PCIe点对点事务，访问通过nvidia_p2p内核模块注册的内存区域。然而，GPUDirect RDMA仅在核函数边界提供一致性保证。GPU内存模型的语义（松散排序、写回缓存）阻止了正在执行的核函数安全地并发访问已注册RDMA的内存，迫使应用程序将计算和通信分开【17，GPU Initiated OpenSHMEM: Correct and Efficient Intra-Kernel Networking for dGPUs, 2020, 2020 IEEE International Parallel and Distributed Processing Symposium (IPDPS)】。

GPUDirect Async的引入。GPUDirect Async（2016年）【18，GPUDirect Async: Exploring GPU Synchronous Communication Techniques for InfiniBand Clusters, 2018, Journal of Parallel and Distributed Computing】引入了部分控制路径卸载：GPU线程通过写入映射到GPU地址空间的NIC门铃寄存器来触发预先配置的网络操作。然而，CPU必须预先构建通信描述符，这限制了操作只能是主机预先配置的那些，从而无法实现完全由设备自主驱动的网络通信。

设备启动通信原语的演进。完全由设备启动的网络通信要求直接在GPU代码中实现网络编程接口。早期的原型如GPUrdma【19，GPUrdma: GPU-Side Library for High Performance Networking from GPU Kernels, 2016, Proceedings of the ACM International Conference on Computing Frontiers】和GIO【17，GPU Initiated OpenSHMEM: Correct and Efficient Intra-Kernel Networking for dGPUs, 2020, 2020 IEEE International Parallel and Distributed Processing Symposium (IPDPS)】将InfiniBand verbs作为设备可调用函数暴露出来，但面临着GPU-NIC内存一致性的挑战。

NVSHMEM和DOCA GPUNetIO的贡献。NVSHMEM【8，NVSHMEM 3.0 Programming Guide, 2023, NVIDIA Corporation】将OpenSHMEM语义扩展到GPU集群，提供了可从CUDA核函数中调用的设备可调用单向操作（put、get、atomics）。这使得设备代码能够在计算和通信之间交错进行，而没有核函数启动开销，其传输后端包括用于节点间传输的IBGDA（带GPUDirect Async的InfiniBand）和用于节点内通信的对称内存机制。DOCA GPUNetIO【20，DOCA GPUNetIO Programming Guide, NVIDIA Corporation】为InfiniBand和RoCE网络（IBGDA）提供了GPU端的RDMA API，暴露了设备函数，使GPU核函数能够直接对NIC进行编程。具体来说，它实现了GPUDirect RDMA（直接GPU数据移动）和GPUDirect Async Kernel-Initiated（GPU控制网络通信）技术。它构成了GIN的GDAKI后端的基石，通过硬件支持的设备verbs实现了直接的GPU到NIC通信。

支持GPU启动通信的网络硬件。GPU启动的网络通信要求网络接口卡通过多种RDMA技术之一（InfiniBand、RoCE或iWARP【21，RDMA Technologies Comparison: InfiniBand, RoCE, iWARP, 2022, NVIDIA Corporation】）支持直接设备访问。

InfiniBand的特性。InfiniBand提供原生的RDMA支持和基于信用的流量控制，实现了约130纳秒的端口到端口延迟，并支持每个子网数万个节点【22，InfiniBand vs RoCE: A Comprehensive Guide, 2025, CloudSwitch】。InfiniBand适配器通过PCIe BAR暴露内存映射的队列对、完成队列和门铃寄存器，当与GPUDirect RDMA结合时，可实现GPU直接访问【13，GPU-Centric Communication on NVIDIA GPU Clusters with InfiniBand: A Case Study with OpenSHMEM, 2017, 24th IEEE International Conference on High Performance Computing】,【15，GPUDirect RDMA, 2013, NVIDIA Corporation】。

RoCE的特性与要求。RoCE在标准以太网上实现RDMA，成本更低，与现有数据中心基础设施的兼容性更广【21，RDMA Technologies Comparison: InfiniBand, RoCE, iWARP, 2022, NVIDIA Corporation】,【22，InfiniBand vs RoCE: A Comprehensive Guide, 2025, CloudSwitch】。流行的RoCEv2变体将InfiniBand传输封装在UDP/IP之上，实现了约400纳秒的端口到端口延迟——虽然高于原生InfiniBand，但对许多工作负载已足够【22，InfiniBand vs RoCE: A Comprehensive Guide, 2025, CloudSwitch】。RoCE需要使用优先级流量控制（PFC）和显式拥塞通知（ECN）的无损以太网配置来防止数据包丢失，这可能使在多租户环境中的部署变得复杂【22，InfiniBand vs RoCE: A Comprehensive Guide, 2025, CloudSwitch】。InfiniBand和RoCE共享相同的用户空间verbs API，从而实现了互连的可移植性【21，RDMA Technologies Comparison: InfiniBand, RoCE, iWARP, 2022, NVIDIA Corporation】。

GPU启动通信的硬件和系统要求。对于GPU启动的通信，硬件要求是NIC支持设备可访问的控制结构。NVIDIA ConnectX系列适配器（ConnectX-6 Dx及更高版本）和BlueField DPU通过DOCA GPUNetIO【20，DOCA GPUNetIO Programming Guide, NVIDIA Corporation】提供了此功能。缺乏此类硬件支持的系统无法启用直接的GPU-NIC通信，必须回退到CPU介导的机制。在GIN的架构中，这一限制催生了双后端设计：GDAKI后端利用DOCA GPUNetIO在支持的硬件上实现直接设备通信，而Proxy后端则通过无锁的GPU到CPU队列和CPU驱动的网络操作，在任意支持RDMA的NIC上提供功能等效的语义。

性能优化的系统配置。为获得最佳性能，需要将GPU和NIC共置在同一个PCIe根复合体上，以最小化点对点延迟并最大化带宽【15，GPUDirect RDMA, 2013, NVIDIA Corporation】,【23，GPUDirect RDMA Requirements and Recommendations, 2021, NVIDIA Corporation】。具有分布式PCIe拓扑的多插槽系统可能会因跨插槽遍历而产生性能损失，从而降低GPUDirect RDMA的效率。此外，GPU启动的网络通信需要nv_peer_mem内核模块（用于GPUDirect RDMA）和适当的驱动程序栈（用于InfiniBand的OFED，用于Mellanox适配器的MOFED）来建立GPU和NIC地址空间之间的内存映射【23，GPUDirect RDMA Requirements and Recommendations, 2021, NVIDIA Corporation】。

NCCL架构及其网络插件。NCCL是用于多GPU机器学习的标准集合通信运行时，为allreduce、allgather、reduce-scatter和broadcast等操作提供拓扑感知的实现。NCCL的架构使用CPU代理线程来协调网络操作，其中GPU核函数将通信描述符入队到主机可见的队列中，然后由CPU线程通过网络插件执行这些操作。虽然这种设计在大规模集合操作中已被证明是稳健的，但NCCL 2.28的Device API【9，Fusing Communication and Compute with New Device API and Copy Engine Collectives in NVIDIA NCCL 2.28, 2025, NVIDIA Developer Blog】通过设备端原语扩展了此架构，使应用程序能够直接从GPU代码实现自定义通信模式，将通信集成到计算核函数中，并为新兴工作负载实现细粒度的计算-通信重叠。

NCCL的生态兼容性。NCCL在生产框架中的广泛采用推动了设备启动通信能力的集成，保留了生态系统的兼容性，同时开启了新的用例。

NCCL网络插件架构。NCCL网络插件架构提供了一个抽象层，将核心库与具体的网络实现解耦。NCCL支持内置于库中的内部插件（例如，Socket和InfiniBand）以及作为共享库（libnccl-net.so）实现NCCL网络API的外部插件。这种设计允许网络供应商和硬件提供商通过专门的传输实现来扩展NCCL，而无需修改NCCL核心。外部插件在运行时动态加载，通过NCCL_NET_PLUGIN环境变量进行选择，从而实现了不同网络技术的无缝集成，并通过版本化的API接口保持版本兼容性。

专门的MoE通信库。LLM中的MoE架构需要动态、负载均衡的all-to-all token路由，其消息大小不可预测，从而产生了传统集合通信【24，DeepSeekMoE: Towards Ultimate Expert Specialization in Mixture-of-Experts Language Models, 2024, arXiv preprint】不擅长处理的非规则通信模式。像DeepEP【25，DeepEP: Efficient Mixture-of-Experts Communication Library, 2025, DeepSeek-AI】和Perplexity的pplx-kernels【26，pplx-kernels: Perplexity GPU Kernels for MoE Communication, 2024, Perplexity AI】这样的专门库，通过CUDA优化的、GPU启动的原语来针对这些工作负载，以实现低延迟的all-to-all传输。这些工作证明了以GPU为中心的通信对MoE工作负载的价值，但它们仍然独立于NCCL的生态系统。

GIN的价值。GIN将GPU启动的RDMA操作引入NCCL，使应用程序能够在统一的运行时中同时利用主机优化的集合操作和设备驱动的点对点通信。

章节过渡。基于这些设备启动的网络技术和NCCL的插件架构，下一节将介绍GIN的设计和实现。

A2 方法细节

本节介绍NCCL GIN的设计和实现。GIN通过将设备启动的单向原语集成到NCCL中，使得GPU线程可以直接从CUDA核函数中发起网络操作，而无需CPU的介入，从而实现了计算与通信的紧密耦合。

GIN的设计保留了NCCL既有的编程模型和生态系统集成，同时为设备驱动的通信增加了一条并行的低延迟路径。这使得像TensorRT-LLM、vLLM和SGLang等生产系统，以及像DeepEP这样的通信库，能够实现传统NCCL无法实现的自定义集合算法和核函数融合模式。

本节分为三部分：III-A介绍三层架构和核心设计原则；III-B介绍设备端API并通过一个实例演示其用法；III-C分析两种插件接口及其后端实现（GDAKI和Proxy），解释其设计理念和性能特点。

A. 核心原则和架构

GIN架构的核心目标。NCCL GIN的架构建立在单向通信语义之上，旨在实现两个关键性能目标：最大化通信-计算重叠和最小化端到端操作延迟。通过使GPU线程能够直接发起RDMA操作——即无需接收方协调即可读写远程内存——GIN消除了主机-设备同步开销和双向握手延迟，允许异步数据传输与计算并发进行。

GIN的三层架构。该架构由三个协同工作的层组成，如Figure 3中的绿色路径（左侧）所示，旨在平衡高性能与广泛的供应商支持：NCCL核心（主机端API）、Device GIN API（GPU可调用原语）和GIN网络插件（可插拔网络后端）。Figure 3中的灰色路径（右侧）显示了现有的基于内置网络插件的双向集合API。我们进一步讨论GIN的三个层次：
i) NCCL核心：主机端的NCCL功能负责管理内存窗口注册、资源分配和通信器初始化，为GIN资源管理提供基础，并通过设备启动的通信能力扩展NCCL的现有基础设施。
ii) Device GIN API：设备端API向GPU核函数暴露统一的接口，使应用程序能够直接从CUDA核函数中调用单向通信操作。它根据底层网络后端，分派到NCCL提供（Proxy）或插件提供（GDAKI）的实现。
iii) GIN网络插件：插件层提供了一种可扩展机制，定义了远程数据移动操作，并支持双重语义——GDAKI和Proxy——以最大化网络覆盖范围。NCCL的InfiniBand传输实现了这两种语义，而外部供应商可以提供自己的实现。在Proxy接口下，NCCL核心拥有控制结构、设备端排队逻辑和设备API实现，而插件仅提供基于CPU的put、signal、test和regMr操作，从而使得没有GPU直通能力的普通网络也能使用GIN，降低了GIN的采用门槛。在GDAKI语义下，插件同时拥有控制路径和设备API：它们通过createContext创建GPU上下文，并提供直接使用内核级API（如DOCA GPUNetIO）对NIC进行编程的设备代码，而NCCL核心则负责协调主机和设备组件之间的结构交换。

关键设计元素。这些组件通过几个关键设计元素实现了设备启动的单向通信：用于单方面数据移动的单向语义、用于零拷贝远程访问的对称内存窗口，以及具有灵活排序语义的异步完成跟踪。

单向通信语义。GIN暴露了单向RDMA原语——用于远程写入的put和用于带远程通知写入的put with signal——使GPU线程能够访问远程内存而无需任何接收方协调。这种单向模型消除了握手协议的开销和接收方参与的需要，允许发起方单方面发出传输请求，并独立控制何时验证完成。对于MoE工作负载中的不规则通信模式（动态令牌路由产生不可预测的流量模式）以及受益于并行、非阻塞对等通信的单次集合实现，单向模型被证明特别有效。

基于窗口的（非）对称内存。通信缓冲区必须在所有rank之间进行集体注册，从而建立在可寻址性上对称的内存窗口，这遵循了MPI RMA窗口模型【27，Remote memory access programming in mpi-3, 2015, ACM Trans. Parallel Comput.】。所有进程都可以访问注册的内存，类似于NVSHMEM的对称堆。GIN窗口设计支持容量上的非对称性：每个rank可以注册不同大小的缓冲区。这种灵活性对于分离式服务架构至关重要，其中预填充（prefill）阶段的rank比解码（decode）阶段的rank需要更大的缓冲区。需要注意的是，NCCL 2.28的当前实现强制要求对称大小，但这一限制将在未来版本中解除。此外，内存分配与注册是解耦的：NCCL-GIN内存窗口允许用户从现有分配中创建窗口。每次注册都会产生包含远程访问元数据的窗口句柄。窗口句柄为后端提供了特定的优化机会。后端可以直接从窗口元数据和使用rank相对偏移的目标地址构建RDMA描述符。

用于网络并行性的GIN上下文。GIN上下文是表达网络并行性的主要抽象。每个上下文抽象了GPU和NIC之间的一个通道，并封装了网络资源和连接（队列对，QPs）。每个通信器可以有多个上下文，使应用程序能够利用跨多个NIC、端口和QP的网络级并行性，从而实现独立的并发通信流。单个上下文可以寻址与通信器关联的每个rank。因此，一个上下文可以向不同的对等方发出多个并发操作。

异步完成跟踪。所有设备启动的操作都异步执行并立即返回，以便其他工作（如计算或通过NVLink的节点内通信）可以并行进行。应用程序通过两种不同的机制使用每个上下文的资源来跟踪操作完成情况。Counters是本地对象，用于在发送方跟踪完成情况，指示源缓冲区何时可以安全重用。与跟踪发布到上下文的所有操作完成情况的flush操作不同，Counters是一个强大的概念，可以按操作跟踪本地完成情况，允许用户高效地描述流水线算法。每个数据移动操作都可以选择性地向用户提供的计数器报告本地完成情况（通过counterID）。

Signals的远程完成跟踪。另一方面，Signals是对称对象，提供远程完成跟踪，确认数据到达目的地并对目的地可见。与OpenSHMEM基于地址的同步不同，GIN使用基于ID的寻址：每个signal（和counter）都由一个整数ID标识，而不是内存地址。这种基于ID的设计简化了资源管理，并能高效地实现硬件完成通知。

排序语义。为了最大化网络效率和吞吐量，GIN操作默认是无序的。然而，GIN仅在同一上下文中对同一对等方的put和signal操作之间提供排序保证。当一个signal操作（无论是独立的signal还是带ncclGin_SignalInc/ncclGin_SignalAdd动作的put）在目的地完成时，它保证了在同一上下文中发往该对等方的所有先前put操作都已完成，并且对远程GPU线程可见。这提供了轻量级的排序，而无需显式的fence操作：应用程序可以批量处理多个put，并将一个signal附加到最后一个操作上，从而确保整个序列的有序远程可见性。相比之下，flush操作仅确保本地完成——所有待处理的操作已被消费，源缓冲区可以安全重用——但对远程可见性不做任何保证。GIN基于signal的排序旨在通过有选择地使用signals而不是全局地提供排序保证来获得最大性能。值得注意的是，GIN不假定或保证GPU线程之间的任何排序。用户有责任使用CUDA同步原语来同步线程。

B. 设备端API和编程模型

API概述。设备端API为GPU核函数提供了对网络操作的直接控制，暴露了可从CUDA设备代码调用而无需CPU干预的方法。编程模型围绕ncclGin对象展开，该对象封装了网络资源（上下文、操作队列和对等连接状态），并提供了数据移动、完成跟踪和同步的方法。后端的选择（DOCA GPUNetIO或Proxy）在通信器初始化时根据硬件能力和用户配置透明地进行，无论底层实现如何，都提供相同的面向设备的接口。

API组织结构。该接口将操作分为四个逻辑类别，反映了通信工作流。数据移动操作（put、putValue、signal）发起对远程对等方的单向传输或通知，提交异步执行的RDMA操作。完成跟踪操作区分了本地完成（flush、readCounter、waitCounter）——表示源缓冲区可以安全重用——和远程完成（readSignal、waitSignal）——确认数据已到达目的地并对远程GPU线程可见。屏障同步（ncclGinBarrierSession）在通信阶段之前协调团队内的所有rank，通过全网络同步确保全局一致性。状态管理操作（resetCounter、resetSignal）重置完成状态以便在多个通信轮次中重用。这种关注点分离实现了对通信-计算重叠的细粒度控制，并支持多样的同步模式。

与GPU内存模型的交互。此外，GIN设备API与GPU内存模型有紧密的交互，并根据用户的提示来优化性能关键的内存排序和一致性任务。例如，put操作从用户那里获取两个提示：所提供数据的可见性范围，以及操作完成后预期的用户可见性。

使用工作流。应用程序通过一个三阶段工作流与GIN交互（如清单1所示）。在初始化阶段，应用程序在创建NCCL设备通信器时启用GIN支持（使用带有适当配置标志的ncclDevCommCreate），这将创建GIN上下文。然后，应用程序使用ncclCommWindowRegister集体地将内存缓冲区注册为窗口，该函数返回用于设备代码的窗口句柄。在核函数执行期间，设备线程实例化一个ncclGin对象，指定所需的上下文索引（通常根据目标对等方或负载均衡需求选择），并发出数据移动操作，可选择性地附加完成动作，如远程signal增量或本地counter更新。最后，核函数通过等待signals或counters进行同步，然后再重用缓冲区或进入后续计算阶段，从而确保通信和计算阶段的正确排序。清单1展示了NCCL GIN接口的简化视图，突出了基本操作。实际的API包括用于高级用例的额外模板参数和选项（例如，协作线程组、内联数据传输），但核心抽象保持一致。

class ncclGin {
  // 构造函数：使用设备通信器和上下文ID进行初始化
  ncclGin(ncclDevComm comm, int contextIndex);

  // 数据移动操作
  void put(team, peer, dstWindow, dstOffset, srcWindow, srcOffset, bytes, ...);
  void putValue(team, peer, dstWindow, dstOffset, value, ...);
  void signal(team, peer, signalId);

  // 本地完成跟踪
  void flush(coop); // 阻塞直到操作完成
  uint64_t readCounter(counterId); // 轮询计数器
  void waitCounter(coop, counterId, expectedValue);
  void resetCounter(counterId); // 重置以备重用

  // 远程完成跟踪
  uint64_t readSignal(signalId); // 轮询信号
  void waitSignal(coop, signalId, expectedValue);
  void resetSignal(signalId); // 重置以备重用
};

// 用于跨rank同步的网络屏障
class ncclGinBarrierSession {
  ncclGinBarrierSession(coop, gin, team, barrierHandle, index);
  void sync(coop); // 全局屏障同步
};

// 可选：将完成动作附加到数据操作上
// 远程signal
put(..., ncclGin_SignalInc{signalId});
// 本地counter
put(..., ncclGin_CounterInc{counterId});

使用示例。清单2演示了使用GIN原语的单向环形交换模式。在这个核函数中，每个rank向其在环形拓扑中的后继者（myRank + 1）发送数据，数据沿环单向流动，这实现了一种流水线通信算法中的常见模式。put操作（第13-16行）将数据从本地的sendWin传输到对等方的recvWin的计算偏移处，并在完成后原子地增加远程signal 0，从而提供数据到达的远程通知。然后，发送rank等待其自己的signal 0被其前驱者增加（第19行），以确保接收到的数据已经到达并且对本地GPU线程可见，然后再进行计算。最后，为后续的通信轮次重置signal（第21行）。这个模式展示了GIN的异步操作、灵活的完成语义和显式同步原语如何实现从设备代码进行高效的重叠点对点通信。

__global__ void ringExchange(
2   ncclDevComm devComm,
3   ncclWindow_t sendWin,
4   ncclWindow_t recvWin,
5   size_t dataSize, int myRank)
6 {
7   // 初始化上下文 0
8   ncclGin gin(devComm, 0);
9   int peer = (myRank + 1) % devComm.nRanks;
10
11  // 向对等方发送数据并发出完成信号
12  // 通过增加对等方的signal
13  gin.put(ncclTeamWorld(devComm), peer,
14          recvWin, myRank * dataSize,
15          sendWin, peer * dataSize, dataSize,
16          ncclGin_SignalInc{0} );
17
18  // 等待前驱节点
19  gin.waitSignal(ncclCoopCta(), 0, 1);
20  // 为下一轮重置
21  gin.resetSignal(0);
22 }

C. 后端（GDAKI和Proxy）实现

双后端抽象。设备API抽象了两种不同的后端实现，实现了前面描述的双插件语义。GDAKI后端通过DOCA GPUNetIO实现GDAKI语义，即通过直接的GPU到NIC通信，插件同时提供设备API和控制路径。Proxy后端通过CPU介导的传输实现Proxy语义，其中NCCL核心提供设备API，而插件仅提供基于CPU的数据路径操作。两个后端都暴露了相同的面向设备的接口，使得在运行时可以透明地选择后端，而无需更改应用程序代码。

GDAKI后端：直接GPU到NIC通信。GDAKI后端以最纯粹的形式实现了设备启动的网络通信，它利用DOCA GPUNetIO使GPU线程能够直接对网络接口卡进行编程，而无需CPU的介入。当一个核函数调用put时，GPU线程在设备内存中构建RDMA工作队列条目（WQEs），用源/目的地址和传输元数据填充它们，并直接写入NIC的门铃寄存器以触发DMA传输。NIC硬件自主管理操作进度：它轮询GPU内存以获取新的WQEs，通过InfiniBand或RoCE执行RDMA事务，并在GPU可见的内存中更新完成队列条目。这种直接的GPU-NIC路径消除了到CPU的PCIe往返，并为小消息实现了低延迟。然而，这种方法需要现代硬件和软件：支持GPU可访问控制结构的ConnectX-6 Dx或更新的NIC，以及所需的CUDA 12.x版本（如【20，DOCA GPUNetIO Programming Guide, NVIDIA Corporation】中所详述）。此外，正确的系统配置——包括GPUDirect RDMA内核模块（nv_peer_mem或dmabuf）和共置的GPU-NIC PCIe拓扑——对于正确操作和最佳性能至关重要。

Proxy后端：CPU辅助通信。Proxy后端通过将CPU引入通信路径作为GPU和NIC之间的中介，牺牲了峰值性能以换取硬件的可移植性。GPU线程将操作描述符——包含源/目的窗口句柄、可能的源内联值、偏移量、大小和完成动作的64字节数据——通过“即发即弃”（fire-and-forget）的存储操作，排入分配在CPU内存中的无锁队列。每个通信器有一个专用的CPU代理线程，该线程被固定在靠近本地rank的GPU和NIC的NUMA节点上，并持续轮询这些队列。一旦检测到新的描述符，代理线程就会提取字段，并通过网络插件的iput/iput_signal接口提交网络操作，该接口映射到标准的InfiniBand verbs或其他网络API。插件负责执行signal并确保所有先前put操作的可见性。完成通知则沿着相反的路径进行：代理线程使用网络插件的test接口轮询完成情况，将完成的操作与其关联的GIN counters匹配，并在GPU可见的内存中更新完成状态（根据GDRCopy的可用性，内存可能位于GPU或CPU上）。虽然CPU的介入引入了比GDAKI更多的延迟，但Proxy后端支持任意CUDA版本、任何支持GPUDirect RDMA的NIC（InfiniBand、RoCE、iWARP）以及Volta或更新的GPU。此外，CPU的介入通过主机端检测简化了调试，并在直接GPU-NIC通信不可用的系统上实现了平滑的性能降级。

后端选择和可移植性。表 I 总结了GDAKI和Proxy后端在架构上的差异。拥有现代NVIDIA端到端网络基础设施（ConnectX-6 Dx或更新的NIC、较新的CUDA版本、正确配置的GPUDirect RDMA）的高性能生产系统倾向于使用GDAKI以获得最小的延迟和零CPU开销。而开发环境、旧硬件部署、多供应商网络结构或GPUDirect支持配置不当的系统则依赖Proxy来实现功能正确性和操作灵活性。运行时在通信器初始化（ncclCommInitRank）期间自动检测可用的后端，通过能力查询探测DOCA GPUNetIO支持，并在必要时回退到Proxy。应用程序可以通过环境变量（NCCL_GIN_BACKEND）覆盖此选择，以进行调试或性能调优。这种设计确保了在不同部署场景下的可移植性，同时在硬件和软件基础设施允许的情况下保留了直接GPU-NIC通信的性能优势。

TABLE I NCCL GIN后端实现的架构比较
表 I NCCL GIN后端实现的架构比较

A7 补充细节

DEEPEP集成

本节介绍了将NCCL GIN集成到DeepEP中的过程，以验证其在需要计算-通信融合和低延迟的工作负载中的有效性。DeepEP是一个专门的MoE通信库，它使用NVSHMEM和IBGDA实现设备启动的稀疏all-to-all通信——即dispatch和combine原语。该库提供了两种风格的dispatch和combine原语：分别用于训练/推理预填充阶段的高吞吐量（HT）核函数和用于推理解码阶段的低延迟（LL）核函数。此次集成展示了如何使用GIN API实现DeepEP的设备启动通信模式，同时保持其性能特性，并与现有的NVSHMEM通信后端共存。

A. 集成需求

DeepEP的通信模式对后端提出多项要求。DeepEP的通信模式提出了几个要求：i) 高QP并行性——HT核函数需要24个QP，而LL核函数需要8-16个QP，以匹配本地专家的数量；ii) 异构拓扑支持——HT使用对称的rank-to-rank RDMA与NVLink转发，而LL使用完全的all-to-all RDMA网状结构；iii) 细粒度同步——需要对循环缓冲区的头/尾指针进行原子更新以实现流控制；iv) 后端共存——NVSHMEM IBGDA和GIN后端必须能够共存，以根据执行环境匹配用户的偏好。

B. 后端集成策略

采用最小抽象层和条件编译。集成采用了一个最小的抽象层来管理生命周期（初始化、内存分配、屏障），同时允许性能关键的操作通过条件编译直接在核函数中使用特定于后端的设备API。这种设计适应了两者之间根本的语义差异：IBGDA使用基于指针的寻址和内存原子操作，而NCCL GIN使用基于窗口的寻址和signal原子操作。

解决四大转换挑战。集成解决了四个关键的转换挑战。首先是多通信器映射：由于NCCL GIN每个通信器提供4个上下文，为了满足DeepEP的QP需求，需要⌈QPs/4⌉个通信器，并通过确定性选择（comm_id = id / 4, ctx_id = id % 4）来分配工作。其次是内存管理：后端将分配的缓冲区注册到所有通信器，并将设备可访问的窗口句柄存储在GPU内存中，使核函数能够将指针算术转换为（窗口，偏移）对。第三是同步：预先分配的结构化signal布局将基于内存的原子操作映射到signal原语（HT：每个通道两个signal，用于头/尾；LL：每个专家一个signal）。第四是语义保留：使用带有原子signal的零字节put操作来模拟释放-获取语义，确保在发出完成信号之前，所有先前的传输都可见。

C. 操作语义映射

编程模型的转换。从NVSHMEM迁移到NCCL GIN需要在这两种截然不同的编程模型之间进行转换，同时保留通信语义。NVSHMEM提供了一个分区全局地址空间（PGAS）的抽象，具有基于指针的寻址和基于内存的同步原语，而NCCL GIN遵循基于窗口的单向模型，使用基于signal的完成跟踪（如前文III-B节所述）。表II将DeepEP的通信模式映射到相应的NVSHMEM和GIN原语，突出了集成的关键语义转换。对于数据传输，核函数动态计算相对于窗口的偏移量，并根据通道或专家ID确定性地选择通信器，从而在QP之间实现负载均衡。对于同步，基于signal的设计将数据移动与完成通知解耦：批量传输使用不带即时信令的put()，随后是显式的signal()操作，这些操作仅在所有先前操作完成后才原子地更新远程计数器。这种模式通过网络原语而非内存排序来实现释放-获取语义，确保在signal到达时数据的可见性。

TABLE II DEEPEP通信核函数中使用的部分DEEPEP定制的NVSHMEM/IBGDA和NCCL GIN API。注意：此表并非NVSHMEM/IBGDA和NCCL GIN API的比较，而是侧重于DeepEP库中使用的相应NVSHMEM和NCCL API。
表II DeepEP通信核函数中使用的部分DEEPEP定制的NVSHMEM/IBGDA和NCCL GIN API

D. 高吞吐量核函数集成

HT核函数的优化策略。HT核函数通过分层通信来优化大批量（4096个token）处理。GPU通过对称的RDMA连接将数据发送到远程节点，然后远程节点通过NVLink将token转发到目标GPU。这在最大化节点内带宽的同时，最小化了节点间流量。RDMA缓冲区包含多个作为QP的通道，每个通道都有发送/接收缓冲区。头指针和尾指针跟踪缓冲区占用情况，提供循环缓冲区的流控制。

HT dispatch核函数的实现。dispatch核函数为SM分配了专门的角色。奇数编号的SM充当发送者（将token传输到远程rank）和NVLink接收者（最终目的地），而偶数编号的SM充当转发者（接收RDMA token并通过NVLink转发）。这种专业化实现了并发的双向通信。为了减少竞争，数据/尾指针更新和头指针更新使用不同的通道，将工作分散到不同的通信器上。

HT核函数中的GIN API使用。遵循操作映射（表II），每个SM角色都使用基于signal的原子操作进行指针管理，并使用基于窗口的put()进行数据传输。远程尾指针signal通过gin.signal(SignalAdd, 1)进行递增，而头指针流控制则依赖于通过gin.readSignal(signal_id)轮询本地头指针signal。数据传输使用单线程的NCCL GIN put()，后跟_syncwarp()以保留warp-collective语义。notify dispatch核函数使用一个协调者SM来刷新所有写入（gin.flush()），在对称RDMA rank之间执行屏障，重置头/尾signal，并在开始主dispatch之前交换元数据。

HT combine核函数的实现。combine核函数镜像了dispatch核函数的专业化分工，每个SM使用25个warp。偶数编号的SM作为NVLink发送者（将输入token分发到本地缓冲区）、RDMA接收者（将远程token与偏置项结合）和监控接收者进度的协调者。奇数编号的SM作为NVLink和RDMA转发者（合并本地token并将其转发到远程rank），并有相应的协调者监控转发者。操作映射与dispatch并行：转发者warp使用单线程put()和__syncwarp()进行数据传输，readSignal()进行头指针轮询，signal()进行尾指针更新。接收者warp使用readSignal()监控尾指针，而协调者warp通过signal()更新头指针。

E. 低延迟核函数集成

LL核函数的优化策略。LL核函数通过全连接的all-to-all RDMA网状连接来优化小批量（1-128个token）处理，实现了直接的GPU到GPU通信。token流式传输嵌入了路由元数据，无需单独的通知阶段，从而最小化了dispatch-combine周期时间。每个专家的signal分配为集群中任何专家对之间提供了直接协调。

LL核函数的SM分配。SM的分配使用 G = ⌈N/S⌉ 个warp组来分发专家，其中N是专家总数，S是可用的SM数量。每个SM通过 expert_idx = sm_id * G + warp_group_id 被分配专家。在每个warp组内，大多数warp处理FP8量化和token发送，而一个计数warp则管理所有分配专家的专家计数和元数据。这种组织方式实现了在数百个专家上的高效并行化（例如，在132个SM上，每个SM 3个warp组，共处理288个专家）。

LL核函数中的混合通信。LL核函数利用了混合的NVLink-RDMA通信。对于每个token传输，核函数通过自定义函数nccl_get_p2p_ptr检查NVLink的可用性。如果可用，则使用warp级别的内存操作直接复制token；否则，使用NCCL GIN的put()执行RDMA传输（表II）。token首先从PyTorch张量复制到RDMA发送缓冲区，在此阶段可选择应用FP8量化。在完成向一个目的地的token传输后，一个计数warp使用带有SignalAdd的零字节put()发送每个专家的token计数，确保所有先前的数据传输都已完成并变得可见，然后才传递计数——这实现了后端集成策略中描述的释放-获取语义。接收方使用gin.readSignal(signal_id)进行轮询，直到token到达。

LL combine核函数的实现。combine核函数将专家的输出路由回源rank并进行加权规约。它使用相同的混合NVLink-RDMA方法，并可选地使用LogFMT压缩来减少数据量。在传输专家输出后，通过带有SignalAdd的零字节put()来通知目的地，确保在接收方开始累加之前完成传输。接收方使用TMA加载warp将专家输出取到共享内存中，然后使用规约warp在FP32精度下应用top-k权重，最后转换为BF16输出。

A4 实验环境与结果

实验环境

本节通过两种互补的方法来评估NCCL GIN和NVSHMEM。首先进行点对点微基准测试（V-A节）以分离协议级别的性能特征，然后评估与生产级MoE通信库DeepEP 1.2.1版的集成。DeepEP评估涵盖了用于训练和推理预填充的高吞吐量（HT）核函数（V-B节），以及用于推理解码的低延迟（LL）核函数（V-C节），测试了混合RDMA+NVLink和纯RDMA两种配置。

所有实验均在NVIDIA的EOS集群上进行，该集群配备H100 GPU（如表III所示），使用NVSHMEM 3.4.5版和NCCL 2.28版。DeepEP基准测试为每个GPU分配24个SM，通信根据自动选择的通道配置分布在NVLink和RDMA上。

TABLE III EOS DGXH100计算节点的硬件规格
表 III EOS DGXH100计算节点的硬件规格

实验结果

A. 点对点微基准测试

实验内容：为了建立基准性能，我们使用ping-pong测试测量了在两个H100 GPU之间，从4字节到4MB不同消息大小下，put with signal操作的延迟。
实验结果：图4展示了NCCL GIN的双后端（GDAKI和Proxy）与NVSHMEM的IBGDA和IBRC传输的性能。
分析结论：
- 对于小消息（4–128字节），NCCL GIN GDAKI实现了16.7 µs的往返延迟，与NVSHMEM IBRC的16.0 µs相当，而NVSHMEM IBGDA为24.3 µs。GDAKI后端的直接GPU-NIC路径消除了CPU代理开销。
- Proxy后端尽管需要经过GPU-CPU队列，但仍达到了18.0 µs的延迟。
- 在较大消息尺寸下，带宽成为瓶颈，所有实现的性能趋于一致。这验证了NCCL GIN用于应用集成的基础性能特征。

图4. 点对点延迟：NVSHMEM IBGDA/IBRC 和 NCCL GIN GDAKI/Proxy 后端。

B. 高吞吐量（HT）核函数

实验内容：HT核函数针对MoE训练和推理预填充中的大批量token（4096个）进行优化，采用分层通信，其中专门的SM角色（发送者、转发者、NVLink接收者）最小化节点间RDMA流量，同时最大化节点内NVLink带宽。
实验结果：图5展示了在2、4和8个节点上，使用FP8和BF16精度时的dispatch和combine带宽，并分别报告了RDMA和NVLink的指标。
分析结论：
- 在所有配置下，两种实现（NCCL GIN和NVSHMEM）都提供了相当的性能。
- 在2节点（16个GPU）、BF16精度下，dispatch操作的RDMA带宽分别为NCCL GIN 84.36 GB/s和NVSHMEM 84.97 GB/s。
- 在8节点（64个GPU）下，两种实现为dispatch操作维持了大约53–54 GB/s的RDMA带宽。
- 在不同规模、精度模式和操作类型下，结果差异在1-2%以内，表明NCCL GIN在保持HT吞吐量的同时，实现了在NCCL基础设施上的标准化。

C. 低延迟（LL）核函数

实验内容：LL核函数针对MoE推理解码中的小批量token（1–128个）进行优化，采用全连接的all-to-all RDMA网状连接，并带有每个专家的signals和混合NVLink-RDMA路径。在BF16精度和7168隐藏维度下，dispatch操作传输14,352字节的消息，combine操作传输14,336字节。实验在启用NVLink（混合RDMA+NVLink）和禁用NVLink（纯RDMA）两种配置下进行评估。

LL核函数 - 启用NVLink (RDMA+NVLink)：
- 实验结果：图6和图7显示了带宽和延迟的比较。
- 分析结论：
- 在1节点（8个GPU）上，NCCL GIN略优于NVSHMEM：dispatch带宽为185.28 GB/s，延迟为40.62 µs，而NVSHMEM为182.15 GB/s和41.43 µs。combine操作几乎相同。
- 在多节点规模下，两者性能相当，但NCCL GIN在不同规模下始终提供更低的延迟（例如，2节点时低9%：142.51 µs vs 157.00 µs）。combine操作的性能差异在1-3%以内。

图6. 启用NVLink时，NCCL GIN和NVSHMEM的LL核函数带宽。
图7. 启用NVLink时，NCCL GIN和NVSHMEM的LL核函数延迟。

LL核函数 - 禁用NVLink (纯RDMA)：

实验结果：图8和图9展示了测量结果。
分析结论：
- 当所有通信都通过RDMA进行时，两种实现都保持了相当的性能，大部分指标差异在1-2%以内。
- 在1节点（8个GPU）上，NCCL GIN的dispatch带宽为47.00 GB/s，延迟为160.82 µs；NVSHMEM为46.79 GB/s和160.67 µs。
- 在8节点（64个GPU）上，两者都维持了大约34–35 GB/s的带宽和219–225 µs的延迟。

图8. 禁用NVLink（纯RDMA）时，NCCL GIN和NVSHMEM的LL核函数带宽。
图9. 禁用NVLink（纯RDMA）时，NCCL GIN和NVSHMEM的LL核函数延迟。

D. 讨论

评估结果表明，NCCL GIN提供的设备启动通信能力具有与NVSHMEM相似的性能特征。在微基准测试和应用工作负载（HT和LL核函数）中，GIN将设备启动的原语与NCCL的拓扑感知集合操作集成在单个运行时中，结合了设备端API的灵活性和NCCL的生产级基础设施。GIN的实现仍在积极开发中，计划的优化包括批处理工作队列条目和摊销多个操作的doorbell成本，以进一步提高性能。

A5 结论

现代AI工作负载，包括MoE推理和编译器生成的融合核函数，要求GPU直接控制网络操作，这些能力超出了NCCL传统的主机启动模型。本文介绍了GIN，作为NCCL 2.28 Device API【9，Fusing Communication and Compute with New Device API and Copy Engine Collectives in NVIDIA NCCL 2.28, 2025, NVIDIA Developer Blog】的一部分，它使GPU线程能够直接从CUDA核函数中发出单向RDMA操作。GIN提供了一个统一的三层架构（主机API、设备API和具有双重语义的可插拔网络后端），并支持通过GDAKI进行直接GPU到NIC通信以及在标准RDMA硬件上进行CPU辅助操作。我们的评估验证了GIN的实际可行性：GDAKI后端为小消息实现了16.7 µs的往返延迟，并且与DeepEP的集成表明，在DeepEP的高吞吐量和低延迟核函数中，只需少量代码更改即可获得具有竞争力的性能。

重要的是，GIN的价值不仅在于原始性能，更在于生态系统的统一，通过与NCCL的生产级基础设施集成，提供了可扩展性和可延展性。应用程序可以访问一套统一的通信抽象——用于NVLink/PCIe的Load/Store Accessible (LSA)、用于NVLink SHARP的Multimem和用于网络RDMA的GIN——从而为每种模式选择合适的原语。关键的是，这种集成保留了NCCL的生产就绪特性：用于多维并行（专家并行、张量并行、流水线并行）的分层通信器、用于弹性大规模训练的容错和弹性机制，以及拓扑感知的优化。这些能力消除了部署多个通信运行时的操作复杂性，同时为新兴工作负载（如MoE推理和编译器生成的融合核函数）提供了所需的灵活性。

未来的工作将集中于在生产应用中更广泛地采用GIN，例如PyTorch分布式训练、TensorRT-LLM推理服务、用于LLM推理的vLLM和SGLang，以及JAX/Triton编译器生成的核函数。我们还计划通过增加额外的单向原语来扩展GIN的API，以支持新兴的通信模式和分布式算法需求。

PaperCache

GPU-Initiated Networking for NCCL

GPU-Initiated Networking for NCCL

A1 主要贡献

A3 背景知识

A2 方法细节

A. 核心原则和架构

B. 设备端API和编程模型

C. 后端（GDAKI和Proxy）实现

A7 补充细节

DEEPEP集成

A. 集成需求

B. 后端集成策略

C. 操作语义映射

D. 高吞吐量核函数集成

E. 低延迟核函数集成

A4 实验环境与结果

实验环境

实验结果

A. 点对点微基准测试

B. 高吞吐量（HT）核函数

C. 低延迟（LL）核函数

D. 讨论

相关工作

A5 结论

💬 评论讨论

GPU-Initiated Networking for NCCL

A1 主要贡献

A3 背景知识

A2 方法细节

A. 核心原则和架构

B. 设备端API和编程模型

C. 后端（GDAKI和Proxy）实现

A7 补充细节

DEEPEP集成

A. 集成需求

B. 后端集成策略

C. 操作语义映射

D. 高吞吐量核函数集成

E. 低延迟核函数集成

A4 实验环境与结果

实验环境

实验结果

A. 点对点微基准测试

B. 高吞吐量（HT）核函数

C. 低延迟（LL）核函数

D. 讨论

相关工作

A5 结论

💬 评论讨论

登录

注册

忘记密码

重发验证邮件