【分布式】NCCL Split Tree kernel内实现情况 - 06

概述

先掠过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在第三位

canmoumou