NVIDIA 集合通信库（NCCL）的最新版本引入了突破性的通信与计算融合技术，可显著提升多 GPU 和多节点系统中的通信吞吐量，降低延迟，并进一步提高 GPU 利用率。
NCCL 2.28 聚焦于 GPU 发起的网络通信、支持通信与计算融合的 Device API、基于拷贝引擎的集合通信，以及新 API 为开发者提供用于构建高效、可扩展的分布式应用。除了在性能方面的多项创新，该版本还通过扩展的 API 功能、优化的开发工具和更清晰的集成路径，显著提升了开发体验，使开发者能够更灵活、高效地编写自定义通信内核，并进一步拓展其应用能力。
版本亮点
支持在性能、监控、可靠性和服务质量方面的持续改进。
-
Device API：支持开发用于通信与计算融合的自定义设备内核（包括支持 GPU 发起网络操作），使内核能够直接执行网络任务。
-
基于拷贝引擎（CE）的集合通信：开发者可利用 CE 驱动 NVIDIA NVLink 传输，有效降低流多处理器（SM）在通信过程中的计算资源争用。
-
NCCL Inspector：提供低开销的分析插件，实现对 NCCL 通信模式的持续可观测性与深入分析。
NCCL Device API
如何实现内核间的直接通信
NCCL 2.28 引入了设备端通信 API，支持在 NVIDIA CUDA 内核中直接发起通信操作。此前，所有 NCCL 操作均由主机端启动，这会带来额外的同步开销。借助这一新 API，内核能够直接启动数据传输，实现通信与计算的深度融合，从而提升整体吞吐量并降低延迟。使用该 API 时，需使用对称内存（Symmetric Memory）窗口的数据缓冲区。
目前支持三种操作模式：
-
Load/Store访问（LSA）：通过 Load/Store 内存操作实现设备间通信，采用 CUDA P2P 技术。
-
Multimem：利用 NVLink SHARP 提供的硬件组播功能，实现设备间的高效通信。
-
GPU Initiated Networking（GIN）：支持 GPU 直接发起网络通信，实现设备之间的数据传输。
如需了解更多详情，请参阅在线文档：
https://docs.nvidia.cn/deeplearning/nccl/user-guide/docs/usage/deviceapi.html
特别值得关注的是，在 NCCL 2.28.7 版本新引入了 GPU-Initiated Networking（GIN），使 GPU 能够直接管理网络操作，无需依赖 CPU 干预。通过该机制，内核可直接对数据传输、信号通知和同步操作进行排队，有效消除传统主机驱动控制路径所带来的性能瓶颈。
通过拷贝引擎卸载来加速 NCCL 性能
要在大规模系统中实现出色的通信性能，需要分配更多的 SM 来充分使用 NVLink 带宽。然而，持续增加的通信任务分配会与计算核心争用资源，从而影响应用程序的整体性能。
CE 集合通信将 NVLink 域内的通信任务从 SM 卸载至专用的硬件 CE 单元。与传统的 NCCL 集合通信不同，基于 CE 的集合通信实现了零 SM 参与的通信操作。该方法适用于仅需数据移动的集合通信，例如 AlltoAll 和 AllGather，从而释放 SM 资源以专注于计算任务，并提升通信与计算的重叠效率。现在，通信与计算可并行执行，不再竞争相同的硬件资源。
基于 CE 的集合通信通过多项优化来提升性能。例如，利用批处理 API（如 cudaMemcpyBatchAsync）将多个 CE 操作合并为单次调用，有效降低 CUDA 驱动开销；同时采用 NVLink 组播机制优化同步信号的广播，提升传输效率。
这些集合通信无需消耗 SM 资源，即可达到与基于 SM 的集合通信相当的性能。下图显示，基于 CE 的 AllGather 实现了高于基于 SM 的 AllGather 的峰值带宽。这一性能优势的另一个原因是，相比由 SM 发起的事务，由 CE 发起的 NVLink 事务采用了更大的事务宽度。
在 NCCL 文档中了解如何启用 CE 集合通信：
https://docs.nvidia.cn/deeplearning/nccl/user-guide/docs/usage/bufferreg.html?highlight=copy%20engine#zero-cta-optimization
使用 NCCL Inspector 轻松实现分析与监控
NCCL Inspector 是一个用于观测、分析与诊断的插件，能够提供针对每个通信器和每个集合操作的详细性能数据及元数据记录。该插件适用于持续运行模式，旨在通过为每次操作生成结构化的 JSON 输出，帮助用户深入分析和调试 NCCL 集合操作，从而在执行分布式工作负载时更好地理解通信行为与性能特征。
NCCL Inspector 通过 NCCL Profiler 插件接口架构集成至 NCCL 应用场景中，主要功能包括：
-
逐个通信器追踪：Inspector 会独立追踪每个 NCCL 通信器，帮助用户分析不同通信上下文的性能特征。这在复杂的分布式应用中尤为关键，因为此类应用中多个通信器可能会用于不同的用途。
-
持续运行：由于插件具备低开销特性，可长期启用，适用于生产环境中的工作负载，在不影响性能的前提下为 NCCL 提供全面的可观察性。
-
事件追踪：插件能够捕获详细的事件记录，包括集合操作的启动与结束事件、内核通道操作以及相关的时间信息。
-
性能指标：NCCL Inspector 会计算并报告多项关键性能指标，如算法带宽、总线带宽、执行耗时、计时来源信息（GPU 或 CPU 计时）、消息大小及集合操作类型。
该插件旨在与其他 NCCL 插件协同工作，可为基于性能反馈优化通信特征的调优插件提供有价值的性能数据。尽管当前设计未实现 NCCL Inspector 与调优器插件之间的上下文直接共享，但 Inspector 生成的详细性能数据可通过外部分析工具进行处理，从而为调优决策提供有力支持。
该插件独立于特定网络运行，因此兼容 NCCL 支持的各类网络技术，确保无论采用哪种底层网络基础设施，Inspector 均可提供有效的洞察。
如需了解更多信息，请参阅 ext-profiler/inspector/ 目录下的 README.md 文件：
https://github.com/NVIDIA/nccl/tree/master/ext-profiler/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 窗口缓冲区注册文档：
https://docs.nvidia.cn/deeplearning/nccl/user-guide/docs/usage/bufferreg.html#window-registration
使用 NCCL 环境插件实现灵活的配置管理
优化通信参数对性能至关重要，NCCL 通过配置文件和环境变量为此提供了可靠的机制。然而，当不同作业需要特定的 NCCL 版本和配置时，目前依赖文件的手动版本匹配方法难以与数据库或云调度系统等现代部署架构有效集成。
新的 NCCL 环境插件 API 提供了一种灵活的编程替代方案，具备以下显著优势：
-
编程版本匹配可为每个 NCCL 版本自动应用正确的配置。
-
与存储无关的配置从文件、数据库或环境变量中获取设置。
-
更精细的控制与更高的灵活性通过以编程方式覆盖部分参数，同时保留其他参数不变。
-
初始化和资源管理仅需启动一次，即可在运行时持续生效，并在释放资源时彻底清除。
-
设计具备前瞻性，支持全新的逐个通信器独立配置。
启用后，插件将与 NCCL 参数子系统集成，覆盖低优先级的配置机制，确保版本感知配置的一致性。该功能简化了大规模、多环境的部署流程，帮助用户突破静态文件系统带来的限制。
插件的共享上下文
随着AI模型训练日益扩展至多个数据中心，通信库面临严峻挑战。重新设计的插件系统支持共享上下文和针对每个通信方的个性化调优，可替代全局初始化机制，并实现跨不同网络环境的上下文感知优化。
以下网络插件 API 已更新：
// Network plugin v11ncclResult_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），以实现隔离和精细化调优。该机制已被应用于调优器和分析器插件，目前可进一步扩展至网络插件。开发者还可通过设置 NCCL_NET_PLUGIN 环境变量，将网络、调优器和分析器插件的代码整合到同一个 .so 库中，从而实现上下文共享与插件间的通信。
NCCL 分析器 API 事件
过去，NCCL 分析器插件接口未能捕获 NCCL API 调用或 CUDA 内核启动事件，因此无法将启动、集合操作或点对点通信事件与对应的 NCCL API 调用关联起来。
分析器现在支持：
-
按用户调用顺序显示 API 事件，以确保操作顺序的正确性，即使在使用 ncclGroupStart/ncclGroupEnd 进行操作分组的情况下也能保持正确的排序，避免因调度顺序变化而影响结果。
-
准确测量 NCCL 操作在 CPU 上的开销，即从用户发起 API 调用到 NCCL 实际调度操作所耗费的时间。
-
将启动的 CUDA 内核与对应的原始 NCCL API 调用直接关联。
-
将 NCCL 调度的集合通信任务和点对点通信任务与其原始 API 调用进行关联。由于 API 事件在图执行启动期间持续存在，因此每次图启动时产生的底层任务均可追溯至用户的原始 API 调用，确保执行过程的可追踪性与一致性。
基于 CMake 的构建系统
NCCL 现已支持适用于 Linux 平台的 CMake，与传统的 Make 方式相比提供了更现代的替代方案。CMake 的引入有助于简化与大型构建流程及跨平台项目的集成，同时保持对传统系统的兼容性。通过标准化的 CMake 支持，项目能够采用 CUDA 和 HPC 生态系统中广泛使用的通用工作流，提升构建的可重现性，并降低维护成本。这将带来更灵活、模块化且更便于开发者使用的构建体验。
开始使用 NCCL 2.28
了解 NCCL 2.28 的新特性，借助增强的可扩展性、优化的集合通信性能以及更高的跨节点效率，进一步突破分布式训练的性能瓶颈。
有关详细文档、源代码及社区支持，请访问 NVIDIA/nccl GitHub 仓库并参考其中提供的示例。若需了解如何针对您的系统架构优化 NCCL，可查阅 NCCL 官方文档：
https://docs.nvidia.cn/deeplearning/nccl/user-guide/docs/index.html
