在过去几个月中，NVIDIA 集合通信库（NCCL）开发者一直在努力开发一系列新的库功能和错误修复。在本文中，我们将讨论 NCCL 2.22 版本的详细信息以及解决的痛点。
版本亮点
NVIDIA Magnum IO NCCL 是一个旨在优化 GPU 之间和多节点通信的库，对于 AI 和 HPC 应用中的高效并行计算至关重要。这次版本的价值在于其新的特点：
- 建立延迟连接以节省 GPU 内存：将连接创建延迟至需要时进行，从而减少 GPU 内存开销。
- 用于成本估算和工作负载均衡的新 API：提供新 API，帮助您优化计算和通信重叠，或研究 NCCL 成本模型。
- 针对
ncclCommInitRank的优化和 Instrumentation：消除冗余拓扑查询，将创建多个 Communicator 的应用的初始化速度提升高达90%。 - 通过 IB 路由器支持多个子网：为跨多个 InfiniBand 子网的作业添加通信支持，从而使 DL 训练作业能够在超过 40,000 个端点的 InfiniBand 网络上运行。
特征
在本节中，我们将深入探讨每个新功能的详细信息：
- 延迟连接建立
- 新的成本模型 API
- 初始化优化和仪器分析
- 新的调谐器插件界面
- 静态插件链接
- 用于中止或销毁的组语义
- IB 路由器支持
延迟连接建立
NCCL 使用一组静态分配的持久性连接和缓冲区，以运行其即时数据传输协议。对于 NCCL 支持的每种给定算法和协议，它都会创建一组单独的连接和缓冲区，每个连接和缓冲区都需要数 MB 的 GPU 内存。
作为参考，算法定义了给定集合的参与者之间的高级数据移动，协议定义了 NCCL 发送数据的方式。根据操作、消息大小、规模和拓扑结构，选择给定的算法和协议，以实现最佳性能。
在 2.22 之前，NCCL 会针对每种组合在对等体之间建立连接，这可能会浪费数 MB 的 GPU 内存来处理永远无法使用的算法和协议。
现在，NCCL 会等待为给定算法建立连接，直到首次需要它。这将大幅减少 NCCL 内存开销，尤其是在窄范围内使用 NCCL 时。例如，如果您反复仅以相同的消息大小运行 ncclAllReduce，则应在给定系统上仅使用一种算法。
该功能默认处于启用状态，但可以通过设置 env NCCL_RUNTIME_CONNECT=0 来禁用它。
在单节点 DGX-H100 的上一个场景中，我们发现仅使用 Ring 算法的 NCCL 的 GPU 内存使用减少了 3.5 倍，仅使用基于 NVSwitch 的归约时的 GPU 内存使用减少了 1.47 倍。
新的成本模型 API
应用程序开发者希望充分利用NVIDIA系统上提供的计算、内存和带宽资源。
理想情况下，计算和通信完全重叠，两者都能完美地完成工作，并将硬件的全部功能发挥到极致。在运行大规模 HPC 和 AI 应用程序时，尤其是在多个平台上运行一个代码库时，很难做到这一点。
为帮助解决此问题，NCCL 添加了一个新的 API，使您能够了解它认为给定操作需要多长时间。此 API 称为 ncclGroupSimulateEnd。其使用方式与 ncclGroupEnd 相同，因此任何熟悉编写 NCCL 代码的人都可以轻松使用。
不同之处在于，它不启动通信操作。相反，NCCL 计算它认为操作需要多长时间，并在提供的 ncclSimInfo_t 结构中设置此操作。
ncclGroupStart()
ncclAllReduce()
ncclGroupSimulateEnd(sim_t)
printf("Estimated completion time=%f microseconds\n", sim.time);
configureComputeAmount(sim.time, &computeIters, &workItemSize);
但是，此 API 返回的值并不完全符合现实。这是基于 NCCL 内部模型的估计值。截至 2.22，此 API 仅返回组中最后一次操作的估计时间。
初始化优化和仪器分析
随着客户工作负载种类繁多且规模不断增加，降低 NCCL 初始化的开销已经成为 NCCL 团队日益优先考虑的事项。
即使是单节点作业，NVIDIA Hopper GPU 上必须单独发现和连接的 NVLink 互连数量也有所增加，这导致初始化时间的大幅增加。
我们希望缩短初始化时间，并且必须研究每个初始化步骤的开销。我们首先分析 ncclCommInitRank 中的每个阶段，并研究不同规模的每个阶段的时间。现在，每当您收集标准 NCCL 日志（NCCL_DEBUG=INFO）时，您都会看到这一点。
此外，还有一个新的 NCCL_PROFILE 调试子系统，如果您不关心 NCCL 初始化日志的其余部分，该子系统将仅提供仪器信息。
之前讨论过的连接建立是一个前景广阔的改进领域。切换到 lazy establishment 可以节省内存并缩短初始化时间。
另一个领域是拓扑发现，这是一个初始化步骤，其中每个 NCCL 排名确定节点上可用的硬件。这包括系统上有哪些 GPU 和 NIC、存在多少 NVLink 互连，以及 PCI 拓扑和 NUMA 亲和力。
事实证明，NCCL 执行 NVLink 发现的方式并不理想，因为每个等级都在自行发现所有链路，从而导致冗余和拥塞。
为解决此问题，我们重复使用了拓扑融合代码，该代码最初在 NCCL 2.21 中引入，是多节点 NVLink (MNNVL) 支持的一部分，其中每个节点上的部分可用信息在使用节点间通信的引导过程中进行组合，从而全面了解 NVLink 拓扑结构。
对于 2.22，我们将此功能扩展到在每个节点内运行。现在，每个等级仅会发现有关其自己的 GPU 的信息，然后使用节点内拓扑融合将这些结果与对应的结果相结合。
在单个 8 个 H100 GPU 系统上，延迟连接建立和节点内拓扑融合相结合可以将 ncclCommInitRank 的执行时间缩短 90%（约 6 秒）。以前大约需要 6.7 秒，现在大约需要 0.7 秒。对于在执行期间创建多个通信器的应用程序，这可以大幅减少初始化时间。
新的调谐器插件界面
借助新的调谐器插件接口（v3），NCCL 为插件提供按集合计算的 2D 成本表，报告执行算法和协议的每种组合所需的估计操作时间。
NCCL 将与检测到的拓扑不兼容的表项设置为 -1，以便向外部调谐器表明这些组合不支持或不允许被覆盖。
为选择特定的组合，外部调谐器会将所需算法或协议组合的值更新为0或整个表的最小值。在插件更新成本表后，NCCL可以使用它为给定集合选择最终配置。
静态插件链接
NCCL 团队提供了一个插件模型，供合作伙伴提供自己的调整或网络后端，以代替 NCCL 内部模型和 InfiniBand 插件。一些合作伙伴希望以静态方式将这些插件与应用程序二进制文件关联起来，以方便起见，并避免加载错误的插件。
如果应用程序已静态链接网络或调优插件，请通过将 NCCL_NET_PLUGIN 或 NCCL_TUNER_PLUGIN 设置为 STATIC_PLUGIN 来指定它。
用于中止或销毁的组语义
之前，ncclCommDestroy 和 ncclCommAbort 将阻塞调用线程，直到完成。
对于多维并行 ML 工作负载，一个进程管理着多个 NCCL 通信器，并且每个通信器最终必须使用这些 API 拆解。我们为这些应用提供了语义，使其能够以分组方式一次销毁多个通信器，从而避免死锁并提供更好的用户体验。
IB 路由器支持
借助此功能，NCCL 可以在由一个或多个路由器连接的不同 InfiniBand 子网中运行。NCCL 会自动检测两个通信端点何时位于同一个 InfiniBand 网络的不同子网上，并交换建立连接和通信所需的 GID 信息。
在子网之间路由时，可使用FLID识别一组要转发的路由器，并在子网之间实现更高性能和自适应路由。NCCL 2.22将自动检测是否存在FLID，并将其用于不同子网上端点之间的连接。
问题修复和次要功能
NCCL 2.22 提供以下额外的更新：
- 增加了对 DGX 上的 Google Cloud
allreduce树算法的支持。 - 在 IB 异步错误中记录了 NIC 名称。
- 修复了聚合集合性能问题。
- 修复了注册发送和接收操作的性能问题。
- 为 NVIDIA 可信计算解决方案添加了基础架构代码。
- 为 IB 和 RoCE 控制消息添加了单独的流量类别，以启用高级 QoS（使用
NCCL_IB_FIFO_TC设置）。 - 增加了对分区的Broadcom PCI交换机子部分的PCI对等通信的支持。
总结
NCCL 2.22 版本引入了一些重要的功能和优化，旨在提高高性能计算（HPC）和 AI 应用的性能和效率。改进还包括新的调优插件界面、插件静态链接支持以及增强的组语义，以防止出现死锁。
有关更多信息，请参阅 Magnum IO 和 NCCL。在 GPU-Accelerated Libraries 论坛上提供反馈。
