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

这篇具有很好参考价值的文章主要介绍了【分布式】小白看Ring算法 - 03。希望对大家有所帮助。如果存在错误或未考虑完全的地方,请大家不吝赐教,您也可以点击"举报违法"按钮提交疑问。

相关系列

【分布式】NCCL部署与测试 - 01
【分布式】入门级NCCL多机并行实践 - 02
【分布式】小白看Ring算法 - 03
【分布式】大模型分布式训练入门与实践 - 04文章来源地址https://www.toymoban.com/news/detail-794064.html

概述

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,第一次循环到第三次循环时:
ring算法,linux,分布式,分布式,算法,NCCL,c++,ring,深度学习

我们将绿色视为这次循环需要传输的数据。
数据ABCD在不同的GPU中流通。
最终达到以下情况,scatter-reduce就完成了:
ring算法,linux,分布式,分布式,算法,NCCL,c++,ring,深度学习
将图中蓝色部分输出,就完成了一次ring算法下的Scatter-Reduce。

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

NCCL代码流程

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 environment
  struct 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 Mode Description Calculation of realChunkSize
NCCL_PROTO_SIMPLE Uses a simple communication protocol. realChunkSize = roundUp(min(chunkSize, divUp(size-gridOffset, nChannels)), (nthreads-WARP_SIZE)*sizeof(uint64_t)/sizeof(T))
NCCL_PROTO_LL Uses a linked list (LL) communication protocol. realChunkSize = size-gridOffset < loopSize ? args->coll.lastChunkSize : chunkSize
NCCL_PROTO_LL128 Uses 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 GPU
      rankDest = ringRanks[nranks-1];
      offset = chunkOffset + rankDest * size;
      prims.send(offset, nelem);

      // k-2 steps: reduce and copy to next GPU
      for (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 result
      rankDest = 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 combinations
  NCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph));

  // Compute nChannels per peer for p2p
  NCCLCHECK(ncclTopoComputeP2pChannels(comm));

  if (ncclParamNvbPreconnect()) {
    // Connect p2p when using NVB path
    int 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 connector
          comm->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 connector
          comm->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

到了这里,关于【分布式】小白看Ring算法 - 03的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • 存储、计算、分布式虚拟化篇(收集整理适合小白)

    超融合 硬件资源的统一管理 虚拟化 计算机资源最大化利用 桌面云 虚拟机上的桌面 虚拟桌面资源池 用户使用的云桌面虚拟机 虚拟应用资源池 应用虚拟化的虚拟机 处理器虚拟化 将一个物理的CPU虚拟成多个逻辑CPU供虚拟机使用 超线程技术 将多线程处理器内部的逻辑内核模拟

    2024年01月16日
    浏览(54)
  • 存储、计算、分布式知识点思维导图(收集整理适合小白)

    IO技术 FC协议 光纤通道协议,为了解决I/O传输瓶颈对于整个存储系统带来的消极影响从而产生的光纤通道标准协议簇 iSCSI技术 一种专门为小型计算机系统设计的I/O技术又被成为小型计算机系统接口,通过网络由专门的服务器提供存储管理,已实现数据的远程存储,便于集中管

    2024年03月11日
    浏览(61)
  • 架构师的36项修炼-03架构核心技术之分布式消息队列

    本课时的主题是分布式消息队列,分布式消息队列的知识结构如下图。 本课时主要介绍以下内容。 同步架构和异步架构的区别。异步架构的主要组成部分:消息生产者、消息消费者、分布式消息队列。异步架构的两种主要模型:点对点模型和发布订阅模型。 分布式消息队列

    2024年01月24日
    浏览(41)
  • Spring Cloud学习(十一)【深入Elasticsearch 分布式搜索引擎03】

    聚合(aggregations)可以实现对文档数据的统计、分析、运算。聚合常见的有三类: 桶(Bucket)聚合:用来对文档做分组 TermAggregation:按照文档字段值分组 Date Histogram:按照日期阶梯分组,例如一周为一组,或者一月为一组 度量(Metric)聚合:用以计算一些值,比如:最大值

    2024年01月23日
    浏览(48)
  • (小白全过程记录)Ubuntu下伪分布式Hadoop环境搭建

    目录 0.准备 1.Hadoop伪分布式环境搭建 2.安装ssh,配置ssh无密码登录 3.通过拖拽的方式将文件从windows传到linux桌面 5.安装hadoop 6.修改hadoop环境变量 7.修改配置文件 core-site.xml 8.修改配置文件hdfs-site.xml文件 9.执行NameNode的格式化 10.开启NameNode和DataNode守护进程 11.访问web页面 12.关闭

    2024年02月03日
    浏览(49)
  • 网络安全---Ring3下动态链接库.so函数劫持

    1.1、原理 Unix操作系统中,程序运行时会按照一定的规则顺序去查找依赖的动态链接库,当查找到指定的so文件时,动态链接器(/lib/ld-linux.so.X)会将程序所依赖的共享对象进行装载和初始化,而为什么可以使用so文件进行函数的劫持呢? 这与LINUX的特性有关,先加载的so中的全局

    2024年02月11日
    浏览(38)
  • elasticsearch(ES)分布式搜索引擎03——(RestClient查询文档,ES旅游案例实战)

    文档的查询同样适用昨天学习的 RestHighLevelClient对象,基本步骤包括: 1)准备Request对象 2)准备请求参数 3)发起请求 4)解析响应 我们以match_all查询为例 3.1.1.发起查询请求 代码解读: 第一步,创建 SearchRequest 对象,指定索引库名 第二步,利用 request.source() 构建DSL,DSL中可

    2024年02月07日
    浏览(53)
  • 【分布式】分布式共识算法 --- RAFT

    CAP原则又称CAP定理,指的是在一个分布式系统中,一致性(Consistency)、可用性(Availability)、分区容错性(Partition tolerance) It states, that though its desirable to have Consistency, High-Availability and Partition-tolerance in every system, unfortunately no system can achieve all three at the same time. 在分布式系

    2024年02月06日
    浏览(51)
  • 高级分布式系统-第15讲 分布式机器学习--分布式机器学习算法

    高级分布式系统汇总:高级分布式系统目录汇总-CSDN博客 按照通信步调,大致可以分为同步算法和异步算法两大类。 同步算法下,通信过程中有一个显式的全局同步状态,称之为同步屏障。当工作节点运行到 同步屏障 ,就会进入等待状态,直到其工作节点均运行到同步屏障

    2024年01月18日
    浏览(46)
  • 算法、语言混编、分布式锁与分布式ID、IO模型

    数据结构和算法是程序的基石。我们使用的所有数据类型就是一种数据结构(数据的组织形式),写的程序逻辑就是算法。 算法是指用来操作数据、解决程序问题的一组方法。 对于同一个问题,使用不同的算法,也许最终得到的结果是一样的,但在过程中消耗的资源(空间

    2024年02月08日
    浏览(46)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包