NVIDIA GPU 集合通信库 NCCL 初始化流程源码级剖析

本文来自博客园，作者：T-BARBARIANS 原创，博文严禁转载，转载必究！
本篇内容将从 NCCL（NVIDIA Collective Communications Library） 源码的角度尽量深入剖析 NCCL 初始化的一些关键业务流程。并且希望通过本篇文章的讲解，能够为大家揭开一些 NCCL 的神秘面纱。
文章内容很长，换言之：信息量很大，一时半会真的说不清楚。
如果尝试深入一些去了解 NCCL，可能需要先了解一些 NCCL 的基本概念，可以参考我写的上一篇有关 NCCL 基础介绍的文章：NVIDIA GPU 集合通信库 NCCL 基本概念简述 。
对于 NCCL 源码，我自己是觉得理解起来挺费劲的。主要是会涉及到很多新知识和新概念，源码里也有很多暴力求解算法，以及触及了大量技术知识盲区，学习起来挺头痛；另外，设备太贵，实验环境比较简陋，只能跟踪到最基础的代码业务流程；再加上才疏学浅，难免有很多撰写不太顺畅，或者理解错误之处，还请各位高手指点一二，共同进步。
NCCL AllReduce 求和操作demo

先举例介绍 NCCL 执行一次 AllReduce Sum 操作的 demo 示例。借助 MPI，使用 NCCL 在多机、多 GPU 实验环境上实现一次 AllReduce Sum 操作，直观的展示 NCCL 完成一次 AllReduce Sum 操作的结果。

运行命令行：

1. mpirun --allow-run-as-root -np 4 --host vm1:2,vm2:2 ./examples/01_communicators/03_one_device_per_process_mpi/one_device_per_process_mpi

复制代码

• 每个 GPU 都有一个  [ global_rank_id，global_rank_id + 1 ]  浮点数组
  
• 通过 NCCL 实现 AllReduce 求和操作，预期是在所有 GPU 上都将得到相同的全局求和结果
  
  
  
  • 每个 GPU 的 result_elements[0] = 6
    
  • 每个 GPU 的 result_elements[1] = 10
    
  
  

下面是 demo 的部分代码，代码框架来源于 NCCL 源码目录：docs/examples/01_communicators/03_one_device_per_process_mpi/one_device_per_process_mpi，并做了一些改动。

1. int main(int argc, char *argv[]) {
2.    
3.     // MPI 初始化与拓扑发现
4.     MPI_Init(&argc, &argv);
5.     MPI_Comm_rank(MPI_COMM_WORLD, &mpi_rank);   // 为每个MPI进程分配一个全局 mpi_rank id
6.     MPI_Comm_size(MPI_COMM_WORLD, &mpi_size);   // 获取MPI总进程数
7.     local_rank = getLocalRank(MPI_COMM_WORLD);  // 获取本节点的 local_rank id
9.     // CUDA 设备绑定
10.     cudaGetDeviceCount(&num_gpus);
11.     cudaSetDevice(local_rank);     // 将每个进程绑定到 local_rank 对应的 GPU
12.     cudaStreamCreate(&stream);
14.     // NCCL 通信器初始化
15.     if (mpi_rank == 0)
16.       ncclGetUniqueId(&nccl_id);  // 生成用于唯一标识一个 NCCL 通信域的 Unique ID
17.     MPI_Bcast(&nccl_id, ...);     // 向其它 MPI 进程广播该 Unique ID
18.     ncclCommInitRank(&comm, mpi_size, nccl_id, mpi_rank);
20.     // 每个 GPU sendbuff 的两个元素设为 [global_rank_id, global_rank_id + 1]
21.     cudaMalloc(sendbuff/recvbuff, size * sizeof(float));
22.     float elements[2] = { (float)comm_rank, (float)(comm_rank + 1) };  
23.     cudaMemcpy(sendbuff, elements, ..., cudaMemcpyHostToDevice);
25.     // 利用 GPU 对 elements[2] 个 float 执行 AllReduce Sum
26.     ncclGroupStart();
27.     ncclAllReduce(sendbuff, recvbuff, size, ncclFloat, ncclSum, comm, stream);
28.     ncclGroupEnd();
29.     cudaStreamSynchronize(stream);
30.     printf("Collective operation completed on rank %d\n", mpi_rank);
32.     // 获取 AllReduce Sum 结果
33.     float result_elements[len] = {0, 0};
34.     CUDACHECK(cudaMemcpy(&result_elements, recvbuff, sizeof(result_elements), cudaMemcpyDeviceToHost));
36.     // 预期 result_elements[0] = 0 + 1 + 2 + 3 = 6
37.     // 预期 result_elements[1] = 1 + 2 + 3 + 4 = 10
38.     for (size_t i = 0; i < len; i++)
39.     {
40.       printf("Device %d correctly received sum: %.0f\n", comm_rank, result_elements[i]);
41.     }
42. }

复制代码

结果日志：

1. # GPU0
2. Collective operation completed on rank 0
3. Device 0 correctly received sum: 6
4. Device 0 correctly received sum: 10
6. # GPU1
7. Collective operation completed on rank 1
8. Device 1 correctly received sum: 6
9. Device 1 correctly received sum: 10
11. # GPU2
12. Collective operation completed on rank 2
13. Device 2 correctly received sum: 6
14. Device 2 correctly received sum: 10
16. # GPU3
17. Collective operation completed on rank 3
18. Device 3 correctly received sum: 6
19. Device 3 correctly received sum: 10

复制代码

生成 NCCL 通信域唯一通信标识符

从上面的代码中可以看到使用了函数：ncclGetUniqueId。原因是使用 NCCL 进行多 GPU 集合通信时，所有参与通信的进程必须属于同一个 NCCL 通信域。通信域定义了一个可以互相通信逻辑组，而创建这个通信组的第一步，就是需要创建一个共享的 Unique Id，作用类似：加群邀请码。
NCCL 通过函数 ncclGetUniqueId 实现如下具体任务：

• 寻找控制面通信接口
  
• 生成加群唯一邀请码 Unique Id
  
• 在控制面通信接口上进行连接监听
  
• 等待其它 MPI 进程连接并协商，并以根 MPI 为中心进行信息传递
  

寻找控制面通信接口

获取控制面通信接口信息（只取一个），如果指定了环境变量，则获取与环境变量匹配的接口信息。

1. ncclResult_t bootstrapNetInit() {
2.   if (bootstrapNetInitDone == 0) {
3.       ......
4.       const char* env = ncclGetEnv("NCCL_COMM_ID");
5.       if (env) {
6.       ......
7.       } else {
8.         NCCLCHECK(ncclFindInterfaces(bootstrapNetIfName, &bootstrapNetIfAddr, MAX_IF_NAME_SIZE, 1, &nIfs));
9.         ......
10.       }
11.       ......
12.       ncclSocketToString(&bootstrapNetIfAddr, line+strlen(line));
13.       INFO(NCCL_BOOTSTRAP, "Bootstrap: Using%s", line);
14.     }
15.   }
16.   return ncclSuccess;
17. }

复制代码

有输出日志：

1. findInterfaces:155 NCCL TRACE Found interface ens3:10.10.10.11<0>
2. bootstrapNetInit:121 NCCL INFO Bootstrap: Using ens3:10.10.10.11<0>  #此时还没有产生端口信息，端口为0

复制代码

生成加群唯一邀请码 Unique Id

生成加群唯一邀请码。唯一邀请码由随机数 magic，和上一步得到的通信接口信息拼接而成

1. struct ncclBootstrapHandle {
2.   uint64_t magic;
3.   union ncclSocketAddress addr;
4. };
6. ncclResult_t bootstrapGetUniqueId(struct ncclBootstrapHandle* handle) {
8.   const char* env = ncclGetEnv("NCCL_COMM_ID");
9.   if (env) {
10.     ......
11.   } else {
12.     NCCLCHECK(getRandomData(&handle->magic, sizeof(handle->magic)));
13.     memcpy(&handle->addr, &bootstrapNetIfAddr, sizeof(union ncclSocketAddress));
14.     NCCLCHECK(bootstrapCreateRoot(handle, false));
15.   }
17.   return ncclSuccess;
18. }

复制代码

在控制面通信接口上进行连接监听

• 在控制面通信接口上进行监听，创建控制面协商引导线程
  
• 通过 ncclSocketInit/ncclSocketListen 得到监听端口，开始在根 MPI 进程进行控制面协商连接监听
  

1. ncclResult_t bootstrapCreateRoot(struct ncclBootstrapHandle* handle, bool idFromEnv) {
3.   ......
4.   NCCLCHECKGOTO(ncclSocketInit(listenSock, &handle->addr, handle->magic, ncclSocketTypeBootstrap, NULL, 0), ret, fail);
5.   NCCLCHECKGOTO(ncclSocketListen(listenSock), ret, fail);
6.   NCCLCHECKGOTO(ncclSocketGetAddr(listenSock, &handle->addr), ret, fail);
8.   args->listenSock = listenSock;
9.   args->magic = handle->magic;
10.   PTHREADCHECKGOTO(pthread_create(&thread, NULL, bootstrapRoot, (void*)args), "pthread_create", ret, fail);
11.   ncclSetThreadName(thread, "NCCL BootstrapR");
12.   ......
13. }

复制代码

1. ncclSocketListen:421 NCCL TRACE Listening on socket 10.10.10.11<36241>

复制代码

等待其它 MPI 进程连接并协商，并以根 MPI 为中心进行信息传递

控制面引导线程传递各个MPI进程的连接信息，为建立控制环逻辑拓扑打基础

1. static void* bootstrapRoot(void* rargs) {
2.   ......
3.   /* Receive addresses from all ranks */
4.   NCCLCHECKGOTO(ncclSocketInit(&sock), res, out);
5.   NCCLCHECKGOTO(ncclSocketAccept(&sock, listenSock), res, out);
6.   NCCLCHECKGOTO(socketRecv(&sock, &info, sizeof(info)), res, out);
7.   NCCLCHECKGOTO(ncclSocketClose(&sock), res, out);
9.   ......
10.   do {
11.     ......
12.     /* 如果当前节点的前驱已存在，MPI 引导线程将当前节点的连接信息发送给当前节点的前驱 */
13.     int prev = (nroots > 1) ? (localId - 1) : BOOTSTRAP_PID(localId - 1, nrecv);
14.     if (prev >= 0 && prev < n2send && memcmp(&zeroAddress, &rankAddressesRoot[prev], sizeof(union ncclSocketAddress)) != 0) {
15.       NCCLCHECKGOTO(rootSend(&rankAddressesRoot[prev], magic, &info.connectInfo), res, out);
16.     } else {
17.       memcpy(&rankInfo[localId], &info.connectInfo, sizeof(union ringConnectInfo));
18.     }
20.     /* 如果当前节点的后继已存在，MPI 引导线程将已缓存的当前节点的后继信息发送给当前的连接节点 */
21.     int next = BOOTSTRAP_PID(localId + 1, nrecv);
22.     if (localId >= 0 && localId < n2send && memcmp(&zeroInfo, &rankInfo[next], sizeof(union ringConnectInfo)) != 0) {
23.       NCCLCHECKGOTO(rootSend(&info.listenRootAddress, magic, &rankInfo[next]), res, out);
24.     } else {
25.       memcpy(rankAddressesRoot + localId, &info.listenRootAddress, sizeof(union ncclSocketAddress));
26.     }
27.     ++c;
28.     TRACE(NCCL_BOOTSTRAP, "Received connect from rank %d total %d/%d", info.rank, c, nrecv);
29.   } while (c < nrecv);
30.   ......
31. }

复制代码

通过 MPI 控制面引导线程的信息中继，参加通信的各个 MPI 进程都能知道自己的后继者是谁，为整个控制面的信息流转打基础。
MPI 根进程将 Unique Id 通过 MPI_Bcast 传递给所有 MPI 进程，里面包含了 MPI 根进程的 magic，IP，port 信息，用作其它 MPI 进程加群的唯一票据。

1.   MPI_Bcast(&nccl_id, NCCL_UNIQUE_ID_BYTES, MPI_CHAR, 0, MPI_COMM_WORLD);

复制代码

通过 MPI_Bcast，所有 rank 上有 unique ID 接收日志：

1. # rank 0
2. Rank 0 received NCCL unique ID
4. # rank 1
5. Rank 1 received NCCL unique ID
7. # rank 2
8. Rank 2 received NCCL unique ID
10. # rank 3
11. Rank 3 received NCCL unique ID

复制代码

以上过程，可总结为如下图所示：

NCCL初始化

这里开始陆续讲解 NCCL 初始化相关流程。
NCCL 通过函数 ncclCommInitRank 实现 NCCL 相关初始化和众多计算逻辑，主要包含如下工作：

• 网络插件初始化。主要涉及 InfiniBand，或者基于 TCP/IP 的 Socket 网络插件的选择。
  
• NCCL 控制环的建立。主要涉及 Unique Id 的产生和用途，以及所有 rank 组成控制面环形拓扑的相关流程。
  
• PCIe 设备发现。记录从 GPU/NIC 出发，直至归属 CPU 这条物理 PCIe 链路上有哪些 PCIe 设备，并生成从 GPU/NIC 至 CPU 链路的 PCIe XML 链路拓扑。
  
• PCIe 设备间建图。PCIe 设备之间的物理链路关系是事实存在的，但是 NCCL 需要将这些链路关系建立到自身的数据结构中，目的是后续 channel 的计算服务。
  
• PCIe 设备间最优可达路径计算。基于 PCIe 建图，NCCL 拥有 PCIe 设备间的连关系。在此基础上，NCCL 使用 BFS 算法，计算任意 PCIe 设备之间的最优可达路径，即跳数最少，带宽最大的可达路径。
  
• 设备内部基于 Ring，Tree 通信算法的 channel 计算。在路径带宽允许的情况下，NCCL 会尽最大可能计算出同一种通信算法的更多通信通道，从而尽可能榨干硬件带宽资源。
  
• 多机之间 Ring，Tree 通道的连接。在多机多卡之间，协商出最终可通信的逻辑算法拓扑。
  

网络插件选择和初始化

ncclNetInit 函数通过 initPluginLibsOnceFunc 实现插件选择。如果没有指定其它插件，NCCL 默认添加两个内部网络插件，分别是：IB（InfiniBand verbs） 和 Socket（TCP/IP）。

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             commAlloc
5.                 ncclNetInit(struct ncclComm* comm)

复制代码

1. ncclResult_t ncclNetInit(struct ncclComm* comm) {
2.   ......
3.   std::call_once(initPluginLibsOnceFlag, initPluginLibsOnceFunc);
4.   ......
5.   for (int pluginIndex = 0; pluginIndex < pluginCount; pluginIndex++) {
7.     ......
8.     if ((netPluginLibs[pluginIndex].ncclNetPluginState >= ncclNetPluginStateInitReady)
9.         && (!comm->config.netName || (strcasecmp(comm->config.netName, netPluginLibs[pluginIndex].ncclNet->name) == 0))) {
10.       // plugin init must be done by all comms to setup the context, therefore we use ">="
11.       NCCLCHECK(ncclNetPluginInit(comm, &netPluginLibs[pluginIndex]));
13.       if (netPluginLibs[pluginIndex].ncclNetPluginState == ncclNetPluginStateEnabled) {
14.         bool isAssigned = false;
15.         NCCLCHECK(ncclNetPluginAssignToComm(comm, pluginIndex, &isAssigned));
16.       }
17.       ......
18.     }
19.   }
20.   ......
21. }

复制代码

什么是 libibverbs？
一个开源用户态库，支持多种 RDMA 硬件。这些 RDMA 硬件可以将相关能力通过 libibverbs 暴露出来，供上层的 NCCL 使用。
NCCL 的 IB 插件是什么？
IB 插件是 NCCL 内置的一个网络插件，可以理解为是对 libibverbs 相关 API 的再次封装。可以实现对 libibverbs 众多 API 的调用，从而最终操作底层的 InfiniBand / RoCE 硬件。NCCL 通过 IB 插件调用这些 libibverbs API，可以在不同 GPU 之间实现 RDMA 的连接建立，以及利用 RDMA 进行数据收发等操作。

1. static void initPluginLibsOnceFunc() {
2.   ......
3.   // Add 2 internal ib and socket plugins
4.   netPluginLibs[pluginCounter].ncclNet = &ncclNetIb;
5.   netPluginLibs[pluginCounter].ncclNetPluginState = ncclNetPluginStateInitReady;
6.   ++pluginCounter;
8.   netPluginLibs[pluginCounter].ncclNet = &ncclNetSocket;
9.   netPluginLibs[pluginCounter++].ncclNetPluginState = ncclNetPluginStateInitReady;
10.   ......
11. }

复制代码

ncclNetInit 函数通过 ncclNetPluginInit 实现插件初始化。通过调用：ncclNet->init 即 ncclIbInit 最终实现 IB 插件初始化。

1. static ncclResult_t ncclNetPluginInit(struct ncclComm* comm, netPluginLib_t* pluginLib) {
2.   ......
3.   if (pluginLib->ncclNetPluginState >= ncclNetPluginStateInitReady && pluginLib->ncclNet) {
4.     ncclNetCommConfig_t commConfig = {};
5.     commConfig.trafficClass = comm->config.trafficClass == NCCL_CONFIG_UNDEF_INT ? NCCL_NET_TRAFFIC_CLASS_UNDEF : comm->config.trafficClass;
6.     if (pluginLib->ncclNet->init(&comm->netContext, comm->commHash, &commConfig, ncclDebugLog, ncclProfilerCallback) != ncclSuccess) goto fail;
7.   }
8.   ......
9. }

复制代码

其中 IB 插件有如下函数定义：

1. ncclNet_t ncclNetIb = {
2.   "IB",
3.   ncclIbInit,
4.   ncclIbDevices,
5.   ncclIbGetProperties,
6.   ncclIbListen,
7.   ncclIbConnect,
8.   ncclIbAccept,
9.   ncclIbRegMr,
10.   ncclIbRegMrDmaBuf,
11.   ncclIbDeregMr,
12.   ncclIbIsend,
13.   ncclIbIrecv,
14.   ncclIbIflush,
15.   ncclIbTest,
16.   ncclIbCloseSend,
17.   ncclIbCloseRecv,
18.   ncclIbCloseListen,
19.   NULL /* getDeviceMr */,
20.   NULL /* irecvConsumed */,
21.   ncclIbMakeVDevice,
22.   ncclIbFinalize,
23.   ncclIbSetNetAttr,
24. };

复制代码

通过 IB 插件定义的上述函数，NCCL 可以实现有关 RDMA 的一系列操作。
设备初始化和信息获取
函数作用说明"IB"插件名称用于日志和 NCCL_NET=IB 匹配ncclIbInit全局初始化初始化 IB 插件一次（如加载驱动、枚举设备、创建上下文）。只在第一个 comm 初始化时调用。ncclIbDevices返回可用设备数返回本机支持的 IB/RoCE 设备数量（如 2 个 mlx5_0/1）。ncclIbGetProperties获取设备属性查询某个设备的详细信息：
• 带宽（speed）
• 端口状态
• 是否支持 GDR（GPUDirect RDMA）
• PCIe 路径（用于 NUMA 亲和性）连接建立
函数作用说明ncclIbListen监听连接在指定设备上创建一个监听 socket（实际是创建 QP 并进入 INIT/SQD 状态），等待对端连接。用于 bootstrap 阶段 root 节点。ncclIbConnect主动连接客户端调用，向对端的 listen 地址发起连接（QP 状态迁移 + 交换 QP 信息）。ncclIbAccept接受连接服务端在收到 connect 请求后，完成连接建立（最终 QP 进入 RTS 状态）。内存注册和数据传输
函数作用说明ncclIbRegMr注册内存区域 (Memory Registration)将一段 CPU/GPU 内存注册到 IB 网卡，使其可被 RDMA 操作（关键！RDMA 必须先注册内存）。ncclIbRegMrDmaBuf注册 DMA-BUF 内存用于 Linux DMA-BUF 框架（如 GPU 显存通过 dmabuf 导出），支持更现代的 GDR 注册方式。ncclIbDeregMr注销内存释放之前注册的内存区域。ncclIbIsend异步发送发起一个 RDMA Write 或 Send 操作（非阻塞），返回请求句柄。ncclIbIrecv异步接收发起一个 RDMA Read 或 Recv 操作（非阻塞）。资源清理相关
函数作用说明ncclIbIflush刷新写操作（可选）确保之前的写操作对远程端可见（某些硬件需要显式 flush）。IB 通常不需要，故可能为空或空操作。ncclIbTest检查操作是否完成查询 Isend/Irecv 返回的请求是否已完成（轮询 CQ 完成队列）。ncclIbCloseSend关闭发送端释放与发送相关的资源（如 QP、CQ）。ncclIbCloseRecv关闭接收端释放接收资源。ncclIbCloseListen关闭监听端释放 listen 创建的 QP。NCCL 成功初始化 IB 插件后，会将 IB 绑定到当前通信域结构体：comm

1. static ncclResult_t ncclNetPluginAssignToComm(struct ncclComm* comm, int pluginIndex, bool* isAssigned) {
2.   ......
3.   if (netPluginLibs[pluginIndex].ncclNetPluginState >= ncclNetPluginStateEnabled) {
4.     comm->ncclNet = netPluginLibs[pluginIndex].ncclNet;
5.     INFO(NCCL_INIT|NCCL_NET, "Assigned NET plugin %s to comm", netPluginLibs[pluginIndex].ncclNet->name);
6.   ......
7. }

复制代码

上述逻辑对应如下日志：

1. ncclIbInitDevices:702 NCCL INFO NCCL_IB_HCA set to mlx5_0,mlx5_1
2. 3045.094624 ncclIbInitDevices:758 NCCL TRACE NET/IB: Device mlx5_0 does not support Data Direct DMA.
3. ncclIbInitDevices:801 NCCL INFO NET/IB: [0] mlx5_0:uverbs0:1/RoCE provider=Mlx5 speed=25000 context=0x5f2592e8ef00 pciPath=/sys/devices/pci0000:00/0000:00:00.0 ar=0
4. ncclIbMakeVDeviceInternal:649 NCCL INFO NET/IB : Made virtual device [0] name=mlx5_0 speed=25000 ndevs=1
5. 3048.209479 ncclIbInitDevices:758 NCCL TRACE NET/IB: Device mlx5_1 does not support Data Direct DMA.
6. ncclIbInitDevices:801 NCCL INFO NET/IB: [1] mlx5_1:uverbs1:1/RoCE provider=Mlx5 speed=25000 context=0x5f2592ed05d0 pciPath=/sys/devices/pci0000:00/0000:00:00.0 ar=0
7. ncclIbMakeVDeviceInternal:649 NCCL INFO NET/IB : Made virtual device [1] name=mlx5_1 speed=25000 ndevs=1
8. ncclIbInitDevices:839 NCCL INFO NET/IB : Using [0]mlx5_0:1/RoCE [1]mlx5_1:1/RoCE [RO]; OOB ens3:10.10.10.11<0>
9. ncclNetPluginInit:187 NCCL INFO Initialized NET plugin IB
10. commAlloc:428 NCCL INFO Using network IB

复制代码

控制环形拓扑的建立

NCCL 使用 bootstrapInit 实现控制环形拓扑的协商和建立。

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             bootstrapInit(int nHandles, void* handles, struct ncclComm* comm)

复制代码

在 NCCL 控制环初始化阶段：

• 每个 MPI 进程创建一个 TCP 监听 socket1，这个 socket 只允许邻居节点的连接，并交换必要的连接信息
  
• 每个 MPI 进程创建第二个 TCP 监听 socket2，这个 socket 只允许 MPI 根进程的连接，用于传递 MPI 邻居信息
  
• 每个 MPI 进程分别向 root MPI 进程发送自身的连接信息（root 节点通过校验 UniqueId（其实就是 UniqueId 里面的 magic），来决定是否接受连接）
  
• 通过 MPI root 进程的 bootstrapRoot 函数逻辑，每个 MPI 进程分别从 root MPI 进程获得自己的后继节点信息（MPI root 进程通过连接每个 MPI 进程的 socket2 发送）
  
• 每个 MPI 进程向自己的后继节点发起主动连接（指向后继节点的socket1），并接受自己前驱节点的连接。到此，每个 rank 已经和下一个 rank 建立了直接连接，形成一个控制逻辑环
  
• 每个 MPI 进程分别再创建 3 个监听 socket，分别是：peerProxyAddresses，peerProxyAddressesUDS，和 peerP2pAddresses
  
  
  
  • peerProxyAddresses。NCCL 使用 代理线程（proxy thread） 来处理一些异步或需要 CPU 参与的任务，例如：
    
    
    
    • 注册/注销 GPU 内存（尤其在 IB/RDMA 场景）
      
    • 处理共享内存（shm）或 CUDA IPC
      
    • 管理连接生命周期
      
    • 每个 rank 启动时会创建一个 本地 proxy 服务线程，并监听一个 TCP 端口
      
    • 其他 rank 若需请求该 rank 的 proxy 服务（如 “请帮我注册这块 GPU 内存”），就连接到 socket：peerProxyAddresses
      
    
    
  • peerProxyAddressesUDS。当多个 rank 运行在 同一台机器时，可以通过：peerProxyAddressesUDS，来加速通信
    
  • peerP2pAddresses。用于 rank 之间建立直接的点对点（P2P）连接，实现任意两个 rank 能直接通信。例如：ncclSend / ncclRecv（点对点通信 API）
    
  
  
• 每个 MPI 进程通过 AllGather 算法，在建立的控制环形拓扑上广播交换 3 个socket：peerProxyAddresses，peerProxyAddressesUDS，和 peerP2pAddresses 的地址信息，使得每个 rank 上都知道其它任意 rank 的 3 个 socket 通信地址
  

1. ncclResult_t bootstrapInit(int nHandles, void* handles, struct ncclComm* comm) {
2.   ······
3.   } else {
4.     // create socket for ring neightbor to contact mee
5.     NCCLCHECK(createListenSocket(comm, comm->magic, &STATE_LISTEN(state, socket), &info.connectInfo.addr, ncclSocketTypeBootstrap));
6.   }
7.   // Create socket for root to contact me using the root's magic
8.   NCCLCHECK(createListenSocket(comm, BOOTSTRAP_HANDLE(handles, curr_root)->magic, &listenSockRoot, &info.listenRootAddress, ncclSocketTypeBootstrap));
10.   ······
11.   NCCLCHECK(sendToRoot(BOOTSTRAP_HANDLE(handles, curr_root), comm, &info));
13.   // get info on my "next" rank in the bootstrap ring from root
14.   NCCLCHECK(ncclSocketInit(&sock));
15.   NCCLCHECK(ncclSocketAccept(&sock, &listenSockRoot));
16.   NCCLCHECK(socketRecv(&sock, &nextPeer, sizeof(nextPeer)));
18.   ······
19.   // accept and connect the ring network
20.   NCCLCHECK(socketRingConnect(&nextPeer.addr, &STATE_RING(state, socket.send), &STATE_LISTEN(state, socket), &STATE_RING(state, socket.recv), comm->magic, state->abortFlag));
22.   ······
23.   // AllGather all listen handlers
24.   // in case of failure, those resources will be free'd when calling bootstrapDestroy, so we can return immediatly
25.   NCCLCHECK(ncclCalloc(&state->peerProxyAddresses, nranks));
26.   NCCLCHECKGOTO(createListenSocket(comm, comm->magic, proxySocket, state->peerProxyAddresses + rank, ncclSocketTypeProxy), result, fail);
28.   NCCLCHECKGOTO(ncclCalloc(&state->peerProxyAddressesUDS, nranks), result, fail);
29.   NCCLCHECKGOTO(getUDS(state->peerProxyAddressesUDS + rank), result, fail);
31.   // Create the service proxy and get the UDS
32.   NCCLCHECKGOTO(ncclProxyInit(comm, proxySocket, state->peerProxyAddresses, state->peerProxyAddressesUDS), result, fail);
34.   // create a socket for others to reach out (P2P)
35.   union ncclSocketAddress peerSocketAddress;
36.   NCCLCHECKGOTO(createListenSocket(comm, comm->magic, &STATE_LISTEN(state, peerSocket), &peerSocketAddress, ncclSocketTypeBootstrap), result, fail);
37.   NCCLCHECKGOTO(ncclCalloc(&state->peerP2pAddresses, nranks), result, fail);
38.   memcpy(state->peerP2pAddresses + rank, &peerSocketAddress, sizeof(union ncclSocketAddress));
40.   NCCLCHECKGOTO(ringAllInfo(comm, state, state->peerP2pAddresses, state->peerProxyAddresses, state->peerProxyAddressesUDS, rasRanks), result, fail);
41.   ······
42. }

复制代码

上述逻辑在 rank0 中对应如下日志：

1. bootstrapInit:656 NCCL TRACE rank 0 nranks 4ncclSocketListen:421 NCCL TRACE Listening on socket 10.10.10.11<36241>ncclSocketListen:421 NCCL TRACE Listening on socket 10.10.10.11<36241>ncclSocketConnect:718 NCCL TRACE Connecting to socket 10.10.10.11ncclSocketConnect:718 NCCL TRACE Connecting to socket 10.10.10.11ncclSocketListen:421 NCCL TRACE Listening on socket 10.10.10.11<36241>ncclSocketListen:421 NCCL TRACE Listening on socket 10.10.10.11<36241>ncclIpcSocketInit:51 NCCL TRACE UDS: Creating socket /tmp/nccl-socket-0-c5b60c2e5e7a9a0cbootstrapAllGather:1074 NCCL TRACE rank 0 nranks 4 size 120 - AllGathersocketRingAllGather:1032 NCCL TRACE socketRingAllGather started: rank=0 nranks=4socketRingAllGather:1034 NCCL TRACE bidirectional bootstrap: totalSteps=2socketRingAllGather:1064 NCCL TRACE socketRingAllGather first message in 0.000115 (1.039933 MB/sec), rest in 0.000017 (21.719457 MB/sec)bootstrapAllGather:1085 NCCL TRACE bootstrapAllGather for 120 B done in 0.000141 sec: 3.404569 MB/secbootstrapAllGather:1086 NCCL TRACE rank 0 nranks 4 size 120 - AllGather DONEbootstrapInit:776 NCCL TRACE rank 0 nranks 4 - DONE

复制代码

以上过程，可总结为如下图所示：

NCCL传输层初始化

NCCL 使用 initTransportsRank 实现传输层初始化。

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             initTransportsRank(struct ncclComm* comm, struct ncclComm* parent, uint64_t timers[TIMERS_INIT_COUNT])

复制代码

控制通道交换邻居信息

交换信息有如下定义，例如 NCCL 版本，主机唯一标识：hostHash，GPU 计算能力：cudaCompCap，等等

1. struct ncclPeerInfo {
2.   int rank;
3.   int cudaDev;
4.   int nvmlDev;
5.   int gdrSupport;
6.   uint64_t hostHash;
7.   uint64_t pidHash;
8.   dev_t shmDev;
9.   int64_t busId;
10.   struct ncclComm* comm;
11.   int cudaCompCap;
12.   size_t totalGlobalMem;
14.   // MNNVL support
15.   nvmlGpuFabricInfoV_t fabricInfo;
16.   int cuMemSupport;
17.   int version;
18. };

复制代码

通过算法：AllGather，在控制环广播。最终每个 rank 都将包含全局邻居信息

1.   NCCLCHECKGOTO(ncclCalloc(&comm->peerInfo, nranks+1), ret, fail); // Extra rank to represent CollNet root
2.   NCCLCHECKGOTO(fillInfo(comm, comm->peerInfo+rank, comm->commHash), ret, fail);
3.   NCCLCHECKGOTO(bootstrapAllGather(comm->bootstrap, comm->peerInfo, sizeof(struct ncclPeerInfo)), ret, fail);

复制代码

版本一致性校验与全局属性推导

做一些基础检查：

• NCCL 版本一致性检查，只有全局版本一致，才允许继续协商
  
• 根据每个 rank 的 hostHash 值，推导物理机节点数
  
• GPU 绑定重复检测。同一物理机上，不能有两个 rank 绑定到同一个 GPU，只能是一个 rank 绑定一个 GPU
  

1.   for (int i = 0; i < nranks; i++) {
3.     if (comm->peerInfo[i].version != comm->peerInfo[rank].version) {
4.       ......
5.       goto fail;
6.     }
8.     if (comm->peerInfo[i].hostHash != comm->peerInfo[rank].hostHash)
9.       nNodes++;
11.     if ((i != rank) && (comm->peerInfo[i].hostHash == comm->peerInfo[rank].hostHash) && (comm->peerInfo[i].busId == comm->peerInfo[rank].busId)) {
12.       ......
13.       goto fail;
14.     }
15.   }

复制代码

同进程rank计算

NCCL 支持多种运行模式，例如：一个进程绑定一个 GPU（One Rank Per Process（ORPP））；一个进程绑定多个 GPU（Single Process Multiple Ranks（SPMR））；一个进程通过多线程绑定多个 GPU 等模式。MPI / PyTorch DDP / DeepSpeed 等分布式深度学习框架通常使用一个进程绑定一个 GPU，即 1 process  1 rank  1 GPU 的模式。
但是如果不是 ORPP 模式，则需要涉及如下代码逻辑，推导多个 rank 与对应进程的关系。 非 ORPP 模式的一些优势：

• 共享虚拟地址空间（可直接访问彼此的 CPU 内存）
  
• 同进程的 rank，集合通信使用的 ring / tree 拓扑会优先使用 Shared Memory (SHM) 或 GPU direct（P2P）进行高效通信
  
• 只有跨节点时才会使用 IB / ROCE 通信
  

1.   do {
2.     // Compute intra-process ranks
3.     int intraProcRank0 = -1, intraProcRank = -1, intraProcRanks = 0;
5.     comm->nvlsRegSupport = 1;
6.     for (int i = 0; i < nranks; i++) {
8.       if ((comm->peerInfo[i].hostHash == comm->peerInfo[rank].hostHash) &&
9.           (comm->peerInfo[i].pidHash == comm->peerInfo[rank].pidHash)) {
10.         // Rank is in same process
11.         if (intraProcRanks == 0)
12.           intraProcRank0 = i;
14.         if (i == rank)
15.           intraProcRank = intraProcRanks;
17.         intraProcRanks++;
19.         if (intraProcRank0 == rank && rank != i) {
20.           comm->peerInfo[i].comm->intraNext = comm->intraNext;
21.           comm->intraNext = comm->peerInfo[i].comm;
22.         }
23.       }
24.     }
26.     comm->intraComm0 = comm0;
27.     comm->intraRank = intraProcRank;    // 当前 rank 在同进程中的序号（0,1,2...）
28.     comm->intraRanks = intraProcRanks;  // 同进程中的总 rank 数
30.   } while(0);

复制代码

NCCL推导PCIe设备链路拓扑

目的是先构建一个有层次关系的全局 XML 拓扑，里面记录了 NCCL 需要使用的硬件信息。

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             initTransportsRank
5.                 ncclTopoGetSystem
6.                     ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct ncclXmlNode** gpuNode)

复制代码

NCCL 使用函数 ncclTopoGetSystem 构建一个描述整机关键硬件拓扑的 XML 树，拓扑将用于后续 NCCL 通信路径的构造。XML 拓扑将包含如下信息：

• 所有 GPU
  
• 所有 NIC
  
• CPU 归属的 NUMA 节点
  
• 体现 PCIe 设备之间的层级结构
  

PCIe 基础介绍

• 每个 CPU 都有自己直连的 Root Complex（简称 RC），RC 会帮助 CPU 和其他组件通信，比如 CPU 和内存，CPU 和 PCIe 系统
  
• PCIe Switch 的作用是扩展 PCIe 端口，下游可以连接 PCIe 设备或者其它 PCIe Switch，上游来的请求将被 PCIe Switch 转发，PCIe 设备可以连在 RC，也可以连在 PCIe Switch
  

假设有如下 PCIe 信息：

1. $ lspci -tv
2. -+-[0000:ff]-+-00.0  Intel Corporation Device 324c
3. |           +-00.1  Intel Corporation Device 324c
4. +-[0000:f2]-+-00.0  Intel Corporation Ice Lake Memory Map/VT-d
5. |           +-00.1  Intel Corporation Ice Lake Mesh 2 PCIe
6. +-[0000:d0]-+-00.0  Intel Corporation Ice Lake Memory Map/VT-d
7. |           +-00.1  Intel Corporation Ice Lake Mesh 2 PCIe
8. |           +-00.2  Intel Corporation Ice Lake RAS
9. |           +-00.4  Intel Corporation Device 0b23
10. |           \-01.0-[d1-d9]----00.0-[d2-d9]--+-00.0-[d3]----00.0  NVIDIA Corporation AD102GL [L20]
11. |                                           +-01.0-[d4]----00.0  NVIDIA Corporation AD102GL [L20]
12. |                                           +-02.0-[d5]----00.0  NVIDIA Corporation AD102GL [L20]
13. |                                           +-03.0-[d6]----00.0  NVIDIA Corporation AD102GL [L20]
14. |                                           +-04.0-[d7]----00.0  Broadcom / LSI Virtual PCIe Placeholder Endpoint
15. |                                           +-05.0-[d8]----00.0  Broadcom / LSI Virtual PCIe Placeholder Endpoint
16. |                                           \-1f.0-[d9]----00.0  Broadcom / LSI PCIe Switch management endpoint
17. +-[0000:95]-+-00.0  Intel Corporation Ice Lake Memory Map/VT-d
18. |           +-00.1  Intel Corporation Ice Lake Mesh 2 PCIe
19. |           +-00.2  Intel Corporation Ice Lake RAS
20. |           +-00.4  Intel Corporation Device 0b23
21. |           \-01.0-[96]--+-00.0  Mellanox Technologies MT27800 Family [ConnectX-5]
22. |                        \-00.1  Mellanox Technologies MT27800 Family [ConnectX-5]

复制代码

GPU 的完整链路信息如下：

1. 0000:d3:00.0 -> ../../../devices/pci0000:d0/0000:d0:01.0/0000:d1:00.0/0000:d2:00.0/0000:d3:00.0
2. 0000:d4:00.0 -> ../../../devices/pci0000:d0/0000:d0:01.0/0000:d1:00.0/0000:d2:01.0/0000:d4:00.0
3. 0000:d5:00.0 -> ../../../devices/pci0000:d0/0000:d0:01.0/0000:d1:00.0/0000:d2:02.0/0000:d5:00.0
4. 0000:d6:00.0 -> ../../../devices/pci0000:d0/0000:d0:01.0/0000:d1:00.0/0000:d2:03.0/0000:d6:00.0

复制代码

NIC 的完整链路信息如下：

1. 0000:96:00.0 -> ../../../devices/pci0000:95/0000:95:01.0/0000:96:00.0
2. 0000:96:00.1 -> ../../../devices/pci0000:95/0000:95:01.0/0000:96:00.1

复制代码

得到信息：

• 四个 GPU 都在相同 Root Complex：pci0000:d0 下面
  
• 0000:d0:01.0 为 Root Complex 的 PCIe Root Port
  
• 0000:d1:00.0 / 0000:d2:01.0 为某个 PCIe Switch 的上游端口和下游端口
  
• 0000:d3:00.0，0000:d4:00.0，0000:d5:00.0，0000:d6:00.0 四个 GPU 分别挂接在上游的 PCIe Switch 上
  

拓扑结构可表示为：CPU → Root Complex → Root Port → PCIe Switch Upstream → PCIe Switch Downstream → GPU。
NCCL 让每个 rank 分别从自身绑定的 GPU 开始，通过 PCIe 路径逐步向上逆向推导，从而获得自身 GPU 在整个 PCIe 链路上的路径信息，每个 rank 只探测自身的 PCIe 链路信息。

1. ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct ncclXmlNode** gpuNode) {
2.   ......
3.   // 向 xml 中添加当前 GPU node
4.   NCCLCHECK(ncclTopoGetPciNode(xml, busId, &node));
6.   ......
7.   // 当前 GPU PCIe 链路信息推导
8.   NCCLCHECK(ncclTopoGetXmlFromSys(node, xml));
9.   ......
10. }

复制代码

其中，ncclTopoGetXmlFromSys 的实现可简略为：

1. ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml* xml) {
2.   ......
3.   NCCLCHECKNOWARN(xmlGetAttrIndex(pciNode, "vendor", &index), NCCL_GRAPH);
4.   if (index == -1) {
5.     if (path) NOWARN(ncclTopoSetAttrFromSys(pciNode, path, "vendor", "vendor"), NCCL_GRAPH);
6.   }
8.   ......
9.   struct ncclXmlNode* parent = pciNode->parent;
10.   if (parent == NULL) {
11.     if (path) {
12.       // Save that for later in case next step is a CPU
13.       NCCLCHECK(ncclTopoGetStrFromSys(path, "numa_node", numaIdStr));
14.       ......
15.       // Go up one level in the PCI tree. Rewind two "/" and follow the upper PCI
16.       // switch, or stop if we reach a CPU root complex.
17.       for (parentOffset = strlen(path)-1; parentOffset>0; parentOffset--) {
18.         if (path[parentOffset] == '/') {
19.           slashCount++;
20.           path[parentOffset] = '\0';
21.           int start = parentOffset - 1;
22.           while (start>0 && path[start] != '/') start--;
23.           // Check whether the parent path looks like "BBBB:BB:DD.F" or not.
24.           if (checkBDFFormat(path+start+1) == 0) {
25.             // This a CPU root complex. Create a CPU tag and stop there.
26.             struct ncclXmlNode* topNode;
27.             NCCLCHECK(xmlFindTag(xml, "system", &topNode));
28.             NCCLCHECK(xmlGetSubKv(topNode, "cpu", &parent, "numaid", numaIdStr));
29.             ......
30.             }
31.           } else if (slashCount == 2) {
32.             // Continue on the upper PCI switch
33.             for (int i = strlen(path)-1; i>0; i--) {
34.               if (path[i] == '/') {
35.                 NCCLCHECK(xmlFindTagKv(xml, "pci", &parent, "busid", path+i+1));
36.                 ......
37.               }
38.             }
39.           }
40.         }
41.         ......
42.       }
43.     } else {
44.       ......
45.     }
46.     pciNode->parent = parent;
48.   ......
49.   if (strcmp(parent->name, "pci") == 0) {
50.     NCCLCHECK(ncclTopoGetXmlFromSys(parent, xml));
51.   } else if (strcmp(parent->name, "cpu") == 0) {
52.     NCCLCHECK(ncclTopoGetXmlFromCpu(parent, xml));
53.   }
54.   ......
55. }

复制代码

算法思想为：逐步向上逆向推导，从而获得 GPU 在整个 PCIe 链路上的硬件信息。算法总结如下：

• 先获取当前 rank 绑定的 GPU Bus ID，在 xml 中添加一个 pci 节点
  
• 调用：ncclTopoGetXmlFromSys，实现向上逐级推导。获取当前 GPU 的父节点，即通过 GPU 的 PCIe Path 推导，得到 GPU：0000:d3:00.0 的父节点为 0000:d2:00.0
  
• PCIe Switch 有上游和下游端口。因此，推导至 PCIe Switch 的上游端口 Bus ID：0000:d1:00.0。判断该 Bus ID 不是 CPU root complex，此时 slashCount = 2，因此将该 ID 作为一个 pci 节点，并添加到 xml 中
  
• 递归调用：ncclTopoGetXmlFromSys，从 PCIe Switch 上游端口 Bus ID：0000:d1:00.0 再次向上推导，重复上面的步骤，推导至 Bus ID：pci0000:d0，发现是一个 CPU root complex，向 xml 中添加该 CPU 节点，并填充 CPU 相关信息，例如 numaid， arch等。
  
• 最后填充 当前 GPU 相关信息，例如 rank ID，sm，gdr 等信息
  

1. ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct ncclXmlNode** gpuNode) {
2.   ......
3.   // 向 xml 中添加当前 GPU 相关信息
4.   NCCLCHECK(ncclTopoGetXmlFromGpu(node, nvmlDev, xml, gpuNode));
5. }

复制代码

NIC 推导也类似 GPU 的推导过程。但是 NIC 是系统级资源，并不绑定某个 rank，每个 rank 都会枚举自己可见的 NIC 资源。
由于每个 rank 只推导自己绑定的 GPU，全局视图需要汇总。因此需要通过 NCCL 控制通道在物理设备内部通过：bootstrapIntraNodeAllGather，实现 XML 信息融合和去重。经过汇总后，前面介绍的 PCIe 链路信息，可以构造出可视化的下列 xml 视图：

1. <system version="1">
2.   <cpu host_hash="0x3bd87a3604afa55d" numaid="0" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="207">
3.   
4.     <pci busid="0000:d1:00.0"  link_speed="16 GT/s" link_width="16">
6.       <pci busid="0000:d3:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
7.         <gpu dev="0" sm="89" rank="0" gdr="1"/>
8.       </pci>
9.           
10.           <pci busid="0000:d4:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
11.         <gpu dev="1" sm="89" rank="1" gdr="1"/>
12.       </pci>
14.           <pci busid="0000:d5:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
15.         <gpu dev="2" sm="89" rank="2" gdr="1"/>
16.       </pci>
18.           <pci busid="0000:d6:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
19.         <gpu dev="3" sm="89" rank="3" gdr="1"/>
20.       </pci>
22.         </pci>
23.        
24.     <pci busid="0000:96:00.0"  link_speed="16 GT/s" link_width="16">
25.       <nic>
26.         <net name="mlx5_0" dev="0" latency="0" speed="25000" port="1" guid="0x7f7afb0003797334" maxconn="131072" gdr="1"/>
27.       </nic>
28.     </pci>
29.        
30.         <pci busid="0000:96:00.1"  link_speed="16 GT/s" link_width="16">
31.       <nic>
32.         <net name="mlx5_1" dev="1" latency="0" speed="25000" port="2" guid="0x7f7afb0003797334" maxconn="131072" gdr="1"/>
33.       </nic>
34.     </pci>
35.        
36.   </cpu>
37. </system>

复制代码

在我们的虚拟机实验环境上，所有设备都挂在同一个 Root Complex 下面，且没有 PCIe Switch 参与，是最简单的拓扑。

1. $ lspci -tv
2. -[0000:00]-+-00.0  Intel Corporation 440FX - 82441FX PMC [Natoma]
3.            +-01.0  Intel Corporation 82371SB PIIX3 ISA [Natoma/Triton II]
4.            +-01.1  Intel Corporation 82371SB PIIX3 IDE [Natoma/Triton II]
5.            +-01.2  Intel Corporation 82371SB PIIX3 USB [Natoma/Triton II]
6.            +-01.3  Intel Corporation 82371AB/EB/MB PIIX4 ACPI
7.            +-02.0  Cirrus Logic GD 5446
8.            +-03.0  Red Hat, Inc. Virtio network device
9.            +-04.0  Red Hat, Inc. Virtio block device
10.            +-05.0  Red Hat, Inc. Virtio memory balloon
11.            +-06.0  Red Hat, Inc. Virtio file system
12.            +-07.0-[01]--
13.            +-08.0  NVIDIA Corporation AD102GL [L20]
14.            +-09.0  NVIDIA Corporation AD102GL [L20]
15.            +-0a.0  Mellanox Technologies MT27800 Family [ConnectX-5]
16.            \-0b.0  Mellanox Technologies MT27800 Family [ConnectX-5]
18. $ ll /sys/bus/pci/devices/
19. 0000:00:00.0 -> ../../../devices/pci0000:00/0000:00:00.0/
20. 0000:00:01.0 -> ../../../devices/pci0000:00/0000:00:01.0/
21. 0000:00:01.1 -> ../../../devices/pci0000:00/0000:00:01.1/
22. 0000:00:01.2 -> ../../../devices/pci0000:00/0000:00:01.2/
23. 0000:00:01.3 -> ../../../devices/pci0000:00/0000:00:01.3/
24. 0000:00:02.0 -> ../../../devices/pci0000:00/0000:00:02.0/
25. 0000:00:03.0 -> ../../../devices/pci0000:00/0000:00:03.0/
26. 0000:00:04.0 -> ../../../devices/pci0000:00/0000:00:04.0/
27. 0000:00:05.0 -> ../../../devices/pci0000:00/0000:00:05.0/
28. 0000:00:06.0 -> ../../../devices/pci0000:00/0000:00:06.0/
29. 0000:00:07.0 -> ../../../devices/pci0000:00/0000:00:07.0/
30. 0000:00:08.0 -> ../../../devices/pci0000:00/0000:00:08.0/
31. 0000:00:09.0 -> ../../../devices/pci0000:00/0000:00:09.0/
32. 0000:00:0a.0 -> ../../../devices/pci0000:00/0000:00:0a.0/
33. 0000:00:0b.0 -> ../../../devices/pci0000:00/0000:00:0b.0/

复制代码

最后得到的 xml 融合信息如下。实验环境中的网卡 mlx5_0 和 mlx5_1 被进行了合并。

1. <system version="1">
2.   <cpu host_hash="0x3bd87a3604afa55d" numaid="-1" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="207">
4.     <pci busid="0000:00:08.0"  vendor="0x10de" device="0x26ba" subsystem_vendor="0x10de" subsystem_device="0x1957" link_speed="16.0 GT/s PCIe" link_width="0">
5.       <gpu dev="0" sm="89" rank="0" gdr="1"/>
6.     </pci>
8.     <pci busid="0000:00:09.0"  vendor="0x10de" device="0x26ba" subsystem_vendor="0x10de" subsystem_device="0x1957" link_speed="16.0 GT/s PCIe" link_width="0">
9.       <gpu dev="1" sm="89" rank="1" gdr="1"/>
10.     </pci>
12.     <pci busid="0000:00:00.0"  vendor="0x8086" device="0x1237" subsystem_vendor="0x1af4" subsystem_device="0x1100" link_speed="" link_width="0">
13.       <nic>
14.         <net name="mlx5_0" dev="0" latency="0" speed="25000" port="1" guid="0x7f7afb0003797334" maxconn="131072" gdr="0"/>
15.         <net name="mlx5_1" dev="1" latency="0" speed="25000" port="2" guid="0x7f7afb0003797334" maxconn="131072" gdr="0"/>
16.       </nic>
17.     </pci>
19.   </cpu>
20. </system>

复制代码

上述逻辑在 rank0 中对应如下日志：

1. // PCIe 链路推导
2. 3068.938120 ncclTopoSetAttrFromSys:472 NCCL TRACE Read from sys /sys/devices/pci0000:00/0000:00:08.0/class -> class=0x030200
3. 3068.949021 ncclTopoSetAttrFromSys:472 NCCL TRACE Read from sys /sys/devices/pci0000:00/0000:00:08.0/vendor -> vendor=0x10de
4. 3068.955660 ncclTopoSetAttrFromSys:472 NCCL TRACE Read from sys /sys/devices/pci0000:00/0000:00:08.0/device -> device=0x26ba
5. ......
7. // 网卡合并
8. 3069.492280 ncclTopoMakeVNics:1314 NCCL TRACE Found physical ncclNet node 0 mlx5_0
9. 3069.494257 ncclTopoMakeVNics:1314 NCCL TRACE Found physical ncclNet node 1 mlx5_1
10. ncclIbMakeVDeviceInternal:649 NCCL INFO NET/IB : Made virtual device [2] name=mlx5_0+mlx5_1 speed=50000 ndevs=2
11. ncclTopoMakeVnic:1029 NCCL INFO TOPO/NET : Made vNic 2
12. 3069.508220 ncclTopoGetVNicParent:1270 NCCL TRACE Re-found physical ncclNet node 0 mlx5_0
13. 3069.510005 ncclTopoGetVNicParent:1270 NCCL TRACE Re-found physical ncclNet node 1 mlx5_1
14. 3069.512078 ncclTopoGetVNicParent:1295 NCCL TRACE Selected parent pci with path 1
15. ncclTopoPopulateNics:1357 NCCL INFO NET/IB : GPU Direct RDMA Disabled for HCA 2 'mlx5_0+mlx5_1'
16. 3069.519502 ncclTopoTrimXmlRec:982 NCCL TRACE Removing node net mlx5_0 (null)
17. 3069.521424 ncclTopoTrimXmlRec:982 NCCL TRACE Removing node net mlx5_1 (null)
19. // 同主机内部 XML 信息融合
20. 3072.011796 bootstrapIntraNodeAllGather:1128 NCCL TRACE rank 0 nranks 2 size 2562056 - ENTER
21. 3072.029868 ncclSocketConnect:718 NCCL TRACE Connecting to socket 10.10.10.11<46225>
22. 3073.529592 socketRingAllGather:1032 NCCL TRACE socketRingAllGather started: rank=0 nranks=2
23. 3073.533829 socketRingAllGather:1034 NCCL TRACE bidirectional bootstrap: totalSteps=1
24. 3074.657969 socketRingAllGather:1064 NCCL TRACE socketRingAllGather first message in 0.001121 (2285.429856 MB/sec), rest in 0.000000 (50236392.156863 MB/sec)
25. 3074.684796 bootstrapIntraNodeAllGather:1142 NCCL TRACE rank 0 nranks 2 size 2562056 - DONE

复制代码

NCCL对PCIe节点建图

通过 PCIe 链路推导得到了 GPU 到 CPU ， NIC 到 CPU，以及 GPU 到 NIC 之间的基于 xml 的拓扑信息，但是还缺乏它们之间的链路连接关系。因此需要通过建图，将 PCIe 设备通过 PCIe 路径连接起来。就比如有多个孤立岛，现在按照路径规划（实际的 PCIe 路径）在它们之间建立桥梁，实现互通，从而为 NCCL 后续最优通信路径的搜索打基础。

NCCL 使用 ncclTopoGetSystemFromXml 在 PCIe 设备之间建图。

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             initTransportsRank
5.                 ncclTopoGetSystem
6.                     ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem, const uint64_t localHostHash)

复制代码

NCCL 定义了六种设备类型，分别是：NCCL_TOPO_NODE_TYPES = [GPU，PCI，NVS，CPU，NIC，NET]，每种设备类型支持最大 NCCL_TOPO_MAX_NODES = 576 个节点，每个节点支持最多 NCCL_TOPO_MAX_LINKS = 576 个 PCIe 连接。其中，struct ncclTopoNode 里面的 struct ncclTopoLink links[NCCL_TOPO_MAX_LINKS]，用于记录当前 PCIe 设备节点的路径信息，也就是记录从它可达其它 PCIe 设备的路径信息。

1. struct ncclTopoNode {
2.   ......
3.   // Type specific data
4.   union {
5.     struct {
6.       int dev; // NVML dev number
7.       int rank;
8.       ......
9.     }gpu;
10.     struct {
11.       int dev; // Plugin dev number
12.       uint64_t pciId;
13.       ......
14.     }net;
15.     struct {
16.       int arch;
17.       int vendor;
18.       ......
19.     }cpu;
20.     struct {
21.       uint64_t device;
22.     }pci;
23.   };
25.   int nlinks;
26.   struct ncclTopoLink links[NCCL_TOPO_MAX_LINKS];
28.   // Pre-computed paths to GPUs and NICs
29.   struct ncclTopoLinkList* paths[NCCL_TOPO_NODE_TYPES];
30.   ......
31. };
33. struct ncclTopoNodeSet {
34.   int count;
35.   struct ncclTopoNode nodes[NCCL_TOPO_MAX_NODES];
36. };
38. struct ncclTopoSystem {
39.   int systemId;
40.   uint64_t hostHashes[NCCL_TOPO_MAX_NODES];
41.   int nHosts;
42.   struct ncclTopoNodeSet nodes[NCCL_TOPO_NODE_TYPES];
43.   float maxBw;
44.   float totalBw;
45.   int inter;
46. };
48. struct ncclTopoLink {
49.   int type;
50.   float bw;
51.   struct ncclTopoNode* remNode;
52. };
54. #define NCCL_TOPO_NODE_TYPES 6
55. #define NCCL_TOPO_MAX_NODES 576
56. #define NCCL_TOPO_MAX_LINKS 576
58. #define GPU 0
59. #define PCI 1
60. #define NVS 2
61. #define CPU 3 // Actually NUMA domains
62. #define NIC 4
63. #define NET 5

复制代码

假设有 XML 视图：

1. <system version="1">
2.   <cpu host_hash="0x3bd87a3604afa55d" numaid="0" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="207">
3.     <pci busid="0000:d1:00.0"  link_speed="16 GT/s" link_width="16">
5.       <pci busid="0000:d3:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
6.         <gpu dev="0" sm="89" rank="0" gdr="1"/>
7.       </pci>
8.           
9.           <pci busid="0000:d4:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
10.         <gpu dev="1" sm="89" rank="1" gdr="1"/>
11.       </pci>
13.         </pci>
14.   </cpu>
17.   <cpu host_hash="0x3bd87a3604afa55d" numaid="1" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="207">
18.     <pci busid="0000:ef:00.0"  link_speed="16 GT/s" link_width="16">
20.       <pci busid="0000:78:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
21.         <gpu dev="0" sm="89" rank="2" gdr="1"/>
22.       </pci>
23.           
24.           <pci busid="0000:79:00.0"  link_speed="16.0 GT/s PCIe" link_width="16">
25.         <gpu dev="1" sm="89" rank="3" gdr="1"/>
26.       </pci>
28.         </pci>
29.   </cpu>
30. </system>

复制代码

通过 ncclTopoAddCpu，从 XML 拓扑的 cpu 开始建图。遍历 XML 的 cpu 节点（如果是多 NUMA 环境，XML 里将是多个 CPU -> PCI 的视图）。

1. ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem, const uint64_t localHostHash) {
2.   ......
3.   for (int s=0; s<topNode->nSubs; s++) {
4.     struct ncclXmlNode* node = topNode->subs[s];
5.     if (strcmp(node->name, "cpu") == 0) NCCLCHECK(ncclTopoAddCpu(node, *topoSystem));
6.   }
8.   ......
9. }

复制代码

从 XML 里遍历当前 cpu 下的 pci 节点。通过 ncclTopoAddPci 去建立 CPU 和 PCIe Switch，以及 PCIe Switch 与 GPU 或者 NIC 的连接。

1. ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* system) {
2.   ......
3.   for (int s=0; s<xmlCpu->nSubs; s++) {
4.     struct ncclXmlNode* node = xmlCpu->subs[s];
5.     if (strcmp(node->name, "pci") == 0) NCCLCHECK(ncclTopoAddPci(node, system, cpu, systemId, numaId));
6.     if (strcmp(node->name, "nic") == 0) {
7.       ......
8.       NCCLCHECK(ncclTopoAddNic(node, system, nic, systemId));
9.     }
10.   }
11.   ......
12. }

复制代码

在 ncclTopoAddPci 里，连接建立的算法思想可以总结如下：

• 如果当前节点就是一个 PCIe 设备，那么通过 ncclTopoCreateNode 将当前节点添加到 struct ncclTopoNodeSet nodes[NCCL_TOPO_NODE_TYPES]，并调用两次 ncclTopoConnectNodes 分别去建立当前 PCIe 设备 与 parent CPU 的连接，以及建立 parent CPU 与 当前 PCIe 设备的连接
  
• 如果当前节点是一个 PCIe Switch，且 Switch 下面还有其它 PCIe 设备，那么继续遍历 Switch 下面的 PCIe 设备，并递归调用 ncclTopoAddPci。当前层次的 PCIe Switch 与 上层 parent CPU 的正反连接在最后递归回溯时再执行
  
• 通过 ncclTopoAddPci 递归，直至达到最底层 PCIe 设备，比如 GPU / NIC，那么将当前节点添加到 struct ncclTopoNodeSet nodes[NCCL_TOPO_NODE_TYPES]，并调用两次 ncclTopoConnectNodes 分别去建立当前 GPU / NIC 与 parent PCIe Switch 的连接，以及建立 parent PCIe Switch 与 当前 GPU / NIC 设备的连接。注意，这里的连接类型是：LINK_PCI，表示通过 PCIe 链路连接。
  

1. ncclResult_t ncclTopoAddPci(struct ncclXmlNode* xmlPci, struct ncclTopoSystem* system, struct ncclTopoNode* parent, int systemId, int numaId) {
2.     ......
3.   } else if (type == PCI) {
4.     NCCLCHECK(ncclTopoCreateNode(system, &node, type, NCCL_TOPO_ID(systemId, busId)));
5.     NCCLCHECK(xmlGetAttr(xmlPci, "vendor", &str));
6.     if (str) node->pci.device += strtol(str, NULL, 0) << 48;
7.     NCCLCHECK(xmlGetAttr(xmlPci, "device", &str));
8.     if (str) node->pci.device += strtol(str, NULL, 0) << 32;
9.     ......
11.     for (int s=0; s<xmlPci->nSubs; s++) {
12.       struct ncclXmlNode* xmlSubPci = xmlPci->subs[s];
13.       if (strcmp(xmlSubPci->name, "pcilink") != 0) { // PCI links will be added later
14.         NCCLCHECK(ncclTopoAddPci(xmlSubPci, system, node, systemId, numaId));
15.       }
16.     }
17.   }
19.   ......
20.   if (node) {
21.     ......
22.     NCCLCHECK(ncclTopoConnectNodes(node, parent, LINK_PCI, width*speed/80.0));
23.     NCCLCHECK(ncclTopoConnectNodes(parent, node, LINK_PCI, width*speed/80.0));
24.   }
25.   return ncclSuccess;
26. }

复制代码

向 node->links 数组里添加 remNode 节点，实现两个节点的连接。调用两次就表示相互互联。

1. ncclResult_t ncclTopoConnectNodes(struct ncclTopoNode* node, struct ncclTopoNode* remNode, int type, float bw) {
2.   
3.   ......
4.   struct ncclTopoLink* link;
5.   for (link = node->links; link - node->links != NCCL_TOPO_MAX_LINKS && link->remNode; link++) {
6.     if (link->remNode == remNode && link->type == type) break;
7.   }
9.   if (link->remNode == NULL) node->nlinks++;
10.   link->type = type;
11.   link->remNode = remNode;
12.   link->bw += bw;
14.   // Sort links in BW descending order
15.   struct ncclTopoLink linkSave;
16.   memcpy(&linkSave, link, sizeof(struct ncclTopoLink));
17.   while (link != node->links) {
18.     if ((link-1)->bw >= linkSave.bw) break;
19.     memcpy(link, link-1, sizeof(struct ncclTopoLink));
20.     link--;
21.   }
22.   memcpy(link, &linkSave, sizeof(struct ncclTopoLink));
23.   ......
24. }

复制代码

回到 ncclTopoGetSystemFromXml 最外层：

• 通过 ncclTopoAddCpu，NCCL 实现了对当前实体节点上众多 GPU，PCI，NVS，CPU，NIC，NET 设备的建图
  
  
  
  • 任意 PCIe 设备都记录了它自身周边通过 PCIe 链路的可达信息
    
  • 每个 PCIe 设备都只知道自己周边有哪些下一跳节点，但是只有一层可达信息
    
  • 每个 PCIe 设备的连接都是双向的，即 parent  child 模式，连接类型为：LINK_PCI
    
  • 每条连接都被赋值了计算出来的带宽值
    
  • 每个 PCIe 设备增加一个连接，就要按照连接带宽进行一次排序，将带宽最大的连接放在 links 数组的最前面
    
  
  
• 通过 ncclTopoAddNvLinks，构建 GPU 和 GPU 之间的 GPU0  GPU1 连接，或者基于 NVSwitch 的 GPU0  NVSwitch  GPU1 连接。连接类型为：LINK_NVL
  
• 通过 ncclTopoAddPciLinks，构建 PCIe 设备之间的 pcilink 连接（如果存在），连接类型为：LINK_LOC。pcilink 不是 PCIe 链路，是一种本地直连技术。
  最终，通过 ncclTopoGetSystemFromXml，。
  
• 通过 ncclTopoConnectCpus，构建 CPU 之间的两两连接，连接类型为：LINK_SYS
  
• 最后通过 ncclTopoSortSystem，将以上所有 PCIe 设备的连接数组 links，按照 NVLink -> PCI-down -> PCI-up -> SYS 的顺序排列。目的是在后续使用 DFS 进行路径搜索时，避免造成路径回退，确保搜索路径是在往目的 PCIe 节点前进。
  

1. ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem, const uint64_t localHostHash) {
2.   ......
3.   for (int s=0; s<topNode->nSubs; s++) {
4.     struct ncclXmlNode* node = topNode->subs[s];
5.     if (strcmp(node->name, "cpu") == 0) NCCLCHECK(ncclTopoAddCpu(node, *topoSystem));
6.   }
8.   ......
9.   NCCLCHECK(ncclTopoAddNvLinks(topNode, *topoSystem, NULL, 0));
10.   NCCLCHECK(ncclTopoAddC2c(topNode, *topoSystem, NULL, 0));
11.   NCCLCHECK(ncclTopoAddPciLinks(topNode, *topoSystem, NULL, 0));
13.   NCCLCHECK(ncclTopoFlattenBcmSwitches(*topoSystem));
14.   NCCLCHECK(ncclTopoConnectCpus(*topoSystem));
15.   NCCLCHECK(ncclTopoSortSystem(*topoSystem));
16. }

复制代码

NCCL 定义的 PCIe 设备间的连接类型

1. #define LINK_LOC 0
2. #define LINK_NVL 1
3. // Skipping 2 for PATH_NVB
4. #define LINK_C2C 3
5. #define LINK_PCI 4
6. // Skipping 5 for PATH_PXB
7. // Skipping 6 for PATH_PXN
8. // Skipping 7 for PATH_P2C
9. // Skipping 8 for PATH_PHB
10. #define LINK_SYS 9
11. #define LINK_NET 10

复制代码

无 NVSwitch，只有 NVLink 直连的拓扑图：

有 NVSwitch 硬件的拓扑图：

但是需要注意，到目前为止，相关 PCIe 设备只有自身到周边一跳设备的连接关系。
NCCL路径计算

NCCL 已经建立了本节点上所有 PCIe 设备的连接图，但是还不知道该如何从某个 PCIe 设备到达另外一个 PCIe 设备。所以就需要计算 GPU, PCI, NVS, CPU, NIC, NET 它们之间的最优互通路径。例如需要计算 GPU 与 NIC 之间的最优路径，拓扑如果足够复杂，GPU 与 NIC 之间的互通路径可能有多条，NCCL 会从路径跳数，或者路径带宽去衡量，从而计算出一条最优路径。通过计算任意设备之间的最优互通路径，为后续 NCCL 进行多通道计算打基础。

NCCL 使用 ncclTopoComputePaths 计算多个 PCIe 设备间的最优可达路径。

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             initTransportsRank
5.                 ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm* comm)

复制代码

ncclTopoComputePaths 的前半部分实现如下：

• 从指定的 CPU 节点出发，计算其它任意 PCIe（GPU, PCI, NVS, NET）设备达到当前 CPU 的最优路径
  
• 从指定的 GPU 节点出发，计算其它任意 PCIe（CPU, PCI, NVS, NET）设备达到当前 GPU 的最优路径
  
• 从指定的 NET 节点出发，计算其它任意 PCIe（GPU, PCI, NVS, CPU）设备达到当前 NET 的最优路径
  
• 从指定的 NVS 节点出发，计算其它任意 PCIe（GPU, PCI, NET, CPU）设备达到当前 NVS 的最优路径
  

1. ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm* comm) {
2.   // Precompute paths between GPUs/NICs.
4.   // Remove everything in case we're re-computing
5.   ncclTopoRemovePaths(system);
7.   // Set direct paths to CPUs. We need them in many cases.
8.   for (int c=0; c<system->nodes[CPU].count; c++) {
9.     NCCLCHECK(ncclTopoSetPaths(system->nodes[CPU].nodes+c, system));
10.   }
12.   // Set direct paths to GPUs.
13.   for (int g=0; g<system->nodes[GPU].count; g++) {
14.     NCCLCHECK(ncclTopoSetPaths(system->nodes[GPU].nodes+g, system));
15.   }
17.   // Set direct paths to NICs.
18.   for (int n=0; n<system->nodes[NET].count; n++) {
19.     NCCLCHECK(ncclTopoSetPaths(system->nodes[NET].nodes+n, system));
20.   }
22.   // Set direct paths to NVSwitches.
23.   for (int n=0; n<system->nodes[NVS].count; n++) {
24.     NCCLCHECK(ncclTopoSetPaths(system->nodes[NVS].nodes+n, system));
25.   }
26.   ......
27. }

复制代码

注意，是从指定设备出发，计算其它任意 PCIe 设备达到当前指定设备的最优路径。意味着在路径探索过程中，会在其它任意中间节点添加到达当前指定的，出发节点的路径信息。
比如有 PCIe 如下拓扑：

• 从 GPU0 出发时，在 PCI_SW1，PCI_SW2，GPU1 中都会添加自身到达 GPU0 的最优路径信息
  
• 从 GPU1 出发时，在 PCI_SW2，PCI_SW1，GPU0 中都会添加自身到达 GPU1 的最优路径信息
  

ncclTopoSetPaths 是实现路径计算的关键函数，通过 BFS 逐层搜索其它 PCIe 设备实现。

1. static ncclResult_t ncclTopoSetPaths(struct ncclTopoNode* baseNode, struct ncclTopoSystem* system) {
3.   ......
4.   nodeList.count = 1;
5.   nodeList.list[0] = baseNode;
7.   // 将出发节点 baseNode 初始化
8.   NCCLCHECK(getPath(system, baseNode, baseNode->type, baseNode->id, &basePath));
9.   basePath->count = 0;
10.   basePath->bw = LOC_BW;
11.   basePath->type = PATH_LOC;
13.   while (nodeList.count) {
14.     nextNodeList.count = 0;
16.     for (int n=0; n<nodeList.count; n++) {
18.       struct ncclTopoNode* node = nodeList.list[n];
19.       NCCLCHECK(getPath(system, node, baseNode->type, baseNode->id, &path));
21.       for (int l=0; l<node->nlinks; l++) {
22.         struct ncclTopoLink* link = node->links+l;
23.         struct ncclTopoNode* remNode = link->remNode;
24.         ......
26.         NCCLCHECK(getPath(system, remNode, baseNode->type, baseNode->id, &remPath));
27.         float bw = std::min(path->bw, link->bw);
28.         ......
30.         // 如果remPath->bw == 0，直接进入
31.         // 如果remPath->bw != 0，比较跳数。当前路径跳数：path->count，小于 remPath->count 时，再比较带宽
32.         // 如果最后当前路径带宽高于之前计算的 remPath->bw，尝试更新
33.         if ((remPath->bw == 0 || remPath->count > path->count) && remPath->bw < bw) {
34.           // Find reverse link
35.           for (int l=0; l<remNode->nlinks; l++) {
36.             if (remNode->links[l].remNode == node && remNode->links[l].type == link->type) {
37.               // 将 remPath->list[0] 的第一个连接指向当前 node
38.               remPath->list[0] = remNode->links+l;
39.               break;
40.             }
41.           }
42.           ......
44.           // 将剩余的，在前面已经得到的路径拷贝到 remPath->list
45.           for (int i=0; i<path->count; i++) remPath->list[i+1] = path->list[i];
47.           // 赋值路径跳数和带宽
48.           remPath->count = path->count + 1;
49.           remPath->bw = bw;
51.           ......
52.           // 赋值路径类型
53.           if (link->type == LINK_PCI && (node->type == CPU || link->remNode->type == CPU)) type = PATH_PHB;
55.           ......
56.           remPath->type = std::max(path->type, type);
58.           // 寻找 nextNodeList 中是否已存在 remNode
59.           // 如果存在，则不会将 remNode 添加至 nextNodeList
60.           // 如果不存在，则将 remNode 添加至 nextNodeList
61.           for (i=0; i<nextNodeList.count; i++) if (nextNodeList.list[i] == remNode) break;
62.           if (i == nextNodeList.count) nextNodeList.list[nextNodeList.count++] = remNode;
63.         }
64.       }
65.     }
67.     // 将 nextNodeList 拷贝至 nodeList ，使用 nodeList 重新遍历
68.     memcpy(&nodeList, &nextNodeList, sizeof(nodeList));
69.   }
70.   ......
71. }

复制代码

节点 GPU0 到 GPU0 的路径为 LOCAL
节点到 GPU0 的路径路由跳数GPU0PATH_LOC0用于循环的 nodeList 只有 GPU0

1. basePath->count = 0
2. basePath->bw = LOC_BW
3. basePath->type = PATH_LOC

复制代码

• 遍历 GPU0
  

GPU0 有两个连接，分别是：PCI_SW1 和 GPU1

1. nodeList = { GPU0 }

复制代码

• baseNode 初始化
  
• 用于循环的节点链表为： nodeList =
  节点到 GPU0 的路径路径跳数GPU0PATH_LOC0
  
• GPU0 有两个连接，分别是：PCI_SW1 和 GPU1
  
  
  
  • PCI_SW1 的 remPath->list[0] 将被赋值为 GPU0
    
  • GPU1 的 remPath->list[0] 也将被赋值为 GPU0
    
  • 最后用于下次循环的节点链表为： nodeList =节点到 GPU0 的路径路径跳数GPU0PATH_LOC0PCI_SW1PCI_SW1 -> GPU01GPU1GPU1 -> GPU01
    
  
  
• 开始下次大循环：while (nodeList = { PCI_SW1, GPU1})
  
  
  
  • 先遍历节点 PCI_SW1，那么它的 remNode 为 PCI_SW2，因此 PCI_SW2 的 remPath->list[0] 将被赋值为 PCI_SW1, remPath->list[1] 将被赋值为 GPU0，此时 remPath->count = 2
    节点到 GPU0 的路径路径跳数GPU0PATH_LOC0PCI_SW1PCI_SW1 -> GPU01GPU1GPU1 -> GPU01PCI_SW2PCI_SW2 -> PCI_SW1 -> GPU02
    
  • 再遍历节点 GPU1，此时 GPU1 的 remNode 同样为 PCI_SW2，但是：
    
    
    
    • remPath->bw != 0，因为遍历 PCI_SW1 时，PCI_SW2 已经被赋值
      
    • 此时 remPath->count = 2，path->count = 1（GPU1 的 path->count 目前还是 1），所以 (remPath->bw == 0 || remPath->count > path->count) 最终为 true
      
    • 因此再看 remPath->bw < bw 是否满足。当前 remPath->bw 也因为遍历 PCI_SW1 时被赋值，但是如果 PCI_SW2 -> PCI_SW1 -> GPU0 的路径带宽确实比 PCI_SW2 -> GPU1 -> GPU0 要小，那么，NCCL 会将 PCI_SW2 的路径信息更新为最优路径：PCI_SW2->GPU1->GPU0。因为 PCI_SW2 已经在遍历节点 PCI_SW1 时被缓存到 nodeList，所以遍历 PCI_SW2 时不再添加。因此，nodeList 只有 PCI_SW2
      
    节点到 GPU0 的路径路径跳数GPU0PATH_LOC0PCI_SW1PCI_SW1 -> GPU01GPU1GPU1 -> GPU01PCI_SW2PCI_SW2 -> GPU1 -> GPU02
    
  
  
• 开始下次大循环：while (nodeList = { PCI_SW2 })，
  
  
  
  • PCI_SW2 有两个连接，分别是：PCI_SW1 和 GPU1
    
  • 遍历 PCI_SW1 和 GPU1 时，都因为  remPath->count(1) > path->count(2)， 不满足（比如 PCI_SW1->count = 1，PCI_SW2->count = 2），导致 nodeList = { NULL }， 因此整个算法流程结束。
    
  
  

ncclTopoComputePaths 的前半部分用于 PCIe 设备之间的最优物理路径计算，获得的是：在硬件能力允许的条件下，怎么走最快的物理路径。但是物理路径可达不代表一些能力实际可用，比如物理 GPU 之间的 P2P PCIe Switch 不支持，或者未开启 PCIe P2P，或者 rank 运行在容器内，共享内存 SHM 方式可能是禁止的。因此，在 ncclTopoComputePaths 后半部分，将标记这些信息，为下一轮进行 PCIe 设备间路径重新计算做铺垫。具体实现可总结为：

• 检查任意 GPU 之间是否支持 P2P 和 SHM，如果都不支持，那么这两个 GPU 之间的路径类型标记为 PATH_NET
  
• 同一物理节点内部，支持通过 GPU 中继，允许比如 GPU0 通过 GPU1 访问 NIC0，但是只允许数据发送方向，即 GPU0 向外发送数据通过 GPU1 中继
  
• 如果 GPU 和 NIC 之间的 GDR（GPU Direct RDMA） 模式不可用，则在它们之间添加中继 CPU 节点，强制数据从 CPU 经过
  
• 提前计算与 NIC 处于同一 CPU 域的 GPU，并缓存
  

1. GPU0 --(PCI)--> PCI_SW1
2. GPU0 --(NVL)--> GPU1

复制代码

NCCL 经过路径计算，如图中每个 PCIe 设备有如下路径信息：

PCIe设备裁剪和路径重新计算

1. ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm* comm) {
2.   ......
3.   // Update path for GPUs when we don't want to / can't use GPU Direct P2P
4.   for (int g=0; g<system->nodes[GPU].count; g++) {
6.     ......
7.     // Remove GPUs we can't (or don't want to) communicate with through P2P or SHM
8.     struct ncclPeerInfo* dstInfo = comm->peerInfo+system->nodes[GPU].nodes[g].gpu.rank;
9.     for (int p=0; p<system->nodes[GPU].count; p++) {
10.       ......
11.       int p2p;
12.       NCCLCHECK(ncclTransports[TRANSPORT_P2P]->canConnect(&p2p, comm, NULL, srcInfo, dstInfo));
13.       if (p2p == 0) {
14.         NCCLCHECK(ncclTransports[TRANSPORT_SHM]->canConnect(&shm, comm, NULL, srcInfo, dstInfo));
15.         if (shm == 0) {
16.           // Mark this peer as inaccessible. We'll trim it later.
17.           system->nodes[GPU].nodes[p].paths[GPU][g].type = PATH_NET;
18.         }
19.       }
20.     }
21.   }
22.   ......
24.   // Update paths for NICs (no GPU Direct, PXN, ...)
25.   for (int n=0; n<system->nodes[NET].count; n++) {
26.     struct ncclTopoNode* netNode = system->nodes[NET].nodes+n;
28.     for (int g=0; g<system->nodes[GPU].count; g++) {
30.       // Check whether we can access the NIC through another NVLink-connected GPU (PXN)
31.       struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
32.       if (ncclPxnDisable(comm) != 1) {
33.         ......
34.         if (localGpuIndex != g && localGpuIndex != -1) {
35.           // PXN = PCI + NVLink.
36.             ......
37.             NCCLCHECK(addInterStep(system, GPU, localGpuIndex, GPU, g, NET, n));
38.         }
39.       }
41.       if (gpu->paths[NET][n].type < PATH_PHB) {
42.         // Update path when we dont want to / can't use GPU Direct RDMA.
43.         ......
44.         if (gdr == 0) {
45.           // We cannot use GPU Direct RDMA, divert all traffic through the CPU local to the GPU
46.           int localCpu;
47.           NCCLCHECK(ncclGetLocalCpu(system, g, &localCpu));
48.           NCCLCHECK(addInterStep(system, CPU, localCpu, NET, n, GPU, g));
49.           NCCLCHECK(addInterStep(system, CPU, localCpu, GPU, g, NET, n));
50.         }
51.       }
52.     }
53.   }
55.   // Pre-compute NET local gpus to accelerate search
56.   for (int n=0; n<system->nodes[NET].count; n++) {
57.     ......
58.     NCCLCHECK(ncclTopoGetLocalGpu(system, net->id, &net->net.localGpu));
59.   }
60.   ......
61. }

复制代码

GPU 之间的裁剪逻辑可总结为：

• 判断 GPU 之间的路径类型是否小于类型 PATH_NET，如果小于则在同一个 domain，否则在不同 domain
  
• 通过 myDomain 记录本 rank 所在 domain
  
• 将与当前 rank 不在同一个 domain 的 GPU 节点都删掉
  
  
  
  • 删除该节点的所有 paths
    
  • 从其它节点的 links 中删除与该节点的连接 link
    
  
  

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             initTransportsRank
5.                 // 路径初始计算
6.                 ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm* comm)           
7.                 // 设备裁剪
8.                 ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* comm)
9.                 // 路径再次计算
10.                 ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm* comm)

复制代码

NCCL 基于裁剪后的拓扑，再次重新计算所有 PCIe 设备的可达路径。从而得到一个运行环境实际可用的拓扑路径，为后续通道计算做铺垫。以上逻辑对应日志如下：

1. ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* comm) {
2.   
3.   int *domains;
4.   int64_t *ids = NULL;
5.   int myDomain = 0;
6.   int ngpus = system->nodes[GPU].count;
7.   ......
8.   for (int g=0; g<system->nodes[GPU].count; g++) {
9.     struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
10.     domains[g] = g;
11.     ids[g] = gpu->id;
12.     for (int p=0; p<g; p++) {
13.       if (gpu->paths[GPU][p].type < PATH_NET) {
14.         domains[g] = std::min(domains[g], domains[p]);
15.       }
16.     }
17.     if (gpu->gpu.rank == comm->rank) myDomain = domains[g];
18.   }
20.   for (int i=0; i<ngpus; i++) {
21.     if (domains[i] == myDomain) continue;
22.     ......
23.     for (g=0; g<system->nodes[GPU].count /* This one varies over the loops */; g++) {
24.       gpu = system->nodes[GPU].nodes+g;
25.       if (gpu->id == ids[i]) break; else gpu=NULL;
26.     }
27.     ......
28.     NCCLCHECKGOTO(ncclTopoRemoveNode(system, GPU, g), ret, fail);
29.   }
30.   ......
31. }
33. ncclResult_t ncclTopoRemoveNode(struct ncclTopoSystem* system, int type, int index) {
34.   ......
35.   for (int t=0; t<NCCL_TOPO_NODE_TYPES; t++) {
36.     free(delNode->paths[t]);
37.    
38.   ......
39.   while (l<node->nlinks && node->links[l].remNode == delNode) {
40.     memmove(...);
41.     node->nlinks--;
42.   }
43. }

复制代码

NCCL通信算法

NCCL 使用 Ring，Tree 等算法在多 GPU 之间进行数据传递。不同的通信算法有对应的适用场景。
NCCL Ring拓扑计算

NCCL Ring 结构顾名思义是将多个 GPU 在逻辑上组织成一个环，从而利用环形拓扑进行数据通信，环形拓扑分为两种场景：

• 单物理节点上，只是将本节点内部的多个 GPU 组织为环，即形成拓扑结构：GPU a -> GPU b -> .. -> GPU x -> GPU a
  
• 多机通信场景，则需要在每个节点构造一个从 NIC 到 GPU-X0，再从 GPU-Xn 到 NIC 的拓扑结构，例如：NET n -> GPU a -> GPU b -> .. -> GPU x -> NET n (or m if crossNic)
  

NCCL 有 channel 的概念，这里的一个 channel 对应一个 ring 环。
NCCL 会尝试 ring 多通道计算，多 channel 的计算逻辑非常复杂。多 channel 的核心作用是可以使得 GPU 利用多 channel 进行数据的并行传输。好比将单车道拓宽为双车道，或者更多车道。但是通道也不是无限扩张，因为每个通道有基础带宽，在通道所经过的路径上，带宽资源也是有限的，NCCL 的目标就是通过多通道无限逼近硬件带宽资源，以求最大化利用。
通道计算由函数 ncclTopoSearchRec 实现

1. # 设备列举
2. ncclTopoPrint:316 NCCL INFO === System : maxBw 6.2 totalBw 24.0 ===
3. ncclTopoPrintRec:289 NCCL INFO CPU/0-ffffffffffffff (1/1/4)
4. ncclTopoPrintRec:289 NCCL INFO + PCI[24.0] - GPU/0-80 (0)
5. ncclTopoPrintRec:289 NCCL INFO + PCI[24.0] - GPU/0-90 (1)
6. ncclTopoPrintRec:289 NCCL INFO + PCI[12.0] - NIC/0-0
7. ncclTopoPrintRec:308 NCCL INFO + NET[6.2] - NET/0-2 (0/7f7afb0003797334/1/6.250000)
9. # 路径展示
10. # 比如：从 GPU/0-80 经过 --PCI(24)->CPU/0-ffffffffffffff，可到达 PCI(24)->GPU/0-90
11. ncclTopoPrint:319 NCCL INFO ==========================================
12. printNodePaths:119 NCCL INFO Paths from GPU/0-80 :
13. printNodePaths:136 NCCL INFO  (5000.000000)
14. printNodePaths:136 NCCL INFO --PCI(24)->CPU/0-ffffffffffffff--PCI(24)->GPU/0-90 (24.000000)
15. printNodePaths:136 NCCL INFO --PCI(24)->CPU/0-ffffffffffffff (24.000000)
16. printNodePaths:136 NCCL INFO --PCI(24)->CPU/0-ffffffffffffff--PCI(12)->NIC/0-0--NET(6.25)->NET/0-2 (6.250000)
18. printNodePaths:119 NCCL INFO Paths from GPU/0-90 :
19. printNodePaths:136 NCCL INFO --PCI(24)->CPU/0-ffffffffffffff--PCI(24)->GPU/0-80 (24.000000)
20. printNodePaths:136 NCCL INFO  (5000.000000)
21. printNodePaths:136 NCCL INFO --PCI(24)->CPU/0-ffffffffffffff (24.000000)
22. printNodePaths:136 NCCL INFO --PCI(24)->CPU/0-ffffffffffffff--PCI(12)->NIC/0-0--NET(6.25)->NET/0-2 (6.250000)
24. printNodePaths:119 NCCL INFO Paths from NET/0-2 :
25. printNodePaths:136 NCCL INFO --NET(6.25)->NIC/0-0--PCI(12)->CPU/0-ffffffffffffff--PCI(24)->GPU/0-80 (6.250000)
26. printNodePaths:136 NCCL INFO --NET(6.25)->NIC/0-0--PCI(12)->CPU/0-ffffffffffffff--PCI(24)->GPU/0-90 (6.250000)
27. printNodePaths:136 NCCL INFO --NET(6.25)->NIC/0-0--PCI(12)->CPU/0-ffffffffffffff (6.250000)
28. printNodePaths:136 NCCL INFO  (5000.000000)

复制代码

如下归纳了 ncclTopoSearchRec 计算通道的算法逻辑（多机场景，只涉及算法主要思想）：

1. ncclCommInitRank
2.     ncclCommInitRankDev
3.         ncclCommInitRankFunc
4.             initTransportsRank
5.                 ncclTopoCompute
6.                     ncclTopoSearchRec(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, struct ncclTopoGraph* saveGraph, int* time)

复制代码

当 ncclTopoSearchRec 完成当前条件设置下的通道搜索，接下来进入如下逻辑：

• 当计算出的通道数 graph->nChannels 等于 NCCL 允许的最大通道数：graph->maxChannels 时，可以暂时结束搜索，跳转到指定标签 done，准备下一次搜索。暗示计算出来的所有通道在途径的链路上，带宽容量都是可以满足的
  

1. (1) 给定一个通道基准带宽值（通道有初始基准带宽，基准带宽刚好小于任意 GPU 到任意 NIC 路径之间的最大带宽），基于基准带宽值去搜索通道。
2. (2) 按照 GPU 的顺序，找出所有 GPU 最合适的网卡，返回网卡列表
3. (3) 按照得到的网卡列表，遍历网卡，准备构造 ring 路径。每一个网卡都要参与路径计算
4.     1. 当前被选择的网卡 Xn，因为 graph->nChannels = 0，强制从 GPU0 开始计算，如果计算成功，得到一条 ring 路径，此时 graph->nChannels 被加 1
5.         a. 如果当前所有通道的总带宽高于通道数 -1 的总带宽，则新增当前通道，保存到 saveGraph
6.         b. 如果带宽相等情况下，但是当前所有通道的总跳数小于之前的通道总跳数，保存到 saveGraph
7.         c. 如果带宽和跳数条件都不满足，则不会将当前 graph 保存到 saveGraph，但是 graph 内容不变
8.     2. 因为 graph->nChannels 被加 1，开始选择另外一块网卡Xn+1，计算公式：(graph->nChannels+i)%netCount，开始计算新通道（graph->nChannels 会影响网卡的选择，此时是递归计算，是网卡 Xn 计算成功后的递归调用，根据通道被加 1 选择了网卡 Xn+1，而不是大循环遍历到网卡 Xn+1，此时最外层大循环还是网卡 Xn），因为 graph->nChannels != 0，从与当前被选择网卡距离最近的 localGpu（与该 NIC 共享 PCIe root complex 或 NUMA node 的 GPU）开始计算，如果计算成功，得到一条 ring 路径，此时 graph->nChannels 再被加 1，保存策略参考 (3).1.abc
9.     3.按照步骤 (3).2，继续迭代，直至无法再迭代，执行回溯。在递归计算的路径上，都会在通道途经的路径上依次扣减 PCIe 设备间的给定基准带宽值，表示已经占用了一部分带宽资源。当回溯时，再按回溯路径依次恢复 PCIe 设备间的带宽。最后回溯至最外层大循环，即第一次选择网卡 Xn 的位置时，带宽恢复完毕，graph->nChannels 也恢复至0
10. (3) 因为回溯至 Xn，接下来选择 localGpu 再次计算，重复过程 (3).1，只不过变为从 localGpu 开始计算，而不再是 GPU0
11. (4) localGpu 计算完毕后，继续尝试所有可达 GPU 路径计算，再次重复过程 (3).1，只不过变为从 GPU 列表循环开始计算，但是会除去 localGpu，因为上一步计算过了
12. (5) 回到网卡大循环，从下一块网卡，比如网卡 Xn+1 开始，再次重复上述整个过程，从步骤 (3) 开始，只是现在变为大循环网卡 Xn+1 出发，graph->nChannels = 0，之前所有 PCIe 设备间被扣减的带宽均已恢复   

复制代码

• 当计算出的通道数小于允许的最大通道数，但是所消耗的总带宽 >= 总带宽时（所有 GPU 中自身link链路的带宽和，取一个最大值），可以暂时结束搜索，跳转到指定标签 done，准备下一次搜索
  

1. // Optimal solution, stop here
2. if (time == -1) goto done;

复制代码

如果前面两个条件都未满足，进入通道搜索第一阶段：NCCL 尝试降低一些搜索条件去继续搜索 ring channel。

• 例如：在基准带宽基础上逐渐减小带宽值，不断尝试是否可以计算出更多的通道。通道带宽也不能无限小，有一个最小经验值去限制。通过减小通道带宽，并尝试增加通道数的方法，去无限逼近通道路径上的带宽上限。例如：假设一个 path 有 64 GB/s 的带宽，按照 18 GB/s 的粒度只能找到 3 个 Channel，总带宽是 54 GB/s，但是按照 16 GB/s 的粒度的话，能够找到 4 个 Channel，刚好利用路径通道的所有带宽
  
• 当降低搜索条件到极限：speedArray[speedIndex+1] / graph->bwInter > .49 不再成立，结束第一个阶段的通道搜索
  

1. // Optimal solution, stop here
2. if (graph->nChannels*graph->bwInter >= system->totalBw) goto done;

复制代码

在以上逻辑完成后，可能会进入通道搜索第二阶段，比如其中有一个条件：(graph->bwIntra && graph->bwInter) >= 25.0，需要基准带宽足够大

• 将当前计算得到的通道数翻倍（需要小于允许的 maxChannels），通道基准带宽减半
  

1. // Decrease bw until we find a solution
2. if ((speedIndex < nspeeds-1) && (graph->nChannels == 0 || (speedArray[speedIndex+1]/graph->bwInter > .49))) {
3.   tmpGraph.bwInter = tmpGraph.bwIntra = speedArray[++speedIndex];
4.   goto search;
5. }

复制代码

• 在带宽减半的基础上，去逐渐反向增加带宽值，不断尝试是否可以计算出带宽更高的总带宽。前提是，在所有通道所经过的路径上，带宽容量都满足。
  

1. ncclTopoDupChannels(graph, ccMin, ngpus);

复制代码

NCCL 通过暴力搜索方法，不断尝试各种条件，计算出基础带宽值合适，通道数更多，总带宽更大的多通道，目标就是尽最大可能榨干硬件带宽资源。上述逻辑对应日志：

1. if (graph->pattern == NCCL_TOPO_PATTERN_RING) {
2.   // increase bw for Ring
3.   tmpGraph.bwIntra = tmpGraph.bwInter = speedArray[--speedIndex];
4.   goto search;
5. }

复制代码

有如下计算得到的 Ring 通道示意图：

NCCL Tree拓扑介绍

NCCL Tree 拓扑也使用 ncclTopoCompute 计算得到，但是 Tree 只能用于 All Reduce 操作。与 ring 一样，一颗 tree 将绑定一个 channel。这里只介绍 tree 拓扑的相关算法概念。NCCL Tree 拓扑也在不断演进，从最初的 Single Binary Tree 演变为 Double Binary Tree，再在 Double Binary Tree 的基础上演变出新的 Split Tree，Balanced Tree。
Single Binary Tree

有如下单二叉树结构：

Single Binary Tree 上执行 All Reduce 操作，规约阶段：

• 叶子节点 1, 3, 5, 7,......，31 分别向自己的父节点：2, 6, 10, 14, 18, 22, 26, 30 发送数据
  
• 非叶子节点 2, 6, 10, 14, 18, 22, 26, 30 把从子节点收到的数据与本地数据进行规约，再将规约后的数据分别发送给自己的父节点
  
• 依此类推，直到 root 节点 0 规约完成所有数据，时间开销 log_N
  

Broadcast 阶段：

• root 节点将规约后的数据广播给自己的子节点 16
  
• 子节点将收到的数据继续广播给自己的子节点 8，24
  
• 依次类推，直到所有的子节点都收到 root 节点的规约数据。时间开销 log_N
  

缺点：

• reduce 和 broadcast 操作在时间线上只能串行执行，需要 root 完成 reduce 后才能进行 broadcast。时间开销比较大
  
• 叶子节点只有一个 parent 节点，没有 child 节点；非叶子节点虽然有两个 child 节点，但是也只有一个 parent 节点。这种结构会导致带宽浪费
  

Double Binary Tree

有如下双二叉树结构：

Double Binary Tree 与 Single Binary Tree 的区别：

• Double Binary Tree 基于 Single Binary Tree 的所有节点生成了第二棵单二叉树
  
• 相对于 Single Binary Tree，在第二棵二叉树中，节点的角色进行了互换，叶子节点变为非叶子节点，非叶子节点转变为叶子节点
  
• 如此，除了 root 点，任意节点都有了两个 parent 节点，两个 child 节点。root 节点也有了一个 child 和一个 parent
  

Double Binary Tree 带来的优势：

• 与 Single Binary Tree 形成互补 Tree，通信更加均衡。基本上所有节点都有两个 parent 连接，两个 child 连接
  
• 增加一个通道，数据可以并行发送和接收，通信时延进一步降低
  

比如执行 All Reduce 操作时，每个节点的数组分别有 2N 个元素，前 N 个元素被映射到 Double Binary Tree 1，剩下的 N 个元素被映射到 Double Binary Tree 2，数据 reduce 操作可以在两棵 Tree 上并行执行，broadcast 阶段也是如此。
Split Tree 和 Balanced Tree

Split Tree 和 Balanced Tree 本质上也都是 Double Binary Tree。但是都在 Double Binary Tree 基础上做了相应变化。
这里需要澄清一个概念，在多机通信时，每个物理节点会选取其中一些节点作为与外部通信的节点；在物理节点内部，GPU 节点之间退化为链路结构。比如现在有 4 个物理节点，每个节点都有 8 个 GPU：

1. ncclTopoPrintGraph:1201 NCCL INFO Pattern 4, crossNic 0, nChannels 1, bw 6.000000/6.000000, type PHB/PHB, sameChannels 1
3. # 当前实验环境只计算出一条channel，从 NET/0-2 出发，经过 GPU/0-80，GPU/0-90，回到 NET/0-2
4. ncclTopoPrintGraph:1224 NCCL INFO  0 : NET/0-2 GPU/0-80 GPU/0-90 NET/0-2

复制代码

那么对于普通 Double Binary Tree，有如下两棵树形结构：

• Tree1 中 GPU-8 和 GPU-24 为叶子节点，GPU-16 和 GPU-0 为非叶子节点
  

1. # node1 的 GPU 编号
2. 7->6->5->4->3->2->1->0
4. # node2 的 GPU 编号
5. 15->14->13->12->11->10->9->8
7. # node3 的 GPU 编号
8. 23->22->21->20->19->18->17->16
10. # node4 的 GPU 编号
11. 31->30->29->28->27->26->25->24

复制代码

• Tree2 中 GPU-0 和 GPU-16 为叶子节点，GPU-8 和 GPU-24 为非叶子节点。相对于 Tree 1，节点角色已经互换
  

1. {8, 24} -> 16 -> 0

复制代码

Double Binary Tree 1 如下图所示：

Double Binary Trees 2 如下图所示：

NCCL 对 Double Binary Tree 的定义：

1. {0, 16} -> 8 -> 24

复制代码

对于 Split Tree，有如下两棵树形结构：

• Tree1 中 GPU-8 和 GPU-24 为子节点，GPU-17 为 parent 节点， GPU-16 为 GPU-1 的子节点
  

1. #define NCCL_TOPO_PATTERN_TREE 3            // All NIC traffic going to/from the same GPU

复制代码

• Tree2 中 GPU-0 和 GPU-16 为子节点，GPU-9 为 parent 节点， GPU-8 为 GPU-25 的子节点
  

1. #define NCCL_TOPO_PATTERN_SPLIT_TREE 2      // Spread NIC traffic between two GPUs (Tree parent on first GPU, tree children on the second GPU)

复制代码

Split Tree 1 如下图所示：

Split Tree 2 如下图所示：

NCCL 对 Split Tree 的定义：

1. {8, 24} -> {17, 16} -> 0

复制代码

带来的优势：

• GPU 不再是单点热点。不再是每个节点的 first GPU 承担全部工作，其它 GPU 也会分担一些计算任务
  
  
  
  • 例如 Tree1 中，GPU-1 承担了子节点的 reduce 计算，而不再是 GPU-0 既要承担全部 reduce，又要执行 broadcast
    
  • 同理，Tree1 中 GPU-17，也为 GPU-16 分担了一些计算任务
    
  
  
• 不再是同一个 GPU 同时处理 RX/TX 流量。NIC 的 ingress 和 egress，落在了不同的 GPU 上，避免带宽下降（虽然是全双工，但实际上 RX/TX同时发生时，容易出现Copy Engine，arbitration争抢，AI说的）
  
  
  
  • 例如 Tree1 中，reduce 操作时，GPU-17 只接收，GPU-16 只发送
    
  
  

对于 Balanced Tree，有如下两棵树形结构：

• Tree1 中 GPU-8 和 GPU-24 为子节点，GPU-17 为 GPU-24 的 parent 节点， GPU-16 为 GPU-8 的 parent 节点
  

1. #define NCCL_TOPO_PATTERN_TREE 3            // All NIC traffic going to/from the same GPU

复制代码

• Tree2 中 GPU-0 和 GPU-16 为子节点，GPU-9 为 GPU-16 的 parent 节点， GPU-8 为 GPU-0 的 parent 节点
  

1. #define NCCL_TOPO_PATTERN_SPLIT_TREE 2      // Spread NIC traffic between two GPUs (Tree parent on first GPU, tree children on the second GPU)

复制代码

Balanced Tree 1 如下图所示：

Balanced Tree 2 如下图所示：

NCCL 对 Balanced Tree 的定义：

1. {0, 16} -> {9, 8} -> 24

复制代码

带来的优势：流量更加负载均衡，NIC 流量分摊到多个 GPU（流量送到多个 GPU，比流量送到同一个 GPU 要更容易吃满 NIC 带宽，AI说的）；两个 parent 节点分别处理不同 child 节点的流量。
NCCL通信算法的选择

NCCL 会根据通信操作类型、数据大小、GPU/节点数量、底层网络拓扑等自动选择最优算法。如下为 Ring 和 Tree 算法的一些对比差异：
特性Ring 算法Tree 算法通信步数O(N)O(log N)带宽效率高（接近最优）中等（根节点易成瓶颈）延迟较高较低适用操作AllReduce（主流）Broadcast, Reduce, AllReduce适用规模小 / 中规模（节点少）大规模（节点多）在链接：https://developer.nvidia.com/blog/massively-scale-deep-learning-training-nccl-2-4/ 中，通过 All Reduce 操作，比较了 Ring 和 Tree 的时延和带宽差异，在大规模集群场景，Tree 的时延是显著低于 Ring 算法的。

NCCL 多节点之间的 Channel 连接

NCCL 在为每个 rank 计算完通道后，接下来就要将多机之间的 channel 进行连接，在多机之间形成逻辑上的通信路径。这里只涉及 Ring channel 的相关逻辑。
首先设置当前 rank 计算得到的通道相关信息，例如通道数，机间带宽等。

1. #define NCCL_TOPO_PATTERN_BALANCED_TREE 1   // Spread NIC traffic between two GPUs (Tree parent + one child on first GPU, second child on second GPU)

复制代码

设置当前 rank 的具体通道信息：

• ringRecv[c] 表示：表示当前 ring channel c 的第一个接收数据的 rank id
  
• ringSend[c] 表示：表示当前 ring channel c 的最后一个 rank id
  
• ringPrev[c] 表示：当前 rank 的前驱 rank id
  
• ringNext[c] 表示：当前 rank 的后继 rank id
  

1.   NCCLCHECKGOTO(ncclCalloc(&allGather3Data, nranks), ret, fail);
3.   for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) {
4.     allGather3Data[rank].graphInfo[a].pattern = graphs[a]->pattern;
5.     allGather3Data[rank].graphInfo[a].nChannels = graphs[a]->nChannels;
6.     allGather3Data[rank].graphInfo[a].sameChannels = graphs[a]->sameChannels;
7.     allGather3Data[rank].graphInfo[a].bwIntra = graphs[a]->bwIntra;
8.     allGather3Data[rank].graphInfo[a].bwInter = graphs[a]->bwInter;
9.     allGather3Data[rank].graphInfo[a].typeIntra = graphs[a]->typeIntra;
10.     allGather3Data[rank].graphInfo[a].typeInter = graphs[a]->typeInter;
11.     allGather3Data[rank].graphInfo[a].crossNic = graphs[a]->crossNic;
12.   }

复制代码

接下来通过控制通道交换所有 rank 的通道信息，目的：

• 对齐通道数和通信带宽
  
• 统一多机通信算法等
  

1. ncclResult_t ncclTopoPreset(struct ncclComm* comm, struct ncclTopoGraph** graphs, struct ncclTopoRanks* topoRanks) {
2.   ......
3.   for (int i=0; i<localRanks; i++) {
4.     if (ringIntra[i] == rank) {
5.       topoRanks->ringRecv[c] = ringIntra[0];
6.       topoRanks->ringSend[c] = ringIntra[localRanks-1];
7.       topoRanks->ringPrev[c] = (i == 0) ? -1 : ringIntra[i-1];
8.       topoRanks->ringNext[c] = (i == localRanks-1) ? -1 : ringIntra[i+1];
9.     }
10.   ......
11. }

复制代码

信息交换并对齐后，接下来就开始跨节点合并通道，实现多机间的逻辑 ring 连接。
这里需要提到一个通信优化： 多机间 ring 通道通信，NCCL 将奇数 node 节点的奇数通道，偶数通道使用的网卡进行交换，避免跨轨通信。

当 NCCL_CROSS_NIC = 2（default），且节点有多块网卡，并计算出多 channel 时，NCCL 会使奇偶通道交替使用 NIC，实现流量负载均衡（NET0 和 NET1都实现收发），例如：

1.   bootstrapAllGather(comm->bootstrap, allGather3Data, sizeof(*allGather3Data))
3.   for (int i=0; i<nranks; i++) {
4.     allTopoRanks[i] = &allGather3Data[i].topoRanks;
5.     // Make sure we align all ranks so that the tuning is consistent across ranks
6.     for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) {
7.       graphs[a]->nChannels = std::min(allGather3Data[i].graphInfo[a].nChannels, graphs[a]->nChannels);
8.       graphs[a]->sameChannels = std::min(allGather3Data[i].graphInfo[a].sameChannels, graphs[a]->sameChannels);
9.       graphs[a]->bwIntra = std::min(allGather3Data[i].graphInfo[a].bwIntra, graphs[a]->bwIntra);
10.       graphs[a]->bwInter = std::min(allGather3Data[i].graphInfo[a].bwInter, graphs[a]->bwInter);
11.       graphs[a]->typeIntra = std::max(allGather3Data[i].graphInfo[a].typeIntra, graphs[a]->typeIntra);
12.       graphs[a]->typeInter = std::max(allGather3Data[i].graphInfo[a].typeInter, graphs[a]->typeInter);
13.       graphs[a]->crossNic = std::max(allGather3Data[i].graphInfo[a].crossNic, graphs[a]->crossNic);
14.     }
15.     comm->maxTreePattern = std::max(comm->maxTreePattern, allGather3Data[i].graphInfo[NCCL_ALGO_TREE].pattern);
16.   }

复制代码

对于 channel 0，NCCL 会在所有节点上，都使用 NET 0 收，NET 1 发，那么流量到达 node-1 时，将会由 node-1 的 NET 0 收，造成了跨轨通信。NCCL 为了避免跨轨通信，将奇数 node 的奇偶通道互换，实现同轨通信。例如 node-0 使用 NET 1 发送时，node-1 实现通过 NET 1 接收数据。

1. # channel 0
2. NET 0 -> GPU a -> GPU b -> .. -> GPU x -> NET 1
4. # channel 1
5. NET 1 -> GPU a -> GPU b -> .. -> GPU x -> NET 0

复制代码

接下来就是通过 connectRings，将多机间的 ring 连接起来，分别设置每个 rank 的前驱和后继，实现多机间 ring 的连接。

1.   // Alternate rings to avoid crossing rails.
2.   // CrossNic values could be not the same on all nodes as it depends on the number of net devs and the NVLink bandwidth.
3.   // Therefore, it's only done if the rank obtained a solution with crossNic=2.
4.   for (int r = 0; r < comm->nRanks; r++) {
5.     if (allTopoRanks[r]->crossNicRing == 2 && (nChannels % 2) == 0 && (comm->rankToNode[r] % 2) == 1) {
6.       // Exchange rings
7.       for (int c=0; c<nChannels; c+=2) {
8.         exchangeValues(allTopoRanks[r]->ringRecv+c, allTopoRanks[r]->ringRecv+(c^1));
9.         exchangeValues(allTopoRanks[r]->ringSend+c, allTopoRanks[r]->ringSend+(c^1));
10.         exchangeValues(allTopoRanks[r]->ringPrev+c, allTopoRanks[r]->ringPrev+(c^1));
11.         exchangeValues(allTopoRanks[r]->ringNext+c, allTopoRanks[r]->ringNext+(c^1));
12.       }
13.     }
14.   }

复制代码

NCCL 初始化的目的，可以引用如下示意图来概括：

初始化总结

基于现有测试环境，以上内容从源码层面分析了 NCCL 的初始化关键流程。包括：

• 网络插件初始化。主要涉及 InfiniBand，或者基于 TCP/IP 的 Socket 网络插件的选择。
  
• NCCL 控制环的建立。主要涉及 Unique Id 的产生和用途，以及所有 rank 组成控制面环形拓扑的相关流程。
  
• PCIe 设备发现。记录从 GPU/NIC 出发，直至归属 CPU 这条物理 PCIe 链路上有哪些 PCIe 设备，并生成从 GPU/NIC 至 CPU 链路的 PCIe XML 链路拓扑。
  
• PCIe 设备间建图。PCIe 设备之间的物理链路关系是事实存在的，但是 NCCL 需要将这些链路关系建立到自身的数据结构中，目的是后续 channel 的计算服务。
  
• PCIe 设备间最优可达路径计算。基于 PCIe 建图，NCCL 拥有 PCIe 设备间的连关系。在此基础上，NCCL 使用 BFS 算法，计算任意 PCIe 设备之间的最优可达路径，即跳数最少，带宽最大的可达路径。
  
• 设备内部基于 Ring，Tree 通信算法的 channel 计算。在路径带宽允许的情况下，NCCL 会尽最大可能计算出同一种通信算法的更多通信通道，从而尽可能榨干硬件带宽资源。
  
• 多机间 Ring，Tree 通道的连接。在多级多卡之间，协商出最终可通信的逻辑算法拓扑。
  
• 也分析了 NCCL tree 由 Double Binary Tree 到 Split Tree，再到 Balanced Tree 的演进和优化思想
  

截至本篇文章发布的时间点，可能是读者们能搜到的有关 NCCL 初始化介绍最全面和详细的文章了。
以上内容是基于源码学习后，再纯手工逐字撰写，包括很多配图的原创。但是，文章也肯定存在很多错误和疏漏之处，也请大家拍砖！
本文来自博客园，作者：T-BARBARIANS 原创，博文严禁转载，转载必究！

来源：程序园用户自行投稿发布，如果侵权，请联系站长删除
免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！