Enabling Fast Inference and Resilient Training with NCCL 2.27

作者: John Bachan
机构: NVIDIA

A1 主要贡献

随着人工智能（AI）工作负载的规模化，快速且可靠的GPU通信变得至关重要，这不仅对训练如此，对规模化推理也日益重要。NVIDIA集合通信库（NCCL）为NVIDIA GPU以及包括PCIe、NVLink、以太网（RoCE）和InfiniBand（IB）在内的多种互连方式提供了高性能、拓扑感知的集合操作，如AllReduce、Broadcast、Reduce、AllGather和ReduceScatter。

本文介绍了最新的 NCCL 2.27 版本，重点展示了其在提升推理延迟、训练弹性和开发者可观测性方面的新特性。核心贡献和创新点如下：

  1. 实现超低延迟的对称内存内核：引入对称内存支持，允许在不同GPU上使用相同虚拟地址的缓冲区，从而通过优化的集合操作内核显著降低所有消息大小的通信延迟。对于小消息，延迟最多可降低9倍。
  2. 通过Direct NIC支持释放全部网络带宽：在特定平台上，通过让CX8 NIC绕过CPU直接与GPU进行PCIe Gen6 x16连接，解决了CPU的PCIe Gen5带宽瓶颈，使GPUDirect RDMA等技术能够实现完整的800 Gb/s网络带宽。
  3. 扩展SHARP以提升大规模训练效率：为NVLink和InfiniBand fabric增加了对SHARP（可扩展分层聚合和归约协议）的支持，并将此支持扩展到AllGather (AG) 和 ReduceScatter (RS) 集合操作。通过将计算密集型任务卸载到网络，显著减少了GPU SM的使用量，从而为模型计算释放更多资源。
  4. 引入Communicator Shrink以增强训练弹性：为应对大规模训练中的设备故障，推出了Communicator Shrink功能。该功能允许在训练过程中动态地排除失败或不再需要的GPU，支持计划内重构（默认模式）和意外故障恢复（错误模式），从而保持训练不中断。
  5. 增强的API和分析工具：
    • 提供了新的对称内存API (ncclCommWindowRegister)，以便开发者注册对称内存，利用低延迟内核。
    • 全面升级了分析基础设施，通过统一代理事件、支持原生GPU时间戳、更新网络插件事件接口等方式，为开发者提供了更精确、更高效的通信性能诊断工具。

A3 背景知识

NCCL的核心价值与功能。随着AI工作负载的扩展，快速可靠的GPU通信对于训练和规模化推理都至关重要。NVIDIA集合通信库（NCCL）为NVIDIA GPU和多种互连技术（包括PCIe、NVLink、以太网（RoCE）和InfiniBand（IB））提供了优化的、拓扑感知的集合操作，例如AllReduce、Broadcast、Reduce、AllGather和ReduceScatter。

NCCL的优势。NCCL通过其通信与计算的单内核实现，确保了低延迟同步，使其成为分布式训练和实时推理场景的理想选择。得益于NCCL的动态拓扑检测和简化的基于C的API，开发者可以在不同节点间扩展应用，而无需为特定的硬件配置进行调优。

本文目标。本文旨在介绍最新的NCCL 2.27版本，展示其在提升推理延迟、训练弹性和开发者可观测性方面的新特性。

A2 方法细节

释放新的性能水平

NCCL 2.27 提供了关键更新，旨在增强跨GPU的集合通信，以应对延迟、带宽效率和扩展性方面的挑战。这些改进同时支持训练和推理，契合了现代AI基础设施不断发展的需求——其中，超低延迟对于实时推理管道至关重要，而强大的容错能力则是保障大规模部署可靠运行的必要条件。该版本的主要亮点包括：带有对称内存的低延迟内核、Direct NIC支持，以及对NVLink和InfiniBand SHARP的支持。

具备对称内存的低延迟内核

对称内存支持与性能提升。此版本引入了对称内存支持，允许跨GPU上具有相同虚拟地址的缓冲区从优化的集合操作中受益。这些内核显著降低了所有消息大小的延迟，如图1所示，对于小消息大小，延迟最多可降低9倍。

图1. NCCL 2.27中使用低延迟内核带来的AllReduce延迟改进

计算精度与确定性。归约操作使用FP32累加器（在NVLink Switch (NVLS)系统上，FP8使用FP16累加器）进行计算，这提高了诸如AllReduce、AllGather和ReduceScatter等操作的准确性和确定性。

支持范围与性能增益。对称内存支持单个NVLink域内的NVLink通信——在NVIDIA GB200和GB300系统中最高可达NVL72（72个GPU），在NVIDIA DGX和HGX系统中最高可达NVL8（8个GPU）。即使在NVL8域上，开发者也能观察到中小消息尺寸下高达2.5倍的性能提升。

Direct NIC 支持

Direct NIC解决带宽瓶颈。NCCL 2.27引入了对Direct NIC配置的支持，旨在为GPU横向扩展通信释放全部网络带宽。在特定的NVIDIA Grace Blackwell平台上，CX8 NIC和NVIDIA Blackwell GPU等组件支持PCIe Gen6 x16，可提供高达800 Gb/s的网络带宽。然而，Grace CPU目前仅支持PCIe Gen5，这将吞吐量限制在400 Gb/s。

Direct NIC架构原理。为了解决这个问题，CX8 NIC暴露了两个虚拟PCIe树：如图2所示，在一个树上，NVIDIA CX8 NIC的数据直接功能（PF）通过PCIe Gen6 x16链路直接连接到GPU PF，从而绕过CPU，避免了带宽瓶颈。在另一个树上，常规的NIC PF连接到CPU的根端口。

图2. Direct NIC架构，GPU与NIC之间通过PCIe Gen6连接，绕过了CPU瓶颈

Direct NIC的优势。这种配置确保了GPUDirect RDMA及相关技术能够达到完整的800 Gb/s带宽，而不会耗尽CPU到GPU的带宽，这在多个GPU共享单个CPU时尤为重要。Direct NIC是为高吞吐量推理和训练工作负载实现全速网络的关键。

NVLink 和 InfiniBand SHARP 支持

扩展SHARP支持至新集合操作。NCCL 2.27增加了对NVLink和IB fabric的SHARP（可扩展分层聚合和归约协议）支持。SHARP允许在网络内部执行归约操作，从而卸载计算密集型任务。这个新版本将SHARP支持引入了AllGather (AG) 和 ReduceScatter (RS) 集合操作，当使用NVLink Sharp加上IB Sharp时，数据可以直接从GPU传输到网络。

对大规模LLM训练的益处。这对于大规模LLM训练尤其有益，因为在这些场景中，AG和RS现在比AllReduce更受青睐，以便更好地重叠计算与通信。传统的基于环的实现可能会消耗16个或更多的SM，但借助NVLink和IB SHARP，这一需求减少到6个SM或更少，从而为模型计算释放了资源，并提升了整体训练效率。其结果是在千卡GPU级别及以上实现了更好的可扩展性和性能。

支持SHARP的集合操作可减少SM使用。

图3. 支持SHARP的集合操作通过将数据聚合卸载到NVLink和InfiniBand网络交换机，减少了SM的使用并提高了可扩展性

使用 NCCL Shrink 增强大规模训练的弹性

引入Communicator Shrink功能。NCCL 2.27引入了Communicator Shrink功能，旨在使分布式训练更加稳健、灵活和高效。在数百或数千个GPU上运行的训练任务容易受到设备故障的影响。Communicator Shrink功能可以在训练期间动态排除发生故障或不再需要的GPU。该功能支持两种操作模式：

  • 默认模式 (Default mode): 用于计划内的重新配置，允许修改设备拓扑，同时确保所有操作都能完成。
  • 错误模式 (Error mode): 自动中止正在进行的操作，以便从意外的设备故障中恢复。

NCCL Shrink 的核心能力。NCCL Shrink 使开发者能够：
* 通过动态重建通信器来维持不间断的训练。
* 通过可配置的资源共享，尽可能地重用资源。
* 以最小的中断优雅地处理设备故障。

NCCL Shrink 使用示例。以下是用于计划内重构和错误恢复场景的NCCL Shrink使用示例代码：
计划内重构：

NCCLCHECK(ncclGroupStart());   
for (int i = 0; i < nGpus; i++) {   
    if (i != excludedRank) {   
        NCCLCHECK(ncclCommShrink(   
            comm[i], &excludeRank, 1,   
            &newcomm[i], NULL, NCCL_SHRINK_DEFAULT)); 
    }
} 
NCCLCHECK(ncclGroupEnd());

错误恢复：

NCCLCHECK(ncclGroupStart()); 
for (int i = 0; i < nGpus; i++) { 
    if (i != excludedRank) { 
        NCCLCHECK(ncclCommShrink( 
            comm[i], &excludeRank, 1, 
            &newcomm[i], NULL, NCCL_SHRINK_ABORT)); 
    } 
} 
NCCLCHECK(ncclGroupEnd());

为开发者提供的更多功能

此版本为开发者提供的其他功能包括对称内存API和增强的分析功能。

对称内存API

对称内存的基础作用。对称内存是NCCL 2.27中的一项基础能力，它支持高性能、低延迟的集合操作。当内存在所有rank上分配到相同的虚拟地址时，NCCL可以执行优化的内核，从而减少同步开销并提高带宽效率。

窗口API介绍。为了支持这一点，NCCL引入了一个用于集合注册对称内存的窗口API：
ncclCommWindowRegister(ncclComm_t comm, void* buff, size_t size, ncclWindow_t* win, int winFlags);
ncclCommWindowDeregister(ncclComm_t comm, ncclWindow_t win);
ncclCommWindowRegister函数将用户分配的内存注册到NCCL通信器。内存必须使用CUDA虚拟内存管理（VMM）API进行分配。winFlags必须包含NCCL_WIN_COLL_SYMMETRIC以启用对称内核优化。所有rank必须提供具有匹配偏移量的缓冲区，以确保对称寻址。取消注册（ncclCommWindowDeregister）是一个本地操作，并且只应在所有相关集合操作完成后进行。

API使用注意事项。ncclCommWindowRegister是集合式和阻塞式的，这意味着当单个线程管理多个GPU时，这些调用必须被包含在ncclGroupStart和ncclGroupEnd之内。如果不需要对称内存，用户可以通过设置NCCL_WIN_ENABLE=0来完全禁用该功能。

对称内存注册流程。图4展示了如何使用NCCL窗口API跨多个GPU注册对称内存。通过对齐虚拟地址，NCCL启用了优化的低延迟内核，从而提高了集合操作的性能。

图4. NCCL中跨GPU的对称内存注册

增强的分析功能

NCCL 2.27对其分析基础设施进行了一系列增强，为开发者和工具提供了更准确、更高效的性能监测手段，以诊断通信性能。

代理事件的统一。之前，NCCL暴露了ncclProfileProxyOp和ncclProfileProxyStep两个事件来跟踪网络代理线程的进度。虽然这些事件提供了不同粒度的信息，但它们也造成了许多监测点的重复。在2.27版本中，NCCL通过移除冗余的ProxyOp状态并引入一个统一的ncclProfilerProxyOpInProgress_v4状态，简化并优化了这一模型。这在不牺牲细节的情况下减少了分析器的开销，并提高了跟踪通信进度时的清晰度。此外，还引入了一个新的ProxyStep事件状态ncclProfilerProxyStepPeerWait_v4，用于反映发送方rank等待接收方发布发送许可信号的时间，整合了以前的功能，同时最大限度地减少了重复。

GPU内核事件的准确性。为了提高计时精度，NCCL现在支持原生GPU时间戳传播。GPU现在使用其内部的全局计时器记录并导出开始和停止时间戳，而不是依赖于通过GPU工作计数器进行的主机端事件计时——这种方法容易受到延迟伪影（例如，内核延迟或合并）的影响。这使得分析工具能够直接从GPU获取精确的内核运行时长，不过开发者在将时间戳转换为主机时间时需要进行校准或插值。

网络插件事件更新。NCCL分析器接口现在支持对网络定义的事件使用recordEventState。这一新机制允许分析器更新正在进行的操作的状态，这对于将实时的网络反馈（如重传信号或拥塞提示）注入性能时间线非常有用。

其他增强功能。
* 分析器初始化：NCCL现在在分析器初始化期间报告通信器元数据，包括名称、ID、节点数、rank数和调试级别。
* 通道报告：报告的通道数量反映的是实际使用情况，而不是理论上的限制。这包括点对点（P2P）操作。
* 通信器标记：ncclConfig_t已扩展，可包含通信器名称，从而改善了被分析操作与特定通信器之间的关联性。

这些更新共同提升了NCCL分析器插件接口的保真度，为开发者提供了对网络动态、GPU计时和操作结构的更深入洞察，这些对于诊断和调优大规模AI工作负载至关重要。

前瞻性支持

前瞻性支持包括：
* 跨数据中心通信：提供初步支持，允许跨地理分布的数据中心执行集合操作。
* 多NIC插件可见性：支持同时利用多种网络配置。

A4 实验环境

本文主要介绍NCCL 2.27的功能，未提供完整的基准测试实验设置，但提及了相关的软硬件配置：

  • 硬件配置:

    • GPU平台: NVIDIA GB200, GB300, DGX, 和 HGX 系统。
    • GPU互联: NVLink, NVLink Switch (NVLS)，支持多达 NVL72 (72 GPUs) 和 NVL8 (8 GPUs) 的NVLink域。
    • CPU: NVIDIA Grace CPU。
    • GPU: NVIDIA Blackwell GPUs。
    • 网卡 (NIC): NVIDIA CX8 NIC。
    • 总线: PCIe Gen5 和 PCIe Gen6 x16。
    • 网络互联: 以太网 (RoCE) 和 InfiniBand (IB)。

  • 软件配置:

    • 通信库: NCCL 2.27。
    • 底层API: CUDA Virtual Memory Management (VMM) API。
    • 测试工具: NCCL-Test repository。

A4 实验结果

本文通过图表展示了NCCL 2.27几个关键特性的性能优势：

  • 低延迟内核性能:

    • 实验内容: 对比NCCL 2.27和NCCL 2.26版本在使用低延迟内核（对称内存）时的AllReduce操作延迟。
    • 实验结果: 如图1所示，NCCL 2.27在所有消息尺寸上都显著降低了延迟。对于小消息（如64B），延迟降低了9倍；对于中等消息（如4KB、128KB），延迟分别降低了7.6倍和5.6倍；对于大消息（8MB），延迟降低了3.3倍。
    • 分析结论: 对称内存和优化的内核能极大减少同步开销，尤其利好于小消息通信频繁的场景，如实时推理。

  • Direct NIC架构优势:

    • 实验内容: 通过架构图对比了有无Direct NIC的系统连接方式。
    • 实验结果: 如图2所示，没有Direct NIC时，GPU与NIC的通信需经过CPU，受限于CPU的PCIe Gen5带宽（400 Gb/s），形成瓶颈。采用Direct NIC后，GPU和NIC通过PCIe Gen6直接连接，绕过CPU，可以达到完整的800 Gb/s吞吐量。
    • 分析结论: Direct NIC是解锁Blackwell平台全部网络带宽的关键，能有效支持高吞吐量的训练和推理任务。

  • SHARP卸载效果:

    • 实验内容: 对比了启用和未启用SHARP时，执行集合操作所需的GPU SM（Streaming Multiprocessor）数量。
    • 实验结果: 如图3所示，传统的基于环的集合操作实现需要16个SM。启用NVLink和IB SHARP后，由于计算聚合任务被卸载到网络交换机，所需的SM数量减少到6个，实现了2.7倍的SM资源节约。
    • 分析结论: SHARP通过卸载数据聚合任务，为模型计算释放了宝贵的GPU计算资源，提高了训练效率和大规模系统的可扩展性。

A5 结论

NCCL 2.27版本通过引入多项关键功能，显著提升了分布式推理和训练工作流的性能、弹性和可观测性。其核心优势包括通过对称内存实现的更低延迟、通过Direct NIC支持实现的更高带宽、通过SHARP卸载实现的更高训练效率，以及通过Communicator Shrink实现的更强容错能力。此外，增强的分析工具为开发者提供了深入洞察通信性能的强大手段。

未来工作展望包括对跨数据中心通信的初步支持和对多NIC插件的可见性，这将进一步扩展NCCL在更复杂和地理上分散的环境中的应用能力。开发者可以访问NVIDIA/nccl GitHub仓库获取详细文档、源代码和支持，以开始使用NCCL 2.27的这些新功能。