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

发表时间: 2024-09

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

文章主图
文章主图
速读

一句话结论 NCCL 2.22 版本通过引入延迟连接建立、节点内拓扑融合以及全新的成本估算 API，大幅降低了 GPU 显存开销与通信器初始化时间，并为计算与通信的重叠优化提供了量化依据。

要解决什么问题 原有做法在三个方面存在卡点。首先是显存浪费严重：NCCL 的高效数据传输协议依赖静态分配的连接和缓冲区，在 2.22 版本之前，初始化阶段会为所有可能的算法和协议组合预先建立连接，每套组合都要占用数兆字节的 GPU 显存，而实际任务往往只用到其中一种，导致大量显存被永远不会使用的组合白白占用。其次是初始化耗时过长：随着 Hopper 架构 GPU 上 NVLink 互连数量的增加，每个 rank（参与通信的独立进程）在拓扑发现阶段都要独立探测所有链路，这种冗余查询在单节点内引发了严重的拥塞，使得创建大量通信器的应用在初始化时面临巨大的时间开销。最后是计算与通信重叠难以完美调度：开发者希望在多平台上让计算和通信充分重叠以榨干硬件性能，但由于缺乏对 NCCL 内部操作耗时的准确预估，很难在代码中动态调整计算工作量来实现两者的负载均衡。此外，在多维并行训练中，多个通信器的阻塞式销毁操作极易引发死锁，且跨 InfiniBand 子网的大规模训练也缺乏原生支持。

怎么做的 针对上述卡点，NCCL 2.22 在核心机制上进行了多项重构。第一，引入延迟连接建立机制。系统不再于初始化时全量建连，而是等到特定算法首次被调用时，才为其建立对应的连接与缓冲区。这一策略直接绕开了预分配机制的显存浪费问题，确保显存只分配给实际执行的算法。第二，重构拓扑发现逻辑以加速初始化。方法是将多节点 NVLink 拓扑融合技术下放到单节点内部：每个 rank 现在只负责探测自身关联的 GPU 信息，随后通过节点内通信将局部信息与对等体合并，拼出完整的拓扑图，从而彻底消除了各 rank 独立扫描全局链路带来的冗余拥塞。第三，提供用于成本估算的新 API `ncclGroupSimulateEnd`。该接口不触发实际通信，而是基于 NCCL 内部模型计算给定集合通信操作的预估耗时 $T_{est}$，并将其写入用户传入的结构体中。开发者可利用该预估时间动态配置计算任务的迭代次数或数据块大小，其核心逻辑可抽象为：$$ \text{ComputeAmount} = f(T_{est}) $$ 第四，新增调优器插件接口 v3，NCCL 会向插件提供一个二维成本表，外部调优器若要强制选择特定的算法与协议组合，只需修改该表：$$ \text{CostTable}[\text{Algorithm}][\text{Protocol}] = 0 \lor \min(\text{CostTable}) $$ NCCL 随后会基于更新后的成本表选择最终配置。最后，在工程实现上增加了多项关键设计：为 `ncclCommDestroy` 和 `ncclCommAbort` 引入组操作语义，允许以分组形式一次性非阻塞地销毁多个通信器，从根本上避免了多维并行中的死锁；支持跨 InfiniBand 子网通信，自动交换 GID 信息并利用 FLID（转发表 LID，用于识别一组转发路由器）实现高性能自适应路由；同时新增了针对初始化阶段的细粒度耗时检测工具，方便开发者定位性能瓶颈。

效果如何 实验在单节点 DGX-H100 系统和单节点 8x H100 GPU 系统上进行，主要对比了开启新特性前后的 NCCL 自身性能。在显存节省方面，通过延迟连接建立机制，当系统仅运行 Ring 算法时，NCCL 的 GPU 显存使用量大幅减少了 3.5 倍；在仅运行基于 NVSwitch 的归约操作时，显存使用量也减少了 1.47 倍。在初始化加速方面，结合延迟连接建立与节点内拓扑融合，单节点 8x H100 系统上的 `ncclCommInitRank` 执行时间从约 6.7 秒骤降至约 0.7 秒，耗时缩减高达 90%，极大缓解了频繁创建通信器场景下的卡顿。不过作者也指出了新 API 的局限性：`ncclGroupSimulateEnd` 返回的耗时纯粹基于 NCCL 内部理论模型估算，与实际物理执行时间无法做到完全一致，且在当前版本中，该 API 仅能返回通信组中最后一个操作的预估时间，尚不支持复杂操作序列的全局精确预估。

A1 主要贡献

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

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 实验

实验环境

实验结果

A7 补充细节

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

A5 结论

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