Fusing Communication and Compute with New Device API and Copy Engine Collectives in NVIDIA NCCL 2.28

标题：在 NVIDIA NCCL 2.28 中利用新的设备 API 和复制引擎集合操作融合通信与计算
作者：Sylvain Jeaugey

主要贡献

NVIDIA Collective Communications Library (NCCL) 的最新版本 2.28 引入了通信与计算的突破性融合，旨在跨多 GPU 和多节点系统实现更高的吞吐量、更低的延迟以及最大化的 GPU 利用率。

NCCL 2.28 的核心创新在于GPU 发起的网络通信（GPU-initiated networking）、用于通信-计算融合的设备 API、基于复制引擎（copy-engine）的集合操作，以及为开发者构建高效、可扩展的分布式应用而设计的新 API。除了性能创新，该版本还通过扩展 API、改进工具和更清晰的集成路径来增强开发者体验，使开发者能够编写自定义通信内核，并以更大的灵活性和效率扩展其应用程序。

发布亮点：
- 设备 API (Device API)：支持开发用于通信/计算融合的自定义设备内核，包括 GPU 发起的网络通信，使内核能直接执行网络操作。
- 基于复制引擎 (CE) 的集合操作：开发者可利用复制引擎（CEs）来驱动 NVIDIA NVLink 传输，从而减少流式多处理器（SM）的计算资源争用。
- NCCL Inspector：提供了一个低开销的性能分析插件，能够对 NCCL 通信模式进行持续的可观测性分析。

方法细节

NCCL 设备 API 如何实现直接内核通信

NCCL 2.28 引入了一个设备侧通信 API，用于在 NVIDIA CUDA 内核内进行直接通信。以往所有的 NCCL 操作都是由主机发起的，这会引入同步开销。通过新的 API，内核可以直接发起数据移动，将通信与计算操作集成在一起，从而带来更高的吞吐量和更低的开销。使用新 API 的前提是，必须使用带有对称内存窗口的数据缓冲区。

目前支持三种操作模式：
1. 加载/存储可访问 (Load/Store Accessible, LSA)：用于可通过内存加载/存储操作访问的设备之间的通信，使用 CUDA P2P 技术。
2. 多内存 (Multimem)：用于利用 NVLink SHARP 提供的硬件多播功能在设备之间进行通信。
3. GPU 发起的网络通信 (GPU Initiated Networking, GIN)：用于由 GPU 发起、通过网络在设备之间进行的通信。

图 1. 设备 API 支持的三种内存通信操作模式及其底层通信机制

值得特别提及的是在 NCCL 2.28.7 中新引入的 GPU 发起的网络通信（GIN）功能。它使 GPU 能够无需 CPU 干预即可管理自身的网络操作。内核可以直接将数据传输、信号和同步步骤加入队列，从而消除了由主机驱动的控制路径所引起的瓶颈。

利用复制引擎卸载加速 NCCL 性能

SM 资源竞争问题。要在规模上实现顶级的通信性能，需要更多的流式多处理器（SMs）来饱和 NVLink 带宽。这种为通信任务增加的资源分配会与计算内核产生资源竞争，可能降低整体应用程序的性能。

CE 集合操作卸载通信任务。CE 集合操作将 NVLink 域内的通信任务从 SM 卸载到专用的硬件复制引擎（CEs）上。与传统的 NCCL 集合操作相比，基于 CE 的集合操作实现了零 SM 占用（zero-SM operation）。这种方法适用于仅需数据移动的集合操作，例如 AlltoAll 和 AllGather。这为计算工作负载释放了 SM 资源，并改善了通信与计算的重叠。现在，两者可以并发执行，而无需争夺相同的硬件资源。

CE 集合操作的优化。基于 CE 的集合操作采用了多种优化措施以增强性能。例如，它们利用批处理 API（如 cudaMemcpyBatchAsync）将多个 CE 操作组合成单次调用，从而减少 CUDA 驱动程序的开销。此外，它们还使用 NVLink 多播优化来高效地广播同步信号。

性能优势。这些集合操作在不占用 SM 资源的情况下，实现了与基于 SM 的集合操作相当的性能。下图显示，基于 CE 的 AllGather 实现了比基于 SM 的 AllGather 更高的峰值带宽。这一性能优势部分源于 CE 发起的 NVLink 事务比 SM 发起的事务使用了更宽的事务宽度。

图 2. 基于 CE 和基于 SM 实现的 NCCL AllGather 性能对比

NCCL Inspector 使性能分析和可观测性变得简单

NCCL Inspector 简介。NCCL Inspector 是一个可观测性、性能分析和分析插件，提供详细的、针对每个通信器和每个集合操作的性能和元数据日志记录。该插件应在持续运行模式下使用。它旨在通过为每个操作生成结构化的 JSON 输出，帮助用户分析和调试 NCCL 集合操作，从而在分布式工作负载运行期间洞察通信模式和性能特征。

核心特性。NCCL Inspector 使用 NCCL Profiler 插件接口架构集成到 NCCL 用例中。其主要特性包括：
- 按通信器跟踪：Inspector 单独跟踪每个 NCCL 通信器，使用户能够分析不同通信上下文中的性能模式。这在复杂的分布式应用中尤其有价值，因为其中可能使用多个通信器用于不同目的。
- 持续运行（Always-on）：该插件的低开销意味着 Inspector 可以在生产工作负载中使用，以提供对 NCCL 的可观测性而不会降低性能。
- 事件跟踪：插件捕获详细的事件跟踪，包括集合操作的开始/停止事件、内核通道操作和计时信息。
- 性能指标：NCCL Inspector 计算并报告关键性能指标，包括算法带宽、总线带宽、执行时间、计时源信息（GPU vs CPU 计时）、消息大小和集合操作类型。

与其他插件的协同。该插件旨在与其他 NCCL 插件协同工作，并可为使用性能反馈来优化通信模式的调优器插件提供有价值的数据。虽然当前设计未在 NCCL Inspector 和调优器插件之间提供直接的共享上下文，但 Inspector 生成的详细性能数据可被外部分析工具用于为调优决策提供信息。

图 3. 使用 Elastic 和 Kibana 仪表板可视化 NCCL Inspector 数据

网络独立性。该插件的运行独立于具体的网络实现，使其与 NCCL 支持的各种网络技术兼容。这确保了无论使用何种底层网络基础设施，Inspector 都能提供洞察。

NCCL 2.28 改进的开发者体验

NCCL 2.28 的改进超越了核心的通信与计算融合，提供了一系列增强功能，以提高灵活性、性能调优和开发者体验。新的 API、内核编排、配置、性能分析和构建系统使开发者能够更好地控制通信工作流，同时在多样化的硬件和环境中获得改进的可观测性和简化的部署。

针对 AllToAll、Gather 和 Scatter 的新主机 API
- 引入了针对 AlltoAll、Gather 和 Scatter 操作的原生主机级 API，这使得 NCCL 能够应用高级优化。例如，NCCL 可以使用复制引擎来减少这些通信模式的 SM 使用量，这也是此版本中引入的一项优化。这些性能增强只有通过专用的原生 API 才可能实现。

对称内核组调用支持
- 单个对称内核在 NVLINK 域中使用时已能提供出色的延迟、带宽和资源利用率。NCCL 2.28 增加了对分组对称内核的支持，以提高性能和资源利用率。用户可以照常注册窗口缓冲区，并在 ncclGroupStart() 和 ncclGroupEnd() 之间调用多个 NCCL 集合操作。在启动期间，NCCL 会自动检测可以利用对称内核的潜在集合操作，并将它们分组调度到单个内核中，以实现更高效率。

通过 NCCL 环境插件实现灵活的配置管理
- 优化通信参数对性能至关重要，NCCL 通过配置文件和环境变量为此提供了强大的机制。然而，当不同作业需要唯一的 NCCL 版本和配置时，当前基于文件的方法需要手动进行版本匹配，这与数据库或云调度器等现代部署系统集成不佳。
- 新的 NCCL 环境插件 API 提供了一种灵活的、程序化的替代方案，具有以下关键优势：
- 程序化版本匹配：为每个 NCCL 版本自动应用正确的配置。
- 存储无关的配置：可从文件、数据库或环境变量中使用设置。
- 增强的灵活性和控制：允许以编程方式覆盖某些参数，同时保留其他参数以进行精细控制。
- 初始化和资源管理：一次性启动并在整个运行时保持活动状态，并能干净地释放资源。
- 面向未来：为新的按通信器配置做好了准备。

• 启用后，该插件会与 NCCL 参数子系统集成，并覆盖优先级较低的配置机制，确保一致的、版本感知的设置。它简化了大规模、多环境的部署，将用户从静态文件系统的限制中解放出来。

插件的共享上下文
- AI 模型训练越来越多地跨越多个数据中心，这给通信库带来了重大挑战。一个重新设计的、支持共享上下文和按通信器调优的插件系统取代了全局初始化，并能在多样化的网络环境中实现上下文感知的优化。
- 以下网络插件 API 已更新：
- ncclResult_t (*init)(void** ctx, uint64_t commId, ncclNetCommConfig_v11_t* config, ...)
- ncclResult_t (*finalize)(void* ctx);
- ncclResult_t (*listen)(void* ctx, ...);
- ncclResult_t (*connect)(void* ctx, int dev, void* handle, ...);

• 现在，每个通信器都使用一个 commId 和网络配置进行初始化，并返回一个用于隔离和精细调优的按通信器上下文句柄（ctx）。该机制已在调优器和性能分析器插件中使用，现已扩展到网络插件。开发者还可以将网络、调优器和性能分析器插件的代码放在一个 .so 库中，通过 NCCL_NET_PLUGIN 变量启用共享上下文和插件间通信。

NCCL Profiler API 事件
- 以前，NCCL Profiler 插件接口无法捕获 NCCL API 调用或 CUDA 内核启动事件，因此无法将启动、集合操作或点对点事件与相应的 NCCL API 调用关联起来。
- Profiler 现在能够：
- 按用户顺序显示 API 事件：即使操作被 ncclGroupStart/ncclGroupEnd 分组（这可能会改变调度顺序），也能确保正确的排序。
- 测量 NCCL 操作在 CPU 上的开销：即用户调用 API 后，NCCL 调度操作所花费的时间。
- 将 CUDA 内核启动与源头的 NCCL API 调用直接关联。
- 将 NCCL 调度的集合和点对点任务与其原始 API 调用链接起来：API 事件在图启动（graph launches）之间保持持久，允许每次图启动的底层任务与用户的原始 API 调用相关联。

基于 CMake 的构建系统

• NCCL 现在也支持使用 CMake 进行 Linux 构建，为 Make 提供了现代化的替代方案。CMake 系统简化了与更大型构建管道和跨平台项目的集成，同时保持了与旧系统的兼容性。通过标准化的 CMake 支持，项目现在可以利用在更广泛的 CUDA 和 HPC 生态系统中使用的熟悉工作流程，从而提高可复现性并减少维护开销。其结果是获得了一个更灵活、模块化和对开发者更友好的构建体验。

实验环境

本文档未提供详细的实验环境配置。从图表的标题中可以提取到部分硬件信息：
- 硬件配置：实验在一台包含 8 个 NVIDIA Blackwell Umbriel GPU 的单节点上进行。
- 其他关于数据集、模型架构、软件配置（除 NCCL 2.28 外）的详细信息均未在文中提及。

实验结果

CE 与 SM 实现的 AllGather 性能对比
- 实验内容：在一台配备 8 个 Blackwell GPU 的节点上，比较了基于复制引擎（CE）和基于流式多处理器（SM）的 AllGather 集合操作的带宽性能。
- 实验结果：如图 2 所示，基于 CE 的 AllGather（特别是 CE Multicast 版本）在处理较大消息时，其峰值带宽显著高于基于 SM 的实现。在消息大小为 4GB 时，CE Multicast 的带宽比 SM Symmetric（NCCL 优化版本）高出约 1.25 倍。
- 结论分析：这一性能优势表明，将通信任务从 SM 卸载到专用的复制引擎，可以有效避免计算资源争用，并利用更宽的 NVLink 事务宽度来提升通信效率。

NCCL Inspector 可视化展示
- 实验内容：展示了 NCCL Inspector 插件收集的数据，并通过 Elastic 和 Kibana 仪表板进行可视化。
- 实验结果：图 3 展示了一个针对 HCA-Only ReduceScatter 操作的集合总线带宽（Collective Bus Bw）随集合操作序列号（Collective Seq No）变化的图表。图中绘制了不同消息大小（如 8.39MB、8.40MB、5.85MB）下的带宽曲线。
- 结论分析：该图表证明了 NCCL Inspector 能够提供持续、低开销的性能监控，并能生成详细的、可供分析和可视化的数据，帮助开发者洞察其分布式工作负载中的通信模式和性能瓶颈。

结论

NCCL 2.28 通过引入一系列新功能，显著提升了分布式训练的可扩展性、优化了集合操作性能并提高了跨节点效率。开发者可以通过访问 NVIDIA/nccl GitHub 仓库获取详细文档、源代码和示例，以利用这些新特性来突破其分布式应用的性能界限。同时，官方文档也提供了关于如何根据特定系统架构调优 NCCL 的指导。

附录

本文档不包含附录部分。