怎么查看网站有没有做推广,网站关键词搜索排名,开发电子商务系统的五个步骤,企业营销型网站建设哪家公司好相关系列
【分布式】NCCL部署与测试 - 01 【分布式】入门级NCCL多机并行实践 - 02 【分布式】小白看Ring算法 - 03 【分布式】大模型分布式训练入门与实践 - 04
概述
NCCL#xff08;NVIDIA Collective Communications Library#xff09;是由NVIDIA开发的一种用于多GPU间…相关系列
【分布式】NCCL部署与测试 - 01 【分布式】入门级NCCL多机并行实践 - 02 【分布式】小白看Ring算法 - 03 【分布式】大模型分布式训练入门与实践 - 04
概述
NCCLNVIDIA Collective Communications Library是由NVIDIA开发的一种用于多GPU间通信的库。NCCL的RING算法是NCCL库中的一种通信算法用于在多个GPU之间进行环形通信。
RING算法的基本思想是将多个GPU连接成一个环形结构每个GPU与相邻的两个GPU进行通信。数据沿着环形结构传递直到到达发送方的位置。这样的环形结构可以有效地利用GPU之间的带宽提高通信的效率。
RING算法的步骤如下 #mermaid-svg-ZOIRjv5qKPKIP8Sq {font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;fill:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .error-icon{fill:#552222;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edge-thickness-normal{stroke-width:2px;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .marker{fill:#333333;stroke:#333333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .marker.cross{stroke:#333333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq svg{font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .label{font-family:"trebuchet ms",verdana,arial,sans-serif;color:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .cluster-label text{fill:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .cluster-label span{color:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .label text,#mermaid-svg-ZOIRjv5qKPKIP8Sq span{fill:#333;color:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .node rect,#mermaid-svg-ZOIRjv5qKPKIP8Sq .node circle,#mermaid-svg-ZOIRjv5qKPKIP8Sq .node ellipse,#mermaid-svg-ZOIRjv5qKPKIP8Sq .node polygon,#mermaid-svg-ZOIRjv5qKPKIP8Sq .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .node .label{text-align:center;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .node.clickable{cursor:pointer;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .arrowheadPath{fill:#333333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edgeLabel{background-color:#e8e8e8;text-align:center;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .edgeLabel rect{opacity:0.5;background-color:#e8e8e8;fill:#e8e8e8;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .cluster text{fill:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq .cluster span{color:#333;}#mermaid-svg-ZOIRjv5qKPKIP8Sq 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-ZOIRjv5qKPKIP8Sq :root{--mermaid-font-family:"trebuchet ms",verdana,arial,sans-serif;} 数据拷贝 数据沿着环形路径传递 传输完成 进行下一轮通信/结束通信过程 初始化 通信缓冲区 等待 接收方 Scatter-Reduce
以Scatter-Reduce为例假设有4张GPURANK_NUM4。 则需要根据RANK_NUM把每张CPU划分为4个chunk。 为什么要这么划分
在 NCCL 中划分 chunk 的数量与 GPU 的数量相关联这是因为 chunk 的目的是将大的消息划分为多个小的数据块以便并行处理和降低通信的延迟。这种划分通常会基于 GPU 的数量以确保每个 GPU 可以处理到一部分数据块从而提高整体的通信效率。
并行性 划分 chunk 可以增加通信的并行性。每个 GPU 处理自己的数据块不同的 GPU 可以并行地执行通信操作从而提高整体的吞吐量。减少延迟 较小的数据块通常可以更快地传输因此通过划分 chunk可以减少每个通信操作的延迟。这对于一些对通信延迟敏感的应用程序是至关重要的。资源分配 NCCL 可能会根据 GPU 的数量来分配适当的资源例如内存等。通过划分 chunk可以更好地管理这些资源。Load Balancing 均衡负载是分布式系统中的一个关键问题。通过根据 GPU 的数量划分 chunk可以更容易地实现负载均衡确保每个 GPU 处理的工作量相对均匀。
划分了chunk以后我们一次RING的通路将会走通4块GPU每次只传输一块chunk的数据。这样需要走很多次通路才能把所有数据传输完。 假如 ringIx0第一次循环到第三次循环时
我们将绿色视为这次循环需要传输的数据。 数据ABCD在不同的GPU中流通。 最终达到以下情况scatter-reduce就完成了 将图中蓝色部分输出就完成了一次ring算法下的Scatter-Reduce。
当然如果要做All-Reduce此时不需要继续按照原来的规则计算类理论上只需要再算一次All-Gather就能把蓝色的块分发给其他几块卡。All-Reduce的相关讲解网络上很多。此处就不讲了。
NCCL代码流程 #mermaid-svg-EN0VkjtBL2AySpsh {font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;fill:#333;}#mermaid-svg-EN0VkjtBL2AySpsh .error-icon{fill:#552222;}#mermaid-svg-EN0VkjtBL2AySpsh .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-EN0VkjtBL2AySpsh .edge-thickness-normal{stroke-width:2px;}#mermaid-svg-EN0VkjtBL2AySpsh .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-EN0VkjtBL2AySpsh .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-EN0VkjtBL2AySpsh .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-EN0VkjtBL2AySpsh .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-EN0VkjtBL2AySpsh .marker{fill:#333333;stroke:#333333;}#mermaid-svg-EN0VkjtBL2AySpsh .marker.cross{stroke:#333333;}#mermaid-svg-EN0VkjtBL2AySpsh svg{font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-EN0VkjtBL2AySpsh .label{font-family:"trebuchet ms",verdana,arial,sans-serif;color:#333;}#mermaid-svg-EN0VkjtBL2AySpsh .cluster-label text{fill:#333;}#mermaid-svg-EN0VkjtBL2AySpsh .cluster-label span{color:#333;}#mermaid-svg-EN0VkjtBL2AySpsh .label text,#mermaid-svg-EN0VkjtBL2AySpsh span{fill:#333;color:#333;}#mermaid-svg-EN0VkjtBL2AySpsh .node rect,#mermaid-svg-EN0VkjtBL2AySpsh .node circle,#mermaid-svg-EN0VkjtBL2AySpsh .node ellipse,#mermaid-svg-EN0VkjtBL2AySpsh .node polygon,#mermaid-svg-EN0VkjtBL2AySpsh .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-EN0VkjtBL2AySpsh .node .label{text-align:center;}#mermaid-svg-EN0VkjtBL2AySpsh .node.clickable{cursor:pointer;}#mermaid-svg-EN0VkjtBL2AySpsh .arrowheadPath{fill:#333333;}#mermaid-svg-EN0VkjtBL2AySpsh .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-EN0VkjtBL2AySpsh .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-EN0VkjtBL2AySpsh .edgeLabel{background-color:#e8e8e8;text-align:center;}#mermaid-svg-EN0VkjtBL2AySpsh .edgeLabel rect{opacity:0.5;background-color:#e8e8e8;fill:#e8e8e8;}#mermaid-svg-EN0VkjtBL2AySpsh .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-EN0VkjtBL2AySpsh .cluster text{fill:#333;}#mermaid-svg-EN0VkjtBL2AySpsh .cluster span{color:#333;}#mermaid-svg-EN0VkjtBL2AySpsh 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-EN0VkjtBL2AySpsh :root{--mermaid-font-family:"trebuchet ms",verdana,arial,sans-serif;} 1 1 1 1 2 2 2 2 4 4 4 4 5 5 5 5 6 6 6 6 7 7 7 7 8 8 8 8 9 9 9 9 10 10 10 10 11 11 11 11 12 12 12 12 13 13 13 13 rank0:fillInfo bootstrap AllGather rank1:fillInfo rank2:fillInfo rank3:fillInfo rank0:getSystem rank1:getSystem rank2:getSystem rank3:getSystem rank0:computePath rank1:computePath rank2:computePath rank3:computePath rank0:search channel rank1:search channel rank2:search channel rank3:search channel bootstrap AllGather rank0:connect rank1:connect rank2:connect rank3:connect rank0:setupChannel rank1:setupChannel rank2:setupChannel rank3:setupChannel rank0:p2pSetup rank1:p2pSetup rank2:p2pSetup rank3:p2pSetup rank0:tuneModel rank1:tuneModel rank2:tuneModel rank3:tuneModel rank0:p2pChannel rank1:p2pChannel rank2:p2pChannel rank3:p2pChannel bootstrap IntraNodeBarrier rank0:NetProxy rank1:NetProxy rank2:NetProxy rank3:NetProxy fillInfo: 这段代码在init.cc中
static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, uint64_t commHash) {info-rank comm-rank;CUDACHECK(cudaGetDevice(info-cudaDev));info-hostHashgetHostHash()commHash;info-pidHashgetPidHash()commHash;// Get the device MAJOR:MINOR of /dev/shm so we can use that// information to decide whether we can use SHM for inter-process// communication in a container environmentstruct stat statbuf;SYSCHECK(stat(/dev/shm, statbuf), stat);info-shmDev statbuf.st_dev;info-busId comm-busId;NCCLCHECK(ncclGpuGdrSupport(info-gdrSupport));return ncclSuccess;
}这段代码的目的是为了获取和存储与通信相关的信息以便在NCCL通信中使用。其中包括设备标识、主机哈希、进程ID哈希、共享内存设备标识、总线ID以及对GDR的支持情况等。
在initTransportsRank中搜索完信息并作第一次AllGather, 收集所有通信节点的信息。 然后再为通信组分配额外的内存以存储每个通信节点的信息包括一个额外的用于表示CollNet root的位置。 遍历节点和复制信息时需要检查是否存在相同主机哈希和总线ID的重复GPU。如果是发出警告并返回ncclInvalidUsage错误。
后面的一系列过程就是计算路径然后这里涉及一些搜索算法通常会将BFS搜索到的路径都存在一个位置选择更优的路径。 搜索时也会根据实际情况判断选择ring算法或者tree算法。 搜索内容可能是无穷的因此NCCL设置了一个超时时间超过该时间则终端搜索。 完成路径的计算后再做一次AllGather。
来到scatter-reduce的实现部分 size_t realChunkSize;if (Proto::Id NCCL_PROTO_SIMPLE) {realChunkSize min(chunkSize, divUp(size-gridOffset, nChannels));realChunkSize roundUp(realChunkSize, (nthreads-WARP_SIZE)*sizeof(uint64_t)/sizeof(T));}else if (Proto::Id NCCL_PROTO_LL)realChunkSize size-gridOffset loopSize ? args-coll.lastChunkSize : chunkSize;else if (Proto::Id NCCL_PROTO_LL128)realChunkSize min(divUp(size-gridOffset, nChannels*minChunkSizeLL128)*minChunkSizeLL128, chunkSize);realChunkSize int(realChunkSize);ssize_t chunkOffset gridOffset bid*int(realChunkSize);这里涉及了NCCL协议的通信模式 一共有三种分别是NCCL_PROTO_SIMPLE、NCCL_PROTO_LL和NCCL_PROTO_LL128。
NCCL_PROTO_SIMPLE
描述 使用简单的通信协议。 差异点 计算realChunkSize时采用了一些特殊的逻辑其中min(chunkSize, divUp(size-gridOffset, nChannels))用于确定实际的块大小并通过roundUp调整为合适的大小。这可能涉及到性能和资源的考虑以及对通信模式的调整。
NCCL_PROTO_LL
描述 使用连续链表Linked ListLL的通信协议。 差异点 在计算realChunkSize时首先检查size-gridOffset loopSize条件如果为真则使用args-coll.lastChunkSize否则使用默认的chunkSize。这可能与LL协议的特性有关具体考虑了循环的情况。 NCCL_PROTO_LL128
描述 使用连续链表的通信协议每次传输128字节。 差异点 计算realChunkSize时采用了min(divUp(size-gridOffset, nChannels*minChunkSizeLL128)*minChunkSizeLL128, chunkSize)的逻辑。这考虑了128字节的限制以及对通信块大小的一些限制。 总体来说这三种协议模式的区别主要体现在计算realChunkSize的逻辑上这可能受到性能、资源利用、通信模式等方面的不同考虑。具体选择哪种协议模式通常取决于系统的特性和应用场景的需求。
Protocol ModeDescriptionCalculation of realChunkSizeNCCL_PROTO_SIMPLEUses a simple communication protocol.realChunkSize roundUp(min(chunkSize, divUp(size-gridOffset, nChannels)), (nthreads-WARP_SIZE)*sizeof(uint64_t)/sizeof(T))NCCL_PROTO_LLUses a linked list (LL) communication protocol.realChunkSize size-gridOffset loopSize ? args-coll.lastChunkSize : chunkSizeNCCL_PROTO_LL128Uses a linked list (LL) communication protocol, with each transfer involving 128 bytes.realChunkSize min(divUp(size-gridOffset, nChannels*minChunkSizeLL128)*minChunkSizeLL128, chunkSize)
最后是正式计算部分 /////////////// begin ReduceScatter steps ///////////////ssize_t offset;int nelem min(realChunkSize, size-chunkOffset);int rankDest;// step 0: push data to next GPUrankDest ringRanks[nranks-1];offset chunkOffset rankDest * size;prims.send(offset, nelem);// k-2 steps: reduce and copy to next GPUfor (int j2; jnranks; j) {rankDest ringRanks[nranks-j];offset chunkOffset rankDest * size;prims.recvReduceSend(offset, nelem);}// step k-1: reduce this buffer and data, which will produce the final resultrankDest ringRanks[0];offset chunkOffset rankDest * size;prims.recvReduceCopy(offset, chunkOffset, nelem, /*postOp*/true);ssize_t offset; int nelem min(realChunkSize, size-chunkOffset); int rankDest;
offset 是一个偏移量变量用于指定数据在通信缓冲区中的位置。 nelem 表示每次操作的元素个数取 realChunkSize 和 size-chunkOffset 的较小值。 rankDest 是目标GPU的排名。
第一步将数据推送到下一个GPU。 计算目标GPU的排名 rankDest 和在通信缓冲区中的偏移量 offset。 调用 prims.send 函数将数据从当前GPU发送到目标GPU。 // k-2 steps: reduce and copy to next GPU
第2到第k-1步 将数据在环形路径上经过各个GPU节点依次进行Reduce操作并将结果复制到下一个GPU。 通过循环依次计算目标GPU的排名 rankDest 和在通信缓冲区中的偏移量 offset。 调用 prims.recvReduceSend 函数接收数据并执行Reduce操作然后将结果发送到下一个GPU。
第k-1步 将最后一个GPU的数据进行Reduce操作得到最终的结果。 计算目标GPU的排名 rankDest 和在通信缓冲区中的偏移量 offset。 调用 prims.recvReduceCopy 函数接收数据并执行Reduce操作然后将结果复制到指定的位置最终产生最终的ReduceScatter结果。
在实际运行中我们在host端的代码只是规定计算流当这些定义好的原子操作加入到stream中去以后就由固定的流来分配实际运行的情况了。
加入Barria在本地intra-node执行一个屏障操作确保同一节点内的所有GPU都达到了同步点。 // Compute time models for algorithm and protocol combinationsNCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, treeGraph, ringGraph, collNetGraph));// Compute nChannels per peer for p2pNCCLCHECK(ncclTopoComputeP2pChannels(comm));if (ncclParamNvbPreconnect()) {// Connect p2p when using NVB pathint nvbNpeers;int* nvbPeers;NCCLCHECK(ncclTopoGetNvbGpus(comm-topo, comm-rank, nvbNpeers, nvbPeers));for (int r0; rnvbNpeers; r) {int peer nvbPeers[r];int delta (comm-nRanks (comm-rank-peer)) % comm-nRanks;for (int c0; ccomm-p2pnChannelsPerPeer; c) {int channelId (deltacomm-p2pChannels[c]) % comm-p2pnChannels;if (comm-channels[channelId].peers[peer].recv[0].connected 0) { // P2P uses only 1 connectorcomm-connectRecv[peer] | (1channelId);}}delta (comm-nRanks - (comm-rank-peer)) % comm-nRanks;for (int c0; ccomm-p2pnChannelsPerPeer; c) {int channelId (deltacomm-p2pChannels[c]) % comm-p2pnChannels;if (comm-channels[channelId].peers[peer].send[0].connected 0) { // P2P uses only 1 connectorcomm-connectSend[peer] | (1channelId);}}}NCCLCHECK(ncclTransportP2pSetup(comm, NULL, 0));free(nvbPeers);}NCCLCHECK(ncclCommSetIntraProc(comm, intraProcRank, intraProcRanks, intraProcRank0Comm));/* Local intra-node barrier */NCCLCHECK(bootstrapBarrier(comm-bootstrap, comm-intraNodeGlobalRanks, intraNodeRank, intraNodeRanks, (int)intraNodeRank0pidHash));if (comm-nNodes) NCCLCHECK(ncclProxyCreate(comm));以上就是整个scatter-reduce的流程。
相关系列
【分布式】NCCL部署与测试 - 01 【分布式】入门级NCCL多机并行实践 - 02 【分布式】小白看Ring算法 - 03 【分布式】大模型分布式训练入门与实践 - 04
