近年来, 机器学习取得了巨大的突破, 在图像处理方面, 微软研究院和谷歌的科学家各自取得了错误率接近3.5%的识别准确度, 远超人类平均水平(5.1%); 在自然语言处理方面, 各大互联网巨头纷纷推出自己的新型机器翻译模型, 微软公司于2018年年初, 宣称在中英新闻翻译领域达到人类水平; 在医疗诊断领域, 由谷歌公司训练的人工智能模型, 对皮肤癌的识别率达到了专业医师水平; 在自动驾驶等其他领域, 同样也得到广泛的应用和突破^[1]. 而在众多机器学习技术中, 以深度神经网络DNNs (Deep Neutral Networks)的发展与应用最为迅速和广泛, 深度神经网络已经成为许多应用领域的核心技术之一. 以深度神经网络为基础的技术占据了如图像识别、自然语言处理、人脸检测, 以及语音处理等诸多机器学习领域. 伴随着NMT (Neural Machine Translation)^[2]等方法的兴起, 面对海量数据, 集中式的机器学习框架在训练时长、处理时效上已经远远不能满足现实的要求^[3].

深度神经网络训练是一个计算密集型问题, 随着深度神经网络中的训练数据及训练规模变得越发庞大, 传统的单机训练模式已不太现实, 过长的训练时间成为困扰众多深度神经网络研究者们的一大问题, 许多该领域的研究提出了充分利用高性能计算HPC (High Performance Computing)丰富计算资源来解决深度神经网络训练所遇到的问题^[4,5]. 为了加速训练过程, 像Caff^[6], Tensorflow^[7], 以及CNTK^[8]这样的深度学习框架, 已经引入了多GPU、多计算节点下的并行或分布式训练功能, 这一系列的变化催生了分布式机器学习的兴起, 分布式机器学习研究的正是如何利用计算机集群训练大规模机器学习模型. 机器学习的学习模型训练, 通常使用迭代式的优化算法, 这使得训练过程中的通信频率很高, 其次分布式机器学习往往处理大数据, 导致通信数据量很大, 因此提高分布式机器学习训练效率的关键是设计高效的通信机制, 对于诸如MPI这样的集合通信库, 需要进行重新的研究和设计, 以求实现低时延、高带宽的通信, 继而缩短训练时间. 对于分布式深度神经网络训练, 要实现高效率的信息通信, 关键是解决每一次训练迭代中, DNN workers之间参数和梯度的交换所带来的通信负载问题. 有几种深度学习框架通过使用MPI通信原理实现分布式训练来解决这个问题, 例如CA-CNTK^[9]使用CUDA-Aware MPI_Bcast进行模型参数的广播通信, 类似于梯度交换中的all-to-all广播形式. 尽管像NCCL^[10]这种专为多GPU集合通信设计的通信库能够解决此类问题, 但却引入了新的问题: NCCL并不支持MPI, 因此需要使用新的API来重新设计应用程序. 为了解决上述问题, 只能重新设计一种更加高效的MPI通信机制, 从而简化深度学习应用程序在多GPU集群中应用时, 代码的修改量, 同时, 充分利用MPI通信机制, 进行分布式深度神经网络训练, 从而缩短训练时间.

遵循这一研究思路, 调查研究了现今应用较多的深度学习框架, 以及数据并行训练的高效广播机制, 最后提出新型广播机制.

1 相关工作

在相关领域论文中, Liu等人提出了一种高效的广播设计^[11], 利用了像硬件多播这样的IB (InfiniBand)特性, 受这一研究工作的启发, 提出了利用新的设计机制, 重新设计和优化MPI_Bcast来应用于新的框架结构, 另一篇论文中, Kandalla等人提出了在 intel MIC (Many Integrated Cores)中, 优化 MPI broardcast 以及 reductions^[12].

此外, Awan等人提出了一种基于IB多播的广播设计^[13], 用于流媒体和DL (Deep Learning)应用. Zhou等人提出了一种对于大消息的广播优化^[14]. Barrett等人提出了利用共享存储和MPI的单边通信特性来优化节点内通信^[15]. Awan等人提出集成NCCL的MPI_Bcast设计^[16], 用于深度学习负载. 集成NCCL MPI_Bcast和提出的新设计的主要区别在于, 后者使得节点间的通信, 不再需要依赖NCCL以及其他外部库, 而可以使用MVAPICH2^[17]运行时库, 实现MPI_Bcast的通信性能达到甚至超过NCCL的水平.

Uber公司提出了一种应用于Tensorflow框架上的开源集合通信库Horovod^[18], 该集合通信库使用环状reduction操作来提高分布式机器学习上多GPU间的通信效率. Anderson等人提出了一种在SPARK上整合MPI集合通信库来提升性能的方法^[19]. Sony公司使用一种2D-Torus all-reduce集合通信来进行分布式DNN训练中梯度的同步^[20], 该集合通信方法已包含在深度学习库NNL (Neural Network Libraries)^[21]中.

2 基于GPU的集合通信

这一部分, 主要讨论基于GPU的集合通信的3个方面: (1) GPU集群中硬件平台配置的多样性. (2) 对于深度学习应用程序的新要求, CUDA-Aware MPI应该如何进行改进. (3) 为特定要求设计的诸如NVIDIA NCCL^[22]的集合通信库如何解决深度学习工作负载问题, 以及这种基于NCCL的MPI设计的限制.

2.1 GPU集群的多样性

HPC集群装载有搭载高速CPU核的高性能计算节点, 配置如无限带宽技术IB的高速传输网络, 但随着科学计算应用对于GPU的需求不断增加, 一些研究中心的HPC集群已经加入了GPU (一个计算节点配置1到2个GPU). 例如美国Cray 公司的超级计算机CS-Storm^[23], 每个计算节点配置16个NVIDIA K-80 GPU, GPU之间采用树状的PCIe拓扑结构连接. 这种系统最初专为计算密集型的应用设计, 包括HPC及DL (Deep Learning)应用. 多GPU系统能够为深度神经网络训练提供绝佳的性能加速^[24,25].

2.2 CUDA-Aware MPI

如前所述, 现代HPC系统^[26]像Cray CS-Storm超级计算机、SUMMIT^[27]等系统都使用 GPU进行计算加速, 提出MPI运行时库支持GPU间的高效通信, 便顺理成章. 在发展初期, 由于缺乏GPU内存直接读取技术, MPI应用程序需要将GPU中的数据复制到主机内存的暂存缓冲区中, 然后再将数据通过网络传输. 在主机通过MPI_Recv 操作接收到数据之后, 会以与上述相似的过程将数据从CPU传输到GPU中, 这种传输方式会极大地影响GPU与CPU之间的通信效率与产出. 而像OpenMPI^[28], MVAPICH2-GDR^[17]这样的MPI库能够提供CUDA-Aware MPI函数, 实现复制数据的透明操作, 从而可以显著提升应用程序的性能与效率. CUDA-Aware MPI运行时库对基于GPU的端对端通信机制作了许多优化, 引入了包括staging, 管道技术, CUDA IPC (Inter-Process Communication)以及RDMA (GPUDirect Remote Direct Memory Access)等技术, 为诸如节点间、节点内、socket间场景下, 提供最佳性能的通信. 对于集合通信, 同样可以利用MPI_Bcast, 来达到基于GPU缓存的直接通信, 实现一种高效的设计.

2.3 基于GPU的集合通信NCCL

NCCL是一种基于GPU的集合通信库, 专门用于深度学习工作负载, NCCL的API十分类似于MPI接口, 提供broadcast, all-gather, reduce, reduce-scatter, all-reduce通信函数, 但NCCL API并不支持MPI, 因此应用程序需要进行大的修改才能使用MPI, 例如NCCL中的广播操作使用ncclBcast(), 而不是MPI中标准的MPI_Bcast(), ncclBcast函数形式如下:

ncclResult _ t

ncclBcast (void*buff, size_t count,

ncclDataType_t datatype,

int root, ncclComm_t comm,

cudaStream_t stream)

关于NCCL API的更多细节可参见文献[29], 准确地说, NCCL的主要目标是为集群中的GPU提供快速的信息通信, 从而显著提高深度学习负载的训练效率, 它是一种优化GPU集合通信的库, 目前已更新到NCCL2版本, NCCL1版本只能用于单节点上的多GPU配置环境, 并且内部的GPU连接使用PCIe或者NVLink, 提供数据交换的wrap级粒度优化, 利用端访问和CUDA (Compute Unified Device Architecture)内核进行数据的复制, 而非标准的CUDA内存复制操作方式. 尽管与MPI十分相似, 但两者的实际目标以及应用平台都迥然相异. MPI是为了实现集群中众多节点间的高效通信而设计, NCCL则为高密度多GPU系统而设计. NCCL库的第一个开源版本为NCCL1, 相关细节及源代码安装包均可在GitHub网站^[30]上得到, NCCL2的相关包目前也已开源, 在本文中, 使用NCCL1库, 以MPI+NCCL1^[16]的设计实现节点内部通信, 使用MPI函数实现节点之间的通信. 从应用程序的角度来说, MPI+NCCL是一种轻量级的方法, 它能提供标准的MPI函数调用(如MPI_Bcast()), 同时在底层, 也能以十分透明的方式利用NCCL, 这种调用方式与NCCL2的ncclBcast()调用完全不同, 后者往往需要对应用程序的代码进行彻底的修改.

2.4 NCCL集成MPI的局限

像MVAPICH2-GDR这样的CUDA-Aware MPI 运行时库十分灵活, 可以集成像NCCL这样的第三方库. 在剩下的章节中, 会对一种基于NCCL的MPI_Bcast^[16]进行评测. 图1是这一方法的大体框架, 图中方形虚线框中的数字圆圈代表节点, 圆形虚线框中的数字圆圈代表节点4中的4个GPU设备, 方形虚线框代表节点1、2、3、4间的通信, 圆形虚线框代表节点4中四个GPU间的通信, 利用MVAPICH2集合通信的分层特性, 节点内部通信只使用NCCL, 两节点之间的通信则使用CUDA-Aware 的端对端通信方式.

MPI+NCCL的通信方式能提供良好的通信性能^[16], 但同时也存在着许多限制, 如流的创建及管理, NCCL communicator的创建及管理, 为获得最优性能可能会创建多个NCCL communicator, 系统中没有到GPU的端对端访问. 通常情况下, NCCL集成MPI 运行时库的设计会使得系统极其复杂. 所提出的方法能够使MPI_Bcast在不依赖NCCL的情况下, 得到相当或者更优的性能, 这一方法的目标, 是研究出最佳的算法及技术设计, 完全利用HPC系统中可用的多样化资源. 因此, 最大的挑战是在不依赖其它库的前提下, 得到处理深度学习工作负载, 性能优异的MPI_Bcast广播设计.

3 现存广播算法性能模型

数十年来, 学术界对于广播算法的研究已十分成熟^[14,31,32], 但类似于GPU这样的加速器的引入, 已经完全改变了该领域的研究, 需要对现存的广播算法重新进行评估, 同时也要探索新的技术设计, 利用GPU特性的硬件和机制来实现这些广播算法的更加高效的应用. 在这一部分, 对几种经典的广播算法进行了分析, 并研究了其性能特点. 在表1中列出了将会用到的一些专业符号.

3.1 基本广播算法

首先分析该领域最常见的广播算法, 并将其应用到基于GPU条件下的广播. 但是请注意, 像那种需要硬件辅助的特殊广播机制^[13,33–35]不在本论文的讨论范围之中.

直接广播算法: 从root进程中直接广播数据到所有其他进程中, 该算法简单利用一个连续回环进行端对端通信调用. 在MPI中, 该算法本质上利用MPI_Send和MPI_Recv的回环调用. 算法时间复杂度为:

${T_{\left( { {\rm {B}}{\rm{cast}}\_{\rm {Direct}}} \right)}} = \left( {n - 1} \right) \times \left( {{t_s} + \frac{M}{B}} \right)$

(1)

由于该算法糟糕的可扩展性, 以及时间复杂度对节点数n的相关性, 在实际应用中, 该算法很少采用.

链/环算法: 与直接广播算法中只有一个root进程发送消息不同, 该算法中, 每一个成功接收到数据的节点, 在环的下一个发送阶段都能充当发送方. 对于像MPI_Bcast这种有根节点的集合通信, 环是由通信进程组成的逻辑链, 并且首节点和尾结点无环绕. 对于非根节点式集合通信, 链的首尾节点是环绕的, 组成一个完整的环. 式(2)描述了该算法的时间复杂度:

${T_{\left( {\rm {Bcast\_Chain}} \right)}} = \left( {n - 1} \right) \times \left( {{t_s} + \frac{M}{B}} \right)$

(2)

K阶/二进制树广播算法: 随着HPC集群规模越来越大, 系统相对于节点数的可扩展性问题也越发重要, 基于二叉树的广播算法已经被提出多年, 并且表现出良好的可扩展性. 该算法将通信进程视作逻辑树状结构, 树的根就是广播操作的根节点, 每一个通信阶段, 根节点将数据传输给它的一个没有接收到该数据的子节点, 在下一阶段, 接收到数据的子节点又将该数据传输给它的没有接收该数据的子节点, 重复这一过程, 直到所有叶节点接收到数据为止. 在二叉树算法中, 一个根节点最多拥有 $\left\lceil {{{\log }_k}n} \right\rceil $ 个孩子, 也即通信的最大次数 , 当K=2时, 即为二叉树算法. 该算法通信开销为:

${T_{\left( {{\rm B}{\rm{cast}}\_{\rm{Knomial}}} \right)}} = \left\lceil {{{\log }_k}n} \right\rceil \times \left( {{t_s} + \frac{M}{B}} \right)$

(3)

式(3)的时间复杂度为 ${\rm{O}}\left( {{{\log }_k}n} \right)$ , 远小于前两个算法的 ${\rm O}\left( n \right)$ . 基于树结构的算法提升了广播操作的可扩展性, 该算法在MPI进行时中广泛使用, 许多集体式操作也都基于该算法.

Scatter-Allgather算法: 对于不了解该算法的人来说, Scatter-Allgather算法有些不太直观, 但仍然掩盖不了它的强大, 该算法可扩展性与消息大小M相关, 而与节点数n无关. Scatter-Allgather机制^[36,37]可以提升大数据量广播操作的性能. 其原理是使用Scatter操作后, 再使用Allgather操作, 从而优化广播带宽, 通常情况下, 会先使用基于二叉树的Scatter操作, 然后进行基于环的Allgather操作来完成广播. 该算法通信开销见式(4), 具体计算过程见^[31]:

$\begin{aligned}[b] & {T_{\left( {\rm {Bcast\_Scatter\_Ring\_Allgather}} \right)}}\\ & = {\log _2}n \times {t_s} + \dfrac{{\left( {n - 1} \right) \times M}}{n} \times \dfrac{1}{B} + \left( {n - 1} \right)\left( {{t_s} + \dfrac{M}{B} \times \dfrac{1}{B}} \right)\\ & = \left( {{{\log }_2}n + n - 1} \right) \times {t_s} + 2 \times \dfrac{{\left( {n - 1} \right) \times M}}{n} \times \dfrac{1}{B} \end{aligned}$

(4)

然而该算法并没有被传统HPC应用广泛使用, 此外CUDA-Aware版本的Scatter-Allgather算法也并不十分常见.

4 CUDA-Aware MPI_Bcast的新型设计

在第3节中讨论了基于现存广播算法的前沿方案与技术, 接下来, 会提出两种新的设计方案: (1) 管道链PC (Pipelined Chain)设计. (2) 拓扑感知管道链TA-PC (Topology-Aware Pipelined Chain)设计. 并详细阐述相关细节. 最后, 将综述MVAPICH2-GDR中的各种设计方案和算法, 从而使新的方案对于所有的消息大小达到最佳的性能.

4.1 大型消息广播的管道方案

从2.1节中讨论几个广播算法的时间开销来看, 除了Scatter-Allgather算法, 一般会以每一个通信步骤为研究粒度, 来考虑整个消息的传输. 然而, 随着近年来互联带宽的提高, 这些算法能够利用分片或管道技术来传输信息. 为了更好利用网络资源和带宽, 管道技术需要进一步的深入探索, 使用MPI_Isend, MPI_Irecv这样的非阻塞式端对端通信, 实现允许通信传输重叠的管道广播, 来达到更佳的带宽资源利用.

传统意义上, 从MPI应用角度来说, 链/环算法被认为是一种低效的算法, 但是随着深度学习应用的出现, 在相对较少的节点或GPU中, 超大信息量通信正在成为MPI运行时库的新应用场景. 因此, 围绕广播算法的传统研究范畴需要重新审视, 接下来的将详细讨论新提出的两种设计方案, 两种方案都引入了管道的概念来充分利用可用带宽.

4.2 管道链设计

提出在MVAPICH2-GDR中的CUDA-Aware 管道链设计实现MPI_Bcast. 上述设计的详细过程见图2, 根进程会将数据分片, 并将这些数据分片轮流发送给逻辑进程链中的右邻居节点进程.

除了链的尾进程, 所有非根进程只接收从左邻居节点发送过来的数据分片. 基于第3节中对链状算法的研究, 该管道链模型的时间开销如下:

$T _{\left({{\rm B}{\rm{cast}}\_{\rm {Chain\_Pipeline}}} \right)} = \left( {\frac{M}{C} + \left( {n - 2} \right)} \right) \times \left( {t_s + \frac{C}{B}} \right)$

(5)

管道链方案理论上实现了更低的通信开销, 从式(5)中可以看出, 对于不同的数据大小、系统结构, 分片大小的选择是十分重要的, 会直接影响通信的性能. 通过对MVAPICH2-GDR运行时库的底层架构进行整体微调, 通过实验证明在一定的消息大小、进程数量、硬件架构中, 能使性能达到最佳的分片大小. 管道链方案设计的主要目标是减少通信开销, 希望整个缓存区的广播开销能够减少至单个分片的广播开销, 加上一些无法用管道技术消除的额外分片的广播开销, 通过式(5)可以发现, 在理论上这是有可能实现的. 但是在实际中, 很难设计出这样的机制, 特别是基于GPU集群的通信. 主要的挑战在于提出一种多样化的设计, 能够处理多类型的CPU、GPU结构, 连接速度(如FDR、EDR等)以及节点拓扑结构. 接下来, 进一步研究了系统节点内多GPU的其他优化可能, 这些会在后续的章节中讨论.

4.3 拓扑感知管道链设计

管道链算法在单GPU系统上性能最佳. 但是, 如果节点内有多个GPU时, 节点的拓扑结构, 以及如何利用PCIe/PLX链路连接GPU到CPU和IB HCA, 都会影响通信的性能, 这将带来新的设计挑战. 比如说, NVIDIA K-80 GPU中含有两个GPU设备, 这两个GPU之间使用PLX互连, 具体拓扑结构见图3, 带有数字的圆圈代表一个节点中的多个GPU. 双向箭头代表闲置的PCIe连接, 用于GPU内存与CPU内存或者两个GPU内存间的数据复制传输, 所谓闲置, 就是该链路可用但无数据传输, 这是由于管道链算法只能在特定的时间点利用某些链路. 指向右侧且带单向箭头的黑色曲线代表从GPU1到GPU2数据的单向移动, 显然, 这对于图3中的多GPU场景来说, 并不是最优方案.

为了解决这一问题, 提出了一种拓扑感知的管道链(TA-PC)设计, 为多GPU节点提供简单却高效的优化设计. 从图3和图4可以看出, 图3只能向右侧传输的PC设计, 变成图4的只向左侧传输的TA-PC设计, 图4对于PCIe链路的利用更加高效. 在图4中, 双向虚线箭头代表可利用链路, 对比于图3中的空闲链路. GPU1正在发送数据分片到它的逻辑链左邻居GPU4, 而非右邻居GPU2. 初始阶段看的话, 这种数据移动方式似乎不合常理(第一次就将数据传输给最远的GPU), 但是由于PCIe链路的双向传输特性, 这种新的数据交换顺序在实现基于GPU的广播时, 表现出优异的性能. 显然, 现在能够实现所有方向的PCIe链路的无冲突利用. 由于数据的传输移动会用到cudaMemcpy(s), 最近发布的NVIDIA驱动能够确保对PCIe链路的完全利用, 此外, 该TA-PC设计不仅能应用在如图4展示的含4个GPU的集群, 也能应用在像Cray CS-Storm这样的集群系统, 该系统每个节点含有16个K-80 GPU. 图4中, 两个K-80与单个CPU之间通过内部的PCIe连接, 组成一个四GPU模块, 对于Cray CS-Storm, 有4个这样的模块, 彼此之间通过PCIe与两个CPU socket之间建立连接. 显然, 如果直接采用PC的设计, 这种拓扑结构中还有许多可用的双向PCIe链路被闲置. 因此, 为了充分利用所有链路, 所有的多GPU节点集群, 包括Cray CS-Storm, 都采用TA-PC设计.

4.4 针对GPU的CUDA-Aware MPI_Bcast优化

如前文所提到的传统广播机制, 基于GPU的广播算法性能不仅与算法本身有关, 也与GPU之间、集群节点间的数据传输移动方式有关.

STG-COLL (Host-staging Scheme): 在文献[38]中提到了一种基于GPU的端对端通信的GPU内存直接读取GPUDirect RDMA技术. 事实上, 对于特定的消息大小, GPU缓存的直接广播会导致性能的下降. 使用STG-COLL技术则能避免这些性能瓶颈. 图5中展示了直接和暂存的设计, 即GDR-COLL与STG-COLL. 本质上说, 根进程首先将GPU内存中的数据移动到主机内存中, 然后再通过主机广播给接收方, 数据就可以通过NVIDIA GDR直接写入GPU内存, 也可以暂存在主机中.

因此, 主机暂存的K阶树设计通信开销为:

$\begin{aligned}[b] & T _{\left({\rm {Bcast\_Knomial\_Staging} }\right) }= \dfrac{M}{{BPCIe}}+ \left\lceil {\log _k n} \right\rceil \times \left( {t_s + \dfrac{M}{B}} \right) \end{aligned}$

(6)

显然, 方程第一项的值为数据复制到主机所用的时间, 只有在第一项的值较小时, 该设计才会表现出更好的性能. 对于超大信息量的通信, 数据暂存的成本会很高. 因此, 对于管道链设计, 当与GPU建立对等访问时, 节点内部的通信不采用主机暂存, 也不使用像CUDA IPC那样的直接机制, 节点之间的数据传输也不使用CUDA GDR, 这样才是合理的.

4.5 MVAPICH2-GDR中的MPI_Bcast设计

像MVAPICH2-GDR这样的MPI运行时库^[17], 设计MPI_Bcast的高效算法优化MVAPICH2-GDR中的端对端通信操作, 图6展示了所提出方案在MVAPICH2-GDR中的MPI_Bcast层次结构图, 完整的MPI_Bcast过程在图中进行了详细的展示, 首先是communi-cator的选择(节点间通信还是节点内通信还是两者兼顾的灵活通信方式), 然后是集合通信传输机制的选择(GDR-COLL或者STG-COLL), 广播算法的选择, P2P的选择(MPI_lsend发送或MPI_lrecv接收数据), 最后是P2P communicator的选择(节点内或节点间的通信).

PC与TA-PC设计都依赖于MPI的Send/Recv操作, 在Send/Recv操作中, 通过许多优化设计, 解决了存储在GPU中的数据的传输问题, 举例来说, 节点间数据传输的管道链以及host-staging机制和节点内部高效数据传输的Loopback/GDR复制机制^[39], 即使使用标准的MPI库进行Send/Recv操作, 对算法(K阶树, Scatter-Allgather, PC等), 以及传输机制(STG-COLL或者GDR-COLL)的选择都是十分重要的. 表2展示了MVAPICH2-GDR中对于信息大小、算法、以及机制的一般选择标准. 但是想要确定确定的消息大小, 使用哪种特定的算法和机制, 则需要进一步的大量微调. 集群配置的差异取决于以下几点: (1) IB HCA模型. (2) CPU架构. (3) 每个节点的进程配置. 基于这些度量参数, 在MVAPICH2代码库中生成并使用了诸如表2的选择标准. 可以直接从MVAPICH2源代码中看到更多细节^[17].

5 实验分析

主要进行下列3种设计方法的性能比较:

(1) 集成NCCL1的 MVAPICH2-GDR (NCCL-MV2-GDR) 或者也称作MPI+NCCL1.

(2) MVAPICH2-GDR中的MPI_Bcast (MV2-GDR-Opt)设计.

(3) NCCL多节点可用版本(NCCL2).

5.2节展示了方法1和方法2的实验比较结果, 5.3节展示了方法2和方法3的实验比较结果. 此外, 还进行了两次实验来验证提出的两种新型设计: (1) PC (Pipeline Chained)广播. (2) TA-PC (Topology-Aware Pipeline Chained)广播.

使用CNTK (Microsoft CognitiveToolkit)进行方法1和方法2应用层的性能对比. 对于方法3, 不太可能使用CNTK进行测评, 因为需要使用ncclBcast API而不是标准的MPI_Bcast API, 所以需要应用层级的代码修改.

5.1 实验平台环境

使用3种不同的集群, 彼此之间存在很大差异, 具体表现在节点拓扑结构, IB HCA型号以及GPU型号的不同, 这样做可以保证集群的多样性, 验证提出的广播设计能适用各种类型的集群, 提高实验结果的可信度.

(1) GPU密集型集群RX1, 由8个节点组成. 每个节点含4个NVIDIA K-80 GPU, 一个K-80含2个GPU设备, 每个节点有8个可用GPU设备, 配置有一个14核dual-socket Intel Xeon E5-2680 (Broadwell) CPU. 每个节点有两个InfiniBand FDR HCA.

(2) 节点密集型集群RX2, 含有16个计算节点, 每个节点配置有单个NVIDIA 帕斯卡P100 GPU, 每个P100含一个GPU设备, 每个GPU节点配置有一个14核dual-socket Intel Xeon E5-2680 (Broadwell) CPU, 以及单个InfiniBand EDR HCA.

(3) 均衡型集群RX3, 每个节点配置单个K-80GPU, 14核dual-socket Intel Xeon E5-2680 (Broadwell) CPU, 同时RX3中也配置了单个的EDR HCA.

5.2 MPI+NCLL1与提出的新型MPI_Bcast

实验中, 将集成NCCL1的MVAPICH2-GDR与提出的MPI_Bcast设计进行了比较, 5.2.1部分提供节点内通信的实验结果, 5.2.2则是节点间通信的实验结果, 本次实验采用OSU Micro-benchmarks^[40].

5.2.1 节点内通信

使用集群中的单个节点, 分别在2、4、8以及16个GPU下比较NCCL和MVAPICH2-GDR的通信性能. 图7展示了实验结果, 在GPU数分别为2和8时, 对应信息大小由1 Byte逐渐增大到8 KBytes, MVAPICH2-GDR的MPI_Bcast广播时延相比于NCCL广播分别对应下降了14倍和9.4倍. 这是由于MVAPICH2-GDR中引入了更为先进的端对端通信设计, 能够以更加灵活的方式应对各种瓶颈问题^[38]. 此外, MVAPICH2-GDR-opt使用CUDA IPC管道设计应对大数据量信息通信, 这种设计能更好的利用可用带宽. 相反, 在NCCL中却没有这些优化机制, 因而在小型、中型消息通信中, NCCL通信性能相比于MVAPICH2-GDR发生了下降, 但在大型与超大型消息通信, NCCL展现了良好的可扩展性能, 而MVAPICH2-GDR也达到了与NCCL相当的性能水平, 这使得不一定要依赖NCCL来提升广播性能^[16], 还可以有其他的选择, 比如说MVAPICH2-GDR.

5.2.2 节点间通信

前面说到, NCCL 1.x版本只适用于单节点上, 所以对于节点间通信, NCCL无法直接实验. 但在Awan等人提出的NCCL广播设计^[16], 解决了这个问题, 使得能够进行节点间的通信性能对比.

我们比较了MVAPICH2-GDR(标记为MV2-GDR-Opt)和集成NCCL的MPI_Bcast^[16](标记为NCCL-MV2-GDR), MVAPICH2-GDR中基于SGL的设计^[39], 利用了IB功能中的Scatter-Gather列表, 能够提升节点间小消息的通信性能. 结果见图8, 对于小、中消息通信, MV2-GDR-Opt在32 GPU集群中, 实现了16.4倍的性能提升, 在64 GPU集群中实现了16.6倍的性能提升. 显然, 对于不同类型的工作负载, 包括深度学习工作负载, MPI_Bcast都能够提供优异的通信性能.

5.3 NCCL2与提出的新型MPI_Bcast

随着NCCL2的开源可用, 实验也能将ncclBroadcast和新提出的MPI_Bcast设计做一个直接的对比. 在5.2的对比中, MPI+NCCL的设计与提出的PC/TA-PC设计仍然使用MPI_Bcast, 因此实验选择标准化的OSU Micro-benchmarks^[40]套件进行实验, 但是, NCCL2中能够直接评估ncclBroadcast API的性能, 即使是在多节点情况下. 使用NCCL2 测试标准^[41]来评估多节点下ncclBroadcast性能.

不太理想的是, 当运行NCCL2测试在每个节点16个GPU的情况下时, 测试程序一直运行在低水平的同步调用状态. 但在每个节点8个GPU的情况下, 程序运行正常, 因此这次只进行了每个节点8个 GPU情况下的测试, 对比结果见图9. 可以看到, 图形趋势与图8中的曲线稍有不同.

在RX2进行实验, 看到了类似于RX1集群中的曲线趋势. 图10中, 在小、中消息通信方面, MVAPICH2-GDR-Opt相对于NCCL2实现了10倍左右的时延性能提升; 从这些结果中可以看出, 提出的MPI_Bcast优化设计具有良好的通用性能, 能够提升深度学习领域以及传统HPC领域应用的通信性能. 此外, 实验结果也证明, 当设计合理时, 对于大多数的信息大小, MPI_Bcast操作也能达到超过ncclBcast的水平.

5.4 PC vs. TA-PC

这一部分会进一步研究提出的PC设计以及TA-PC设计, 并理解TA-PC设计相对于PC设计的优势. 为了更好凸显TA-PC设计, 节点内至少要有2个GPU. 因此, 实验在RX3和RX1集群(2-8GPU/node)上进行.

TA-PC的实验结果见图11 (a), 在集群RX3上, 相比于PC设计实现了27%的性能提升, 同样, 在RX1集群上8个节点的64个GPU上也进行了实验, 实验结果见图11 (b), TA-PC设计也实现了50%的性能提升.

5.5 应用程序性能测试

使用CA-CNTK^[9]的训练框架实验, CA-CNTK使用CUDA-Aware MPI_Bcast用于训练进程中参数的交换, 图12(a)中, 展示了在ImageNet-1K^[42]数据库中, NCCL-integrated MVAPICH2 (图中标记为NCCL-MV2-GDR)和提出的优化版本MVAPICH2-GDR (图中标记为MV2-GDR-Opt), 使用VGG^[43]模型后训练时间的比较结果, 从图中可以看到. MV2-GDR-Opt在32GPU集群中, 训练时间的缩短比例达到7%, 在其他各种情形下, 也跟NCCL-MV2-GDR不相上下, 甚至表现更好. 这是由于训练中VGG模型的消息大小混合, 其中也有数据量大的消息, 这表明, 对于应用程序负载中的超大消息, 优化后的MPI运行时库也能够保证良好的性能表现.

5.5.1 VGG训练

为了更好地分析优化设计的性能优势, 我们特意展示了两张数据图, 对VGG训练进行更加详细的介绍, 见图12(b). 信息的大小往往取决于进程数, 举例来说, 图中的趋势说明, 随着GPU数量的增加(4–16 GPU), 大消息所占的比例越来越少, 这是由于CNTK会根据参与DNN训练的进程数量, 来对张量进行域分解, 而像Caffe这样的深度学习框架, 张量在进程交换过程中, 大小和模式保持不变. 因此, 对于给定的像MVAPICH2-GDR这样的通信运行时库, 为了实现可伸缩的DNN训练性能, 需要优化所有消息的大小.

6 总结与展望

在一个以深度学习和深度神经网络为驱动的新AI系统时代下, 人们围绕通信运行时库进行着不断地探索, 然而传统的集合通信库如NCCL也存在着不支持MPI等诸多的问题, 这也促使了像MVAPICH2和其他类型为特殊需求而设计的通信库的兴起. 这也使得研究者们能够探索支持深度学习负载的新型MPI运行时库设计. 在本文中, 我们深入研究了MPI_Bcast, 这对于并行深度神经网络训练来说是十分重要的, 使用micro-benchmark和深度学习框架进行实验, 对MPI_Bcast进行有效的设计与调整, 最高可以取得16.6倍的加速性能, 在64 GPU集群上使用Microsoft CNTK框架训练VGG网络, 实现最高7%的性能提升, 同时, 也提供了NCCL2和提出的MPI_Bcast设计的性能对比, 新设计在小、中消息通信中实现了最高10倍的时延缩短, 此外, 也看到提出的TA-PC设计相比于PC设计, 在RX1系统上实现了最高50%的性能提升. 实验结果充分说明, 提出的解决方案在可移植性以及性能方面有较大的优势. 提出的设计方案已在MVAPICH2-GDR 2.3rc1版本中公开发布^[17]. 目前只是针对MPI_Bcast进行优化设计, 而对MPI_Reduce和MPI_Allreduce等其他集合通信的支持还有待进一步的研究, 以期在未来能支持全方位的并行DNN训练.