S5111: Scaling Deep Learning Training: Fast Inter-GPU Communication with NCCL

Sylvain Jeaugey | GTC 2023

NCCL：多GPU计算的关键通信库

NCCL (NVIDIA Collective Communications Library) 是专为多GPU计算设计的关键通信库，针对从桌面到DGX Superpod的所有平台进行了优化。它支持PCI、NVLink以及NVSwitch + Network等多种互连方式，以实现高效的GPU间通信。

下载链接: https://developer.nvidia.com/nccl，可用于NGC容器。
源代码: https://github.com/nvidia/nccl。

NCCL：多GPU计算的关键通信库
深度学习：从数据到AI
深度学习训练
NCCL 概述
NCCL API
集合通信带宽
点对点通信带宽
集合操作 (Collective Operations
拓扑检测 (Topology Detection
- 将算法映射到硬件
- 并行性策略示例
节点间通信 (Inter-node Communication
- 轨优化设计 (Rail-optimized design
点对点通信 (Point-to-point Communication
新特性与未来规划 (New and Future
总结 (Summary

深度学习：从数据到AI

深度学习流程涉及将海量数据转化为可部署的AI模型：

数据阶段: PB级数据。
训练阶段: 利用8到数千个GPU进行训练，需要高功率、大内存。训练通常耗时1小时到1周。NCCL旨在通过高效的多GPU训练来缩短训练时间。
模型部署: 训练好的模型部署到嵌入式设备或数据中心，用于推理，需要低功率、低内存。推理时间通常为1毫秒到1秒。

深度学习训练

单GPU训练

在单GPU上进行深度学习训练时，数据经过GPU处理，计算梯度并更新模型参数以最小化误差。

数据并行

数据并行是深度学习训练中最常见的并行方法。在这种模式下，数据被分成多个批次，每个GPU处理一个批次的数据并计算局部梯度。然后，NCCL AllReduce操作将这些局部梯度聚合为全局梯度，用于更新全局参数。

特点: 随着GPU数量的增加，需要增加批次大小。
挑战: 批次大小过大可能会导致准确性下降。

流水线并行

流水线并行允许训练无法单独放入单个GPU内存中的大型模型。模型被分割成多个部分，每个部分由不同的GPU处理。NCCL Send/Receive用于在模型层之间传递数据，而NCCL AllReduce用于同步全局梯度。

优势: 支持更大模型的训练。
挑战: 发送/接收流水线中可能存在“气泡”，对于小批次大小效率较低。

张量并行

张量并行允许将单个张量操作（例如矩阵乘法）分布到多个GPU上。在图示中，矩阵B被分成B0、B1、B2、B3四个部分，每个GPU处理一部分计算，然后NCCL Allgather将结果C聚合起来。

Megatron / GPT 训练

目标是训练一个包含1万亿参数的GPT模型。为了实现这一目标，采用了多种并行策略的组合：

在6个节点上进行流水线并行。
在每个8个GPU的节点上进行张量并行。
在64组6个节点上进行数据并行。

通过这种组合方式，实现了502 Petaflops的性能，使用了3072个GPU。

NCCL 概述

NCCL 提供了GPU间通信的原语，并针对所有平台进行了优化。它作为一个库实现，在GPU上启动CUDA核函数，并通过CUDA流语义与计算核函数交错执行。

NCCL 的核心功能包括：

初始化:
- CPU端通信。
- 基于Socket的带外引导。
- 容错处理。 (GTC '21)
硬件拓扑检测:
- 构建包含GPU、CPU、NIC、PCI交换机、NVLinks和NVSwitches的拓扑图。
- 通过图搜索发现适合集合通信算法的拓扑结构。 (GTC '20)
集合通信算法:
- 支持环形(Ring)、树形(Tree)、CollNet (链式/直接) 等算法。
- 支持NVLink SHARP。 (GTC '22)
点对点通信:
- 发送/接收语义。
- 通信调度。
- 聚合。 (GTC '22)

NCCL API

NCCL 提供了一系列丰富的API，用于管理通信上下文、执行集合操作和点对点通信等。

初始化相关API:
- ncclGetUniqueId、ncclCommInitRank、ncclCommFinalize 等，用于通信环境的设置和销毁。
- 支持版本从2.0到2.18 (Planned)。
集合操作API (自2.0版本):
- ncclReduce: 归约操作。
- ncclBroadcast: 广播操作。
- ncclAllReduce: 全归约操作。
- ncclReduceScatter: 归约散射操作。
- ncclAllGather: 全收集操作。
点对点通信API (自2.0版本):
- ncclSend: 发送数据。
- ncclRecv: 接收数据。
融合操作API (自2.0版本):
- ncclGroupStart 和 ncclGroupEnd: 用于将多个NCCL操作融合到一个组中执行。
自定义缩放因子API (自2.11版本):
- ncclRedOpCreatePreMulSum 和 ncclRedOpDestroy: 创建和销毁自定义归约操作。
其他实用API:
- ncclGetVersion、ncclGetLastError 等，用于获取版本信息和错误诊断。

集合通信带宽

NCCL在不同互连技术和配置下的Allreduce总线带宽（GB/s）表现如下：

技术/配置	多GPU (GB/s)	多节点 (GB/s)
PCI Gen4	24	10 (100GbE, TCP/IP Sockets)
PCI Gen5	48	12 (100Gb RDMA (IB/RoCE))
NVLink2 (V100)	122	82 (8x 100Gb RDMA (DGX-2))
NVLink3 (A100)	232	192 (8x 200Gb RDMA (DGX A100))
NVLink4 (H100)	360	330 (8x 400Gb RDMA (DGX H100))
NVLink4 SHARP (H100)	480	480 (8x 400Gb RDMA SHARP (DGX H100))

点对点通信带宽

NCCL在不同互连技术和配置下的Alltoall总线带宽（GB/s）表现如下：

技术/配置	多GPU (GB/s)	多节点 (GB/s)
PCI Gen4, 2x4 GPUs	6	1.25 (100Gb RDMA (IB/RoCE))
PCI Gen5, 2x4 GPUs	12	2.5 (200Gb RDMA (IB/RoCE))
NVLink2 (V100)	122	6 (8x 100Gb RDMA (DGX-2))
NVLink3 (A100)	230	24 (8x 200Gb RDMA (DGX A100))
NVLink4 (H100)	340	48 (8x 400Gb RDMA (DGX H100))

集合操作 (Collective Operations)

环形算法 (Ring Algorithm)

环形算法用于 allreduce 操作，其特点是：

算法简单，适用于所有拓扑，计算均衡。
但延迟随 GPU 数量线性增加，在大规模场景下性能会迅速下降。

在 DGX 上实现多环形算法时，会构建复杂的互联结构以支持多个节点间的环形通信。

树形算法 (Tree Algorithm)

树形算法也用于 allreduce 操作，其特点是：
* 对数级别的延迟扩展优于环形算法。
* 但计算不均衡通常会导致峰值带宽达不到最佳状态。

在 DGX 上实现多树形算法时，同样会构建复杂的互联结构以支持多个节点间的树形通信。

Collnet 算法 (链式版本)

Collnet 算法的链式版本具有以下特点：
* 显著减少网络流量，如果网络是瓶颈，可提高峰值带宽。
* 在大规模部署时，延迟几乎恒定，因此性能稳定。

在 DGX 上，Collnet/链式算法通过多节点间链条实现互联通信。

COLLNET 算法 (直接版本)

Collnet 算法的直接版本具有以下特点：
* 当每个 GPU 都有专用的 NIC 时，效果更好。
* 需要 nvswitch 或 alltoall 连接。
* 比链式版本具有更好的节点内延迟，尤其是在使用 CUDA 图注册缓冲区时。
* alltoall 阶段可能无法达到最佳带宽。

H100 SHARP (NVLink SHARP)

H100 SHARP 利用 NVLink SHARP 技术优化通信：
* 通过 NVLink 减少通信量，实现更高的带宽。

H100 SHARP NVLink SHARP示意图
H100 SHARP NVLink SHARP通信减少示意图

H100 SHARP (NVLink 和 IB)

H100 SHARP 结合了所有 SHARP 技术，包括 NVLink 和 InfiniBand (IB) 进行网络内归约 (in-network reductions)，以达到新的性能水平。

拓扑检测 (Topology Detection)

将算法映射到硬件

拓扑检测是将算法映射到硬件的关键过程，主要分为三个阶段：

拓扑检测 (Topology detection)
- 构建包含所有 GPU、NIC、CPU、PCI 交换机、NVLink、NVSwitch 的图。
- 支持为虚拟机注入拓扑。
图搜索 (Graph search)
- 在节点内部为环形、树形和链形结构找到最佳路径集合。
- 为每种算法建立性能模型并进行调优。
CUDA 内核 (CUDA Kernels)
- 优化归约和复制操作，以最小化 SM (流多处理器) 使用。
- 使用 CPU 线程进行网络通信。

并行性策略示例

拓扑检测支持多种并行性策略，例如：

纯数据并行 (Pure data parallelism)：在一个全局通信器中，每个节点包含 8 个 GPU。
在 4 个 GPU 上实现管道并行 (Pipeline parallelism on 4 GPUs)：在 4 个 NCCL 通信器中，每个通信器包含 2 个 GPU。
在 8 个 GPU 上实现管道并行 (Pipeline parallelism on 8 GPUs)：在 8 个 NCCL 通信器中，每个通信器包含 1 个 GPU。

节点间通信 (Inter-node Communication)

轨优化设计 (Rail-optimized design)

节点间通信的轨优化设计与经典架构相比，具有显著优势：
* 经典结构设计 (Classic fabric design)：路由必须完美以确保所有流使用不同的链路。
* 轨优化设计 (Rail-optimized design)：所有流量都局限于叶交换机。路由冲突是不可能发生的。

点对点通信 (Point-to-point Communication)

点对点语义：分组和 Alltoall 示例 (Page 32)

NCCL 定义了两个用于点对点通信的函数：ncclSend 和 ncclRecv。这些函数通过 ncclGroupStart/ncclGroupEnd 分组，可以轻松编写任何通信模式，包括 alltoall(v,w)、scatter(v)、gather(v)、邻居集体操作等。

Alltoall 示例：

ncclGroupStart();
for (int i=0; i<nranks; i++) {
  ncclSend(sendbuffs[i], count, type, i, comm);
  ncclRecv(recvbuffs[i], count, type, i, comm);
}
ncclGroupEnd();

组调用中的操作顺序无关紧要，所有操作将同时进行且无死锁。
如果不进行分组，ncclSend 和 ncclRecv 操作可能会阻塞，直到远程 GPU 上的匹配操作执行完毕。

点对点通信：旋转环 (Page 33)

点对点通信调度采用环形原则：同时从 rank - delta 接收数据，并发送到 rank + delta。这保证了无死锁操作。

为确保双向使用 NIC，我们交替使用 +delta 和 -delta。
为确保 NIC 和 NVLink 使用的重叠，我们交替使用圆的两端（delta = 0 和 delta = N/2）。
最后，为了更好的延迟和重叠，我们并行执行 X 步，X 最大为 256。

Alltoall 通信：PXN 优化 (Page 34)

Alltoall 通信的大多数传输通过顶级交换机以轨优化设计进行。

PXN 还允许将所有消息聚合到同一目的地。
PXN (PCI x NVLink) 允许 GPU 使用中间 GPU 向远程 NIC 发送数据。

新特性与未来规划 (New and Future)

容错 (Fault tolerance) (Page 36)

在 NCCL 2.14 之前，NCCL 提供了 ncclCommAbort 函数来停止在 GPU 上运行的核函数。
NCCL 2.14 扩展了 ncclCommAbort 的范围到 CPU 调用，这些调用也可能阻塞。

容错流程对比：

旧流程 (Init -> Communication -> Destroy)	新流程 (Init -> Poll -> Communication -> Poll -> Finalize -> Poll -> Destroy)
Init：如果缺少 rank 可能阻塞	`config.blocking = 0; ncclCommInitRankConfig(..., config);`
Communication：由于动态连接建立可能阻塞
Destroy：可能需要与其他 rank 同步	Call `ncclCommGetAsyncError` until the return code is not `ncclInProgress`. Call `ncclCommAbort` if needed.
	`ncclCommFinalize(comm);`
	Destroy：只释放资源

未来展望 (Future) (Page 38)

更多 H100 支持
- NVLink SHARP + IB SHARP
- NVLink SHARP + Ring (用于非 IB 网络)
通信器拆分
- 创建子通信器，可能共享资源。
CUDA 网络
- 添加设备代码以优化网络通信（CUDA 发起的通信）。

总结 (Summary) (Page 39)

NCCL 是用于多 GPU 计算的关键通信库。
已针对所有平台进行优化，从桌面到 DGX Superpod。
可从 https://developer.nvidia.com/nccl 和 NGC 容器下载。
源代码位于 https://github.com/nvidia/nccl。

PaperCache

S5111: Scaling Deep Learning Training: Fast Inter-GPU Communication with NCCL

S5111: Scaling Deep Learning Training: Fast Inter-GPU Communication with NCCL

NCCL：多GPU计算的关键通信库

目录

深度学习：从数据到AI