Offloading communication control logic in GPU accelerated applications
Offloading communication control logic in GPU accelerated applications
发表时间: 2017-05
Elena Agostini
NVIDIA and National Research Council of Italy
Santa Clara, CA, USA
Email: [email protected]
Davide Rossetti
NVIDIA
Santa Clara, CA, USA
Email: [email protected]
Sreeram Potluri
NVIDIA
Santa Clara, CA, USA
Email: [email protected]
速读
一句话结论 本文提出了 GPUDirect Async 技术，通过让 GPU 直接触发和同步网络通信操作，将 CPU 从计算与通信交替的关键控制路径中移除，在多个科学计算基准测试中实现了最高 40% 的通信时间缩减与整体应用性能提升。
要解决什么问题 现有的多 GPU 加速应用在处理计算与通信交替的任务时，存在严重的控制路径开销问题。虽然传统的 GPUDirect 技术优化了数据在底层移动的效率，但通信的控制权依然完全依赖 CPU。具体而言，在典型的“计算并发送”工作流中，CPU 需要不断循环执行以下操作：将计算任务排队到 GPU、同步等待 GPU 内核执行完毕、在 InfiniBand 主机通道适配器（HCA）上触发网络通信，最后还要不断轮询 HCA 的完成队列以确认通信结束。这种机制导致 CPU 必须始终以 100% 的峰值功耗状态运行，以保证计算到通信的状态转换尽可能快。更致命的是，当应用进行强扩展导致分配给单个 GPU 的计算规模减小时，CPU 介入控制所带来的延迟和转换开销会占据整体运行时间的极大比例。这不仅使得 GPU 无法及时获取新任务，网络接口卡也处于闲置等待状态，最终造成 GPU 和网络硬件的利用率双双低下，成为限制系统扩展性的核心卡点。
怎么做的 核心思路是将 GPU 与网络适配器之间的控制依赖关系直接卸载到 GPU 自身，让 GPU 绕过 CPU 直接访问并操作网卡的控制寄存器。为了实现这一目标，该方法首先利用 GPU 对等映射技术，通过底层驱动函数将驻留在主机内存中的 HCA 队列对门铃寄存器和完成队列映射到 GPU 的地址空间中，并获取 GPU 可寻址的指针。在此基础上，系统构建了包含底层设备驱动 libmlx5、Verbs API 扩展 libibverbs、同步库 LibGDSync 以及顶层消息传递库 LibMP 的完整软件栈。通过 LibMP，开发者可以采用不同的执行模型来重构应用。最核心的是流异步模型（SA），它利用 CUDA 引入的流控制函数，将通信任务与计算内核一起排队到同一个 CUDA 流中。具体而言，使用向内存地址写入值的指令来敲响 HCA 门铃以触发发送，使用在特定内存位置排队等待条件满足的指令来轮询完成队列。这样一来，CPU 只需要一次性准备并排队一长串的计算和通信任务即可进入低功耗状态，由 GPU 按照流的顺序自动触发网络发送并等待接收完成。此外，系统还设计了内核发起模型（KI），允许负责执行计算的流式多处理器直接在 CUDA 内核线程内部访问网卡寄存器，通过简单的值赋值和比较来触发通信。这种方式能向高度并行的 GPU 硬件暴露更多并发性，甚至支持内核融合技术，将计算与通信控制逻辑写在同一个内核中。
效果如何 实验在两套硬件环境下展开：一是配备双 Tesla K40m GPU 和 Mellanox FDR Connect-IB 网卡的单节点服务器，二是配备 256 张 Tesla K20c GPU 和对应 InfiniBand 网卡的 Wilkes 集群。对比的基线方法是 NO-SA 模型，该模型使用默认的 IB Verbs 实现同步函数，代表了依赖 CPU 调度的传统 MPI 通信路线。在 2DStencil 模板计算任务中，当网格尺寸较小时，内核发起模型相比同步基线实现了 40% 的性能增益，流异步模型也获得了约 20% 的增益；但随着网格尺寸增大，计算耗时足以掩盖通信时间，增益会显著收窄。在 HPGMG-FV 多网格求解器任务中，流异步模型在 16 进程下最高带来 24% 的时间增益，而内核发起模型通过单内核重叠操作最高实现了 26% 的增益。在分子动力学代理应用 CoMD-CUDA 中，虽然通信仅占总时间的 7%，但流异步模型在通信阶段实现了 25% 到 35% 的提速，使整体运行时间改善约 5%。在 LULESH2-CUDA 代理应用中，27 进程下流异步模型平均缩短了 13% 的总执行时间。作者也指出了该方法的局限与代价：系统中的流异步混合模型由于需要 CPU 和 GPU 频繁通过主机内存协调，性能最慢且不被推荐；内核发起模型虽然上限高，但编程极其复杂，且在持久内核场景下面临 CUDA 模型无法保证内存一致性的问题；此外，当前实现仅绑定特定厂商的 InfiniBand 硬件，且为了规避内存一致性问题并排除底层数据直通技术的干扰，实验中的通信缓冲区均被强制放置在主机内存中。
主要贡献
本文从引言中提取的核心问题是：在多GPU加速的应用中，传统的GPUDirect技术主要关注数据移动的效率，但控制路径仍依赖CPU，导致CPU在计算和通信阶段的转换中产生开销，尤其在强扩展应用中，GPU和网络接口卡利用率低下。研究目标是介绍GPUDirect Async技术，该技术利用CUDA 8.0引入的功能，优化GPU与网络适配器（如InfiniBand）之间的控制路径，使GPU直接触发和同步通信操作，从而减少CPU在关键路径中的参与。创新点包括：构建GPUDirect Async技术（第II-C节）；设计并开发简单消息传递库（libMP）以演示GPUDirect Async在应用中的使用（第II-D节）；呈现三种使用GPUDirect Async构建的通信模型（第II-E节）；移植并评估代表不同科学领域的迷你应用和基准，展示GPUDirect Async在实际应用中的益处（第III节）。这是第一篇解释GPUDirect Async构建块、其启用的通信模型以及对科学应用性能益处的论文。
背景知识/关键Observation/设计原则
GPUDirect Async概述。 科学应用通常在计算和通信阶段之间交替。在多节点GPU应用中，从计算到通信的转换涉及在GPU上启动计算内核、等待其完成并通过网络发送数据。使用GPUDirect RDMA和InfiniBand互连的工作流程如下：1) CPU将计算任务排队到GPU并同步等待完成；2) CPU在InfiniBand HCA上排队通信任务；3) HCA直接从GPU内存获取数据；4) HCA发送数据。GPUDirect Async通过使GPU触发HCA上的通信并使HCA解锁CUDA任务来消除对CPU的依赖。CPU只需准备并排队计算和通信任务。带有GPUDirect Async的计算并发送工作流程如下：1) CPU准备并将计算和通信任务排队到GPU；2) GPU完成计算任务并在HCA上触发通信；3) HCA直接从主机内存或GPU内存获取数据；4) HCA发送数据。这种将GPU-HCA依赖关系卸载到GPU的能力允许CPU排队比以前更长的操作链。CPU随后可以处于低功耗状态或执行其他有用工作。
动机。 图3显示了典型的多GPU应用在使用InfiniBand时的计算并发送期间的时间线。CPU负责迭代地将工作调度到GPU、等待GPU内核完成、在HCA上触发通信并轮询HCA以完成通信。CPU必须以峰值功耗状态100%运行，以确保完成检测和转换尽可能快。当应用强扩展且GPU计算规模减小时，这些转换占运行时的相当一部分，导致GPU和NIC利用不足。通过利用GPUDirect Async，可以将多个计算和发送迭代卸载到CUDA流中，从关键路径中移除CPU，如图4所示。这不仅释放了CPU，还隐藏了GPU内核启动的延迟。这也适用于接收并计算工作流程。
InfiniBand队列。 目前，GPUDirect Async仅针对Mellanox InfiniBand互连[9，Infiniband. http://www.mellanox.com/pdf/whitepapers/IB Intro WP 190.pdf]实现。通信操作通过将命令发布到称为队列对（QPs）的发送/接收队列对来向IB HCA（主机通道适配器）发出，然后写入与QP关联的门铃寄存器，以告知HCA关于要处理的新请求。当请求完成（即数据已发送或接收）时，HCA将条目添加到完成队列（CQ）。应用需要轮询CQ以了解请求何时完成（图5）。
方法细节
实现。 GPUDirect Async技术涉及使GPU访问HCA上的QPs门铃寄存器和CQ（在我们的案例中驻留在主机内存中）。这通过两个CUDA驱动函数的组合实现：cuMemHostRegister()，它页锁定现有主机内存范围并将其映射到GPU的地址空间；以及cuMemHostGetDevicePointer()，它检索GPU可寻址的指向注册内存范围的指针。具体来说，cuMemHostRegister()在主机内存的情况下需要CU_MEMHOSTREGISTER_DEVICEMAP标志，在属于第三方PCIe设备（我们的案例中是InfiniBand HCA）的内存映射I/O空间的情况下需要CU_MEMHOSTREGISTER_IOMEMORY标志。后一种过程称为GPU对等映射。所有NVIDIA GPU卡，除了基于新“Pascal”架构的那些，支持仅有40位的PCIe总线地址。因此，Mellanox发布了HCA固件，该固件强制HCA PCIe BAR（基地址寄存器）分配在适当的地址范围内。有了映射到GPU的门铃寄存器和CQ，就可以使用CUDA内核线程访问它们（我们在第II-E节中将其称为内核发起通信模型），或者我们可以指示CUDA流使用最新CUDA 8.0版本中引入的以下CUDA驱动函数等待（轮询CQ）或写入（响铃门铃）这些位置。我们将其称为流异步通信模型。
1) cuStreamWaitValue32(stream, value, address, condition)：在给定内存位置上排队CUDA流的同步。在操作之后排序的工作将阻塞，直到内存上的给定条件满足。
2) cuStreamWriteValue32(stream, value, address)：将值写入内存地址。
3) cuStreamBatchMemOp(stream, list operations)：这是前述函数的批量版本，输入操作列表（等待或写入）及其参数，按列表中出现的顺序排队。
在算法1中，有一个CUDA流将整数写入映射到GPU内存的PCIe设备内存的伪代码。GPUDirect Async将GPU对等映射与cuStreamWriteValue32() API结合，以映射并“响铃”HCA门铃寄存器。要轮询CQ等待CQE，它注册CQ主机内存并使用cuStreamWaitValue32() API。
// Algorithm 1 GPU peer mapping and CUDA stream write
struct device_bar {
void *ptr;
CUdeviceptr dPtr;
size_t len;
} db;
// ....
device_driver_get_bar(&db.ptr, &db.len);
int flags = CU_MEMHOSTREGISTER_IOMEMORY;
cuMemHostRegister(db.ptr, db.len, flags);
cuMemHostGetDevicePointer(&db.dPtr, db.ptr, 0);
// ....
cuStreamWriteValue32(stream, db.dPtr + offset, someValue, 0);
使用GPUDirect Async，CPU任务减少：
- 启动用于计算的CUDA内核。
- 分配并注册通信缓冲区（设备或主机固定内存）。
- 如前所述将HCA特定数据结构映射到GPU。
- 准备发送/接收请求描述符并将它们转换为基本操作序列，例如由负责CUDA流抽象的GPU前端单元执行。
- 清理CUDA流成功读取的CQE。
相反，GPU任务增加：
- 执行用于计算的CUDA内核。发送WQE响铃HCA门铃。
- 等待与发送或接收WQE相关的CQE，在CQ上轮询。
带有GPUDirect Async的InfiniBand工作流程如图6所示。
软件栈。 为了利用GPUDirect Async技术，我们在软件栈的不同级别实现了或修改了库，如图7所示：
1) libmlx5：（供应商/设备特定）Mellanox Connect-IB InfiniBand HCA的低级设备驱动程序，允许用户空间进程直接访问Mellanox HCA硬件，具有低延迟和低开销。标准实现已扩展以满足GPUDirect Async的需求。
2) libibverbs：（Verbs API）OpenFabrics Infiniband Verbs API的实现。标准实现已扩展了特定于GPUDirect Async的新Verbs。
3) LibGDSync：（NVIDIA开源）由我们的NVIDIA团队实现，它由一组混合API组成，其中IB Verbs和CUDA GPU流操作合并。它负责创建尊重GPUDirect Async约束的IB跟踪数据结构，在需要时注册主机内存，直接在GPU流上发布发送指令和完成轮询。
4) LibMP：（NVIDIA开源）由我们的NVIDIA团队实现，它位于栈的顶层，是一个消息传递库（类似于MPI），作为技术演示器开发，以轻松在MPI应用中部署GPUDirect Async技术：初始化MPI环境（即通信器、秩、拓扑等）后，可以用LibMP相应的调用替换标准MPI通信调用：mp_isend_on_stream()代替MPI_Isend()，mp_wait_on_stream()代替MPI_Wait()等。它利用LibGDSync API并提供基本的点对点和单边通信调用。目前，集体调用不可用。
在算法2中，我们呈现了使用LibMP函数的典型GPUDirect Async应用的结构，其中进程使用CUDA流交换数据，在GPU上交替通信和计算期。
GPUDirect Async模型。 GPUDirect Async有不同的执行模型，由LibMP库实现：
1) 流异步模型（SA）：异步。这是前述部分中描述的模型，其中通信任务通过CUDA驱动函数与计算任务（如CUDA内核、CUDA内存拷贝等）一起排队到CUDA流中；因此，通信相对于CPU是异步的，但相对于CUDA流是同步的。它是最简单的模型，因为它只需用SA LibMP调用替换MPI调用。该模型要求：
- 移除所有CUDA同步原语。
- 用异步版本替换所有CUDA同步函数。
- 用LibMP调用替换所有MPI通信调用。
在算法3中，有一个使用GPU的MPI应用的伪代码示例。要应用SA模型，必须如算法4中修改算法3。
// Algorithm 3 MPI generic example C-pseudocode
MPI_Irecv(recvBuffer, ...);
cudaMemset(computeBuffer, ...);
first_cuda_kernel(computeBuffer, ...);
cudaMemcpy(sendBuffer, computeBuffer, ....);
second_cuda_kernel(sendBuffer, ...);
cudaDeviceSynchronize();
MPI_Isend(sendBuffer, ...);
// Algorithm 4 SA model generic example C-pseudocode
mp_irecv(recvBuffer, ...);
cudaMemsetAsync(computeBuffer, ...);
first_cuda_kernel(computeBuffer, ...);
cudaMemcpyAsync(sendBuffer, computeBuffer, ....);
second_cuda_kernel(sendBuffer, ...);
mp_isend_on_stream(sendBuffer, ...);
2) 内核发起模型（KI）：异步。流式多处理器（SM），负责执行CUDA内核，可以直接访问QPs门铃寄存器和CQ，以发送消息或等待完成。这可以与GPUrdma[6，F. Daoud, A. Watad, M. Silberstein GPUrdma: GPU-side library for high performance networking from GPU kernels, Proceedings of the 6th International Workshop on Runtime and Operating Systems for Supercomputers, Article No. 6]启用的GPU发起IB通信比较。GPUrdma在GPU上实现IB verbs，可能占用更多资源（如CUDA内核寄存器）来执行这些复杂函数。这可能影响占用率从而影响计算效率。另一方面，KI模型在CPU上准备发送和接收请求，并让GPU内核线程仅触发通信和轮询完成。这转化为GPU映射内存上的简单值赋值和比较。GPUrdma要求CUDA内核边界内的 consistency 点，这不是CUDA模型保证的。这是一个问题，KI方法在持久内核的情况下也面临，如[24，GPUDirect RDMA Considerations. http://docs.nvidia.com/cuda/gpudirectrdma/#design-considerations]所述。
3) 流异步混合模型（SAH）：异步。类似于SA，但CPU和GPU使用固定主机内存区域协调通信。GPU触发CPU响铃门铃或轮询CQE操作，更新内存区域；当任务成功完成时，CPU进行另一个内存区域写入以告知GPU。该模型是最慢的，不推荐用于性能目的。
4) 无异步模型（NO-SA）：同步。为了比较目的，LibMP有相同异步函数的同步版本，使用默认IB Verbs实现，即mp_send()、mp_isend()、mp_wait()等。
在以下部分，我们将展示内核发起可以比流异步更快，即通过允许内核融合技术，从而向高度并行的GPU HW单元暴露更多并发性。缺点是它更复杂，主要因为程序员需要手动将不同子任务（发送、接收或计算）调度到单独的CUDA内核线程块，尊重算法的约束。因此，我们能够在所有实际应用中使用SA模型，而KI模型仅应用于HPGMG-FV。我们还将展示流异步混合模型是最慢的解决方案。
实验环境
实验使用两个环境。第一个环境（E1）由两个标准2U Xeon服务器组成，每个服务器配备Mellanox双端口FDR Connect-IB HCA和单个Tesla K40m（提升时钟设置为875 MHz），运行RHEL 6.6、CUDA 8.0 RC的预发布GPU显示驱动程序和OpenMPI 1.10.3。第二个环境（E2）是英国剑桥大学的Wilkes集群[16，Wilkes cluster Cambridge, UK. www.hpc.cam.ac.uk]，由128个Dell T620服务器、256个NVIDIA Tesla K20c GPU和256个Mellanox Connect IB卡互连，配备CUDA 8.0 Toolkit、显示驱动程序367.44和OpenMPI 1.10.3。为了避免内存一致性问题（如[24]所述）并更好地评估Async性能而不受GPUDirect RDMA干扰，所有通信缓冲区驻留在主机内存中。基准和迷你应用包括Pingpong、2DStencil、HPGMG-FV、CoMD-CUDA和LULESH2-CUDA。HPGMG-FV使用几何多网格求解器，输入精细级盒子数量和大小的对数；CoMD-CUDA使用EAM势，数据集为2,048,000个原子；LULESH2-CUDA使用结构化网格大小60。
实验结果
Pingpong实验。 该实验是两个MPI进程（rank0和rank1）之间的简单乒乓，交换数据并执行固定时间CUDA内核。在E1环境中2个进程的情况下，根据图8，KI模型比所有其他模型更快（时间范围31.5-32.5 μs），而SAH模型是最慢的（时间范围36.5-38.5 μs）。
2DStencil实验。 2DStencil基准模拟5点模板计算（例如一阶导数）在单精度浮点数2D矩阵上，通过GPU的2D网格域分解并行化，带有一层鬼细胞。比较NO-SA、SA和KI模型。图9显示在E1环境（2进程）中异步模型相对于同步模型的增益；图10显示在E2环境（4进程）中的增益。对于小网格大小，KI模型更快（40%增益）比SA（约20%增益）。仅对于大尺寸，计算足够大以隐藏通信时间，增益显著减少，甚至在KI模型中低于零。在SA情况下，基准从单个（网格大小小于512）切换到双CUDA流实现，这在图中显示为性能突变。
HPGMG-FV实验。 HPGMG-FV[10，HPGMG https://hpgmg.org]是LBNL开发的HPC基准，使用有限体积（FV）[12，Finite Volume method. https://en.wikipedia.org/wiki/Finite volume method]和全多网格（FMG）[13，Full MultiGrid method. https://en.wikipedia.org/wiki/Multigrid method]方法求解椭圆问题。集群实现中，工作负载在进程间公平分布，每个问题级别分为相同大小的盒子。NVIDIA在[11，N. Sakharnykh High-Performance Geometric Multi-Grid with GPU Acceleration. https://devblogs.nvidia.com/parallelforall/highperformance-geometric-multi-grid-gpu-acceleration]中改进了混合解决方案，确保每个级别在合适架构上执行（阈值10000元素以上用GPU，否则CPU，如图11）。使用SA和KI模型实现异步通信。在E1环境中2进程，LibMP模型比MPI更快，KI达到13%时间增益（图14，仅GPU级别）。SA实现如算法6，移除pack和send之间的cudaDeviceSynchronize()，在E2环境中最多24%增益（图16，2进程，log2(size)=4）。KI实现如算法7，使用单个CUDA内核重叠操作（如图17，使用原子操作和inter-block barrier如图18），在E2环境中最多26%增益（图19）。
CoMD-CUDA实验。 CoMD[17，ExMaTex. http://www.exmatex.org/comd.html]是ExMatEx项目的代理应用，实现经典分子动力学，使用短程原子势（LJ和EAM）。NVIDIA的CUDA版本[19，CoMD-CUDA Code. https://github.com/NVIDIA/CoMD-CUDA]在E2环境中使用SA模型重新实现EAM势的force exchange和atoms exchange（4、8、16节点，2,048,000原子数据集）。通信仅占总时间的7%（16进程），但SA模型在通信期中相对于MPI有25%-35%增益（图20），总时间改善约5%。
LULESH2-CUDA实验。 LULESH[20，Lulesh Website. https://codesign.llnl.gov/lulesh.php]是LLNL的代理应用，离散分区空间域。CUDA实现在E2环境中使用SA模型全异步执行（27节点，网格大小60，增加迭代）。相对于MPI的增益如图21所示，平均13%总执行时间改善。
结论
GPUDirect Async是一种使GPU与第三方设备直接控制路径的技术，已随CUDA 8.0发布。InfiniBand网络对Async的支持以新实验OFED verbs集（Mellanox OFED 3.4加上[22，GPUDirect libmlx5. https://github.com/gpudirect/libmlx5]中的更新）和中层抽象库[23，GPUDirect libgdsync. https://github.com/gpudirect/libgdsync]的形式出现，作为BSD。GPUDirect Async仍在开发中：我们正在评估用例和新功能，如在GPU内存上分配InfiniBand元素（即CQ和QP），并测试与GPUDirect RDMA的交互。此外，我们希望使用更多节点测试异步应用以更好地评估其可扩展性，并扩展软件栈以使用不同于InfiniBand的网络协议。
💬 评论讨论
欢迎在这里分享您的想法和见解！