NCCL中channel的概念表示一个通信路径，为了更好地利用带宽和网卡，以及同一块数据可以通过多个channel并发通信，另外后续可以看到一个channel对应了一个GPU SM，所以基于这些原因，NCCL会使用多channel，搜索的过程就是搜索出来一组channel。如上节所述，单机的情况下会在ncclTopoTrimSystem函数里删除网卡，因此我们先看下...
NCCL中channel的概念表示一个通信路径,为了更好地利用带宽和网卡,以及同一块数据可以通过多个channel并发通信,另外后续可以看到一个channel对应了一个GPU SM,所以基于这些原因,NCCL会使用多channel,搜索的过程就是搜索出来一组channel。 如上节所述,单机的情况下会在ncclTopoTrimSystem函数里删除网卡,因此我们先看下单机...
struct ncclConnect { char data[CONNECT_SIZE];}; struct ncclConnect *connect; NCCLCHECKGOTO(ncclCalloc(&connect, 2), ret, affinity_restore); for (int c=0; cnChannels; c++) { struct ncclChannel* channel = comm->channels+c; NCCLCHECKGOTO(setupChannel(comm, c, rank, nranks, rings+c*nr...
这里将所有channel的prev,next,send,recv信息打平到数组中,例如recv[0]表示第一个ring中rank0的recv是哪个rank,然后开始计算当前机器第一个rank的prev和最后一个rank的next。 static ncclResult_t connectRings(struct ncclComm* comm,int* ringRecv,int* ringSend,int* ringPrev,int* ringNext,int* firstRanks)...
NCCL中用Channel的概念表示一个通信路径，在初始化的过程会自动感知拓扑并计算最佳的通信路径。为了更好的利用带宽和网卡实现并发通信，NCCL会使用多channel。NCCL-test运行日志里列出了16组channel如下：### ChannelNum:16 bm-2204kzq:252978:253054 [0] NCCL INFO Channel 00/16 : 0 7 5 6 ...
### ChannelNum:16 bm-2204kzq:252978:253054 [0] NCCL INFO Channel 00/16 : 0 7 5 6 4 3 1 2 8 15 13 14 12 11 9 10 bm-2204kzq:252978:253054 [0] NCCL INFO Channel 01/16 : 0 7 5 6 4 3 1 10 8 15 13 14 12 11 9 2 ...
### ChannelNum:16 bm-2204kzq:252978:253054 [0] NCCL INFO Channel 00/16 : 0 7 5 6 4 3 1 2 8 15 13 14 12 11 9 10 bm-2204kzq:252978:253054 [0] NCCL INFO Channel 01/16 : 0 7 5 6 4 3 1 10 8 15 13 14 12 11 9 2 ...
channel获取 channel的概念: nccl中channel的概念表示一个通信路径,为了更好的利用带宽和网卡,以及同一块数据可以通过多个channel并发通信,nccl会使用多channel,搜索的过程就是搜索出来一组channel。 具体一点可以参考以下文章: 如何理解Nvidia英伟达的Multi-GPU多卡通信框架NCCL? - Connolly的回答 - 知乎https://www.zhih...
实际运行过程中,NCCL每个channel会做intra-node reduce -> inter-node all-reduce -> intra-node broadcast,这三个步骤是可以流水线处理的。 通过构建Double Binary Tree,每个节点至多是一个叶子和一个非叶子结点。因此每个节点至多只需要收发两倍的数据量,这点跟Ring的实现也是一样的。因此Double Binary Tree也是一...
第1章 多机多卡运行nccl-tests 和channel获取 第2章 多机多卡nccl-tests 对比分析 前言 NCCL（NVIDIA ...