NCCL 介绍

NCLL github：https://github.com/NVIDIA/nccl

模型规模和单一算例下GPU集群数量的指数化增长，对于大规模的神经网络的训练和推理，需要对计算任务进行分割并协调GPU集群来执行单个同步计算。

常用的 GPU 并行计算技术：

• 数据并行（DP）：在不同的GPU上运行同一批数据的不同子集
• 流水线并行（PP）：在不同的GPU上运行模型的不同层
• 张量并行（TP）：将单个数学运算（如矩阵乘法）拆分到不同的GPU上运行
• 专家混合（Mixture-of-Experts），只用模型每一层中的一小部分来处理数据
• 专家并行（EP）：每个 GPU 负责某个或者某几个专家的计算

无论使用哪一种并行运算技术，始终无法避免GPU与GPU之间的相互通讯。

1 集合通信

GPU与GPU之间的通信包含了点对点通信（单发单收）和集合通信（多发/多收）。在并行计算环境中，多路GPU节点同时工作以加速计算过程和处理大规模的数据集。为了实现这一目标，需要高效的通信机制来协调计算节点间的信息数据交换。在此过程中，集合通信起到至关重要的作用。

集合通信（Collective Communication）是并行计算和分布式系统中用于在多个进程或计算节点间进行数据交换的一种通信模式。与点对点通信（在两个节点之间直接发送和接收消息）不同，集合通信涉及到一组节点之间的协调通信。常见的集合通信有如下几种：

• Broadcast：从一个节点向其他所有节点发送相同的数据。
• Scatter：一个节点将不同的数据块分发给组内的其他节点，每个节点收取唯一的一块数据。
• Gather：与Scatter相反，组内的其它节点将数据发送给指定的某一个节点。
• Reduce：一个节点接受组内其他节点发出的数据，并生成汇总结果。
• All-Gather：每个节点都发送自身的数据，也接受组内其他节点发出的数据。
• All-Reduce：每个节点都发送自身的数据，也接受组内其他节点发出的数据，并生成汇总结果。
• All-to-All：每个节点将不同的数据块分发给组内的其他节点，同时也接收其他节点分发的数据块。

集合通信对于需要广泛数据交换的算法（例如矩阵操作、多维数据分布）至关重要。其在实现上也非常复杂。这种通讯模式的高实时性、同步性对网络的带宽、延时都提出了苛刻的要求。

2 GPU 拓扑

集合通信仅仅是并行计算中的一种通讯方法。其底层仍然受传输层以下的各种类型通信信道的限制，以及节点间复杂拓扑结构的限制。

集合通信是否得以高性能得实现，常依赖于通信节点间拓扑的识别。服务器技术、网络技术、板间通信技术的发展，使得GPU之间的互联方式也变得纷繁复杂。

DGX-1系统为例，GPU与GPU之间的通信包含：PCIe、NVLink、Infiniband、QPI等通信方式。每一种通信方式都会经过多样的芯片、电路、光路，使用各种通信协议，进而拥有不同的带宽、时延特性。

NCCL是一个Nvidia专门为多GPU之间提供集合通讯的通讯库，或者说是一个多GPU卡通讯的框架，它具有一定程度拓扑感知的能力，提供了包括 AllReduce、Broadcast、Reduce、AllGather、ReduceScatter 等集合通讯 API，也支持用户去使用 ncclSend()、ncclRecv() 来实现各种复杂的点对点通讯，如One-to-all、All-to-one、All-to-all等，在绝大多数情况下都可以通过服务器内的PCIe、NVLink、NVSwitch等和服务器间的RoCEv2、IB、TCP网络实现高带宽和低延迟。

它解决了前面提到的GPU间拓扑识别、GPU间集合通信优化的问题。NCCL屏蔽了底层复杂的细节，向上提供API供训练框架调用，向下连接机内机间的GPU以完成模型参数的高效传输。

Caffe2、Chainer、MxNet、PyTorch和TensorFlow等领先的深度学习框架已将NCCL集成到多GPU多节点系统上加速深度学习训练。在NCCL工作流中，主要包含了节点和设备的拓扑探测、根据探测出的结构构建最优的ring或tree形式的通信结构。

与MPI相比而言，NCCL主要用于优化在 NVIDIA GPU 上的多GPU并行计算，特别是深度学习和机器学习任务。它专注于加速GPU之间的数据通信，以提高训练深度学习模型时的性能。而MPI则相对更通用，不仅仅局限于对Nvidia GPU的通信优化。MPI更多的聚焦在支持在大规模计算集群上的并行计算，特别是在超级计算机上进行科学计算。不过二者很多API接口相似度高，迁移和兼容相对简便；NCCL也可以轻松与MPI结合使用。

MPI专注于CPU间通信，NCCL聚焦在GPU间通信。不过由于NCCL是Nvidia针对自家产品推出的专用的通信库，其性能优化方面有得天独厚的优势。

3 NCCL Test

NCCL 安装：

wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/x86_64/cuda-keyring_1.0-1_all.deb
sudo dpkg -i cuda-keyring_1.0-1_all.deb
sudo apt-get update
sudo apt install libnccl2=2.19.3-1+cuda12.3 libnccl-dev=2.19.3-1+cuda12.3

NCCL Test用于再NCCL部署后，简单地测试环境地正确性以及环境地性能。NCCL Test github: https://github.com/NVIDIA/nccl-tests/tree/master

make MPI=1 MPI_HOME=/path/to/mpi/gcc/openmpi-4.1.5rc2 CUDA_HOME=/path/to/cuda NCCL_HOME=/lib/x86_64-linux-gnu/

# MPI、CUDA、NCCL路径可以通过下述方式查找
whereis cuda
ldconfig -p | grep libnccl
which mpirun

4 NCCL Ring & Double binary trees

nvidia 官方 blog：https://developer.nvidia.com/blog/massively-scale-deep-learning-training-nccl-2-4

在集合通讯库中，频率最高、用途最广的操作无非是All Reduce操作。在NCCL 2.4版本之前，通常使用Ring实现All Reduce，在2.4版本后，NCCL提供了双二叉树用于增强All Reduce操作的性能。

对于Ring结构而言，其缺点是自身的线性结构导致通信的延迟与节点数量成线性增长关系（即时延取决于Ring的长度），在GPU节点数量庞大时延时明显。而Tree结构，改善了这一缺点，其通信延时与节点数量成对数关系（即时延取决于Tree的深度）。

Tree中的每个节点都代表一个GPU，灰色的叶子节点需执行直接计算，绿色的节点不参与直接计算而参与集合通信。

任意一个节点如果在左侧树中是绿色节点，那么在右侧必然是灰色节点。（这用于保证每个节点都会参与直接计算，不存在划水的节点）除根节点以外的每个节点都有2个子节点和2个父节点，这种结构保证了每个节点都能从2个子节点接收和发送数据，也都能向2个父节点发送和接收数据（注意收发关系是对等的、但收发的内容并不是一致的）。

建两棵树的目的是为了方便将数据分成2组，展开数据并行计算。

以All Reduce为例，每棵树以leaf→root方向开始reduce（2对1和3进行reduce，4对2和6进行reduce，……），最终数据到达root后，由root→leaf方向开始broadcast（16把对8和24reduce后的结果broadcast给8和24，8再把结果broadcast至4和12，……）。这样保证每个树的每个节点都实现了本树拥有数据的All Reduce操作，对于两棵树各自而言，实现了各自持有数据集的All Reduce，每个节点都拥有两棵树All Reduce后的结果，每个节点再对两个数据集的结果进行合并，实现整个数据集在所有节点上的All Reduce。