当前位置：首页 > news >正文

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

news 2025/8/6 3:46:00

概述

先掠过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