相关系列

【分布式】NCCL部署与测试 - 01
【分布式】入门级NCCL多机并行实践 - 02
【分布式】小白看Ring算法 - 03
【分布式】大模型分布式训练入门与实践 - 04

概述

先掠过Tree算法在拓扑方面以及树的生成方面是如何实现的,本期主要讲kernel内部的情况。
先放上2.11.4部分的tree,后续增添2.18版本中nccl的改动,以及rccl的处理。

如果你看过其他的一些文档,应该知道double binary tree的一些构造。
即我们可以将tree分为三类,朴素的tree、double binary tree和split tree、balanced tree。

1.1 Tree

首先最朴素的tree,存在浪费带宽的情况。

1.2 double binary tree

因此引入double binary tree
把allreduce可以拆分为reduce和broadcast两个过程,reduce是自下而上(Tree1),broadcast自上而下(Tree2),这样构造两棵树,第一棵树的叶子节点在第二棵树中是中间节点,这样就能更好的做到流水并行。

1.3 SplitTree
但是这样又有个问题,根节点要向Tree2所有的中间节点发送消息,同时还要从Tree1的中间节点接收消息,Root会不会太忙碌了?因此又提出了SplitTree
再多出一个接收节点,用来平衡。通过切分的方式,把所有向上的父节点放到GPU0,向下的传输放到GPU1上,把上行和下行切分开来。
代价:例如broadcast的时候,多了一个GPU0到GPU1的传输操作。在代码中注释为,Spread NIC traffic between two GPUs。

1.4 balanced tree
所有的父节点放到同一张GPU0上,但子节点放到GPU0和GPU1两张显卡上。因此向上传递的时候具有一些不确定性。英伟达推荐CUDA sm>70所有的Tree更推荐Balanced Tree。不过nccl 2.11.4的版本这次先不提及balanced tree的细节。

初始化和拓扑

2.1 Tree的初始化与差异

初始化的部分,节点内要确定backToNIC是哪张卡。但节点内搜索Tree本质上和Ring没有区别。

Tree的代码出现差异的部分在:ncclTopoPreset

NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &allGather3Data[rank].topoRanks));

在connect.cc文件中:

 int parentIndex = 0;
int child0Index = treeGraph->pattern == NCCL_TOPO_PATTERN_TREE ? 0 : 1;
int child1Index =
    treeGraph->pattern == NCCL_TOPO_PATTERN_SPLIT_TREE ? 1 : 0;

topoRanks->treeToParent[c] = treeIntra[parentIndex];
topoRanks->treeToChild0[c] = treeIntra[child0Index];
topoRanks->treeToChild1[c] = treeIntra[child1Index];
channel->tree.up = i == 0 ? -1 : treeIntra[i - 1];
channel->tree.down[0] = i == localRanks - 1 ? -1 : treeIntra[i + 1];

这一段是判断Tree的类型,然后判断哪个GPU来连接Child。然后进行赋值。

最后两行是完成在节点内的ring环,并确定出入节点。

2.2 ncclGetBtree

 * Illustration :
 * 0---------------8
 *          ______/ \______
 *         4               12
 *       /   \            /  \
 *     2       6       10     \
 *    / \     / \     /  \     \
 *   1   3   5   7   9   11    13

二进制化以后,根据从右往左位出现的第一个1的位数来判断是哪一层,例如没有1就是root,
叶子节点在最下层,所以1在第四位
第二层的节点,1在第二位
第三层的节点,1在第三位

03-11 10:17