FAST INTER-GPU COMMUNICATION WITH NCCL FOR DEEP LEARNING TRAINING, AND MORE

SYLVAIN JEAUGEY, DAVID ADDISON

目录

• 多GPU计算 (MULTI-GPU COMPUTING

• 深度学习训练 (DEEP LEARNING TRAINING

  • 从数据到AI (From data to AI
  • 单GPU上的深度学习训练 (On a single GPU
  • 多GPU上的深度学习训练 (On multiple GPUs

• AllReduce 性能 (ALLREDUCE PERFORMANCE

  • 深度学习训练：NCCL性能为何关键 (Why NCCL performance is critical

• AllReduce 算法 (ALLREDUCE ALGORITHMS

  • 环形算法 (RING ALGORITHM
  • 树形算法 (TREE ALGORITHM

• COLLNET算法与Allreduce性能分析

  • COLLNET算法

    • 链式版本
    • 直接版本

  • 算法总结：优缺点

• 硬件拓扑

  • Allreduce性能
  • 拓扑检测：将环形拓扑映射到硬件
  • 拓扑检测：将树形拓扑映射到硬件
  • 节点间通信：轨优化设计
  • 硬件拓扑：性能考量
  • GPU Direct考量：ACS和ATS

• 点对点通信

  • Alltoall通信：示例与实现
  • 点对点通信：网络代理
  • Alltoall带宽
  • ALLTOALL 通信优化
  • 初始化 / 引导
  • 未来发展

• 总结

---

多GPU计算 (MULTI-GPU COMPUTING)

NCCL是多GPU计算的关键通信库，针对所有平台进行了优化，从桌面级到DGX Superpod。

NCCL代码可在https://developer.nvidia.com/nccl下载，并在NGC容器中可用。源代码位于https://github.com/nvidia/nccl。

图示展示了不同GPU互联方式的硬件配置：
- PCI: 多个GPU通过PCIe连接。
- NVLink: 多个GPU通过NVLink连接。
- NVLink + Infiniband: 多个GPU通过NVLink连接，并与Infiniband网络接口卡连接，用于多节点通信。

深度学习训练 (DEEP LEARNING TRAINING)

从数据到AI (From data to AI)

深度学习过程包括从数PB级数据到AI的转化，涉及训练和推理两个主要阶段。

- 数据: 兆字节级别的数据。
- 训练: 需要8到数千个GPU，高功耗、高内存，训练时间从1小时到1周。NCCL旨在使多GPU训练尽可能高效，以减少训练时间。
- 推理: 在嵌入式设备或数据中心进行，低功耗、低内存，推理时间从1毫秒到1秒。

单GPU上的深度学习训练 (On a single GPU)

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

图示：数据流入GPU进行计算，产生梯度，并更新模型参数以减少误差。

多GPU上的深度学习训练 (On multiple GPUs)

在多GPU训练中，每个GPU处理一部分数据并计算局部梯度。然后，通过NCCL AllReduce操作聚合这些局部梯度，形成全局梯度，并用于更新所有GPU上的模型参数。

图示：多个GPU并行处理数据，每个GPU计算局部梯度。NCCL AllReduce操作将局部梯度汇集并分发全局梯度，然后用于更新每个GPU上的模型参数。

AllReduce 性能 (ALLREDUCE PERFORMANCE)

AllReduce的性能对于多GPU和多节点深度学习训练至关重要。

该图表展示了NCCL在不同互连技术和硬件配置下的AllReduce总线带宽（GB/s）：
- 多GPU（PCI/NVLink）:
- PCI Gen4: 24 GB/s
- PCI Gen5: 48 GB/s
- NVLink2 (V100): 122 GB/s
- NVLink3 (A100): 232 GB/s
- NVLink4 (H100): 330-350 GB/s
- NVLink4 SHARP (H100): 500+ GB/s

• 多节点（IB/RoCE/TCP/IP）:
  • 100GbE, TCP/IP Sockets: 10 GB/s
  • 100Gb IB/RoCE: 12 GB/s
  • 8x 100Gb (DGX-2): 82 GB/s
  • 8x 200Gb (DGX A100): 192 GB/s
  • 8x 400Gb (DGX H100): 330-350 GB/s
  • 8x 400Gb SHARP (DGX H100): 500+ GB/s

深度学习训练：NCCL性能为何关键 (Why NCCL performance is critical)

NCCL的性能对于实现高效的深度学习训练扩展至关重要。

该图表比较了两种不同互联方式下的深度学习训练性能：

1. DL训练规模扩展，PCI Gen 3 + 100G:

• 效率: 77%
• 最大加速比: 640x
• 图表显示，随着GPU数量的增加（从1/256/256到2048/64K/32），训练时间（尤其是Allreduce部分）显著增加，导致整体加速比低于理想情况。

2. DL训练规模扩展，NVLink + 8x200G:
- 效率: 97%
- 最大加速比: 4800x
- 图表显示，在NVLink + 8x200G互联下，随着GPU数量的增加，训练时间（尤其是Allreduce部分）保持在较低水平，实现了接近理想的线性扩展。

理论模型参数: 模型大小 500MB，计算时间 300ms / 256个样本。
图示横轴表示：GPU数量 / 总批次大小 / 每个GPU的批次大小。

AllReduce 算法 (ALLREDUCE ALGORITHMS)

环形算法 (RING ALGORITHM)

环形算法是AllReduce的一种实现方式，数据在环形拓扑结构中的节点之间传递。

图示：AllReduce环形算法的初始状态，每个Input数据（Input0, Input1, Input2, Input3）对应一个GPU和最终的Output。

图示：展示了环形算法中数据如何在一个逻辑环中进行通信，例如，Output0从Input0和Input1接收数据，并传递给Output1。

树形算法 (TREE ALGORITHM)

树形算法是AllReduce的另一种实现方式，它利用树形结构进行数据的归约和广播。

图示：树形算法的详细过程。
- 阶段1: 归约 (Reduce)
- 阶段2: 广播 (Broadcast)
- 关键特点: 每棵树用于处理总数据的一半。每个机器在一棵树上是节点，在另一棵树上是叶子。每台机器接收2倍的半数据，并发送2倍的半数据。
- 图中展示了两棵树，左侧用于将数据归约到节点0，右侧用于将数据归约到节点31，然后通过广播将结果分发给所有节点。

图示：展示了树形算法在多GPU环境下的应用，其中涉及两个树结构（Tree 1 #7, Tree 2 #2）以及它们之间的连接（#3）。每个GPU处理各自的Input（Input0, Input1, Input2, Input3），并最终产生Output（Output0, Output1, Output2, Output3）。

COLLNET算法与Allreduce性能分析

COLLNET算法

链式版本

COLLNET算法的链式版本展示了网络内归约如何连接到一系列GPU。每个GPU（Input0到Input3）处理其各自的输入并生成对应的输出。

直接版本

COLLNET算法的直接版本描绘了四个GPU的互连，每个GPU都连接到网络内归约，并直接与其他GPU连接。这种配置在每个GPU拥有专用网卡时效果更好，并且需要nvswitch或alltoall连接。

算法总结：优缺点

本节总结了Rings、Trees和Collnet算法在延迟、带宽、计算和网络模式方面的优缺点。

Algorithm	Latency	Bandwidth	Computing	Network Pattern
Rings	Linear	Perfect	Uniform	Few flows
Trees	Log	Close to perfect	Imbalanced	Many flows
Collnet	Constant	Close to 2x	Almost uniform	Minimal flows

Collnet算法的延迟为常数，带宽接近2倍（可能受NVLink限制），计算分布几乎均匀，并且网络流最少。

下方的图表进一步展示了Allreduce的延迟和带宽表现：
* Allreduce延迟：Collnet显示出恒定的低延迟，而Rings和Trees的延迟随GPU数量线性或对数增长。
* Allreduce带宽：Collnet展现出最高的带宽，其次是Rings，最后是Trees（最佳和最差情况）。

硬件拓扑

Allreduce性能

本节展示了不同硬件配置下Allreduce的性能，包括多GPU和多节点配置，测量单位为GB/s。

在多GPU配置中，NVLink提供了显著的带宽提升，例如NVLink4 SHARP (H100) 可以达到500+ GB/s。
在多节点配置中，使用IB/RoCE和NVLink4 SHARP (DGX H100) 也能达到500+ GB/s的性能。

拓扑检测：将环形拓扑映射到硬件

幻灯片展示了大型GPU集群中将环形通信模式映射到物理硬件上的复杂互连方式。

不同硬件平台（如PCI平台、NVLink Cubemesh (DGX-1) 和 NVswitch (DGX A100)）对环形拓扑的映射方式各不相同。

拓扑检测：将树形拓扑映射到硬件

与环形拓扑类似，本节展示了PCI平台、NVLink Cubemesh (DGX-1) 和 NVswitch (DGX A100) 如何映射树形通信模式。

节点间通信：轨优化设计

本节对比了经典的互连结构设计与轨优化设计。经典的结构设计要求完美的路由以确保所有流使用不同的链路。而轨优化设计则能确保所有流量都本地化到叶子交换机，从而从根本上避免路由冲突。

硬件拓扑：性能考量

本节讨论了GPU之间以及GPU与网卡（NIC）之间的通信性能。
* GPU到GPU：通过PCI交换机可以获得完美性能，而CPU通常只能达到峰值性能的50%到80%。
* GPU到网卡：需要GPU Direct RDMA以直接通过PCI交换机进行通信。

GPU Direct考量：ACS和ATS

幻灯片解释了在GPU Direct通信中，ACS（Access Control Services）和ATS（Address Translation Services）的作用。
* 如果ACS未进行正确的CPU端配置，会导致GPU Direct中断。目前大多数系统默认启用ACS，因为它在BIOS中启用，而Linux在没有安装VM管理程序（例如KVM）时不会配置ACS。
* ACS会将所有PCI到PCI的事务转发到CPU根复合体进行访问控制检查，这至少会将双向带宽减半。
* ATS通过缓存ACS结果，允许大部分流量直接从GPU流向网卡，从而提高效率。

点对点通信

Alltoall通信：示例与实现

Alltoall通信的NCCL用法示例代码：

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

通过8个p2p通道（即SMs/CTAs），可以并行处理64个对等体。图示进一步说明了每个SM/CTA的8个warp如何与8个不同的对等体进行发送和接收。

点对点通信：网络代理

点对点通信的网络代理机制描绘了CPU代理如何协调GPU（通过Send FIFO发送缓冲区）与网络链路之间的通信，并将数据通过Recv FIFO接收缓冲区传递给目标GPU。

Alltoall带宽

本节展示了不同硬件配置下Alltoall的有效带宽性能，单位为GB/s。
* 在多GPU配置中，DGX H100 (NVswitch) 实现了330-350 GB/s的带宽。
* 在多节点或PCI平台配置中，DGX H100 (NVLink4) 也能达到330-350 GB/s的最高带宽。

ALLTOALL 通信优化

ALLTOALL COMMUNICATION
PXN 优化

在全互联（Alltoall）通信中，大多数数据传输通过顶层交换机进行，尤其是在针对机架优化的设计中。
PXN（PCI x NVLink）允许 GPU 使用远端 NIC 发送数据，通过中间 GPU 进行传输。
PXN 还支持将所有消息聚合到相同的目的地。

(Page 31)

ALLTOALL LATENCY
NCCL Alltoall 延迟
DGX A100, InfiniBand, NCCL 2.12

下图展示了 NCCL Alltoall 延迟在不同节点/GPU 数量配置下的表现，并对比了 PXN 开启和关闭的情况。图表显示，随着节点和 GPU 数量的增加，PXN 开启时（灰色柱）的延迟显著低于 PXN 关闭时（深灰色柱），并且性能增益（绿色柱）也随之提高。在 128/1024 节点/GPU 配置下，PXN 提供了接近 1.25 的性能增益因子。

(Page 32)

初始化 / 引导

INIT / BOOTSTRAP
(Page 33)

NCCL INIT
引导与操作

NCCL 初始化过程分为两个主要阶段：

• 一次性操作（通常在 worker0 上执行）：

  • 生成一个唯一的 ID：ncclUniqueId id;
  • 获取该唯一 ID：ncclGetUniqueId(&id);
  • 将该 ID 广播给所有其他 rank：broadcast_to_all_ranks(&id);
  • 此步骤使用 CPU 侧通信（MPI 或其他方式）。

• 在所有并行 worker 上执行的操作：

  • 初始化：

    • 获取唯一的 ID：get_unique_id(&id);
    • 初始化 NCCL 通信器：ncclCommInitRank(&comm, nranks, &id, rank);

  • 通信：

    • 执行 AllReduce 操作：ncclAllReduce(..., comm);

  • 清理：

    • 销毁 NCCL 通信器：ncclCommDestroy(comm);

该过程的引导线程/套接字通过 CPU 与各个 GPU 节点进行通信。

(Page 34)

NCCL BOOTSTRAP
网络考虑

• 引导（Bootstrap）通过简单的、未加密的 TCP/IP 套接字进行。
• 每个 rank 需要能够通过 TCP/IP 与其他每个 rank 进行通信。这可能需要防火墙配置或将接口设置为“网桥”。
• NCCL_SOCKET_IFNAME 可用于选择用于引导的接口。
• 如果关注安全性，NCCL_SOCKET_IFNAME 应指向一个位于私有网络或 VLAN 上的接口。

未来发展

FUTURE
(Page 36)

H100 SHARP
In-nvswitch 和 In-network

H100 SHARP 利用 NVLink 交换机和网络内（in-network）的规约功能，优化了数据传输和聚合。该图展示了 H100 GPU 之间如何通过 NVLink 交换机（中心矩形）连接，并利用网络内规约能力进行高效的通信。四个象限代表不同的 GPU 组，它们之间可以进行网络内规约操作。

(Page 37)

USER BUFFER REGISTRATION
改进 alltoall 延迟

用户缓冲区注册通过优化数据传输路径来改进 alltoall 延迟。发送缓冲区（sendbuff）中的数据通过发送 FIFO 传输至 GPU，再通过互联结构（可能是 NVLink 和网络交换机）到达目标 GPU。接收端通过接收 FIFO 将数据写入接收缓冲区（recvbuff）。这个过程通过直接注册用户缓冲区来减少数据复制和提高传输效率。

(Page 38)

FUTURE
其他改进

• 持续优化 allreduce 和 alltoall 算法

  • 更多 allreduce 算法（direct+tree/ring, collnet+chain, ...）
  • 更高效的 CUDA 图支持
  • 更高效的聚合
  • 提高非 alltoall 情况下的点对点性能

• 改进可用性

  • 更好地与性能分析器集成
  • 更好的错误报告

• 改进容错性

  • 能够中止初始化/销毁操作

总结

SUMMARY
(Page 40)

SUMMARY

• NCCL: 针对多 GPU 计算的关键通信库。
• 针对所有平台进行了优化，从桌面级到 DGX Superpod。
• 可从 https://developer.nvidia.com/nccl 下载，并在 NGC 容器中提供。
• 源代码可在 https://github.com/nvidia/nccl 获取。

NCCL 旨在为多 GPU 系统提供高性能的集合通信原语。

(Page 41)

(Page 42)