【分布式】小白看Ring算法 - 03

相关系列

【分布式】NCCL部署与测试 - 01
【分布式】入门级NCCL多机并行实践 - 02
【分布式】小白看Ring算法 - 03
【分布式】大模型分布式训练入门与实践 - 04

概述

NCCL（NVIDIA Collective Communications Library）是由NVIDIA开发的一种用于多GPU间通信的库。NCCL的RING算法是NCCL库中的一种通信算法，用于在多个GPU之间进行环形通信。

RING算法的基本思想是将多个GPU连接成一个环形结构，每个GPU与相邻的两个GPU进行通信。数据沿着环形结构传递，直到到达发送方的位置。这样的环形结构可以有效地利用GPU之间的带宽，提高通信的效率。

RING算法的步骤如下：

数据拷贝
数据沿着环形路径传递
传输完成
进行下一轮通信/结束通信过程
初始化
通信缓冲区
等待
接收方

Scatter-Reduce

以Scatter-Reduce为例，假设有4张GPU，RANK_NUM=4。
则需要根据RANK_NUM把每张CPU划分为4个chunk。
为什么要这么划分？

在 NCCL 中，划分 chunk 的数量与 GPU 的数量相关联，这是因为 chunk 的目的是将大的消息划分为多个小的数据块，以便并行处理和降低通信的延迟。这种划分通常会基于 GPU 的数量，以确保每个 GPU 可以处理到一部分数据块，从而提高整体的通信效率。

  1. 并行性： 划分 chunk 可以增加通信的并行性。每个 GPU 处理自己的数据块，不同的 GPU 可以并行地执行通信操作，从而提高整体的吞吐量。
  2. 减少延迟： 较小的数据块通常可以更快地传输，因此通过划分 chunk，可以减少每个通信操作的延迟。这对于一些对通信延迟敏感的应用程序是至关重要的。
  3. 资源分配： NCCL 可能会根据 GPU 的数量来分配适当的资源，例如内存等。通过划分 chunk，可以更好地管理这些资源。
  4. Load Balancing： 均衡负载是分布式系统中的一个关键问题。通过根据 GPU 的数量划分 chunk，可以更容易地实现负载均衡，确保每个 GPU 处理的工作量相对均匀。

划分了chunk以后，我们一次RING的通路将会走通4块GPU，每次只传输一块chunk的数据。这样需要走很多次通路才能把所有数据传输完。
假如 ringIx=0，第一次循环到第三次循环时：
在这里插入图片描述

我们将绿色视为这次循环需要传输的数据。
数据ABCD在不同的GPU中流通。
最终达到以下情况，scatter-reduce就完成了：
在这里插入图片描述
将图中蓝色部分输出，就完成了一次ring算法下的Scatter-Reduce。

当然，如果要做All-Reduce，此时不需要继续按照原来的规则计算类，理论上只需要再算一次All-Gather，就能把蓝色的块分发给其他几块卡。All-Reduce的相关讲解网络上很多。此处就不讲了。

NCCL代码流程

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->hostHash=getHostHash()+commHash;info->pidHash=getPidHash()+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 List，LL）的通信协议。
差异点： 在计算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 realChunkSize
NCCL_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 : chunkSize
NCCL_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 j=2; j<nranks; ++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 r=0; r<nvbNpeers; r++) {int peer = nvbPeers[r];int delta = (comm->nRanks + (comm->rank-peer)) % comm->nRanks;for (int c=0; c<comm->p2pnChannelsPerPeer; c++) {int channelId = (delta+comm->p2pChannels[c]) % comm->p2pnChannels;if (comm->channels[channelId].peers[peer].recv[0].connected == 0) { // P2P uses only 1 connectorcomm->connectRecv[peer] |= (1<<channelId);}}delta = (comm->nRanks - (comm->rank-peer)) % comm->nRanks;for (int c=0; c<comm->p2pnChannelsPerPeer; c++) {int channelId = (delta+comm->p2pChannels[c]) % comm->p2pnChannels;if (comm->channels[channelId].peers[peer].send[0].connected == 0) { // P2P uses only 1 connectorcomm->connectSend[peer] |= (1<<channelId);}}}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

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/165141.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

GoLand 2023.2.5(GO语言集成开发工具环境)

GoLand 2023.2.5(GO语言集成开发工具环境)

GoLand是一款专门为Go语言开发者打造的集成开发环境&#xff08;IDE&#xff09;。它能够提供一系列功能&#xff0c;如代码自动完成、语法高亮、代码格式化、代码重构、代码调试等等&#xff0c;使编写代码更加高效和舒适。 GoLand的特点包括&#xff1a; 1. 智能代码补全&a…
阅读更多...
Ubuntu安装CUDA驱动

Ubuntu安装CUDA驱动

Ubuntu安装CUDA驱动 前言官网安装确认安装版本安装CUDA Toolkit 前言 CUDA驱动一般指CUDA Toolkit&#xff0c;可通过Nvidia官网下载安装。本文介绍安装方法。 官网 CUDA Toolkit 最新版&#xff1a;CUDA Toolkit Downloads | NVIDIA Developer CUDA Toolkit 最新版文档&…
阅读更多...
NX二次开发UF_CAM_update_list_object_customization 函数介绍

NX二次开发UF_CAM_update_list_object_customization 函数介绍

文章作者&#xff1a;里海 来源网站&#xff1a;https://blog.csdn.net/WangPaiFeiXingYuan UF_CAM_update_list_object_customization Defined in: uf_cam.h int UF_CAM_update_list_object_customization(tag_t * object_tags ) overview 概述 This function provids the…
阅读更多...
UDP客户端使用connect与UDP服务器使用send函数和recv函数收发数据

UDP客户端使用connect与UDP服务器使用send函数和recv函数收发数据

服务器代码编译运行 服务器udpconnectToServer.c的代码如下&#xff1a; #include<stdio.h> #include<stdlib.h> #include<string.h> #include<unistd.h> #include<arpa/inet.h> #include<sys/socket.h> #include<errno.h> #inclu…
阅读更多...
Okhttp 浅析

Okhttp 浅析

安全的连接 OkHttpClient: OkHttpClient: 1.线程调度 2.连接池,有则复用,没有就创建 3.interceptor 4.interceptor 5.监听工厂 6.是否失败重试 7.自动修正访问,如果没有权限或认证 8是否重定向 followRedirects 9.协议切换时候是否继续重定向 10.Cookie jar 容器 默认…
阅读更多...
pycharm 创建的django目录和命令行创建的django再使用pycharm打开的目录对比截图 及相关

pycharm 创建的django目录和命令行创建的django再使用pycharm打开的目录对比截图 及相关

pytcharm创建django的项目 命令行创建的django 命令行创建项目时 不带路径时 (.venv) D:\gbCode>django-admin startproject gbCode 命令行创建项目时 带路径时 -- 所以如果有目录就指定路径好 (.venv) D:\gbCode>django-admin startproject gbCode d:\gbCode\
阅读更多...
洛谷P1219 [USACO1.5] 八皇后【n皇后问题】【深搜+回溯 经典题】【附O(1)方法】

洛谷P1219 [USACO1.5] 八皇后【n皇后问题】【深搜+回溯 经典题】【附O(1)方法】

P1219 [USACO1.5] 八皇后 Checker Challenge 前言题目题目描述输入格式输出格式样例 #1样例输入 #1样例输出 #1 提示题目分析注意事项 代码深搜回溯打表 后话额外测试用例样例输入 #2样例输出 #2 王婆卖瓜 题目来源 前言 也是说到做到&#xff0c;来做搜索的题&#xff08;虽…
阅读更多...
微机原理_2

微机原理_2

一、单项选择题(本大题共15小题,每小题3分,共45分。在每小题给出的四个备选项中,选出一个正确的答案&#xff0c;请将选定的答案填涂在答题纸的相应位置上。&#xff09; 下列数中最大的数为&#xff08;&#xff09; A. 10010101B B. (126)8 C. 96H D. 100 CPU 执行 OUT 60H,…
阅读更多...
西门子（Siemens）仿真PLC启动报错处理

西门子（Siemens）仿真PLC启动报错处理

目录 一、背景&#xff1a; 二、卸载软件 三、安装软件 三、启动软件 四、下载PORTAL项目 五、测试 一、背景&#xff1a; 在启动S7-PLCSIM Advanced V3.0仿真PLC时报错&#xff0c;报错信息为&#xff1a;>>Siemens PLCSIM Virtual Switch<<is misconfigu…
阅读更多...
Ubuntu 23.10 服务器版本 ifconfig 查不到网卡 ip(已解决)

Ubuntu 23.10 服务器版本 ifconfig 查不到网卡 ip(已解决)

文章目录 1、问题描述2、 解决方案 1、问题描述 服务器&#xff1a;ubuntu 23.10 经常会遇到虚拟机添加仅主机网卡后&#xff0c;通过 ifconfig 无法获取其网卡 ip 2、 解决方案 修改网卡配置文件&#xff1a; # 进入网卡配置文件目录 cd /etc/netplan # 备份原始文件 cp …
阅读更多...
ArgoWorkflow教程(一)---DevOps 另一选择？云原生 CICD: ArgoWorkflow 初体验

ArgoWorkflow教程(一)---DevOps 另一选择？云原生 CICD: ArgoWorkflow 初体验

来自&#xff1a;探索云原生 https://www.lixueduan.com 原文&#xff1a;https://www.lixueduan.com/posts/devops/argo-workflow/01-deploy-argo-workflows/ 本文主要记录了如何在 k8s 上快速部署云原生的工作流引擎 ArgoWorkflow。 ArgoWorkflow 是什么 Argo Workflows 是…
阅读更多...
网络安全如何自学？

网络安全如何自学？

1.网络安全是什么 网络安全可以基于攻击和防御视角来分类&#xff0c;我们经常听到的 “红队”、“渗透测试” 等就是研究攻击技术&#xff0c;而“蓝队”、“安全运营”、“安全运维”则研究防御技术。 2.网络安全市场 一、是市场需求量高&#xff1b; 二、则是发展相对成熟…
阅读更多...
Android设计模式--装饰模式

Android设计模式--装饰模式

千淘万漉虽辛苦&#xff0c;吹尽黄沙始到金 一&#xff0c;定义 动态地给一个对象添加一些额外的职责。就增加功能来说&#xff0c;装饰模式相比生成子类更为灵活。 装饰模式也叫包装模式&#xff0c;结构型设计模式之一&#xff0c;其使用一种对客户端透明的方式来动态地扩展…
阅读更多...
＜JavaEE＞ 什么是线程（Thread）？进程和线程有什么区别？

＜JavaEE＞ 什么是线程（Thread）？进程和线程有什么区别？

目录 一、线程&#xff08;Thread&#xff09;的概念 二、线程存在的意义 2.1 并发编程 2.2 比进程更“轻量” 三、使用线程时应该注意 四、进程和线程的区别 五、Java中的线程和操作系统中的线程是不同的概念 六、多线程编程 一、线程&#xff08;Thread&#xff09;的…
阅读更多...
图神经网络的数学原理总结

图神经网络的数学原理总结

图深度学习(Graph Deep Learning) 多年来一直在加速发展。许多现实生活问题使GDL成为万能工具&#xff1a;在社交媒体、药物发现、芯片植入、预测、生物信息学等方面都显示出了很大的前景。 本文将流行的图神经网络及其数学细微差别的进行详细的梳理和解释&#xff0c;图深度学…
阅读更多...
Linux中flask项目开启https访问

Linux中flask项目开启https访问

1.下载阿里云免费证书 2.项目添加https配置 3.服务器开启https访问 3.1 重新安装OpenSSL 3.2.重新安装Python 上一次已经讲过Linux安装部署Python: Linux安装Python3.10与部署flask项目实战详细记录,今天记录一下Python项目如何支持https访问…
阅读更多...
性能相关的闪存特性

性能相关的闪存特性

一、多Plane操作 上章提到若干个Plane组成Die或者叫LUN,即一个Die上有多个Plane 每次进行写操作时&#xff0c;控制器先将数据写入页缓存中&#xff0c;等同一个Die上另一个Plane也写数据的时候&#xff0c;再同时写入&#xff0c;原来单独操作一个Plane的时间变成了可以同时做…
阅读更多...
Springmvc实现增删改差

Springmvc实现增删改差

一、包结构 二、各层代码 (1)数据User public class User {private Integer id;private String userName;private String note;public User() {super();}public User(Integer i, String userName, String note) {super();this.id i;this.userName userName;this.note note;…
阅读更多...
Qt实现自定义IP地址输入控件（百分百还原Windows 10网络地址输入框）

Qt实现自定义IP地址输入控件（百分百还原Windows 10网络地址输入框）

在开发网络相关的程序时，我们经常需要输入IP地址，例如源地址和目标地址。Qt提供了一些基础的控件，如QLineEdit，但是它们并不能满足我们对IP地址输入的要求，例如限制输入的格式、自动跳转到下一个输入框、处理回车和退格键等。因此，我们需要自己编写一个自定义的IP地址输入…
阅读更多...
【赠书第7期】从零基础到精通Flutter开发

【赠书第7期】从零基础到精通Flutter开发

文章目录 前言 1 安装Flutter和Dart 2 了解Flutter的基础概念 2.1 Widget 2.2 MaterialApp和Scaffold 2.3 Hot Reload 3 编写你的第一个Flutter应用 3.1 创建一个Flutter项目 3.2 修改默认页面 3.3 添加交互 4 深入学习Flutter高级特性 4.1 路由和导航 4.2 状态管…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023