GPU集合通信库在B站的应用和改进

1. 背景

上篇文章 万字长文解析：大模型需要怎样的硬件算力 深入探讨了大型语言模型（LLMs）在硬件资源方面的需求和面临的挑战，详尽地阐述了如何进行大模型的硬件选型，以及在实际工作中如何根据模型的特定需求来优化硬件资源配置。继此话题之后，本篇文章将重点介绍支撑大模型运作的核心组件——集合通信库，介绍其在大模型架构中的关键作用和实现机制，以及B站是如何应用和改进它的。

随着模型规模的不断增长，单块显卡已经无法满足模型对于显存的需求，分布式训练逐渐成为主流，其中通信库负责了拓扑感知、集合通信原语实现、数据传输等工作，扮演着至关重要的角色。在分布式训练集群逐步普及和规模化的过程中，各个厂商，尤其是云和GPU硬件制造商，对于整个集群的性能和效率不断提出更高的要求，也因此涌现了一批xCCLs(x Collective Communication Libraries)，例如HCCL、ACCL、oneCCL和TCCL等，从侧面也反映了通信库的重要性。

鉴于通信库的原理和实现都异曲同工，本文只针对开源的NCCL通信库来进行讲解，结合B站大模型训练的落地实践经验，拆分解析AI基础软件中通信库的实现，以及如何在工程化中，借助通信库来发现和解决问题。

2. 什么是集合通信库

集合通信库是实现高性能多GPU计算的核心组件，它允许GPU之间进行高效的数据交换和协同工作。

通信库提供了一系列基础的通信原语，这些原语是构建复杂通信模式和算法的基础，以下是一些集合通信的基本概念：

• 集合通信（Collective Communication）：指的是多个节点以特定的规则进行信息交互；

• 节点：指的是参与集群通信的成员，它们可以是CPU进程也可以是GPU卡、NPU卡。如果是CPU进程之间的集合通信，就是CPU并行计算，如果都是GPU(or NPU)就是GPU并行计算；

• Ring：Ring是一种通信拓扑结构，它将参与通信的GPU组织成一个逻辑环。在这个环中，每个GPU既是发送者也是接收者，数据在环中的GPU间以特定的顺序按顺序传递；

• Channel：实现通信的具体路径，一个Channel表示一个通信路径，它允许数据在两个GPU之间或者在GPU与其他设备之间流动。为了更好地利用带宽和网卡，以及使得同一块数据可以通过多个Channel并发通信，通信库一般会使用多个Channel；Ring是描述GPU间通信的拓扑结构，而Channel是实现通信的具体路径；

• world size：参与整个训练过程的总设备数，这些设备可能包括GPU、TPU或其他类型的处理器；

• 通信域（Communicators）：用于表示一组GPU之间的通信上下文。创建通信域时，每个GPU都会被分配一个唯一的标识符，即rank，这些rank标识符用于在集体通信操作中确定数据传输的方向和顺序；通信域的建立是进行集体通信的前提，它允许不同的GPU之间相互识别，并开始数据的发送和接收；

• 集合通信原语（Collective Primitives）: 通信原语是通信域中ranks之间信息交互的特定规则；以班级学生的例子来说明：每个学生有语文、数学、英语等几门学科作业。在每个“群聊组”中约定把哪些人哪些学科的作业发给谁的规则，就是通信原语；

3. 挑战和问题

（图1）

如图1，由于通信库位于框架以下，对用户是无感知的，因此在训练过程中呈现出来的一些问题，如训练中断或性能不佳等，除了模型本身的实现和超参影响外，还可能是通信库的拓扑感知或数据传输出了问题。为了能够快速定位这些问题，我们需要把通信库隐藏起来的，或未显示输出的一些关键信息和指标，呈现给用户，用来辅助模型的设计和训练者，更好的利用底层硬件资源。以NCCL为例，我们总结目前所面临的挑战和问题包括：

1. 通信拓扑中实际物理GPU和网卡互联信息的缺失；

2. 各个通信原语所处理完成数据量的缺失；

3. 机内和机间传输数据量的缺失；

4. 通信域和框架并行设计间映射的缺失；

5. 真实网络拓扑信息的缺失；

本文将主要介绍B站针对上述第1至第3问题的解决方案。对于第4和第5问题，将在后续的文章中进行详细阐述。

为解决问题1，我们基于NCCL的开源版本强化了拓扑信息的解析能力；对于问题2和问题3，我们引入了带宽数据统计功能，可以实时监控和分析通信原语处理的数据量以及机内和机间传输的数据量，为性能调优提供数据支持。

4. 什么是拓扑感知

NCCL支持了Flat Ring 和Double Binary Tree两种图算法，以适应不同规模的集群需求。在小型至中等规模的集群中，即GPU数量从数十到数百的范围内，Flat Ring算法因其通信模式简单、延迟低而展现出优势。然而，当扩展到更大规模的集群，比如数千甚至数万GPU时，Double Binary Tree算法因其卓越的扩展性和流水线处理能力，成为更佳的选择。本文主要以简单的Flat Ring算法为例，来介绍NCCL通信拓扑感知流程。

通过扫描本地机器的PCI硬件信息，NCCL可以获取到GPU和网卡是如何互联的，他们可以是直接连接在一个PCI桥下，也可以是跨数个桥甚至Socket才能连接，另外还有NVLink这样的高速连接方式，可参考v100的硬件架构：

（图2）

每个机器节点内部根据这些硬件信息，结合框架层面Rank，World Size等参数，来决定哪些GPU和网卡以什么方式进行连接，形成本地小环，多机情况下环的首尾都会是网卡，多台机器之间的本地小环直接按顺序对应相连，形成多节点的大环（Ring），本地GPU间以机内传输，机器之间以大环对应的网卡进行网络传输，实现封闭的数据传输拓扑。

单机情况，World Size参数为4时，机内搜索到的第一条环路示例如下：

•  GPU/0 GPU/1 GPU/2 GPU/3

数据在环中的GPU间以特定的顺序传递，每个GPU既是数据发送者也是接收者，以此环路为例，最终形成环路如图3：

（图3）

如上图3，GPU0总是固定发送数据到GPU1，接收来自GPU3的数据。

多机情况，World Size参数为8时，机内搜索到的第一条环路示例如下：

        服务器节点0(node0)：

• NET/0 GPU/0 GPU/1 GPU/2 GPU/3 NET/0 

        服务器节点1(node1)：

• NET/0 GPU/4 GPU/5 GPU/6 GPU/7 NET/0

        两台服务器节点内部都搜索到相应环路，环路的首尾节点通常都是一张网卡，不同机器节点间通过网络传输实现小环的相连形成了更大的环路，以上述两台机器的第一条环路为例，最终形成机间环路如图4：

（图4）

如上图4，GPU0总是固定发送数据到GPU1，通过机间网络传输接收来自GPU7的数据；GPU3总是固定通过机间网络发送数据到GPU4，接收来自GPU2的数据；

上述机内和机间的环路示例只选取了探测到的第一条环路，在实际应用中，通信库会根据机器的PCI、QPI、网卡等硬件的带宽特性，探测并建立多条环路。例如在多机环境下，服务器节点中有多少个网卡，通信库就会依次使用不同的本地网卡作为环路的起始节点进行探测，一般情况下会形成与网卡数量一致的环路数量。

通信库支持同一数据块在多个环路进行并行传输，从而实现资源的充分利用。为了达到这个目的，NCCL通常会把最初探测到的环路数量翻倍，这意味着，最终形成的环路数量将是最初探测到的环路数量的两倍。这种环路复制的方法显著提升了数据传输的并行度，从而增强了整体的通信效率。为了提供更多的灵活性，NCCL还允许通过设置环境变量来控制最终形成的环路数量，但这个数量不会超过32个，这样既保证了通信的高效性，又允许用户根据不同的应用场景和硬件配置进行适当的调整。

拓扑探测流程实现在initTransportsRank函数中，其主要参数comm→topo描述了一个图结构，包括顶点和边。GPU，CPU，NET等设备节点都属于图的顶点，顶点到其他相邻顶点的连接作为边，边存放了两个设备间通信的路径信息，主要为带宽，路径类型等。本地拓扑检测主要是对comm→topo参数的填充，再通过它来搜索找到带宽最大化的路径，形成本地环路，大致流程如图5：

（图5）

每块GPU对应的训练任务都会通过这个流程完成完整拓扑的汇聚，所有的分布式任务都会存有一份环路信息，来完成之后的集合通信操作。通常框架都会创建多个环来实现各种并发技术，这种情况下会创建多个大小不同的环。感知本地硬件拓扑并根据用户需求完成建环，这就是通信库拓扑感知的作用。

通常情况下，物理网络拓扑会采用rail-optimized的方式，来确保所有机器的同号网卡都连接到同一个交换机。这是因为NCCL在构建拓扑时，侧重于本地环带宽的最大化，并不会考虑网卡在网络拓扑中的连接方式，如果不采用rail-optimized的话就会出现跨交换机传输的情况，形成性能瓶颈。

但是当参与训练的机器规模足够大时，单个交换机受端口数量的限制，无法再满足这一需求，此时就需要把网卡的物理拓扑信息传递给通信库，并优化网卡的选择来避免跨交换机传输的情况，我们计划在未来的工作中实现这一能力。

5. 如何强化拓扑解析

NCCL能输出通信域对应的Ring信息到日志文件，但由于各种问题无法满足工程化需要，主要为：

• 并发原因导致的日志割裂，无法有效解析；

• 仅包含通信域内rank连接信息，无法显示推演出物理GPU的拓扑连接；

• 不同NCCL版本可能存在日志差异，导致解析错误；

因此我们需要实现统一的信息输出方式，来支持平台化输出拓扑信息，通过在NCCL初始化后收集各个环对应的物理设备信息，统一格式并输出到新的日志，可以得到：

（图6）服务器节点0拓扑信息

（图7）服务器节点1拓扑信息

如图6、7，自上而下分别是通信域id，总的rank数量，本地rank数量，本地rank和GPU的映射，初始的通道数量及环信息，最终的通道数量及环信息。

可清晰知道通信域id为0xbcbd26d283a4efaf的通信域信息，此通信域是跨服务器的，总共包含12GPU,本服务器有4个，全局rank编号与本地GPU设备编号映射关系，由此就可清晰知道服务器节点内GPU互联信息和节点间网卡互联信息。

正如上一节介绍的，NCCL会根据全局环的探测结果来决定如何创建Channel，这里存在的隐患是，如果集群中某台服务器节点的初始环路数量低于其他节点（部分网卡down），那为了匹配这个短板，其他机器也不得不降低最终生成的Channel数量，这类瓶颈往往导致训练效率低下又无明显报错信息难以排查。

拓扑解析强化后输出的初始和最终通道数量，可以帮助我们快速定位到这类短板，如果InitChannelsNum * 2 < FinalChannelsNum，就意味着该机器节点由于其他短板节点，没有发挥出该有的效率。

将拓扑解析与性能数据获应用到监控平台，能可视化一些分布式训练的拓扑及性能信息，如下图显示了4个服务器节点的训练任务的拓扑信息：

（图8）

平台化展示这些信息，可以帮助我们直观的发现通信域物理拓扑连接的健康情况，及时发现硬件故障，快速定位问题机器。

6. 通信原语的实现

 6.1 什么是通信原语

通信原语是指一组基本的、不可分割的通信操作，它们是构建更复杂通信模式和算法的基础，常见的通信原语规则如下图：

（图9）

NCCL提供了众多集合通信原语api来进行集合间通信，包括Broadcast、Reduce、AllGather、ReduceScatter、AllReduce，还有点对点的send/recv；

• Broadcast（广播）：操作涉及一个发送者（root）GPU和多个接收者GPU，发送者GPU上的数据被复制并发送到所有接收者GPU；当需要将相同数据分发到多个GPU上进行计算时使用；

• Reduce（归约）：操作涉及多个发送者GPU和一个接收者GPU，所有发送者GPU上的数据被某种方式（如求和、求最大值等）合并，然后发送到接收者GPU；在多GPU上进行并行计算后，需要将结果汇总到一个GPU上进行进一步处理；

• AllGather（全部收集）：允许每个GPU发送数据，并且所有GPU都会收到所有其他GPU发送的数据副本；当每个GPU需要其他所有GPU的数据进行后续独立计算时使用；

• ReduceScatter（归约散布）：结合了Reduce和Scatter操作，多个发送者GPU上的数据首先被Reduce，然后Reduce的结果散布到每个GPU上；适用于需要将归约后的数据分发到各个GPU进行进一步操作的场景；

• AllReduce（全部归约）：Reduce和Broadcast的组合。所有GPU发送数据，数据被Reduce后，将结果广播给所有GPU；在并行计算中，当所有GPU都需要获得归约操作的最终结果时使用；

• Send/Recv（发送/接收）：点对点通信操作，一个GPU发送数据给另一个特定的GPU，数据传输是一对一的；用于GPU之间的直接数据交换；

6.2 通信原语的实现

以AllReduce为例，Ring AllReduce(SIMPLE协议)实现流程如下：

• 根据建立的拓扑，已确定数据在多个GPU间的传递顺序；

• 框架调用NCCL的ncclAllReduce api，代入数据源/目的地址、操作数量及操作类型等；

• launch kernel后GPU进入ncclKernel入口开始执行，最终进入到Fn参数为ncclFuncAllReduce的runRing()函数；

namespace {
  template<typename T, typename RedOp, typename Proto>
  __device__ __forceinline__ void runRing(ncclWorkElem *args) {
    ......
    int ringIx = ring->index;                                /* myrank在Ring中的索引 */
    const ssize_t loopSize = nChannels*nranks*chunkSize;    /* 每次循环处理的数据量 */
    const ssize_t size = args->count;                        /* 待处理的总数据量 */
    ......
    /* 创建Primitives对象来做实际的send/recv等操作 */
    Primitives<T, RedOp, FanSymmetric<1>, 1, Proto, 0> prims
      (tid, nthreads, &ring->prev, &ring->next, args->sendbuff, args->recvbuff, args->redOpArg);
      (tid, nthreads, &ring->prev, &ring->next, args->sendbuff, args->recvbuff, args->redOpArg);
  
    /* 一个大循环，多个Channel运行处理的总量为size，单次循环迭代处理总数据量为gridOffset，其中单个Channel处理数据量nranks*chunkSize */
    for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) {
      ssize_t realChunkSize;
      ......
      /* 根据当前线程块对应Channelid和chunk块号找到send/recv缓冲区偏移 */
      auto calcOffset = [&]__device__(int chunk)->ssize_t
            return gridOffset + bid*nranks*realChunkSize + chunk*realChunkSize;
   
     /* 单次迭代的数据被分为n块（n = nranks），每块数据大小为realChunkSize */
      ssize_t offset;
      int nelem;
      int chunk;
  
      /* 第一步：myrank发送自己第ringIx-1块数据给下一rank */
      chunk = modRanks(ringIx + nranks-1);
      offset = calcOffset(chunk);
      nelem = min(realChunkSize, size-offset);
      prims.send(offset, nelem);
   
      /* 第二步：myrank分别接收第ringIx-j块数据并与自身的ringIx-j块数据reduce后发给下一rank */
       for (int j=2; j<nranks; ++j) {
        chunk = modRanks(ringIx + nranks-j);
        offset = calcOffset(chunk);
        nelem = min(realChunkSize, size-offset);
        prims.recvReduceSend(offset, nelem);
      }
           
      /* 第三步：myrank接收第ringIx块数据并与自身第ringIx块数据reduce，到此这块数据已被所有rank reduce过，是最终结果数据，将这块结果数据存储并发送给下一rank */
      chunk = ringIx + 0;
      offset = calcOffset(chunk);
      nelem = min(realChunkSize, size-offset);
      prims.directRecvReduceCopySend(offset, offset, offset, nelem, /*postOp=*/true);
          
      /* 第四步：myrank依次接收第ringIx-j块结果数据，存储并发送给下一rank */
      for (int j=1; j<nranks-1; ++j) {
        chunk = modRanks(ringIx + nranks-j);
        offset = calcOffset(chunk);
        nelem = min(realChunkSize, size-offset);
        prims.directRecvCopySend(offset, offset, nelem);
      }
  
      /* 第五步：myrank接收最后的第ringIx+1块结果数据，到此myrank接收完n块结果数据 */
      chunk = modRanks(ringIx + 1);
      offset = calcOffset(chunk);
      nelem = min(realChunkSize, size-offset);
      prims.directRecv(offset, nelem);
    }
    /* myrank处理完nranks*chunkSize大小的数据，开始下一轮迭代 */
  }

这里的数据处理主要分为：每块数据的reduce，及结果数据的分发。

• 需处理的数据总量为args→count，即调用ncclAllReduce api时的入参count，单次处理 nChannels*nranks*chunkSize 的数据量，多次循环迭代直到处理完；

• 单次处理中，数据按Ring的rank数量（假设为n）被划分为各个chunk，若当前rank在Ring中索引为x，则它负责数据chunk x的最后一次reduce；

• 对应上述代码第一步，rank x发送chunk x-1给rank x+1(这里若x - 1 < 0，则取 x + n -1，后面皆如此)， 并在第二步中陷入等待接收rank x -1的数据，如图10：

（图10）NCCL Ring allreduce 流程展示

• 以此类推，当rank x - 1完成第一步后，rank x会收到chunk x - 2；

• 继续代码的第二步，并和自己的chunk x -2 进行reduce操作后发送给rank x + 1，继续陷入等待rank x -1的chunk x -3；

• 以此类推，直到rank x收到rank x - 1发送的chunk x + 1，完成reduce操作后发送给rank x + 1，进入第三步；

• 第三步中，rank x收到来自rank x -1的chunk x，此时这块数据chunk x已经在rank x + 1 到rank x -1上都完成了reduce，最后和自己的chunk x进行reduce后则完成全部的reduce；

• rank x把完成的chunk x拷贝到目的地址后，发送给rank x + 1，随后进入第四步，如图11：

（图11）NCCL Ring allreduce 流程展示

第四步中，rank x等待rank x -1发送完成reduce后的chunk x-1，拷贝到目的地址后，发送给rank x + 1；

以此类推，直到rank x收到来自rank x -1的chunk x +2，拷贝到目的地址后，发送给rank x + 1，随后进入第五步；

第五步，rank x收到rank x-1发来的chunk x + 1，至此rank x的chunk 0 - n-1 全部完成了reduce操作，如图12：

（图12）NCCL Ring allreduce 流程展示

• 以此类推当所有rank完成第五步后，所有rank的chunk 0 - n-1 全部完成了reduce操作并将结果数据存到了相应目的地址；

7. 如何获取带宽数据

结合allreduce代码流程，获取各个通信域内流量数据的方法可参考图13：

（图13）

• 在集合通信api， 如ncclAllReduce处埋点，获取对应通信原语下发数据量；

• launch kernel后，GPU执行完对应集合通信的数据处理和传输后，再埋点获取通信完成量；

• 参照图13的实现方式，可以在网络send/recv回调结束时埋点，获取跨机传输流量；

一含2GPU卡的通信域内，各种集合原语的通信完成数据量获取示例如下：

（图14）

（图15）

完成量终端图14中，自上而下分别是对应rank编号不同原语类型的传输字节数和原语调用次数。

通过周期性采集变化值，就可以得到对应原语在各个rank上的下发和完成带宽，平台化后就可以直观地发现机器之间是否存在带宽明显低于其他训练的现象。同时通过分析下发和完成带宽的匹配度，例如下发带宽经常大幅高于完成带宽的情况，就可能是存在网络拥塞之类的问题，导致完成效率产生波动。

后续计划将各种通信原语进一步细化，按框架并行设计来将各通信域和并行类型进行映射，获取不同并行的各种通信原语数据，有助于优化大规模分布式系统中的并行设计和通信模式，从而提升整体的计算效率和性能。

训练过程中每个通信域的各种集合通信完成量的平台显示示例图：

（图16）

通过平台化展示这些信息，可以帮助我们直观的发现数据传输的健康情况，及时发现瓶颈的存在，快速定位问题训练任务。

8. 总结与展望

目前，各大厂商虽然都开发了自己的性能诊断和监控工具，但各种工具或方法缺乏统一标准，难以按一定标准界定通信库的使用是否达到了最优化。在这个背景下，B站也在不断探索和尝试，当前我们主要强化了拓扑信息的解析能力并引入了带宽数据统计功能来改进通信库的应用；未来，我们计划针对通信库在通信域和框架并行设计间映射的不足、以及真实网络拓扑信息的缺失等问题，进行进一步的优化；在性能统计方面，我们也将通过更精细地获取性能数据，对集合通信原语的计算过程进行深入分析，分别统计其在单个机器内部以及跨机器间的性能指标，帮助我们识别和解决潜在的性能瓶颈，从而进一步提升通信效率。

随着训练模型的规模日益庞大和计算需求的持续增长，集合通信的性能和效率变得日益关键，在这样的技术发展浪潮中，国内一些厂商也积极投身于高性能通信库的研发，以满足日益增长的计算需求。例如，基于国产AI处理器平台的HCCL就是国产通信库的一个代表。虽然HCCL在底层的通信原理上和NCCL有相似之处，但在具体的实现方式、优化策略以及对特定硬件平台的适配上，两者之间存在着较大的差异。为了更好地掌握通信库的工作状态，并进一步优化性能，如何去实现这些国产通信库的性能监控也在后续计划中。

9.参考

大语言模型--Tensor Parallel（张量并行）（https://cloud.tencent.com/developer/article/2378690）

使用 NVIDIA Collective Communication Library 2.12 将所有 all2all 性能翻倍（https://developer.nvidia.com/zh-cn/blog/doubling-all2all-performance-with-nvidia-collective-communication-library-2-12/）

分布式训练技术总结（https://orion-wyc.github.io/2020/09/19/2020-09-19-%E5%88%86%E5%B8%83%E5%BC%8F%E8%AE%AD%E7%BB%83%E6%8A%80%E6%9C%AF%E6%80%BB%E7%BB%93/）

-End-

作者丨SYS团队