作者: Sylvain Jeaugey
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。