Zhiyi Hu1∗, Siyuan Shen1∗, Tommaso Bonato1, Sylvain Jeaugey2, Cedell Alexander3, Eric Spada3, James Dinan2, Jeff Hammond2, Torsten Hoefler1
1ETH Zurich, Switzerland; 2NVIDIA Corporation; 3Broadcom Inc.
* These authors contributed equally to this work

主要贡献

本文针对NVIDIA Collective Communication Library (NCCL)的内部设计进行全面分析，核心问题在于NCCL尽管开源并有文档化API，但其内部机制如通信通道编排、协议选择以及跨设备和节点的内存移动处理缺乏透明度，导致性能分析和瓶颈识别困难。研究目标是通过系统探索NCCL的四个主要方面，包括API结构和通信通道管理、通信协议（Simple、LL和LL128）的详细检查、数据传输模型分析以及集体通信算法的全面剖析，提供对NCCL架构的深入理解。创新点在于基于这些洞见开发ATLAHS，一个应用跟踪驱动的网络模拟工具链，能够准确再现大规模AI训练负载中的NCCL通信模式，从而为系统研究者和性能工程师优化或模拟大规模集体通信提供指导。本分析基于NCCL版本2.19.1，其核心机制预计在未来版本中保持一致。

背景知识与设计原则

NCCL API概述。 NCCL专为GPU集群提供高度优化的集体通信操作，强调低延迟和高带宽。其核心通过清晰高效的API管理GPU到GPU通信，抽象复杂技术细节。NCCL主要提供四类函数：通信器管理、集体通信、点对点通信和组调用。通信器管理类似于MPI，所有操作在通信器上下文中进行，每个参与GPU维护一个通信器对象。用户首先初始化通信器并定义参与GPU集合。在单一进程或线程管理所有设备时，使用ncclCommInitAll集体创建通信器；在多进程或多线程环境中，每个进程使用共享唯一标识符调用ncclCommInitRank建立通信器。通信任务完成后，通过ncclCommDestroy安全销毁通信器，确保所有挂起操作完成前清理，或使用ncclCommAbort立即终止并取消操作，用于错误恢复避免死锁。集体通信提供5个操作：ncclAllReduce、ncclBroadcast、ncclReduce、ncclAllGather和ncclReduceScatter。历史上，ncclBcast作为ncclBroadcast的就地变体模拟MPI_Bcast行为，但现在ncclBroadcast支持独立发送和接收缓冲区，ncclBcast主要为兼容性保留。点对点通信通过ncclSend和ncclRecv支持。组调用使用ncclGroupStart和ncclGroupEnd聚合操作，延迟执行直到组结束，减少启动开销和延迟，支持模拟SendRecv、All-to-One、One-to-All或All-to-All模式。

启动策略。 NCCL支持三种常见执行模型，每种有不同权衡。每GPU一个CPU进程模型提供更大控制，通过绑定每个GPU到单独进程，将关联CPU代码调度到本地NUMA域，提高数据局部性和减少内存访问延迟。每GPU一个CPU线程模型在单一CPU进程通过多线程管理多个GPU，实现高效进程内内存共享，直接访问跨rank的内存包括GPU缓冲区，减少通信期间内存复制开销。多个GPU一个CPU线程模型虽因顺序内核启动和降低并发性而性能受限，但提供简单性、最小CPU开销和确定性执行，适合小规模部署或原型环境优先易实现而非最高性能。

通信通道。 NCCL通过GPU、CPU和网络接口三个硬件组件编排通信。GPU执行归约并移动缓冲区数据，CPU启动内核并管理主机侧协调，NIC跨节点传输数据包。当单一流多处理器（SM）处理GPU工作时，大消息可能过载该SM，未充分利用其他SM，并无法饱和NVLink或InfiniBand等链路【7，NCCL GitHub Issue #578: Ring AllReduce performance discrepancy, 未知年份, GitHub, https://github.com/NVIDIA/nccl/issues/578】【8，NCCL GitHub Issue #1302: Broadcast logic and topology insights, 未知年份, GitHub, https://github.com/NVIDIA/nccl/issues/1302】。为避免瓶颈，NCCL将每个集体操作细分为通信通道，每个通道作为单独CUDA块在自身SM上运行，库分区输入缓冲区使通道并行操作不相交块。这种细粒度并行提高聚合吞吐，尤其对大负载否则会序列化于一个SM。跨通道分布工作还帮助在NVLink平台平衡流量跨多个NIC，每个通道可独立通过不同NIC离开节点，提高链路利用率、减少空闲时间并平衡NVLink、PCIe和InfiniBand等互连负载。然而，激进使用多通道可能负面影响网络效率。当每通道块大小小于NIC传输使用的512 KiB FIFO缓冲区大小时，代理线程发送部分填充缓冲区。这种未充分利用可降低PCIe和网络吞吐，尤其当多个队列对（QP）激活以启用等价多路径路由（ECMP）负载平衡时。NCCL通过启发式减少小消息的nChannels解决此问题（参考enqueue.cc中的calcP2pChunkSize函数）。尽管如此，选择最优通道数仍是GPU侧并行与网络利用效率的权衡。通道管理在通信器级别协调，每个GPU接收0到n-1的唯一rank，其中n是参与GPU总数。在通信器初始化期间，NCCL建立初始通道结构集，其总数主要由系统拓扑和架构默认引导。当调用集体操作时，NCCL动态选择算法和协议。根据运行时选择，NCCL内部调优模型考虑所选策略、当前消息大小、可用带宽和每通道线程配置，确定使用多少预建立通道。虽然早期版本允许用户通过环境变量如NCCL_NTHREADS影响通道行为，但现在不鼓励手动调优，最近版本通常忽略这些设置并可能导致错误行为。分配给每个通道的逻辑通信拓扑直接塑造操作期间GPU间数据流动。在环拓扑中，每个GPU识别其直接前驱和后继形成单向通信环。在树拓扑中，每个GPU跟踪其父和子rank，建立逻辑通信树。为增加带宽利用，NCCL采用双二叉树结构【9，Full bandwidth broadcast, reduction and scan with only two trees, 2007, European Conference on Recent Advances in Parallel Virtual Machine and Message Passing Interface, 无URL】【10，Energy, Memory, and Runtime Tradeoffs for Implementing Collective Communication Operations, 2014, Journal of Supercomputing Frontiers and Innovations, 无URL】：无节点在两树中均为非叶，至多一节点在两树中为叶。第二树在节点数偶数时镜像第一树，奇数时一位置移位。这些拓扑在通信器初始化时建立，并在所有集体操作中重用。对于使用ncclGroupStart和ncclGroupEnd的分组点对点操作，NCCL尽可能将每个传输分配到单独通道，实现跨传输的任务级并行。

通信协议概述。 NCCL采用多种通信协议优化集体操作期间数据传输效率。三种协议Simple、LL（Low Latency）和LL128设计实现带宽与延迟的不同权衡。本节概述每种协议机制。表I总结三种协议关键特性。

Simple协议。 Simple协议设计最大化带宽利用，用于大消息传输。它通过将数据分为相对大块并跨通信通道分派，确保充分利用网络接口和GPU内存系统的高吞吐。为保持内存一致性，协议使用内存栅栏强制正确顺序和数据可见性。接收者必须等待完整块传输后访问。虽然有效确保正确性，但内存栅栏引入显著开销。此开销对小消息成为限制因素，其中同步成本主导整体传输时间。因此，虽然Simple协议对大消息实现接近峰值带宽，但处理小负载时遭受高延迟。

LL (Low Latency)协议。 为解决Simple协议的延迟问题，NCCL包含LL协议，优化小消息大小，其中带宽通常未充分利用。LL协议不依赖内存栅栏，而是使用轻量标志基同步。小标志与数据一起传输信号其有效性，使接收者一数据可用即继续，无需昂贵内存屏障。每传输由4字节数据后跟4字节标志组成，使用8字节原子操作一起发送。此方法显著减少同步开销，提高延迟敏感负载响应性。LL强制中间缓冲区驻留主机内存，以便CPU轮询标志并检测数据何时准备通过NIC发送。这必要因为通过PCIe轮询GPU内存远慢于DRAM访问，并需显式同步确保主机数据可见性。虽然此设计实现低延迟，但阻止使用GPU Direct远程直接内存访问（RDMA），严重限制带宽。因此，LL通常仅实现峰值带宽的25–50%，取决于互连。因此，仅在延迟关键且带宽次要的小传输中首选。

LL128协议。 LL128协议改进LL，保持其低延迟属性，同时显著增加带宽效率，尤其在NVLink等高性能互连上。像LL一样，使用标志基同步消除内存栅栏，但以128字节单位传输数据。其中120字节专用数据，8字节保留标志，允许协议利用约95%峰值带宽。在网络路径上，LL128类似于Simple协议，发送GPU聚合相对大数据块前通知CPU准备发送。虽然此块基聚合限制跨节点流水线，但LL128仍受益于节点内细粒度流水线，由于其较小传输粒度。此低延迟和高吞吐组合使LL128适合广泛消息大小。然而，LL128有更严格硬件要求。它依赖原子128字节写入，必须不被内存系统或互连拆分或重排序。在此类操作不保证的系统中，由于PCIe限制或其他架构约束，NCCL禁用LL128避免数据损坏。协议选择因此不仅受消息大小影响，还受系统级能力影响。
表II NCCL通信特性和传输

协议选择与比较。 NCCL在运行时基于用户设置（即NCCL_PROTO）、集体算法和内部性能启发式动态选择Simple、LL和LL128协议。若未显式指定，NCCL使用调优模型考虑系统拓扑、GPU架构、消息大小和预定义性能指标选择最佳算法-协议对。此选择受资源可用性约束，如协议特定缓冲区内存。通常，小消息选择LL/LL128减少延迟，大消息使用Simple最大化吞吐。

方法细节

数据传输方法和传输层概述。 高效数据移动是NCCL通信性能核心，尤其在多GPU和多节点环境中。如表II总结，NCCL根据通信发生在单一节点内（intra-node）或跨多个节点（inter-node）采用不同数据传输策略和传输机制，每种传输针对特定硬件和互连类型优化，支持可扩展集体。

图1. NCCL中节点内数据传输路径示意图。每路径颜色编码指示所选传输和硬件支持。

节点内数据传输。 NCCL采用复杂分层方法处理节点内通信，优先同一物理机器上GPU间最低延迟和最高带宽路径（见图1）。此策略 heavily 利用NVIDIA的GPUDirect Peer-to-Peer (P2P)技术，使GPU直接访问彼此内存而不通过CPU系统内存暂存。核心是P2P传输，主要在src/transport/p2p.cc管理。当GPU通过NVIDIA NVLink互连时，NCCL优先此路径，实现GPUDirect P2P over NVLink利用这些专用高速直接GPU到GPU链路。若NVLink不可用，NCCL可利用PCIe总线上的GPUDirect P2P通信，也由P2P传输层管理。此提供通常远优于使用cudaMemcpy主机内存暂存的回退。NCCL P2P传输的关键优化是P2P_DIRECT模式，当通信rank属于同一进程时启用。虽然单进程和多进程通信均利用无CPU参与的GPU到GPU传输，但P2P_DIRECT模式通过两种方式显著提高效率。首先，它绕过IPC句柄需求，使用同一地址空间内的直接GPU内存指针。更重要的是，它通过使用directSend和directRecv原语消除中间数据复制，这些原语直接在源和目标缓冲区间传输数据，而非通过中间FIFO缓冲区路由。尽管此优化数据路径，NCCL仍使用共享结构（如ncclSendMem和ncclRecvMem）中的原子头尾计数器保持正确同步，确保适当顺序并防止数据竞争。因此，P2P_DIRECT通过简化内存寻址和更直接数据传输路径提供实质性能益处，建立在基础GPUDirect P2P能力上。NCCL可能利用Shared Memory (SHM)传输，不仅当直接GPU到GPU P2P不可用时，还当P2P次优时。特别地，跨插槽PCIe上的P2P常生成CPU处理差的P2P数据包，导致性能下降。SHM通过路由流量经系统内存避免此，使用PCIe到内存和内存到PCIe传输，CPU通常更好优化处理。在SHM模式中，一个GPU的控制进程写入共享内存段，然后由另一GPU进程读取。注意，在一些多插槽系统中，当每个GPU驻留单独CPU插槽并有支持GPUDirect RDMA的本地NIC时，NCCL可能使用NIC进行节点内GPU间通信。而非遍历CPU互连，NCCL可能路由数据通过GPU–NIC–NIC–GPU路径，利用PCIe带宽避免CPU瓶颈。此行为由NCCL的拓扑感知逻辑确定，并可使用环境变量如NCCL_CROSS_NIC控制。

图2. NCCL中节点内数据传输路径示意图。每路径颜色编码指示所选传输和硬件支持。

节点间数据传输。 NCCL中的节点间通信编排位于不同物理节点GPU间数据交换。此过程涉及GPU执行NCCL内核、CPU上运行代理线程管理网络操作以及底层网络结构。如图2所示，NCCL基于可用硬件在标准TCP Socket传输或高性能InfiniBand (IB) Verbs传输间选择。

基于Socket的通信。 当网络接口不支持RDMA时，NCCL采用socket传输，在transport/net_socket.cc实现。在此模式中，中间缓冲区分配为CUDA pinned主机内存。在发送侧，数据从GPU复制到此缓冲区前使用标准socket调用发送过网络。在接收侧，数据接收到主机缓冲区然后复制到GPU。此依赖主机内存作为暂存区招致跨PCIe总线的额外内存复制开销。发送和接收均遵循rendezvous协议，其中发送者和接收者在实际数据传输前协调缓冲区就绪。

IB Verbs传输。 对于InfiniBand或RoCE等高性能网络，NCCL使用IB传输，在net_ib.cc实现。IB传输利用RDMA能力实现最小CPU干预的节点间直接数据移动。与socket传输一样，所有传输经中间缓冲区暂存，但缓冲区位置取决于硬件支持和配置。默认若NIC无法直接访问GPU内存，中间缓冲区分配在主机内存。GPU内核复制数据到此缓冲区，代理线程发布RDMA操作使用RDMA write从主机内存移动数据到远程节点【12，NCCL GitHub Issue # 609: why uses rdma write for default ib traffic, 2021, GitHub, https://github.com/NVIDIA/nccl/issues/609】。在接收侧，过程反转：NIC写入传入数据到主机缓冲区，代理线程协调从主机到设备内存的复制。代理线程角色是管理这些DMA和RDMA操作。与socket传输一样，使用rendezvous协议同步发送者和接收者前数据传输。在以下段落中，我们突出IB传输中的一些关键特性和优化。

GPUDirect RDMA优化。 IB传输的关键优化是GPUDirect RDMA (GDRDMA)，使NIC直接访问GPU内存，消除主机内存暂存需求。GDRDMA仅当NIC和GPU连接到同一PCIe开关时使用。在此情况下，中间缓冲区分配在GPU内存。CPU代理线程使用机制如nv_peer_mem【13，Developing a Linux Kernel Module using GPUDirect RDMA, 2025, 无会议, 无URL】或Linux DMA-BUF子系统【14，Buffer Sharing and Synchronization (dma-buf), 2025, 无会议, kernel.org】注册此GPU内存到RDMA-capable NIC，允许NIC映射并直接访问GPU内存。然后NIC的DMA引擎直接执行RDMA读或写到/从GPU，绕过CPU和主机内存完全。

每对等体多通道连接。 作为提高带宽利用和减少拥塞的优化，IB传输默认为每个远程GPU和每个NIC实例化2个逻辑通道（由NCHANNELS_PER_NET_PEER参数控制）。每个逻辑通道维护自己的ncclIbSendComm结构，嵌入独立InfiniBand QP束。在执行期间，主机侧网络代理在发出ncclNet->isend()调用时交替使用两个sendComm句柄，从而跨QP集拆分流量。此轮询策略增加有效每QP块大小，为ECMP感知结构引入路径多样性，并提升整体互连效率——所有无需额外协调开销。

QP布局。 对于每对rank，RDMA插件建立两个可靠连接（RC）QP，每个方向一个。正向QP负责批量数据流：代理发出一个或多个RDMA_WRITE工作请求直接推送用户数据到对等缓冲区，后跟RDMA_WRITE_WITH_IMM用于完成通知。此最终请求行为取决于通信模式：小消息将数据和大小信息组合单一操作最小化开销；大消息使用自适应路由发送零字节写入以传输大小作为立即数据保证顺序，同时允许批量数据使用自适应路由；聚合多消息传输写入精确大小信息到远程数组，同时使用立即数据完成信号。反向QP仅携带微小clear-to-send (CTS)消息，即单一RDMA_WRITE广告远程缓冲区地址、rkeys和标签信息。虽然理论上同一QP可复用相同功能，但将CTS分离到自身通道隔离延迟关键控制流量与带宽饥饿数据流，允许网络以最小头阻塞交付。

使用循环回RDMA_READ的本地刷新。 当启用GPUDirect RDMA时，发送者必须确保所有未决PCIe写入在内核消耗数据前到达GPU内存。NCCL通过在最后接收完成后发出虚拟RDMA_READ实现此。专用“flush” QP连接到自身，即其ready-to-receive (RTR)阶段使用自身本地QP号作为目的地。因此，读从未离开主机，但verbs层仍等待先前写入的PCIe完成，提供廉价顺序屏障。
表III NCCL集体操作支持的算法和协议

图例：✓=支持，✗=不支持。

NCCL集体算法概述。 集体算法是NCCL核心，使GPU间高效同步通信。它们管理数据移动和依赖、优化通信路径，并随GPU数增加扩展。NCCL通过将每个集体操作分解为低级通信原语并跨多个并行通道分布实现这些算法。算法选择，通常环或树，取决于特定集体操作和相关执行参数如消息大小和拓扑【15，Optimization of collective reduction operations, 2004, ICCS, 无URL】【16，Bandwidth optimal all-reduce algorithms for clusters of workstations, 2009, Journal of Parallel and Distributed Computing, 无URL】【17，Optimization of collective communication operations in mpich, 2005, The International Journal of High Performance Computing Applications, 无URL】。本节概述NCCL集体算法的设计和主要特征。

算法和协议支持概述。 虽然NCCL提供六种协议，但并非所有适用于每个算法，其可用性可能基于硬件特征和运行时约束而变。表III总结NCCL版本2.19中5个集体操作支持的算法和通信协议。此信息从src/device目录相应头文件提取。除了常用Ring和Tree算法，表还突出对专用算法CollNet和NVLS的支持。NVLS和CollNet主要设计优化AllReduce性能，NVLS还通过利用特定硬件能力支持ReduceScatter和AllGather。CollNet算法针对网络基础设施可参与集体操作场景，如使用NVIDIA SHARP (Scalable Hierarchical Aggregation and Reduction Protocol)技术，允许归约或其他部分集体计算卸载到网络交换机，从而减少数据移动和延迟【18，NCCL Issue #320: NVLS and CollNet Support, 2021, GitHub, https://github.com/NVIDIA/nccl/issues/320】。CollNet Direct启用节点内全对全通信。相反，CollNet Chain线性排列GPU，并执行链上归约和链下广播【19，Nccl github issue #919: Question: Nccl tree algorithm behaviour, 2023, GitHub, https://github.com/NVIDIA/nccl/issues/919】。NVLS算法设计利用NVIDIA的NVLink Switch (NVSwitch)系统，提供多GPU服务器或NVSwitch结构内高带宽直接GPU到GPU通信路径，实现更高效集体操作【2，NVIDIA Collective Communications Library (NCCL) Documentation, 2025, 无会议, 无URL】。纯NVLS和NVLS Tree算法均使用NVLink SHARP进行节点内归约，但节点间处理不同：NVLS通过CollNet和SHARP启用交换机继续归约，而NVLS Tree使用树基扇出【19，Nccl github issue #919: Question: Nccl tree algorithm behaviour, 2023, GitHub, https://github.com/NVIDIA/nccl/issues/919】。然而，本文不包括NVLS和CollNet分析，因为其实现高度依赖特定硬件（NVSwitch和SHARP启用网络），使其代表性较低。我们承认NCCL持续演进，最近在版本2.23引入额外算法如Parallel Aggregated Trees (PAT)【20，New scaling algorithm and initialization with nvidia collective communications library 2.23, 2025, 无会议, 无URL】【21，Pat: a new algorithm for all-gather and reduce-scatter operations at scale, 2025, 无会议, 无URL】。尽管如此，由于新算法尚未广泛采用，我们后续讨论仍集中在Ring和Tree算法。

通信原语。 NCCL通过从低级通信原语集组合实现高级集体操作。这些原语形成NCCL集体算法基础，封装基本操作如发送、接收、归约和跨GPU复制数据。常见原语包括send、recv、recvReduceSend、recvCopySend和recvReduceCopySend，以及其“direct”变体如IV-A节讨论。每原语代表独特数据移动或计算模式，命名约定清晰指示操作序列。例如，recvReduceSend表示GPU从对等接收数据、与本地缓冲区执行归约并发送结果到下一GPU的步骤。在执行期间，NCCL运行时跨循环步骤迭代分派这些原语，实现不同算法、拓扑和传输层的灵活协调。每NCCL原语行为进一步由所选通信协议塑造。同步、缓冲区管理和传输粒度取决于使用Simple、LL或LL128协议而变。注意，这些低级原语高度优化于具有固定小源和目的地数的集体，如环和树，通常涉及一源和一目的地（或树拓扑某些至多三）。虽然此方法为许多标准集体算法实现高效率，但对如全对全需要处理N源和N目的地的模式较不有效。
表IV 默认配置下每协议的NCCL通道缓冲区大小

NCCL集体的迭代执行。 NCCL通过首先将用户输入数据分配到可用通信通道处理集体操作，实现通道级并行。每通道负责输入的连续段，由总元素数（count）和通道数确定。此分区如图3可视化，总数据拆分使每个通道如Channel 0和Channel 1独立操作其分配区域。每通道工作起始索引由workOffset给出，大小由channelCount给出。为促进高效数据传输和计算，NCCL为每个通道分配固定大小缓冲区，其容量取决于所选通信协议（Simple、LL或LL128，如表IV所示）。若通道数据区域大于其总缓冲区，NCCL将数据分解为几个外循环迭代。每迭代处理至缓冲区大小的段（每迭代loopCount元素），通道循环所需循环覆盖所有分配元素。在每个外循环迭代内，NCCL通过将通道缓冲区分为固定数段（称为槽）实现流水线。通常为8，由NCCL_STEPS参数设置。每槽可独立推进通信和计算不同阶段，允许流水线重叠数据传输与归约或复制。在每个基本步骤中，跟随V-B节描述的通信原语，处理数据块（chunkCount元素，或循环中最终的lastChunkCount），并映射到缓冲区槽。此分块机制允许NCCL保持通信通道繁忙，重叠新块与进行中操作最大化吞吐。

图3. NCCL跨通信通道和循环迭代的数据分区策略可视化。

在NCCL中，数据移动基本单位称为元素，其含义取决于集体操作。对于ncclAllGather和ncclBroadcast，每个元素是一个字节，因为这些操作聚焦高效移动和连接数据。此字节级粒度赋予NCCL打包和传输数据的灵活性，独立于底层类型。对于ncclAllReduce、ncclReduceScatter和ncclReduce，每个元素对应用户定义数据类型（如float或int），因为这些操作要求仅在数据类型级有意义的算术归约。图3展示此过程在行动中。图中每个单元代表sendBuff中一个数据元素。为说明目的，此示例假设channelCount等于2，chunkCount等于2，loopCount等于4。Channel 0从其workOffset开始，并在loopCount的循环迭代中处理元素，进一步分解为chunkCount的块。Channel 1对其自身区域遵循相同逻辑。通过协调此分区和流水线，NCCL实现跨所有参与GPU的高效、并行和可扩展集体操作。

将通信流水线映射到CUDA层次。 为完全理解NCCL性能特性，考察这些通信通道和数据流水线如何映射到GPU并行执行模型至关重要。NCCL通过仔细结构化层次编排跨多个SM的数十万线程，与CUDA线程组织对齐。NCCL内核以网格维度(nChannels, 1, 1)启动，其中nChannels表示操作的活跃通信通道数。此一对一映射确保每个CUDA块(blockIdx.x)对应 точно 一个通信通道。在每个块内，NCCL使用从NCCL_MIN_NTHREADS到NCCL_MAX_NTHREADS的可变线程数。确切线程数由NCCL自动调优系统确定，并在内核启动准备期间存储在plan->threadPerBlock中。块Idx.x和通道ID间映射通过位掩码方法管理。每内核接收指定当前操作活跃通道的channelMask。设备代码使用位填充计数计算块索引和通道ID对应：对于给定blockIdx.x，通道ID由找到channelMask中第blockIdx.x个设置位确定。在每个块内，NCCL在warp级组织线程到专用角色。前两个warp处理初始化任务：warp 0加载通信器元数据(ncclDevComm)到共享内存，而warp 1加载通道特定数据(ncclDevChannel)。剩余warp执行实际通信和计算工作。对于集体操作，工作warp数由工作描述符中的nWarps字段控制。不同算法不同分配这些warp。例如，在NVLS AllReduce中，warp细分为不同阶段组。对于点对点操作，warp在发送和接收操作间分区，分布基于并发传输数动态计算。在每个通道缓冲区内，NCCL_STEPS槽启用线程级细粒度流水线。warp内线程协作以循环方式移动数据通过这些槽。每槽包含ncclConnFifo结构，具有模式、偏移、大小和数据指针字段，允许数据同时存在不同流水线状态：计算中、排队传输、网络中飞行或准备消耗。在最细粒度，NCCL在每个warp内单个线程间分布工作。对于批量数据移动，线程在展开方式每迭代处理多个数据元素。确切量取决于协议。线程分歧通过warp统一操作仔细管理。例如，warp中所有线程执行相同操作序列（send、reduce、copy）但在不同数据元素或内存地址上。此确保GPU SIMT架构的最大利用。关键地，NCCL从不执行“一个通信任务”如高层描述建议。相反，它同时维护数十并发数据移动流水线。多个通道跨不同SM并行执行，每个通道内多个槽流水线数据流动不同阶段，每个通道内多个warp处理通信不同阶段。此多级并行对实现高带宽利用至关重要。
表V NCCL环AllReduce一个循环迭代中的步骤


图4. NCCL中环拓扑连接4 GPU的环AllReduce算法示意图，突出单一循环迭代内GPU通信原语序列。

定性算法分析。 既然我们确立所有常见NCCL集体算法遵循迭代处理模型，一个重要区别在于GPU是否能流水线连续循环迭代。基于此特性，算法可分组为流水线和非流水线两类。在以下节中，我们据此组织集体算法并提供每个的定性分析。对于每个算法，我们描述每个循环迭代内执行的基本步骤特定序列。虽然我们最初考虑量化复杂度分析以参数如数据大小和alpha-beta模型界定算法运行时，但发现此方法不切实际，因为影响性能的因素众多。变量如GPU跨节点分布有重大影响。例如，单一节点上4 GPU体验与置于单独节点的4 GPU非常不同的带宽和延迟。包含所有这些变量会使模型太复杂，并矛盾于保持复杂度界限简单有用的目标。因此，我们分析保持定性，并聚焦本质行为而非尝试提供详细理论运行时估计。

非流水线模式。 在非流水线模式中，每个GPU必须完成一迭代所有任务前开始下一。环AllReduce、环AllGather和环ReduceScatter遵循此模式。在以下分析中，k表示参与集体的GPU数。

环AllReduce。 NCCL中的环AllReduce算法结合分布式归约阶段与数据传播阶段，确保所有k参与GPU接收完整元素级归约结果。操作每个循环分为2k − 1步骤，如表V详细。从ReduceScatter-like阶段开始，如图4上部示意。最初，在步骤0，每个GPU发送其本地数据一段到邻居。在下一k − 2步骤，每个GPU重复执行recvReduceSend操作：从前邻接收数据段，与其本地数据对应段执行元素级归约，并转发归约结果到环中后续GPU。此迭代归约继续至步骤k − 1。在此步骤，每个GPU接收数据段、执行最终归约，从而产生完全归约段，并复制结果到输出缓冲区指定位置前发送此段。在此步骤，每个GPU接收数据段、执行最终归约，并复制结果到输出缓冲区指定位置前发送完全归约段。在下一k − 2步骤，每个GPU执行一系列recvCopySend操作。在每个步骤，GPU从前邻接收完全归约段、直接复制到输出缓冲区适当位置，并不变转发此段到下一GPU。环AllReduce操作在步骤2k − 2结束，每个GPU执行最终recv完成完全归约数据收集。
表VI NCCL环AllGather一个循环迭代中的步骤

环AllGather。 环AllGather算法使k参与GPU每个收集所有rank贡献的完整数据块集。算法在k − 1通信步骤上进行，使用连接GPU的逻辑环拓扑。在初始步骤（表VI中步骤0），每个GPU i准备其本地数据块。若操作就地，块已在输出缓冲区第i段。否则，GPU使用copySend原语从输入缓冲区复制数据到该段。此设置后，每个GPU发送其本地块到右手邻居。在下一k − 2步骤，每个GPU执行recvCopySend操作序列。在每个步骤，GPU从左手邻居接收块、存储到输出缓冲区正确段，并转发到右手邻居。最终步骤是recv操作交付最后缺失块。此步骤后，所有GPU持有集体数据的完整有序副本。
表VII NCCL环ReduceScatter一个循环迭代中的步骤

环ReduceScatter。 环ReduceScatter算法跨初始分布在k GPU上的数据块执行元素级归约，后跟将完全归约结果独特段散布回每个GPU。开始时，每个GPU的sendbuff包含k个独特数据块，随着它们围绕逻辑环拓扑移动逐步归约。表VII总结环ReduceScatter单一循环迭代每个步骤执行的原语。在初始步骤，每个GPU i发送其本地数据块之一到立即邻居GPU (i+ 1)%k，启动环绕数据移动。在后续k − 2步骤，每个GPU执行一系列recvReduceSend操作：从左手邻居(GPU (i − 1)%k)接收部分归约数据块、元素级结合此块与其sendbuff中对应本地块，并发送不同部分归约块到右手邻居。在最终步骤，每个GPU从左手邻居接收最后数据块、应用最终归约操作，并直接复制完全归约结果到其自身recvbuff。

图5. NCCL中树拓扑连接4 GPU的树AllReduce算法示意图，突出单一循环迭代内GPU通信原语序列。

流水线模式。 NCCL中的树AllReduce、环Broadcast和环Reduce算法遵循流水线执行模式。

树AllReduce。 树AllReduce算法在每个循环迭代内进行两个独特阶段：归约阶段后跟广播阶段。数据移动由涉及4 GPU的示例示意，如图5。虽然示意图显示4 rank上的完整树，但注意分支结构仅跨节点构建。每个节点内，NCCL简单链链接本地GPU。在NCCL备选实现中，这些两阶段常通过将SM分区为两个不均组并发执行。一组处理向根归约，而另一同时从根执行广播。此不对称分配为带宽密集归约阶段分配更多线程，实现更好资源利用。在归约阶段，叶GPU通过使用send操作向上发送其本地数据到父启动归约。中GPU使用recvReduceSend原语从一个或多个子接收数据、与自身数据执行元素级归约，并向上传递结果。最后，根GPU执行recvReduceCopySend，通过结合传入数据与其本地缓冲区完成归约，并复制完全归约结果到用户提供输出缓冲区。在广播阶段，完全归约结果向下树传播。根使用recvCopySend操作发送结果到其子。中GPU从父接收数据、复制到自身输出缓冲区，并使用相同recvCopySend原语转发到其子。叶GPU使用简单recv接收数据并复制到其输出缓冲区。每类GPU角色使用的设备原语序列总结在表VIII。
表VIII NCCL树AllReduce一个循环迭代中的步骤

表IX NCCL环Broadcast一个循环迭代中的步骤

环Broadcast。 NCCL环Broadcast算法从用户指定根GPU传播数据到通信器所有其他GPU。虽然使用环拓扑，但通信模式有效形成定向链，从根开始顺序通过每个GPU直到最后接收数据。操作从根GPU开始。如表IX所示，根若其发送缓冲区也是接收缓冲区则执行就地send操作，或copySend其中数据从其独特发送缓冲区先复制到接收缓冲区然后传输。在任一情况下，根发送其数据块到环中立即后继。链中后续每个GPU执行recvCopySend原语：从前驱接收数据块、复制到自身接收缓冲区，然后转发数据块到后继。此过程继续直到数据到达链中最后GPU。此最后GPU简单执行recv操作，复制传入数据到其接收缓冲区，并不进一步发送，因为逻辑链中所有GPU现已接收广播数据。
表X NCCL环Reduce一个循环迭代中的步骤


图6. 环和树AllReduce在节点间和节点内运行时协议运行时比较。每数据点包括20次运行带预热阶段。对于节点内通信为可读性仅报告中值，因为方差很低。

环Reduce。 NCCL环Reduce算法执行跨多个GPU分布数据的元素级归约，将最终结果聚合到用户定义根GPU。像环Broadcast一样，操作利用从环拓扑派生的逻辑链，沿此数据流动并向根累积。如表X所示，链从第一GPU发送其本地数据块到环中下一GPU开始。中间GPU执行recvReduceSend原语：每个接收部分归约块、使用其自身对应数据应用元素级归约，并转发更新结果到下一GPU。此过程重复直到数据到达目的地根。根GPU以recvReduceCopy完成操作：接收最终部分结果、与其本地数据归约，并存储完全归约输出到其接收缓冲区。

方法细节中的引用汇总。
- [7]：NCCL GitHub Issue #578: Ring AllReduce performance discrepancy, 未知年份, GitHub, https://github.com/NVIDIA/nccl/issues/578。在通信通道段落中，描述为大消息可能过载单一SM并无法饱和链路的问题。
- [8]：NCCL GitHub Issue #1302: Broadcast logic and topology insights, 未知年份, GitHub, https://github.com/NVIDIA/nccl/issues/1302。在通信通道段落中，描述为类似问题。
- [9]：Full bandwidth broadcast, reduction and scan with only two trees, 2007, European Conference on Recent Advances in Parallel Virtual Machine and Message Passing Interface, 无URL。在通信通道段落中，描述为双二叉树结构。
- [10]：Energy, Memory, and Runtime Tradeoffs for Implementing Collective Communication Operations, 2014, Journal of Supercomputing Frontiers and Innovations, 无URL。在通信通道段落中，描述为双二叉树结构。
- [12]：NCCL GitHub Issue # 609: why uses rdma write for default ib traffic, 2021, GitHub, https://github.com/NVIDIA/nccl/issues/609。在IB Verbs传输段落中，描述为使用RDMA write移动数据。
- [13]：Developing a Linux Kernel Module using GPUDirect RDMA, 2025, 无会议, 无URL。在GPUDirect RDMA优化段落中，描述为注册GPU内存机制。
- [14]：Buffer Sharing and Synchronization (dma-buf), 2025, 无会议, kernel.org。在GPUDirect RDMA优化段落中，描述为Linux DMA-BUF子系统。
- [15]：Optimization of collective reduction operations, 2004, ICCS, 无URL。在NCCL集体算法概述段落中，描述为算法选择依赖。
- [16]：Bandwidth optimal all-reduce algorithms for clusters of workstations, 2009, Journal of Parallel and Distributed Computing, 无URL。在NCCL集体算法概述段落中，描述为类似。
- [17]：Optimization of collective communication operations in mpich, 2005, The International Journal of High Performance Computing Applications, 无URL。在NCCL集体算法概述段落中，描述为类似。
- [18]：NCCL Issue #320: NVLS and CollNet Support, 2021, GitHub, https://github.com/NVIDIA/nccl/issues/320。在算法和协议支持概述段落中，描述为SHARP技术。
- [19]：Nccl github issue #919: Question: Nccl tree algorithm behaviour, 2023, GitHub, https://github.com/NVIDIA/nccl/issues/919。在算法和协议支持概述段落中，描述为CollNet和NVLS细节。
- [2]：NVIDIA Collective Communications Library (NCCL) Documentation, 2025, 无会议, 无URL。在算法和协议支持概述段落中，描述为NVLS设计。
- [20]：New scaling algorithm and initialization with nvidia collective communications library 2.23, 2025, 无会议, 无URL。在算法和协议支持概述段落中，描述为PAT算法。
- [21]：Pat: a new algorithm for all-gather and reduce-scatter operations at scale, 2025, 无会议, 无URL。在算法和协议支持概述段落中，描述为PAT算法。

实验环境

实验使用Alps超级计算系统，在瑞士国家超级计算中心（CSCS），配置16节点配备NVIDIA Grace Hopper Superchips (GH200)。每个节点提供150GB/s高带宽节点内互连，并通过25GB/s每方向网络链路连接到Cray Slingshot互连。无明确数据集或模型架构细节；基准聚焦NCCL集体操作性能。硬件配置包括GH200 GPU、节点内NVLink和Slingshot网络。软件配置基于NCCL版本2.19.1，无指定代码实现语言或OS，但隐含使用CUDA和NCCL库。

实验结果

AllReduce协议运行时比较（图6）。 实验内容：在节点间和节点内设置下基准环和树AllReduce的三种协议（Simple、LL、LL128），消息大小从小型到GB级，每数据点20次运行带预热。结果：在节点间，对于树和环算法，LL和LL128对小于64 KiB小消息性能最佳，但随消息大小增加到GB级，性能急剧下降于Simple协议，主要由于LL和LL128细粒度标志基同步开销，需要处理数百万小同步操作（每8或128字节一个）跨网络。LL128虽受益于更大缓冲区并在NVLink高效，但RoCE上大传输的累积同步成本超过优势，甚至落后于LL。Simple使用更大传输和更少同步事件，对网络延迟较不敏感，对极大消息维持高吞吐更好。在节点内，LL128一致性能优异跨所有消息大小，利用NVLink优势；对小消息，LL128与LL相当或略差，而对大消息几乎匹配Simple（预期5%慢，如表I）。LL和Simple在相反极端最佳，Simple大消息，LL小消息。在两设置中，环算法对大消息卓越，树算法对小消息最佳。分析结论：确认LL/LL128适合小消息尤其节点间，Simple优于大分布式传输；考虑节点内/间差异重要，LL128性能显著不同；通常依赖NCCL自动调优提供稳健性能。

其他集体算法运行时比较（图7）。 实验内容：基准各种NCCL集体（AllGather、Broadcast、Reduce、ReduceScatter）协议在节点间和节点内，每数据点20次运行带预热，节点内报告中值以可读性。结果：行为遵循AllReduce相同趋势，无具体数字，但确认类似协议性能模式。分析结论：LL/LL128小消息优势，Simple大消息；环大消息好，树小消息好。

图7. 各种NCCL集体在节点间和节点内运行时协议运行时比较。每数据点包括20次运行带预热。对于节点内通信为可读性仅报告中值，因为方差很低。