Offloading communication control logic in GPU accelerated applications

Elena Agostini
NVIDIA and National Research Council of Italy
Santa Clara, CA, USA
Email: elena.ago@gmail.com

Davide Rossetti
NVIDIA
Santa Clara, CA, USA
Email: drossetti@nvidia.com

Sreeram Potluri
NVIDIA
Santa Clara, CA, USA
Email: spotluri@nvidia.com

主要贡献

本文从引言中提取的核心问题是：在多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的网络协议。