熊掌号接合网站做seo个人网站有哪些站
相关系列
【分布式】NCCL部署与测试 - 01
 【分布式】入门级NCCL多机并行实践 - 02
 【分布式】小白看Ring算法 - 03
 【分布式】大模型分布式训练入门与实践 - 04
目录
- 相关系列
 - 概述
 - 1.1 Tree
 - 1.2 double binary tree
 
- 初始化和拓扑
 - 2.1 Tree的初始化与差异
 - 2.2 ncclGetBtree
 
- Kernel内部实现
 - 初始化
 - 原语差异
 - 流程
 - RunTreeUpDown
 - RunTreeSplit
 
概述
先掠过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在第三位
