理解NCCL源码必看：详细论述NCCL源码做了哪些事

Happy_Enger

已于 2024-04-22 13:08:09 修改

阅读量2.8k

点赞数 25

分类专栏： NCCL 文章标签： nvidia NCCL 源码

于 2024-03-21 11:21:44 首次发布

本文链接：https://blog.csdn.net/u014443578/article/details/136902252

版权

NCCL 专栏收录该内容

6 篇文章 16 订阅

订阅专栏

本文详细解析了NCCL源码的初始化过程，包括获取UniqueID、Bootstrap网络的构建和AllGather同步数据实现，以及机器拓扑、XML转无向图、路径计算和Channel搜索的算法。文章通过实例说明了如何通过Transport建立P2P通信和利用共享内存的设计原理。

摘要由CSDN通过智能技术生成

前言

NCCL 源码解析总目录

本文主要通过图示的方式来讲解 NCCL 源码初始化API做了哪些事，对NCCL源码实现的功能有一个整体直观的认识。
对于即将研究源码的同学来说，可以从整体上先了解一下，与源码一同食用，更健康。

我的思维可能比较跳脱，有的内容没有提到的，留言或者请自行查阅资料理解。

如有问题，请留言指正。

有些遗漏之处，还没涉及，后面补；
闲话后面再补。

NCCL 源码具体实现概览

ncclCommInitRank() 主要实现如下几个功能：
请添加图片描述

1. 初始化及获取 UniqueID

初始化主要是软件运行需要的一些准备，以及所需要的网络初始化；
需要关注的是，获取UniqueID, 本质是一个数组，保存当前机器的IP和Port，什么时候用到，后文会提到。
请添加图片描述

2. Bootstrap 网络

2.1 目的

建立Bootstrap网络的目的是，线程或进程通过TCP网络建立一个环，所有线程或进程通过 ALLGather 方式同步一些信息
该功能完成之后结果如下图：
在这里插入图片描述

2.2 AllGather 同步数据实现步骤

首先，我前面提到建立了 bootstrap 网络之后，就是用来做 allgather 同步信息的，本节介绍一下 AllGather是如何操作的。
假设有3个GPU，对应有三个线程或进程，如下图，每个长方形代表一个线程，每个线程有一个独有的相关数据，要实现所有线程gather所有信息，即每个线程都要拥有全部数据，最快速的方法就是通过AllGather的方式，传输数据的通过是上图中通过TCP建立的环。
在这里插入图片描述

第一步，前一个向后一个发送自己的数据
继续重复

最终，所有的线程都拥有了所有的数据。

2.3 Bootstrap 网络如何建立

首先，同样假设我们的硬件组成中有三个GPU，每个GPU对应一个线程，Rank 也是唯一标识一个GPU的ID。启动NCCL初始化中会遍历主机中的网卡，对应3个GPU，创建三个线程。现在我们要将他们通过TCP连成一个环。
虽然线程是资源共享的，但是在NCCL中，他们之间的数据不是通过全局变量进行同步的，也就是说，为了软件结构上的统一，NCCL 中线程和进程的处理基本一致，不会因为线程之间可以资源共享而采用不同的方式，即，你可以认为这里三个线程和进程一样，是独立的，互相无法感知的。所以，现在就有一个问题，这三个线程如何建立一个环，要连接TCP，必须知道服务端的IP和端口，但是现在每个线程独立，并不知道如何联系。
NCCL 另外创建了一个 Root 线程，这个线程只有一个，是 Rank0 线程创建的，所以一个初始化中只有一个。创建这个为了啥，看后文。
继续操作，每个Rank线程都创建两个TCP服务监听端口，端口1是为了让其他Rank线程可以联系到本地，端口2是为了让root线程可以联系我。
下一步，root 线程也建立一个监听端口，等待所有 rank 线程链接。
但是，其他rank线程链接root会面临同样的问题，线程相互独立，没有ip和端口不知道怎么连接。
还记得前面提到的UniquelID吗，保存的是ip和端口，创建所有线程的时候，这个信息就都传递进来了，所以，root 线程使用UniquelID建立Socket监听，其他线程通过这个链接root 线程。
OK, 既然知道了，那就链接TCP服务器，连接的目的是发送本地线程的服务器监听端口1给Root线程进行一个汇总，统一管理。
执行完上一步，Root线程就会知道了所有Rank线程的监听端口1的socket了。
下一步，在root线程中，按照rank顺序建立 ring, 具体实现如图，环是 0->1->2, 那么就向Rank0发送Rank1 的监听Socket, 向 Rank1 发送 Rank2 的监听Socket。。。
到此，所有Rank线程就知道了要链接的监听Socket，Root 线程使命结束。
下一步就简单了，直接连接就可以了，围成一个环。因为是网络通信，所以无论是线程还是进程，无论是单机内还是多机，线程或进程都可以建立这个环。

3. 机器拓扑

把服务器内相关的硬件信息抽象出来，让软件知道。获取的设备：GPU、PCI Bus、网卡；

3.1 目的

软件获取硬件连接信息，为了后续软件建立环等操作

3.2 流程

在这里插入图片描述

3.3 例程

假设现在硬件连接如下：
4个GPU, 通过NVLINK 相互连接，有一根的也有两根的，GPU4 单独拿出来；

拓扑流程走完，软件就可以生成如下文件：
在这里插入图片描述

4. XML 转无向图

XML用于算法对软件不友好，因此需要把XML格式转为数据结构中的无向图结构，无向图可视化就是如下图：

请添加图请添加图片描述
片描述

其中，有几个重点需要注意:
节点的类型： 节点的类型分为 GPU、PCI、NVS(nv switch)、CPU、NIC、NET
节点的信息： 节点的信息最主要的是 nlinks, 表示与该节点相连的设备数量，包括自己，比如与GPU0相连的设备数量有5个
边的信息： 边的信息则要关注以下内容：
在这里插入图片描述

Link.remNode 边的对端节点就是对端节点的对象
Link.bw 边的带宽是理论计算的带宽，且是累计的，如上，两个NVLINK则是40，一个是20
Links 是一个数组，保存到其他设备的所有边
边的类型定义如下
LINK_LOC 跟自己相连的边
LINK_NVL NVLink 连接的边
LINK_PCI PCIe总线连接的
LINK_SYS CPU之间连接的边
LINK_NET 网络连接的边

在这里插入图片描述
计算NVLINK的理论带宽，是根据SM选择的，定义如下。

5. 路径计算

5.1 目的

搜索所有设备到GPU、NIC、NVSwitch的通路，并计算相应信息。

5.2 结果

路径计算完毕之后，就会得到路径信息，如下
count 为路径的数量
bw 为路径的带宽
type 为路径的类型
在这里插入图片描述

如GPU0 ，所有设备到GPU0的路径如下
在这里插入图片描述
请添加图片描述
其中路径类型的定义，从上到下，速率依次降低：

5.3 实现步骤

记住：

计算的原则是使用广度优先搜索
最优的路径是路径最短且带宽最大
按照带宽大小遍历对端节点

5.3.1 广度优先搜索

以GPU0为例，要计算所有设备到GPU0的路径，使用的算法是广度优先搜索，先简单介绍一下该算法。

假设这是一个树，从GPU0开始，先访问GPU0
GPU0一层的节点就一个，把GPU0的子节点GPU3 GPU1 GPU2放入队列；
GPU3 出队, 处理GPU3，将GPU3 的子节点放入队列
继续出队直到结束。

总之，如下图，从上到下，从左到右遍历节点。
在这里插入图片描述

5.3.2 例程

以GPU0 为例，计算所有设备到GPU0的路径。

从GPU0 开始遍历节点，GPU0 到GPU0没啥好说的，自己跟自己带宽最大，计算一遍，就不会再更新了。
假设程序遍历到GPU3
针对GPU3，再遍历GPU3 相连的节点，有5个
同样假设，我们遍历到 GPU3 的对端节点GPU2。
这个意思就是现在要计算GPU2 到GPU0的路径，通过GPU3计算，也就是说要看一下当前GPU2记录的到GPU0的路径跟GPU2走GPU3到GPU0的路径相比较，有点绕，仔细看这句话。
比较的指标是路径最短以及带宽最大，如果要走GPU3，那么要选择 min(GPU3到GPU0的bw, GPU3到GPU2的bw)，这个好理解，就是选最小的带宽
如下例子，假设GPU2到GPU3的带宽是40，GPU3 到GPU0的带宽是40，那么GPU2->GPU3->GPU0 的带宽是40，GPU2->GPU0还有两条路带宽是20, 一个是直连，一个是过GPU1。
这里会选择直连作为最优路径，因为虽然带宽小，但是路径短，即边的数量只有1，比其他两个的2条边少。
通过第四步，可以更新GPU2到GPU0 的路径了。通过暴力搜索，所有的节点都能计算出来。

5.3.3 结果

GPU0 的路径结果如下。有连接类型，设备类型，还有带宽。
在这里插入图片描述

6. 搜索 Channel

以上步骤，都是为了本节的搜索 channel 做准备。channel 搜索的结果不止一个，也可能是多个。
同样需要记住的关键点：

使用的是递归暴力搜索
根据设置的条件，最严格到一步一步放松，条件包括路径类型和带宽
满足的要求，即怎么才算成功搜索到一个channel：
1. 路径要通
1. path 路径类型要满足
1. 链路带宽 > 要求的条件带宽
1. channel * bw 要尽量大

6.1 目的

在这里插入图片描述

6.2 例程

为了方便起见，以以下拓扑链接为例，搜索 ring channel，链接路径类型都是 nvlink, 相应贷款也列出来了。
根据GPU的SM计算能力，设置带宽条件，先从60开始，路径类型也从nvlink开始，从设备GPU0开始暴力搜索，判断GPU0到GPU1或者GPU3的路径类型满足，但是带宽不满足60，那么就失败了；先降低路径类型，从nvlink 一步一步降低路径类型要求，但是无论如何带宽60是不满足的，所以，恢复路径类型的要求到nvlink, 将带宽要求降低到40，40也不满足，因为GPU0开始，到GPU3是通的，但是转一圈回来，GPU2、GPU1到GPU0的带宽都是小于40的，所以成不了环。
下一个节点的选择有两种策略，一个是按照PCI顺序，即GPU0的下一个节点是GPU1,GPU2,GPU3，这些都是要遍历的，针对GPU1, 下一个几点就是GPU2, GPU3，这样递归遍历；另一个策略是按照带宽大小寻找下一个节点；
当路径类型要求为nvlink，路径带宽条件为20的时候，就有路可以通了。下一个节点是按照带宽大小寻找的，即GPU0下一个节点是GPU3, GPU3到GPU2,GPU2到GPU1。两个条件都满足，可以构成 0->1->2->3，也可以构成channel, 0->3->2->1, 这是一个循环。同时，链路上的带宽是减掉当前遍历的带宽的，如下图，好理解一点，本来是40的变为20，本来是20的变为0。
还要在上面的基础上进行遍历，要充分利用带宽。流程还是一样的，就不细讲了，我们可以直接看图，按照上面的思路直接选，会花圈就行，以下的信息可以画出2个 channel 来，即 0->3->1->2, 0->2->1->3.
所以，最终，搜索出4个channel。还可以从另一个角度去看 channel, 站在GPU3的角度，带宽为20的话，有两条nvlink到GPU0, 这个四个channel里只看GPU3的下一个节点，就会看到2个0, 1个1，1个2。这是最终的一个效果，即充分利用了所有的带宽。