在了解NCCL之前，我们需要先了解一下GPU之间的通信方式。

单机：

1 GPU Shared Memory

GPU和GPU之间的数据传输需要经过Host memory，即需要先把GPU上的数据拷贝至CPU host memory再拷贝至另一块GPU 。

 

2  GPU Direct P2P（Peer-to-Peer)

在相同的RC下，GPU可以通过PCIe直接访问目标GPU的显存，避免了通过拷贝到CPU host memory作为中转，降低延迟。

 

3 Nvlink

GPU和GPU之间的数据传输不再通过PCIe，直接走NVlink，具有更大的带宽。

 

4 NVswitch

GPU和GPU之间直接通过NVswitch相连。

 

NCCL的基本概念：

NCCL，MPI这类集合通信库描述的都是多进程通信。一个进程叫做rank，每个进程有自己的rank号。进程的集合叫做comunicator。

一个进程可以属于一个communicator，也可以属于多个communicators，一个进程在不同的communicator中有不同的rank号。要使用NCCL进行通信，每个设备上都要有一个NCCL Communicator object。通过ncclCommInitRank()和ncclCommInitAll()初始化创建communicator 。

创建communicator前root rank（rank号为0）需要通过ncclGetUniqueId() 创建ID，ID需要broadcast到所有进程上去。这个ID相当于是一个unique object，这样每个进程有了这个ID后就知道自己是communicator的一部分，知道自己是communicator的一部分后就可以开始集合通信了。 

 

创建Communicator：

ncclCommInitRank   ncclResult_t  ncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank)     

创建单个communicator，这个函数用在多线程、多进程程序中，即存在多个rank。每个rank都属于该 communicator，

使用ncclCommInitRank时，需要配合ncclGetUniqueId和MPI_Bcast使用。                                                  

ncclGetUniqueId产生id，MPI_Bcast将id广播到communicator的其他rank上，这样所有的rank都知道自己属于某个communicator了，然后就可以做集合通信了。 

 

ncclCommInitAll   ncclResult_t  ncclCommInitAll(ncclComm_t* comms, int ndev, const int* devlist)

创建一个或多个communicator（comms数组），用在单进程单线程程序（也就是只有一个rank）中。

由于是单进程单线程只有一个rank，所以不需要 ncclGetUniqueId和MPI_Bcast广播unique_id，

也不需要MPI_Comm_rank和MPI_Comm_size知道各自rank的id和一共有多少个rank（rank的size）。 

 

NCCL初始化流程梳理：

1 initTransportsRank函数

initTransportsRank函数处于ncclCommInitRank的流程中，是NCCL初始化的一部分，它是一个重要函数。

每一个rank都会执行initTransportsRank，它主要做了以下几件事：

1  检测系统里的设备和设备之间的拓扑结构

2  计算当前系统中的RING、TREE、COLLNET结构

3  建立设备之间的连接（p2p，shared memory，跨主机net） 

 

initTransportsRank流程梳理：

initTransportsRank流程梳理 - ncclTopoGetSystem ：

1  ncclTopoGetSystem首先创建了一个xmlNode，即XML的根节点。

并设置XML根节点的属性， "system" ["version"] = NCCL_TOPO_XML_VERSION.

2  通过ncclTopoFillGpu将GPU作为一个xmlNode加入到这颗XML树中。 

 

NCCL中以一颗XML树来代表系统中的拓扑结构，每一个设备（CPU，GPU，NIC，PCIe switch）都表现为一个XML结点。

在NCCL中，XML结点数据结构表示如下：

ncclXmlNode表示一个节点，记录了父节点和所有子节点，节点有name和attr，通过xmlSetAttr函数进行设置属性。

 

initTransportsRank流程梳理 -> ncclTopoGetSystem -> ncclTopoFillGpu：

经过ncclTopoFillGpu建立好的包含GPU的XML树如下：

上图中的每一个方框就代表一个xmlNode，每个xmlNode都有对应的属性，比如CPU就有对应的arch，affinity，numaid等属性.

ncclTopoFillGpu中ncclTopoGetXmlFromSys函数通过sysfs中获取gpu节点到cpu的path，通过这个路径转成xml树，并读取该路径下相关属性设置到xml里。

path就是如/sys/devices/pci0000:10/0000:10:00.0/0000:11:00.0/0000:12:00.0/0000:13:00.0/0000:14:00.0/0000:15:00.0/0000:16:00.0/0000:17:00.0，其中gpu的busId是0000:17:00.0 然后读取path下的属性，获取class（PCI设备类型），link_speed，link_width等设置到xml pciNode中，ncclTopoGetStrFromSys就是通过读取path下的内核文件，获得link_speed，link_width等相关属性并保存到strValue中。ncclTopoGetStrFromSys函数在ncclTopoGetXmlFromSys中被调用。 

 

NCCL中可通过设置环境变量NCCL_TOPO_DUMP_FILE来书输出XML文件，并通过该XML文件来查看机器的拓扑结构。

最终生成的XML文件如下所示：

NCCL的XML拓扑文件中，使用上游端口的PCI Bridge的busId来表示这个PCI switch。

一个PCI switch由上游端口的PCI Bridge和下游端口的PCI Bridge组成。所以在NCCL输出的拓扑文件中，PCI的busId看起来像是隔了一跳。 

附PCIe switch结构图：

 

initTransportsRank流程梳理  拓扑图相关数据结构 

ncclTopoSystem中记录的是全局所有类型的节点，有GPU，PCI，NVS，CPU，NIC，NET这几种类型的节点（共NCCL_TOPO_NODE_TYPES（7）种类型）。

其中ncclTopoNodeSet表示某种类型的所有节点，同种类型最多可以有NCCL_TOPO_MAX_NODES（256）个。 

具体的某个节点（GPU，NIC，CPU，PCI ...）的数据结构由struct ncclTopoNode来表示：

type表示节点的类型，是NIC，CPU，PCI还是GPU

nlinks表示了具体有几条边

links存储了具体连接的边

paths存储了和该结点到其他结点的路径，比如paths[0]就存储了所有GPU到该结点的路径，

paths[1]就存储了所有PCI结点到该结点的路径。paths[type][id]就是表示到type类型的第id个node的路径， node->paths[GPU][n] 就表示节点node到第n个GPU的路径。 

 

ncclTopoLink数据结构如下：

图中的边由ncclTopoLink表示，bw表示带宽；remNode表示当前边连接的对端节点,type区分边的类型，具体有以下几种类型： 

最后计算出来的结点和结点之间的路径关系由ncclTopoLinkList表示。 count表示到节点一共有几条边（即有几跳，hops）。这个路径的带宽是width，即count条边中带宽最小为width。list为具体的边。type为结点之间的最终连接类型。比如PATH_NVL，PATH_PIX，PATH_SYS等。 

 

initTransportsRank流程梳理  拓扑图相关数据结构  结点与结点之间的连接类型：

#define PATH_LOC 0     // PATH_LOC为节点到自己

#define PATH_NVL 1     // PATH_NVL表示结点和结点之间只经过NVLink

#define PATH_NVB 2     // PATH_NVB表示通过GPU中转的Nvlink

#define PATH_PIX 3      // PATH_PIX表示结点间经过最多一个PCIe switch，即在通过一个PCI switch下，典型的就是GPU p2p

#define PATH_PXB 4     // PATH_PXB表示结点间经过了多个PCIe witch，但是没有经过CPU

#define PATH_PXN 5    // PATH_PXN表示GPU和NIC之间通过PCI+NVLINK中转 （PXN即PCI * NVlink

#define PATH_PHB 6    // PATH_PHB表示节点间经过了CPU

#define PATH_SYS 7    // PATH_SYS表示结点间经过不同numa之间的路径，即需要经过QPI总线

#define PATH_NET 8   //  PATH_NET表示跨网络，比如GPU direct RDMA

#define PATH_DIS 9    //  无连接 

 

PXN介绍 ：

一个GPU如果要和一个跨NUMA的NIC通信的话，需要经过CPU，跨QPI总线。有了PXN之后，GPU可以直接通过NVlink与另一块GPU通信，然后再将数据转到local NIC上，这能够充分利用带宽降低时延。

PXN是NCCL 2.12的新特性，即NVlink + PCI，先走Nvlink然后再走PCI。 需要CPU proxy线程配合，告诉NIC数据已经ready。

PXN可以结合rail topology进行优化，可以减小交换机的跳数（Node0上的GPU0和Node1上的GPU3通信，3跳变1跳），减小时延。 

 

initTransportsRank流程梳理  路径计算：

在完成拓扑计算后，接下来就是计算路径，主要是ncclTopoComputePaths，ncclTopoTrimSystem，ncclTopoComputePaths三个函数。 主要是通过ncclTopoComputePaths计算系统中GPU到其他所有结点的路径，NIC到其他所有结点的路径，并将路径信息存入ncclTopoNode，ncclLinkLists等结点相关的数据结构中 。

 

 

 

initTransportsRank流程梳理  通信channel的建立：

       

在计算完节点间各设备的拓扑结构，就要开始计算通信路径了，也就是channel。为了更好的利用带宽和网卡，同一块数据可以通过多个channel并发通信，所以nccl在作数据发送流程时会使用多个channel发送数据。 channel的建立是为了后续作数据通信准备的，nccl中建立了基于ring，tree，CollNet等结构的channel，并将计算出的channel存入相关的数据结构中。 

Channel：

执行nccl_test程序：./build/all_reduce_perf -b 8 -e 256M -f 2 -g 4 

通过该日志可以看到，GPU之间建立的channel图如下：

单机上执行nccl-test指定GPU的个数，可以看到NCCL将4块GPU（id分别为0，1，2，3）构成了一个ring，ring环中一共有两个channel，用于数据的传输。其中GPU和GPU之间通过p2p方式通信（无NVlink）

可通过在/etc/nccl.conf中设置NCCL_MAX_NCHANNELS来指定最大channel的数量. 

 

 

initTransportsRank流程梳理  通信channel的数据结构：

通信路径channel搜索完毕的结果会存储在结构体ncclTopoGraph中

nChannels表示一共有几条路径

bwIntra表示节点内单个channel的带宽

bwInter表示节点间单个channel的带宽

typeIntra表示节点内channel的路径类型

typeInter表示节点间channel的路径类型

intra数组存储了节点内每个路径的类型， 比如[0,1,2,3,4,5,6,7]，就是由8个GPU构成的一个ring，存储在intra数组中

inter数组存储了节点间每个channel的路径

比如两台机器，每台机器有8块GPU ，那么

第一台机器的ring为： graph->intra: GPU/0 GPU/7 GPU/6 GPU/3 GPU/2 GPU/5 GPU/4 GPU/1

                                         graph->inter: NET/0 NET/0

第二台机器的ring为： graph->intra: GPU/10 GPU/9 GPU/8 GPU/13 GPU/12 GPU/15 GPU/14 GPU/11

                                  graph->inter: NET/0 NET/0 

 

 

通信channel的搜索：ncclTopoCompute流程分析 ：

channel搜索的目标：搜索出尽可能多的channel，每条channel的带宽尽可能大

channel搜索的流程：先设置一系列较为苛刻的条件（比如先把带宽设置得很高(LOC_BW=5000GB/s)，节点内类型graph->typeIntra设置为PATH_NVL等），不符合条件则降低标准（比如降低带宽，降低类型）继续搜索。

搜索的核心是函数是ncclTopoSearchRec，本质是都搜到ring环和可用的NIC，并填到相关数据结构中。 

如上图所示： 满足条件则跳到done，搜索结束。如不满足条件，则减小带宽值等，返回ncclTopoSearchRec对重新对channel进行搜索。 

在建立好机器间多个rank的大环后，执行ncclTransportP2pSetup建立当前rank和prev，next的之间的通信链路， NCCL中针对不同的All Reduce算法建立了不同的通信链路，包含了ring ,Tree, nvls（nvlink+sharp，将Allreduce卸载到NVswitch上去做）

 

本文着重分析了NCCL中的初始化相关的一些流程，最后总结一下NCCL的主要feature

NCCL feature：

1 结点内的各个设备（NIC，GPU，CPU，PCIe switch）的拓扑发现与检测；

2 利用 SHARPV2 的网络内 all reduce 操作，充分利用交换机性能，将峰值带宽提升 2 倍；

3 通过图形搜索，找到更佳的高带宽、低延迟的环和树集合；

4 支持多线程和多进程模型，与MPI联合编程；

5 节点间通信支持IB，ROCE，TCP等多种网络协议；

6 使用 Infiniband 动态路由重新路由流量，缓解端口拥塞