数据中心/云端

在 NVIDIA NCCL 2.28 中将通信和计算与新的设备 API 和复制引擎集合融合

NVIDIA 集合通信库(NCCL)的最新版本引入了突破性的通信与计算融合技术,可显著提升多 GPU 和多节点系统中的通信吞吐量,降低延迟,并进一步提高 GPU 利用率。

NCCL 2.28 聚焦于 GPU 发起的网络通信、支持通信与计算融合的设备级 API、基于复制引擎的集群架构,以及 新 API 为开发者提供的用于构建高效、可扩展分布式应用。除了在性能方面的多项创新,该版本还通过扩展的 API 功能、优化的开发工具和更清晰的集成路径,显著提升了开发体验,使开发者能够更灵活、高效地编写自定义通信内核,并进一步拓展其应用能力。

版本亮点 

支持在性能、监控、可靠性和服务质量方面的持续改进。

  • 设备 API: 支持开发用于通信与计算融合的自定义设备内核(包括支持 GPU 发起网络操作的内核),使内核能够直接执行网络任务。
  • 基于复制引擎(CE)的集群: 开发者可利用 CE 驱动 NVIDIA NVLink 传输,有效降低流多处理器(SM)在通信过程中的计算资源争用。
  • NCCL Inspector: 提供低开销的分析插件,实现对 NCCL 通信模式的持续可观测性与深入分析。

NCCL 设备 API 如何实现内核间的直接通信

NCCL 2.28 引入了设备端通信 API,支持在 NVIDIA CUDA 内核中直接发起通信操作。此前,所有 NCCL 操作均由主机端启动,这会带来额外的同步开销。借助这一新 API,内核能够直接启动数据传输,实现通信与计算的深度融合,从而提升整体吞吐量并降低延迟。使用该 API 时,需采用具备对称内存窗口的数据缓冲区。

目前支持三种操作模式:

  • 可加载/存储访问(LSA): 通过内存加载/存储操作实现设备间通信,采用 CUDA P2P 技术。
  • Multimem: 利用 NVLink SHARP 提供的硬件组播功能,实现设备间的高效通信。
  • GPU Initiated Networking(GIN): 支持 GPU 主动发起网络通信,实现设备之间的数据传输。
Flowchart showing three NCCL Device API operation modes for GPU memory communication. “LSA” (Load/Store Accessible Memory) enables access to remote memory over PCI for load/store operations. “Multimem” (NVLink Sharp) provides multicast load/reduce operations over NVLink. “GIN” (GPU-Initiated Networking) supports put, wait, and signal operations over network mechanisms such as GDAKI and Proxy. The diagram maps each mode to its corresponding interconnect technology. Flowchart showing three NCCL Device API operation modes for GPU memory communication. “LSA” (Load/Store Accessible Memory) enables access to remote memory over PCI for load/store operations. “Multimem” (NVLink Sharp) provides multicast load/reduce operations over NVLink. “GIN” (GPU-Initiated Networking) supports put, wait, and signal operations over network mechanisms such as GDAKI and Proxy. The diagram maps each mode to its corresponding interconnect technology.
图1: 设备API支持三种操作模式,用于内存通信及底层通信机制。

如需了解更多详情,请参阅在线文档

特别值得关注的是,GPU-Initiated Networking(GIN),新引入了 NCCL 2.28.7,使 GPU 能够自主管理网络操作,无需依赖 CPU 干预。通过该机制,内核可直接对数据传输、信号通知和同步操作进行排队,有效消除传统主机驱动控制路径所带来的性能瓶颈。

通过复制引擎卸载来加速 NCCL 性能

要大规模实现出色的通信性能,需要更多的SM来充分使用NVLink带宽。然而,持续增加的通信任务分配会与计算核心争用资源,从而影响应用程序的整体性能。

CE 集合将 NVLink 域内的通信任务从 SM 卸载至专用的硬件 CE 单元。与传统的 NCCL 集合不同,基于 CE 的集合实现了零 SM 参与的通信操作。该方法适用于仅需数据移动的集合操作,例如 AlltoAll 和 AllGather,从而释放 SM 资源以专注于计算任务,并提升通信与计算的重叠效率。现在,通信与计算可并行执行,不再竞争相同的硬件资源。

基于 CE 的集合通过多项优化来提升性能。例如,利用批处理 API(如 cudaMemcpyBatchAsync)将多个 CE 操作合并为单次调用,有效降低 CUDA 驱动开销;同时采用 NVLink 组播机制优化同步信号的广播,提升传输效率。

这些群集无需消耗 SM 资源,即可达到与基于 SM 的群集相当的性能。下图显示,基于 CE 的 AllGather 实现了高于基于 SM 的 AllGather 的峰值带宽。这一性能优势的另一个原因是,相比由 SM 发起的事务,由 CE 发起的 NVLink 事务采用了更大的事务宽度。

Line chart titled “Allgather Bandwidth – Size (1 Node – 8 GPUs Blackwell Umbriel)” comparing CE Multicast, CE Unicast, SM Symmetric (NCCL Optimized), and SM NonSymmetric (NCCL Default). The x-axis shows message sizes from 4 MB to 4 GB, and the y-axis shows bandwidth in GBps. CE Multicast achieves the highest bandwidth at larger message sizes, peaking around 780 GBps, while SM NonSymmetric peaks near 620 GBps—indicating a 1.25× bandwidth gain for CE Multicast over SM Symmetric at 4 GB. Line chart titled “Allgather Bandwidth – Size (1 Node – 8 GPUs Blackwell Umbriel)” comparing CE Multicast, CE Unicast, SM Symmetric (NCCL Optimized), and SM NonSymmetric (NCCL Default). The x-axis shows message sizes from 4 MB to 4 GB, and the y-axis shows bandwidth in GBps. CE Multicast achieves the highest bandwidth at larger message sizes, peaking around 780 GBps, while SM NonSymmetric peaks near 620 GBps—indicating a 1.25× bandwidth gain for CE Multicast over SM Symmetric at 4 GB.
图2. NCCL All-to-All 性能对比:基于CE与基于SM的实现

查看 NCCL 文档中的要求,了解如何启用 CE 群集,在 NCCL 文档中。

使用 NCCL Inspector 轻松实现分析与监控

NCCL Inspector 是一个用于可观测性、分析与诊断的插件,能够提供针对每个通信器和每个集合操作的详细性能数据及元数据记录。该插件适用于持续运行模式,旨在通过为每次操作生成结构化的 JSON 输出,帮助用户深入分析和调试 NCCL 集合操作,从而在执行分布式工作负载时更好地理解通信行为与性能特征。

NCCL Inspector 通过 NCCL Profiler 插件接口架构集成至 NCCL 应用场景中,主要功能包括:

  • 每通信者追踪: Inspector 会独立追踪每个 NCCL 通信者,帮助用户分析不同通信上下文中的性能表现。这在复杂的分布式应用中尤为关键,因为此类应用通常会为不同用途创建多个通信者。
  • 持续运行: 由于插件具备低开销特性,可长期启用,适用于生产环境中的工作负载,在不影响性能的前提下为 NCCL 提供全面的可观察性。
  • 事件追踪:插件能够捕获详细的事件记录,包括集合操作的启动与结束事件、内核通道活动以及相关的时间信息。
  • 性能指标:NCCL Inspector 会计算并报告多项关键性能指标,如算法带宽、总线带宽、执行耗时、计时来源信息(GPU 或 CPU 计时)、消息大小及集合操作类型。

该插件旨在与其他 NCCL 插件协同工作,可为基于性能反馈优化通信模式的调优插件提供有价值的性能数据。尽管当前设计未实现 NCCL Inspector 与调优器插件之间的上下文直接共享,但 Inspector 生成的详细性能数据可通过外部分析工具进行处理,从而为调优决策提供有力支持。

Line chart titled “HCA-Only ReduceScatter,” showing collective bus bandwidth versus collective sequence number for three NCCL Inspector runs. The x-axis spans sequence numbers from 0 to about 2,800, and the y-axis shows collective bus bandwidth. Three lines represent individual runs: green for 8.39MB_r_32 (highest), blue for 8.40MB_r_32, and red for 5.85MB_r_32 (lowest), each remaining mostly stable across the range. Line chart titled “HCA-Only ReduceScatter,” showing collective bus bandwidth versus collective sequence number for three NCCL Inspector runs. The x-axis spans sequence numbers from 0 to about 2,800, and the y-axis shows collective bus bandwidth. Three lines represent individual runs: green for 8.39MB_r_32 (highest), blue for 8.40MB_r_32, and red for 5.85MB_r_32 (lowest), each remaining mostly stable across the range.
图3:NCCL Inspector 数据的弹性及 Kibana 仪表板可视化

该插件独立于特定网络实现运行,因此兼容 NCCL 支持的各类网络技术,能够确保无论底层网络基础设施如何,Inspector 均可提供有效的洞察。

如需了解更多信息,请参阅 ext-profiler/inspector/ 目录下的 README.md 文件.

使用 NCCL 2.28 改进开发者体验

NCCL 2.28 在核心通信与计算融合的基础上,进一步通过一系列增强功能提升了灵活性、性能调优能力以及开发体验。新增的 API、内核编排机制、配置选项、分析工具和构建系统,使开发者能够更精细地控制通信工作流,增强在不同硬件和环境下的可观测性,同时简化部署流程。以下将简要介绍这些关键改进。

适用于 AllToAll、Gather 和 Scatter 的新主机 API

通过引入用于 AlltoAll、Gather 和 Scatter 操作的原生主机级 API,NCCL 能够实现更高级的优化。例如,在此版本中,NCCL 可利用复制引擎降低这些通信模式对 SM 的占用率。此类性能提升仅在使用专用的原生 API 时方可实现。

对称核函数组调用支持

在使用 NVLINK 域时,单个对称内核即可实现优异的延迟表现、带宽性能和资源利用率。NCCL 2.28 新增了对分组对称内核的支持,进一步提升了性能与资源利用效率。用户可按照常规方式注册窗口缓冲区,并在 ncclGroupStart()ncclGroupEnd() 之间发起多个 NCCL 群集操作。在启动过程中,NCCL 会自动识别可应用对称内核及分组优化的潜在群集,并将其调度至单一内核中执行,从而提升整体运行效率。有关窗口缓冲区注册的详细信息,请参考 NCCL 窗口缓冲区注册文档

使用 NCCL 环境插件实现灵活的配置管理

优化通信参数对性能至关重要,NCCL 通过配置文件和环境变量为此提供了可靠的机制。然而,当不同作业需要特定的 NCCL 版本和配置时,目前依赖文件的手动版本匹配方法难以与数据库或云调度系统等现代部署架构有效集成。

新的 NCCL 环境插件 API 提供了一种灵活的编程替代方案,具备以下显著优势:

  • 编程版本匹配可为每个 NCCL 版本自动应用相应的配置。
  • 与存储方式无关的配置则从文件、数据库或环境变量中获取设置。
  • 更精细的控制与更高的灵活性通过以编程方式覆盖部分参数,同时保留其他参数不变实现。
  • 初始化和资源管理 仅需启动一次,即可在运行时持续生效,并支持资源的有序释放。
  • 设计具备前瞻性,支持全新的每通信设备独立配置。

启用后,插件将与 NCCL 参数子系统集成,覆盖低优先级的配置机制,确保版本感知配置的一致性。该功能简化了大规模、多环境的部署流程,帮助用户突破静态文件系统带来的限制。

插件的共享上下文

随着AI模型训练日益扩展至多个数据中心,通信库面临严峻挑战。重新设计的插件系统支持共享上下文和针对每个通信方的个性化调优,可替代全局初始化机制,并实现跨不同网络环境的上下文感知优化。

以下网络插件 API 已更新:

// Network plugin v11
ncclResult_t (*init)(void** ctx, uint64_t commId, ncclNetCommConfig_v11_t* config, ...)
ncclResult_t (*finalize)(void* ctx);
ncclResult_t (*listen)(void* ctx, ...);
ncclResult_t (*connect)(void* ctx, int dev, void* handle, ...);

每个通信器现在均使用 commId 和网络配置进行初始化,并返回对应的上下文句柄(ctx),以实现资源隔离和精细化调优。该机制已被应用于调优器和分析器插件,目前可进一步扩展至网络插件。开发者还可通过设置 NCCL_NET_PLUGIN 环境变量,将网络、调优器和分析器插件的代码整合到同一个 .so 库中,从而实现上下文共享与插件间的通信。

NCCL 分析器 API 事件

过去,NCCL 分析器插件接口未能捕获 NCCL API 调用或 CUDA 内核启动事件,因此无法将启动、集合操作或点对点通信事件与对应的 NCCL API 调用关联起来。

分析器现在支持:

  • 按用户调用顺序显示 API 事件,以确保操作顺序的正确性,即使在使用 ncclGroupStart/ncclGroupEnd 进行操作分组的情况下也能保持正确的排序,避免因调度顺序变化而影响结果。
  • 准确测量 NCCL 操作在 CPU 上的开销,即从用户发起 API 调用到 NCCL 实际调度操作所耗费的时间。
  • 将启动的 CUDA 核函数与对应的原始 NCCL API 调用直接关联,建立清晰的执行路径。
  • 将 NCCL 调度的集合通信任务和点对点通信任务与其原始 API 调用进行关联。由于 API 事件在图执行启动期间持续存在,因此每次图启动时产生的底层任务均可追溯至用户的原始 API 调用,确保执行过程的可追踪性与一致性。

基于 CMake 的构建系统

NCCL 现已支持适用于 Linux 平台的 CMake,为传统的 Make 构建方式提供了更现代化的替代方案。CMake 的引入有助于简化与大型构建流程及跨平台项目的集成,同时保持对现有构建系统的兼容性。通过标准化的 CMake 支持,项目能够采用 CUDA 和 HPC 生态系统中广泛使用的通用工作流,提升构建的可重现性,并降低维护成本。这将带来更灵活、模块化且更便于开发者使用的构建体验。

开始使用 NCCL 2.28

了解 NCCL 2.28 的新特性,借助增强的可扩展性、优化的集合操作性能以及更高的跨节点效率,进一步突破分布式训练的性能瓶颈。

有关详细文档、源代码及社区支持,请访问 NVIDIA/nccl GitHub 仓库并参考其中提供的 示例。若需了解如何针对您的系统架构优化 NCCL,可查阅 NCCL 官方文档

 

标签