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为例，我们总结目前所面临的挑战和问题包括：通信拓扑中实际物理GPU和网卡互联信息的缺失；各个通信原语所处理完成数据量的缺失；机内和机间传输数据量的缺失；通信域和框架并行设计间映射的缺失；真实网络拓扑信息的缺失；本文将主要介绍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 ?__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, 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 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