网站上传后,西安网站推广招聘,中国建设银行app官方下载,凡科网站怎么做外链相关系列
【分布式】NCCL部署与测试 - 01 【分布式】入门级NCCL多机并行实践 - 02 【分布式】小白看Ring算法 - 03 【分布式】大模型分布式训练入门与实践 - 04 目录 相关系列概述1.1 Tree1.2 double binary tree初始化和拓扑2.1 Tree的初始化与差异2.2 ncclGetBtreeKernel内部…相关系列
【分布式】NCCL部署与测试 - 01 【分布式】入门级NCCL多机并行实践 - 02 【分布式】小白看Ring算法 - 03 【分布式】大模型分布式训练入门与实践 - 04 目录 相关系列概述1.1 Tree1.2 double binary tree 初始化和拓扑2.1 Tree的初始化与差异2.2 ncclGetBtree Kernel内部实现初始化原语差异流程RunTreeUpDownRunTreeSplit 概述
先掠过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 sm70所有的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在第三位 #mermaid-svg-jxTbd7mVEE7H65tP {font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;fill:#333;}#mermaid-svg-jxTbd7mVEE7H65tP .error-icon{fill:#552222;}#mermaid-svg-jxTbd7mVEE7H65tP .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-jxTbd7mVEE7H65tP .edge-thickness-normal{stroke-width:2px;}#mermaid-svg-jxTbd7mVEE7H65tP .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-jxTbd7mVEE7H65tP .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-jxTbd7mVEE7H65tP .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-jxTbd7mVEE7H65tP .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-jxTbd7mVEE7H65tP .marker{fill:#333333;stroke:#333333;}#mermaid-svg-jxTbd7mVEE7H65tP .marker.cross{stroke:#333333;}#mermaid-svg-jxTbd7mVEE7H65tP svg{font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-jxTbd7mVEE7H65tP .label{font-family:"trebuchet ms",verdana,arial,sans-serif;color:#333;}#mermaid-svg-jxTbd7mVEE7H65tP .cluster-label text{fill:#333;}#mermaid-svg-jxTbd7mVEE7H65tP .cluster-label span{color:#333;}#mermaid-svg-jxTbd7mVEE7H65tP .label text,#mermaid-svg-jxTbd7mVEE7H65tP span{fill:#333;color:#333;}#mermaid-svg-jxTbd7mVEE7H65tP .node rect,#mermaid-svg-jxTbd7mVEE7H65tP .node circle,#mermaid-svg-jxTbd7mVEE7H65tP .node ellipse,#mermaid-svg-jxTbd7mVEE7H65tP .node polygon,#mermaid-svg-jxTbd7mVEE7H65tP .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-jxTbd7mVEE7H65tP .node .label{text-align:center;}#mermaid-svg-jxTbd7mVEE7H65tP .node.clickable{cursor:pointer;}#mermaid-svg-jxTbd7mVEE7H65tP .arrowheadPath{fill:#333333;}#mermaid-svg-jxTbd7mVEE7H65tP .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-jxTbd7mVEE7H65tP .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-jxTbd7mVEE7H65tP .edgeLabel{background-color:#e8e8e8;text-align:center;}#mermaid-svg-jxTbd7mVEE7H65tP .edgeLabel rect{opacity:0.5;background-color:#e8e8e8;fill:#e8e8e8;}#mermaid-svg-jxTbd7mVEE7H65tP .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-jxTbd7mVEE7H65tP .cluster text{fill:#333;}#mermaid-svg-jxTbd7mVEE7H65tP .cluster span{color:#333;}#mermaid-svg-jxTbd7mVEE7H65tP div.mermaidTooltip{position:absolute;text-align:center;max-width:200px;padding:2px;font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:12px;background:hsl(80, 100%, 96.2745098039%);border:1px solid #aaaa33;border-radius:2px;pointer-events:none;z-index:100;}#mermaid-svg-jxTbd7mVEE7H65tP :root{--mermaid-font-family:"trebuchet ms",verdana,arial,sans-serif;}