发表时间: 2025-01
作者: Sylvain Jeaugey
一句话结论 NVIDIA 在 NCCL 2.23 版本中通过引入对数级扩展的 PAT 算法、多 ID 负载均衡的初始化机制以及节点内用户缓冲区注册功能,解决了大规模集群下小数据量通信和初始化的扩展性瓶颈,并提供了细粒度的性能分析接口。
要解决什么问题 随着深度学习模型和 GPU 集群规模的扩大,多 GPU 及多节点间的通信面临四个具体的底层机制卡点。首先,在大规模集群中进行小数据量的 ReduceScatter 和 AllGather 操作时,传统的通信算法(如递归加倍)网络步骤数过多,扩展性极差。其次,在创建通信器时,NCCL 的引导步骤需要获取所有 rank 的地址,这依赖所有 rank 向同一个已知的唯一 ID 交换地址,这种“多对一”的通信模式在大规模场景下会造成严重的初始化时间瓶颈。第三,NCCL 的默认设计允许用户无需注册和维护任何持久性缓冲区即可工作,这极大地简化了可用性,但也带来了性能妥协:如果没有直接访问权限,NCCL 在传输数据时必须进行更多的控制流操作和缓冲处理,这不仅消耗更多 GPU 资源,还会在移动相同数据量时产生更高开销和多余拷贝,增加了内存子系统压力。最后,随着集群规模增加,性能异常变得越来越难以检测和定位,现有的系统缺乏能够以极低开销为正在运行的作业收集细粒度遥测数据的领域特定诊断工具。
怎么做的 针对上述卡点,NCCL 2.23 在算法、初始化、内存管理和监控四个维度进行了重构。第一,针对 ReduceScatter 和 AllGather 引入了基于 Bruck 算法变体的并行聚合树(PAT)算法。该算法为每个 rank 执行一个移位的二项树,其核心优势在于网络步骤数随节点数呈对数级 $O(\log N)$ 增长,且适用于任意数量的 rank 而不仅限于 2 的幂次方。随着数据量增加,PAT 会逐步增加网络传输次数以保持最小的缓冲需求。目前该算法专为每节点单 GPU 场景设计,完美契合大语言模型训练中张量并行与节点内 NVLink 对齐、其他并行维度跨节点单 GPU 的拓扑。第二,引入全新的 `ncclCommInitRankScalable` 初始化 API。该机制允许在创建通信器时传入多个唯一 ID,NCCL 会将地址交换负载均匀分散到这些 ID 上,只要提供的 ID 数量与通信器规模成比例增长,就能实现恒定的引导时间;同时支持通过 `NCCL_OOB_NET_ENABLE=1` 启用 IB/RoCE 等高速网络进行带外引导通信,并可通过 `NCCL_OOB_NET_IFNAME` 指定网络接口。第三,实现了节点内用户缓冲区注册。用户可通过 `ncclCommRegister` 显式注册或利用 CUDA Graphs 自动捕获(此版本将图分配的 `cudaMalloc` 等调用变为异步以加快捕获),使 NCCL 获得缓冲区的直接访问权限,从而彻底消除节点内节点间的额外拷贝,提升计算与通信的重叠率。该功能覆盖了多节点 NVLink 系统中每个域内的通信,以及除 `ncclReduce` 和 `ncclReduceScatter` 外的所有集合操作与 P2P 操作。第四,设计了层级化的分析器插件 API。通过 32 位事件激活掩码控制监控粒度,将事件按组、集合操作或 P2P 操作、代理进度通道、代理进度步骤等层级严格组织。API 暴露了五个生命周期回调:`init` 接收掩码并返回隔离的上下文对象;`startEvent` 分配并初始化事件返回句柄;`recordEventState` 允许第三方工具(如 PyTorch Kineto)在代理与 GPU 及网络协调传输数据发生状态转换时更新状态;`stopEvent` 标记完成并回收对象;`finalize` 释放资源。
效果如何 本文档为纯工程发布说明与算法设计阐述,作者并未在文中提供具体的数据集、模型规模、硬件配置等实验设置,也未提供任何定量的性能对比基线和图表数据。关于方法的有效性,作者仅基于算法复杂度给出了定性预期:PAT 算法在处理中小型消息时性能更优,且这种优势会随着工作负载规模的扩大而愈发明显。在局限性与代价方面,作者明确指出 PAT 算法当前仅支持每节点一个 GPU 的场景;节点内用户缓冲区注册功能虽然能提升性能,但无法使 `ncclReduce` 和 `ncclReduceScatter` 两个操作从中受益;此外,利用高速网络加速带外通信的功能由于在拓扑检测之前执行,若设备配置错误可能导致问题,因此默认处于禁用状态。此外,更新还包含边缘场景的修复:将默认 IB 超时从 18 增至 20,利用致命 IB 异步事件捕获链路断开错误,在每节点超两个 GPU 的 AMD CPU 上将 P2P 级别设为 PXB,并修复了混合大小操作的性能回归。
NVIDIA集合通信库(NCCL)是为NVIDIA GPU和网络优化的多GPU及多节点通信原语实现。它是多GPU深度学习训练的核心软件组件,负责处理各种GPU间通信,无论是通过PCI、NVLink还是网络。NCCL利用先进的拓扑检测、优化的通信图和调优模型,在NVIDIA GPU平台上实现开箱即用的最佳性能。本文将讨论NCCL 2.23版本中发布的新功能和修复。
NVIDIA Magnum IO NCCL库旨在优化GPU间和多节点通信,这对AI和高性能计算(HPC)应用中的高效并行计算至关重要。此版本的价值在于其新功能:
* 针对ReduceScatter和AllGather的全新PAT算法:引入了基于Brucks算法的并行聚合树(PAT)算法,用于AllGather和ReduceScatter,实现了对数级扩展性。
* 加速初始化:提升了初始化性能,包括能够使用带内网络进行引导通信。
* ncclCommInitRankScalable:一个新的初始化API,用于在创建通信器时使用多个ncclUniqueId,以加速大规模场景下的初始化。
* 节点内用户缓冲区注册:利用已注册的用户缓冲区进行节点内操作,提升性能。
* 新的分析器插件API:提供API钩子以测量细粒度的NCCL性能。
ncclCommInitRankScalable APIncclCommInitRankScalable,它允许在创建通信器期间利用多个唯一ID。这一新增功能避免了初始化过程中的“多对一”通信模式,从而提供了更具扩展性的初始化性能。ncclCommInitRankScalable的工作原理:通过ncclCommInitRankScalable,用户现在可以自由地提供多个唯一ID用于引导过程。为了获得最大增益,NCCL会将负载分散到多个唯一ID上。如果提供的唯一ID数量与通信器的规模成比例增长,就可以实现恒定的引导时间。ncclNet设备的使用发生在拓扑检测之前)。可以通过设置 NCCL_OOB_NET_ENABLE=1 来启用它。NCCL_OOB_NET_IFNAME 来指定应使用的网络接口。默认情况下,NCCL将使用在该网络上找到的第一个 ncclNet 设备。ncclCommRegister 注册他们的缓冲区,以便NCCL能够利用所有可用的优化。NCCL团队一直在努力为注册的用户缓冲区增加更多的使用场景。2.23版本为NvLink和PCIe P2P传输实现了节点内用户缓冲区(UB)注册支持。ncclReduce 和 ncclReduceScatter(因为它们无法从中受益),所有NCCL集合操作以及基于send/recv的操作都支持此功能。ncclCommRegister 显式注册缓冲区,缓冲区仅在调用相应的NCCL集合操作时才被注册。第二种方式是通过CUDA Graphs捕获NCCL操作,所有用户缓冲区将在图捕获期间自动注册。更多指南和要求,请参考NCCL文档。插件加载与事件激活:新的 NCCL_PROFILER_PLUGIN 环境变量控制分析器插件的加载和初始化,其方式与其他NCCL插件相同。加载后,分析器插件可以通过设置事件激活掩码来启用NCCL事件分析。NCCL在初始化期间将此掩码暴露给分析器。事件激活掩码是一个32位整数,其中每一位代表一个NCCL分析器事件。目前,NCCL支持以下事件:
ncclProfileGroup (bit-0): 组事件ncclProfileColl (bit-1): 集合操作事件ncclProfileP2p (bit-2): 点对点事件ncclProfileProxyOp (bit-3): 代理进度通道事件ncclProfileProxyStep (bit-4): 代理进度步骤事件ncclProfileProxyCtrl (bit-5): 代理进度内部状态事件事件的层次结构:NCCL以分层形式表示事件。例如,集合操作可以组合在一起,而代理操作则协助GPU通过可用的网络通信通道进行单个数据块的点对点传输。因此,NCCL向分析器呈现相应的事件时会保留这种关系。NCCL事件层次结构图如下所示:
ncclProfileGroup
+- ncclProfileColl
| +- ncclProfileProxyOp
| +- ncclProfileProxyStep
+- ncclProfileP2p
+- ncclProfileProxyOp
+- ncclProfileProxyStep
# ncclProfileProxyCtrl
ext-profiler/example 目录中提供了一个示例分析器插件,可用作开发第三方分析器插件的模板。API回调函数定义:分析器插件接口总共定义了以下五个函数回调:
ncclResult_t (*init)( void** context, int* eActivationMask);ncclResult_t (*startEvent)( void* context, void** eHandle, ncclProfilerEventDescr_t* eDescr);ncclResult_t (*stopEvent)( void* eHandle);ncclResult_t (*recordEventState)( void* eHandle, ncclProfilerEventState_t eState, NcclProfilerEventStateArgs_t* eStateArgs);ncclResult_t (*finalize)(void* context);init 函数:分析器的 init 函数接收一个事件激活掩码指针,并向NCCL返回一个不透明的上下文对象。该上下文为分析器实例之间提供了隔离,而事件激活掩码则由分析器用来通知NCCL应分析哪些事件;例如,设置 *eActivationMask = ncclProfileColl | ncclProfileProxyOp。
startEvent 函数:分析器的 startEvent 函数接收一个分析器上下文和一个事件描述符。分析器使用描述符信息来分配一个新的事件对象并对其进行初始化。之后,分析器返回一个不透明的句柄,NCCL可以用它来对事件执行进一步的操作,例如记录状态更新。stopEvent 函数:分析器的 stopEvent 函数接收一个事件句柄,并将该事件标记为完成。此后,该事件句柄将不能再被使用(分析器可能会在内部回收相应的对象以用于未来的事件)。recordEventState 函数:分析器的 recordEventState 函数接收一个事件句柄、一个事件状态以及(可选的)一个事件状态参数对象。此函数使分析器能够更新在NCCL中可能经历不同状态转换的事件。一个例子是代理事件,其中代理在传输数据时需要与GPU和网络进行协调,从而在此过程中从一个状态转换到另一个状态。finalize 函数:分析器的 finalize 函数接收分析器上下文,并释放与其相关的所有资源。cudaMalloc 和 cudaMemcpy 在图分配期间的调用变为异步,从而显著加快图捕获速度。ncclCommInitRank 还是 ncclCommSplit。NCCL_CONF_FILE 变量。NCCL_CROSS_NIC 设置为1时的树图搜索问题。文章未提供详细的实验环境配置信息,如数据集、模型、硬件(GPU、CPU、网络)或软件(操作系统、代码库)的具体细节。
文章主要对新功能进行了定性描述,并未提供定量的实验结果、图表或性能对比数据。文中提到的一些预期性能提升(如PAT算法对中小型消息的性能改善)是基于算法设计原理的推断。
NVIDIA NCCL 2.23版本为优化AI和HPC应用中至关重要的GPU间和多节点通信引入了多项新功能和改进。关键增强功能包括:新的PAT算法、大规模场景下的加速初始化、节点内用户缓冲区注册以及新的分析器插件API。