Memory Efficiency, Faster Initialization, and Cost Estimation with NVIDIA Collective Communications Library 2.22

文章标题：通过 NVIDIA Collective Communications Library 2.22 实现内存效率、更快的初始化和成本估算
作者：Giuseppe Congiu
机构：NVIDIA

A1 主要贡献

NVIDIA Magnum IO NCCL 是一个为优化 GPU 间和多节点通信而设计的库，对于 AI 和 HPC 应用中的高效并行计算至关重要。NCCL 2.22 版本通过引入以下新特性，解决了用户的关键痛点：

• 延迟连接建立以节省 GPU 内存：该功能将连接的创建推迟到实际需要时，从而显著减少了 GPU 内存开销。
• 用于成本估算和工作负载平衡的新 API：提供了一个新的 API，帮助用户优化计算与通信的重叠，或研究 NCCL 的成本模型。
• ncclCommInitRank 的优化和检测工具：通过消除冗余的拓扑查询，为创建大量通信器的应用程序将初始化速度提高了高达 90%。
• 支持 IB Router 的多个子网：增加了对跨多个 InfiniBand 子网的作业进行通信的支持，使得深度学习训练作业能够在超过 40K 个端点的 InfiniBand 网络上运行。

A3 背景知识

NCCL 的连接与缓冲机制。NCCL 的高效数据传输协议（eager data transfer protocol）依赖于一组持久且静态分配的连接和缓冲区。NCCL 为其支持的每一种算法和协议组合都会创建一套独立的连接和缓冲区，而每一套都需要占用数兆字节的 GPU 内存。在此背景下，算法定义了在给定集合通信操作中数据在参与者之间的高级移动方式，而协议则定义了 NCCL 发送数据的具体方法。NCCL 会根据操作类型、消息大小、规模和拓扑结构来选择特定的算法和协议组合，以实现最佳性能。

A2 方法细节

延迟连接建立 (Lazy connection establishment)

解决内存浪费问题。在 2.22 版本之前，NCCL 会在初始化阶段为所有可能的算法和协议组合预先建立对等体之间的连接。这种方式可能导致数兆字节的 GPU 内存被浪费在那些永远不会被使用的算法和协议上。

实现延迟连接。现在，NCCL 会等到某个特定算法首次被需要时，才为其建立连接。这种“延迟”建立机制能够大幅降低 NCCL 的内存开销，尤其是在 NCCL 使用范围较窄的场景中。例如，如果一个应用在特定系统上只反复运行相同消息大小的 ncclAllReduce 操作，那么它应该只需要使用一种算法。

控制与效果。该功能默认启用，但可以通过设置环境变量 NCCL_RUNTIME_CONNECT=0 来禁用。在一个单节点 DGX-H100 系统上，当仅使用 Ring 算法时，该功能使 NCCL 的 GPU 内存使用量减少了 3.5 倍；当仅运行基于 NVSwitch 的归约操作时，内存使用量减少了 1.47 倍。

新的成本模型 API (New cost model API)

优化计算与通信重叠的需求。应用程序开发者希望充分利用 NVIDIA 系统上可用的计算、内存和带宽资源。理想情况下，计算和通信应该完美重叠，两者都有充足的工作负载，从而将硬件性能发挥到极致。然而，在运行大规模 HPC 和 AI 应用时，尤其是在多平台上运行同一代码库时，实现这一点非常困难。

ncclGroupSimulateEnd API 的引入。为了解决这个问题，NCCL 增加了一个名为 ncclGroupSimulateEnd 的新 API，它能够让用户了解 NCCL 预估的给定操作所需的时间。该 API 的使用方式与 ncclGroupEnd 相同，便于熟悉 NCCL 编程的用户上手。

API 工作原理与使用。与 ncclGroupEnd 不同，ncclGroupSimulateEnd 不会启动任何通信操作。相反，NCCL 会计算它认为操作将花费的时间，并将此估算值设置在用户提供的 ncclSimInfo_t 结构中。用户可以利用这个预估时间来调整计算量，以更好地平衡计算与通信。

ncclGroupStart();
ncclAllReduce();
ncclGroupSimulateEnd(sim_t);
printf("Estimated completion time=%f microseconds
", sim.time);
configureComputeAmount(sim.time, &computeIters, &workItemSize);

API 的局限性。需要注意的是，此 API 返回的值是基于 NCCL 内部模型的估算，并不与实际执行时间完全一致。截至 NCCL 2.22 版本，该 API 仅返回组中最后一个操作的预估时间。

初始化优化和检测工具 (Initialization optimizations and instrumentation)

降低初始化开销的必要性。随着客户工作负载的规模和多样性不断增加，降低 NCCL 的初始化开销已成为 NCCL 团队日益重要的优先事项。即使是单节点作业，NVIDIA Hopper GPU 上增加的 NVLink 互连数量也导致了初始化时间的显著增加，因为每个互连都必须被单独发现和连接。

引入初始化阶段的检测工具。为了改进初始化时间，团队首先需要研究每个初始化步骤的开销。他们对 ncclCommInitRank 内部的每个阶段进行了检测，并研究了在不同规模下各阶段的耗时。现在，当用户收集标准的 NCCL 日志时（NCCL_DEBUG=INFO），可以看到这些计时信息。此外，还新增了一个 NCCL_PROFILE 调试子系统（NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=PROFILE），它只提供检测信息，适合那些不关心其他初始化日志的用户。

优化点一：延迟连接建立。之前讨论的延迟连接建立不仅节省了内存，也同时减少了初始化时间，是一个重要的优化方向。

优化点二：拓扑发现。拓扑发现是初始化的一个步骤，其中每个 NCCL rank 确定节点上可用的硬件资源，包括系统中的 GPU 和 NIC、NVLink 互连数量，以及 PCI 拓扑和 NUMA 亲和性。研究发现，NCCL 执行 NVLink 发现的方式是次优的，因为每个 rank 都会独立发现所有链路，导致了冗余和拥塞。为了解决这个问题，团队重用了最初在 NCCL 2.21 中为多节点 NVLink (MNNVL) 支持而引入的拓扑融合代码。该代码原本用于在引导过程中通过节点间通信合并每个节点上的部分信息，从而获得完整的 NVLink 拓扑图。在 2.22 版本中，该功能被扩展到在每个节点内部使用。现在，每个 rank 只发现其自身 GPU 的信息，然后通过节点内拓扑融合与对等体合并这些结果。

优化的综合效果。通过结合延迟连接建立和节点内拓扑融合，在一个单节点的 8x H100 GPU 系统上，ncclCommInitRank 的执行时间可以减少 90%（约 6 秒），从之前的约 6.7 秒缩短到约 0.7 秒。对于那些在执行期间创建大量通信器的应用程序来说，这可以极大地减少总初始化时间。

新的调优器插件接口 (New tuner plugin interface)

v3 接口的工作方式。通过新的调优器插件接口（v3），NCCL 为插件提供了一个针对每个集合操作的二维成本表。该表报告了对于每种算法和协议的组合，执行该操作所需的预估时间。

兼容性与选择机制。NCCL 会将表中与检测到的拓扑不兼容的条目设置为 -1，以告知外部调优器这些组合不受支持或不允许被覆盖。为了选择一个特定的组合，外部调优器需要将所需算法/协议组合对应的值更新为 0 或整个表中的最小值。在插件更新成本表之后，NCCL 会利用这个表来选择给定集合操作的最终配置。

静态插件链接 (Static plugin linking)

满足合作伙伴需求。NCCL 团队提供了一个插件模型，允许合作伙伴提供自己的调优或网络后端，以替代 NCCL 的内部模型和 InfiniBand 插件。一些合作伙伴希望将这些插件静态链接到他们的应用程序二进制文件中，以方便使用并避免加载错误版本的插件。

指定静态插件。如果应用程序已经静态链接了网络或调优器插件，可以通过将 NCCL_NET_PLUGIN 或 NCCL_TUNER_PLUGIN 设置为 STATIC_PLUGIN 来指定它。

终止或销毁的组语义 (Group semantics for abort or destroy)

解决多维并行中的死锁问题。在之前的版本中，ncclCommDestroy 和 ncclCommAbort 会阻塞调用线程直至操作完成。对于多维并行机器学习工作负载，一个进程可能管理多个 NCCL 通信器，并且最终必须使用这些 API 来拆除它们。这种阻塞行为可能导致死锁。

提供组操作语义。为了提供更好的用户体验并避免死锁，新版本为这些应用程序提供了组语义，允许它们以分组的方式一次性销毁多个通信器。

IB Router 支持

跨 InfiniBand 子网通信。通过此功能，NCCL 可以在由一个或多个路由器连接的不同 InfiniBand 子网之间进行操作。NCCL 能够自动检测到两个通信端点位于 InfiniBand 网络的不同子网上，并交换建立连接和通信所需的 GID 信息。

利用 FLID 提升性能。当在子网之间进行路由时，可以使用 FLID (Forwarding table LID) 来识别一组用于转发的路由器，从而实现子网间更高性能的自适应路由。NCCL 2.22 会自动检测 FLID 的存在，并在不同子网的端点之间建立连接时使用它。

A4 实验

实验环境

• 硬件配置:
  • 单节点 DGX-H100 系统
  • 单节点 8x H100 GPU 系统

实验结果

• 延迟连接建立的内存节省：

  • 在单节点 DGX-H100 上，当仅使用 Ring 算法时，NCCL 的 GPU 内存使用量减少了 3.5 倍。
  • 当仅运行基于 NVSwitch 的归约操作时，内存使用量减少了 1.47 倍。

• 初始化时间优化：

  • 在单节点 8x H100 GPU 系统上，结合延迟连接建立和节点内拓扑融合，ncclCommInitRank 的执行时间减少了 90%，从约 6.7 秒降至约 0.7 秒。

A7 补充细节

NCCL 2.22 版本还提供了以下错误修复和次要功能更新：

• 在 DGX Google Cloud 上增加了对 allreduce 树算法的支持。
• 在 IB 异步错误日志中记录了 NIC 的名称。
• 修复了聚合集合操作的性能问题。
• 修复了已注册发送和接收操作的性能问题。
• 为 NVIDIA Trusted Computing Solutions 添加了基础设施代码。
• 为 IB 和 RoCE 控制消息添加了独立的流量类别，以支持高级 QoS 服务质量（通过 NCCL_IB_FIFO_TC 设置）。
• 增加了对跨分区 Broadcom PCI 交换机子部分进行 PCI 点对点通信的支持。

A5 结论

NCCL 2.22 版本引入了多项重要的功能和优化，旨在提高高性能计算（HPC）和 AI 应用程序的性能与效率。改进还包括一个新的调优器插件接口、对插件静态链接的支持，以及增强的组语义以防止死锁。如需更多信息，请参阅 Magnum IO 和 NCCL 的相关文档，并通过 GPU-Accelerated Libraries 论坛提供反馈。