E. Agostini *, D. Rossetti, S. Potluri

NVIDIA, Santa Clara, CA, United States

主要贡献

本文探讨了GPUDirect Async技术的动机和构建块，该技术是CUDA 8.0引入的新功能，允许GPU直接与第三方设备同步，例如直接触发和轮询排队到InfiniBand Connect-IB网络适配器的通信操作，而无需CPU参与GPU应用的通信关键路径。核心问题是传统GPU加速并行应用中，CPU作为GPU计算和网络通信任务的协调者，导致同步开销（如cudaStreamSynchronize和MPI_Wait），从而限制性能，尤其在强缩放时计算和通信时长缩短时。研究目标是通过微基准测试和性能模型展示GPUDirect Async在多GPU MPI应用中的潜在益处，包括HPGMG-FV（几何多网格应用的代理）和CoMD-CUDA（经典分子动力学代码的代理），并报告一个不提供优势的测试案例，即大规模图的广度优先搜索算法实现。创新点包括引入两种异步通信模型：Stream Asynchronous（SA）模型，通过CUDA流混合通信和计算；Kernel-Initiated（KI）模型，由CUDA内核触发通信；开发LibMP作为简单消息传递库来演示GPUDirect Async的使用；提出性能模型来解释领域分解数值应用的结果；使用代表不同科学领域的基准测试展示益处。

背景知识/关键Observation/设计原则

相关工作概述。自GPU作为通用加速器引入以来，许多论文研究了优化GPU和NIC之间数据路径的方法，例如[2, Ammendola et al., GPU peer-to-peer techniques applied to a cluster interconnect, 2013, IEEE International Parallel & Distributed Processing Symposium, IPDPS]和[4, Ammendola et al., NaNet: a flexible and configurable low-latency NIC for real-time trigger systems based on GPUs, 2014, Journal of Instrumentation]，这些工作使用自定义FPGA-based NIC原生实现NVIDIA P2P和RDMA协议；类似地，[27, Oden and Froning, GGAS: Global GPU address spaces for efficient communication in heterogeneous clusters, 2013, IEEE International Conference on Cluster Computing, CLUSTER]实验了使用EXTOLL互连的GPUDirect RDMA。到目前为止，只有少数论文探讨了将通信控制路径卸载到GPU的方法，如表1所示，该表基于通信阶段（通信描述符准备、触发）的位置（CPU、GPU流、GPU SMs）、NIC控制结构的位置（主机内存、GPU内存）、基准应用以及相关数据缓冲区的位置（使用GPUDirect RDMA通信位于GPU内存的数据缓冲区）来分类这些方法。

表1 比较引用的不同方法与本文提出的方法（最后两行）
表1 比较引用的不同方法与本文提出的方法（最后两行）

S. Kim 等人的工作。[23, Kim et al., GPUnet: Networking Abstractions for GPU Programs, 2014, USENIX Symposium on Operating Systems Design and Implementation]描述了一个原生GPU网络层，提供类似于BSD的套接字抽象和高层次网络API给GPU程序。虽然这些API可以直接由CUDA线程调用，但它们实际上由运行在CPU上的代理代理执行。

Lena Oden 等人的工作。[28, Oden et al., Infiniband-Verbs on GPU: A case study of controlling an Infiniband network device from the GPU, 2014, IEEE International Parallel & Distributed Processing Symposium Workshops, IPDPSW]探索了从CUDA内核生成和发布InfiniBand发送和接收操作的不同方法。在一个实验中，他们实现了GPU侧的IB Verbs子集，修改了NVIDIA内核模式驱动程序的开源部分和Mellanox用户空间驱动程序，以将一些关键InfiniBand HCA资源（例如内存队列和硬件门铃寄存器）映射到GPU。这些GPU Verbs API使用临界区以单个CUDA线程块粒度序列化对IB QP的访问。他们显示了不令人满意的结果，导致作者得出结论：GPU原生设计不如传统的CPU控制网络传输。

F. Daud 等人的工作。[12, Daoud et al., GPUrdma: GPU-side library for high performance networking from GPU kernels, 2016, International Workshop on Runtime and Operating Systems for Supercomputers]实现了GPU侧的全局地址空间编程接口（GPI）库子集，使GPU代码能够直接进行高性能RDMA通信，即完全绕过CPU，CUDA内核线程既准备又触发与通信相关的NIC命令。与[28]类似，他们将InfiniBand资源映射到GPU。他们还实验了将其中一些资源（QP、CQ）备份到GPU内存而不是主机内存。在最后两篇论文中，他们在GPU侧重新实现了通信栈的一部分。此外，他们必须黑入GPU和/或HCA驱动程序，以允许GPU访问NIC门铃并将控制结构放置在GPU内存。这种方法有两个缺点：GPU侧栈使用更多GPU资源（例如寄存器），可能降低占用率，从而降低计算内核的性能，其中使用了通信函数。此外，据我们所知，它们受到与在持久CUDA内核中使用通过GPUDirect RDMA更新的接收数据缓冲区相关的GPU内存一致性问题[19, NVIDIA, GPUDirect RDMA design considerations, http://docs.nvidia.com/cuda/gpudirect-rdma/#design-considerations]的影响。GPUDirect Async正式引入了一种机制来围栏针对GPU内存缓冲区的传入流量。该机制暴露为新的FLUSH MemOP，可以在等待通信完成通知后和释放预启动GPU内核之前排队。

图1 GPUDirect RDMA 计算并发送工作流
图1 GPUDirect RDMA 计算并发送工作流

Ohio State University团队的工作。[34, Venkatesh et al., MPI-GDS: High performance MPI designs with GPUDirect-aSync for CPU–GPU control flow decoupling, 2017, International Conference on Parallel Processing]与本文作者合作，呈现了使用GPUDirect Async技术在MVAPICH2（MPI-GDS）中的早期结果。该论文的范围是探索利用GPUDirect Async的设计，同时尊重MPI规范的要求。在这方面，本文构成了该工作的前提。更具体地说，MPI-GDS提供与CUDA流同步的MPI点对点原语。支持MPI标签匹配和会合协议，并使用混合方法实现，其中CPU主动推进协议的一部分，代价是额外的开销。在本文中，我们测量GPUDirect Async本身的潜在性能，而不管其他GPUDirect技术的使用。此外，我们还探索单边通信原语和CUDA内核发起的通信。最后，他们关注微基准，而这里我们呈现应用基准。

与之前工作的比较。与我们之前的工作[1, Agostini et al., Offloading communication control logic in GPU accelerated applications, 2017, IEEE/ACM International Symposium on Cluster, Cloud and Grid Computing]相比，本文引入了几项改进：对Async技术和其软件栈的更详细描述；捕获领域分解多GPU应用通信模式的通用性能模型。该性能模型应用于GPUDirect Async模型，以澄清每个异步模型达到性能增益所需的要求；这里讨论了新的微基准和负面测试案例，澄清了GPUDirect Async的一些限制。

GPUDirect Async的动机。科学应用通常在计算和通信阶段之间交替。在多节点GPU应用中，从计算到通信的转换涉及在GPU上启动计算内核、等待其完成，然后通过网络发送数据。使用GPUDirect RDMA启用的InfiniBand HCA的工作流如图1所示：CPU将一些计算任务排队到GPU（内核启动）并同步等待其完成；CPU将通信任务排队到InfiniBand HCA；HCA直接从GPU内存获取数据，得益于GPUDirect RDMA；HCA通过网络注入相关消息；CPU通过等待完成与HCA同步（图中未显示）。注意，当不使用GPUDirect RDMA时，合适的GPU到主机复制将包含在通信任务中。GPUDirect Async通过启用GPU在HCA上触发通信和HCA解锁CUDA任务来消除对CPU的依赖。CPU只需准备和排队计算和通信任务。在GPUDirect Async存在下的计算并发送工作流如图2所示：CPU准备并将计算和通信任务排队到GPU；GPU完成计算任务并直接在HCA上触发挂起的通信；HCA直接从主机内存获取数据，或在使用GPUDirect RDMA时从GPU内存获取；HCA发送数据。在后一种情况下，CPU工作负载发生变化。例如，在将所有必要任务排队到GPU后，CPU可以返回并做其他有用工作。我们在这里注意到GPUDirect Async独立于GPUDirect RDMA，因此我们可以隔离地实验它，即使用前者而不启用后者。[3, Ammendola et al., GPU peer-to-peer techniques applied to a cluster interconnect, 2013, CASS 2013 workshop at IEEE International Parallel & Distributed Processing Symposium, IPDPS]注意到GPUDirect RDMA性能高度依赖于PCIe结构，即连接NIC到GPU的PCIe桥和交换机的类型和数量，以及特定的GPU架构[32, Rossetti et al., State of GPUDirect technologies, 2016, GPU Technology Conference]。例如，对于会合协议中的大消息大小，在发送侧通过主机内存的流水线暂存比使用GPUDirect RDMA更有效率。这里做同样的事情，如在MPI-GDS[34]中，将要求CPU推进通信，这违背了我们隔离基准GPUDirect Async的目标。考虑到我们的基准平台之一中GPUDirect RDMA效率低下，我们决定不使用它。

图2 GPUDirect Async 计算并发送工作流
图2 GPUDirect Async 计算并发送工作流

图3描绘了典型多GPU应用的时序，包括计算和通信阶段。CPU迭代地将工作调度到GPU，等待GPU内核完成，在HCA上触发通信，并最终轮询HCA以完成通信。CPU必须大多数时间以峰值性能运行，以确保响应性，即不同阶段尽可能快速调度。通常，当应用强缩放时，由于几何和/或物理属性，每个计算节点的GPU计算和网络通信长度都会减少。此外，计算和通信的减少速率不同，例如前者与体积缩放，而后者与表面缩放，如领域分解方法。因此，即使算法允许它们重叠，也越来越难以将通信隐藏在计算后面。超过某个点，应用不再缩放。不缩放体制的开始可以提前——即应用在较小的GPU数量停止缩放——如果CPU在启动计算和通信任务时产生的开销与分别在GPU和NIC上执行它们所需的时间相同。在这些情况下，启动GPU计算可能需要数十微秒，这可能与执行该任务的时间相同，或通过网络交换几千字节数据的时间相同。类似地，一些应用，如稍后介绍的HPGMG，可能经历阶段——粗粒度级别——其中将计算移回CPU更方便，以避免启动GPU工作的开销。通过利用GPUDirect Async，整个并行计算阶段可以卸载到CUDA流。这反过来允许重叠——从而支付成本——迭代i的工作提交，同时迭代i-1由GPU协调，有效地将CPU从关键路径移除。如图4所示定性显示，有时候CPU变得空闲，可能持续较长时间，在此期间它可以做有用工作。当不可能时，可以允许CPU进入更深的睡眠状态，从而降低应用功率剖面。

图3 通信期 MPI 多GPU应用时序
图3 通信期 MPI 多GPU应用时序

图4 具有GPUDirect Async的通信期多GPU应用时序
图4 具有GPUDirect Async的通信期多GPU应用时序

方法细节

GPUDirect Async实施。目前，对GPUDirect Async的支持需要MLNX OFED 4.0中的扩展IB Verbs API，并限于最新一代Mellanox InfiniBand主机通道适配器（HCA）[22, Mellanox, InfiniBand Standard, http://www.mellanox.com/pdf/whitepapers/IB_Intro_WP_190.pdf]，这些由libmlx5用户空间提供程序库支持。传统上，CPU通过填充数据结构（工作请求或WQE）到发送或接收内存队列，然后更新某种门铃寄存器（两者都与特定队列对（QP）关联）来向IB HCA发布通信操作。门铃更新需要通知HCA有新请求准备好处理。在最近Mellanox HCA的特定情况下，触发发送操作需要两个不同的门铃更新：一个到主机内存中的32位字（DBREC），另一个到HCA PCIe资源之一（基地址寄存器或BAR）中特定偏移的硬件寄存器。当使用内核旁路时，用户空间进程通过使用HCA BAR页面的非缓存内存映射I/O（MMIO）映射直接更新DB。当请求完成——即数据已发送或接收——HCA将新完成队列条目（CQE）添加到与该QP创建时关联的发送或接收完成队列（CQ）。应用需要轮询相应的CQ以检测请求是否已完成（图5）。

图5 InfiniBand HCA 发送/接收请求处理
图5 InfiniBand HCA 发送/接收请求处理

虽然CPU仍负责准备命令，但GPUDirect Async要求GPU直接访问HCA门铃寄存器和CQ（在我们的情况下驻留在主机内存），使用两个CUDA驱动函数的组合：cuMemHostRegister()来页锁定现有主机内存范围并将其映射到GPU的地址空间；以及cuMemHostGetDevicePointer()来检索相应的设备指针。特别是，当注册属于第三方PCIe设备（在我们情况下是InfiniBand HCA）的MMIO地址范围时，使用CU_MEMHOSTREGISTER_IOMEMORY标志。后者对应于创建所谓的GPU对等映射，即GPU到对等PCIe设备的映射。注意，在当前实现中，整个MMIO范围必须物理上连续，并标记为CPU缓存抑制。由于Pascal架构之前的NVIDIA GPU中的硬件限制，需要特殊的Mellanox HCA固件来让HCA PCIe资源（BAR）放置在适当的地址范围。一旦门铃寄存器和CQ映射到GPU，就可以（a）在CUDA流上或（b）从CUDA内核线程访问它们。我们将前者称为Stream Asynchronous（SA）通信模型——见Section 4.2——后者称为Kernel-Initiated（KI）通信模型——Section 4.3。在SA模型中，我们广泛使用CUDA内存操作API，如下所述，以等待（轮询）CQE或写入（ ringing）门铃寄存器。1. cuStreamWaitValue32(stream, value, address, condition)：将CUDA流的同步排队到给定内存位置。在操作后有序的工作将阻塞，直到给定条件（EQual, Greater-or-EQual, AND）满足。例如，这允许阻塞CUDA流，直到NIC信号特定完成事件（CQE）。2. cuStreamWriteValue32(stream, value, address)：将传递的值写入由设备地址标识的内存。该API用于 ringing QP门铃寄存器。3. cuStreamBatchMemOp(stream, count, mem_ops[])：前述函数的批量版本，输入内存操作向量（等待或写入）。

图6 具有GPUDirect Async的InfiniBand HCA 发送/接收请求处理
图6 具有GPUDirect Async的InfiniBand HCA 发送/接收请求处理

当使用GPUDirect Async时——见图6中的交互图——CPU仍需要：分配和注册通信缓冲区（设备或主机固定内存）；如上所述将HCA特定数据结构映射到GPU；准备并在QP上发布WQE；准备发送和接收请求描述符并将它们转换为内存操作序列；一旦CUDA流成功读取CQE，就轮询它们。相反，GPU有更多任务：通过 ringing 门铃触发准备好的WQE；在发送或接收WQE相关的CQE上等待，轮询CQ。鉴于CUDA内存操作API提供的少量功能，即CUDA流只能阻塞在内存位置，我们不能在那里实现完整的CQE解析器和调度器，如在CPU上。因此，如果我们想字面上保持CPU脱离关键路径，GPU流需要CQE和发送/接收操作严格关联，即第i个WQE的完成将放置在第i个可用CQE中。例如，如果在QP创建时使用不同的发送和接收CQ，以及放弃共享接收队列（SRQ），就是这种情况。请注意，错误处理和恢复仍由CPU完成。当轮询CQE时，CPU可能观察到带有错误的完成。在这种情况下，它负责中止GPU和HCA的所有未完成工作。

软件栈。为了利用GPUDirect Async技术，我们在图7所示软件栈的不同级别实现了或修改了库。

图7 GPUDirect Async 软件栈
图7 GPUDirect Async 软件栈

libibverbs。libibverbs实现了OpenFabrics InfiniBand Verbs API规范。在版本4.0中，Mellanox引入了新的Peer-Direct Async API（例如见peer_ops.h头文件），针对NVIDIA GPUDirect Async技术。

libmlx5。它是管理最近Mellanox InfiniBand HCA的供应商特定低级提供程序库。它允许用户空间进程以低延迟和低开销直接访问Mellanox HCA硬件（内核旁路）。

LibGDSync。由作者开发，它概念上在InfiniBand Verbs上实现了GPUDirect Async支持，通过桥接CUDA和Verbs API之间的差距。它由一组低级API组成，这些API仍非常类似于IB Verbs，尽管在CUDA流上操作。LibGDSync负责创建Verbs对象，即队列对（QP）、完成队列（CQ）、尊重GPUDirect Async约束的结构，当需要时注册主机内存，直接在GPU流上发布发送指令和完成等待。像gds_stream_queue_send或gds_stream_wait_cq这样的函数，在内部使用如前一节3.2所述的CUDA流MemOp API。

LibMP。由作者实现，它是一个基于LibGDSync API构建的轻量级消息传递库，作为技术演示器开发，以轻松在应用中部署GPUDirect Async技术。一旦MPI环境初始化（即通信器、秩、拓扑等），就可以用相应的LibMP原语替换标准MPI通信原语，例如用mp_isend_on_stream()替换MPI_Isend()，用mp_wait_on_stream()替换MPI_Wait()等。LibMP的功能和设计权衡是：使用IB verbs的发送/接收语义的点对点通信原语：接收缓冲区按照它们在特定QP上发布的顺序消耗；单边异步通信，例如在远程内存地址上的put和get；不支持MPI风格的标签匹配；没有集体通信原语。如前所述，每个QP都有自己的CQ。WQ和CQ的深度可以在运行时设置；在我们的基准中，我们使用了512个条目的默认深度。在我们的实验中，WQ、CQ和DBREC驻留在主机内存；在未来版本中，我们计划启用使用GPU内存用于CQ和DBREC。通信原语的参数（即目标/源对等秩、消息大小、缓冲区指针）在CPU发布WQE时使用，在收集描述符并将它们转换为CUDA API调用之前。因此，它们必须在WQE发布时已知，并且不能例如是GPU计算的结果，这在一些应用中可能增加复杂性，如下所示。虽然原则上可以通过在GPU工作内直接修改WQE来更改其中一些参数，例如在触发它们之前，但这会带来如[12, Daoud et al., GPUrdma: GPU-side library for high performance networking from GPU kernels, 2016, International Workshop on Runtime and Operating Systems for Supercomputers]和[28, Oden et al., Infiniband-Verbs on GPU: A case study of controlling an Infiniband network device from the GPU, 2014, IEEE International Parallel & Distributed Processing Symposium Workshops, IPDPSW]中讨论的知名挑战。

系统要求。Async的要求是：Mellanox Connect-IB或更高版本的HCA，可能带有特殊固件版本；MLNX OFED 4.0用于Peer-Direct Async Verbs API；CUDA 8.0用于第3.2节所述的流内存操作API；NVIDIA显示驱动程序版本384或更高；LibGDSync库，可在[17, GPUDirect LibGDSync, http://github.com/gpudirect/libgdsync]获得；需要特殊的NVIDIA内核驱动程序注册表键来启用GPU对等映射；nvidia_peer_memory内核模块；GDRcopy库[15, GDRcopy, https://github.com/NVIDIA/gdrcopy]。在算法1中，我们呈现了典型GPUDirect Async应用的结构，使用LibMP函数，其中两个进程使用Stream Asynchronous模型交换一些数据，混合通信和计算任务。

GPUDirect Async模型。如前述部分所述，LibMP呈现两种不同的执行模型：Stream Asynchronous模型（SA），其中通信相对于主机是异步的，相对于CUDA流是同步的；Kernel-Initiated模型（KI），其中通信由内核内的CUDA线程触发。在本节中，通过抽象性能模型的帮助，我们将我们的Async模型的行为与标准MPI通信模型进行比较。我们考虑GPU加速MPI应用的典型执行流，其中每个MPI秩在计算和其他对等体的通信之间交替。后来，在第6节中，我们将使用我们的性能模型来探索我们期望GPUDirect Async优于MPI模型的条件。

算法1 LibMP 示例 C-伪代码。

1: numRanks=2, Nreq=1;   
2: ▷ Initialize MPI and CUDA environment   
3: initialize_MPI_environment();   
4: cuda_init();   
5: myRank = get_MPI_rank();   
6: ...   
7: ▷ Initialize LibMP environment   
8: mp_init(MPI_COMM_WORLD, !myRank, numRanks);   
9: ...   
10: ▷ Create mp requests descriptors   
11: mp_request_t * sreq, rreq;   
12: host_memory_alloc_request(sreq, Nreq);   
13: host_memory_alloc_request(rreq, Nreq);   
14: ...   
15: ▷ Allocate send/receive buffers   
16: memory_alloc_buffer(sendBuffer, sizeS);   
17: memory_alloc_buffer(recvBuffer, sizeR);   
18: ...   
19: ▷ Register related memory regions   
20: mp_reg_t sreg, rreg;   
21: mp_register(sendBuffer, sizeS, &sreg);   
22: mp_register(recvBuffer, sizeR, &rreg);   
23: ...   
24: ▷ Post a Receive WQE   
25: mp_irecv(recvBuffer, sizeR, !myRank, &rreg, &rreq));   
26:   
27: ▷ Start a CUDA kernel to prepare send buffers   
28: launch_cuda_kernel(sendBuffer, ...., stream);   
29:   
30: ▷ Trigger HCA for Send WQE   
31: mp_isend_on_stream(sendBuffer, sizeS, !myRank, &sreg, &sreq, stream);   
32:   
33: ▷ Wait (poll) for Receive CQE   
34: mp_wait_on_stream(&rreq, stream);   
35:   
36: ▷ Start a CUDA kernel to work on received data   
37: launch_cuda_kernel(recvBuffer, ...., stream);   
38: ...   
39: ▷ Cleanup CQEs   
40: mp_wait(&rreq);   
41: mp_wait(&sreq);   
42: ...   
43: ▷ Synchronize and cleanup   
44: cudaDeviceSynchronize();   
45: mp_deregister(&rreg);   
46: mp_deregister(&sreg);   
47: cleanup_MPI_environment();

CPU同步模型。作为常规多GPU MPI应用的示例，我们考虑D维迭代模板计算的内核，使用领域分解方法并行化。可以识别三个独立阶段：1. 计算并发送：X次，启动（LAi时间）一些CUDA任务（运行Ai时间）如内核或内存传输，在主机上执行一些操作，如与CUDA流的同步（TH时间），然后发送计算数据（Si时间）。2. 内部计算：Y次，在主机上执行一些操作（TH时间）并启动（LBj时间）一些CUDA任务（Bj时间），处理内部数据元素，即不依赖于来自邻居节点的数据。3. 接收并计算：Z次，等待从其他进程接收（Wk），在主机上执行一些操作（TH时间）并启动（LCk时间）CUDA任务（Ck时间），处理接收数据。考虑上述模式的R次迭代，如图8所示，方程(1)分别表示在CPU（TCPUS）、GPU（TGPUS）和整个应用（TS）上花费的时间。Tidle是GPU在等待CPU工作时花费的空闲时间。

图8 迭代多GPU领域分解MPI应用的典型时序
图8 迭代多GPU领域分解MPI应用的典型时序

总时间TS将等于CPU时间，因为CPU总是忙碌，即最坏情况下等待GPU任务的完成，由THsync参数表示：

$TS = TCPUS = THsync + \sum (LA + A + TH + S + LB + B + TH + LC + C + W + TH)$

在以下部分，我们检查LibMP通信模型的情况，给出它们应用的一些示例。

流同步、CPU异步模型（SA）。如前所述，在此模型中，通信与CUDA任务（如内核、内存传输等）一起排队到CUDA流。通常，此模型相对容易使用，因为它需要对MPI应用进行很少更改（即用mp_isend_on_stream修改MPI_Isend，忽略CUDA同步原语）。计算和通信任务相对于主机代码是异步执行的，但与CUDA流同步。如果可能更改原始算法以与以下公式(2)一致（由图9表示），则第4.1节引入的应用类可以使用SA模型利用：

图9 多GPU SA模型应用的一般通信模式，公式(2)表示
图9 多GPU SA模型应用的一般通信模式，公式(2)表示

其中LSi和LWk分别是CPU在CUDA流上排队发送或等待接收操作所花费的时间。在此模型中，Tidle时间可以忽略，因为由于异步行为，CPU在不等待它们完成的情况下在CUDA流上排队大量顺序任务。为了确保异步行为，在通信期需要：移除所有CUDA同步原语；用相应的CUDA异步原语替换所有非异步CUDA原语；通信参数必须在发布时已知（例如发送或接收缓冲区大小、目标秩、指针等）；所有MPI函数必须替换为LibMP函数。一个明显的副作用是CPU有更少的工作要做，因为主机代码既不做同步也不做通信，因此在异步上下文中不相关；因此，我们可以认为公式(2)中的TH参数可以忽略。与公式(2)一致的算法相对于同步版本表示改进，如果以下3个条件得到验证。

C1条件：异步性。在公式(2)中，总执行时间等于GPU时间，如果：

$\sum (LA + LS + LB + LW + LC) < \sum (A + S + B + W + C)$

即CPU在CUDA流上排队任务所需的时间（启动时间）必须小于GPU执行这些任务的时间（C1条件）。没有这个条件，异步无法发生，因为CPU启动时间大于GPU执行时间。

C2条件：时间增益。SA模型（公式(2)）比同步模型（公式(1)）更快，如果：

$TGPUSA < TS$

这总是用更强的条件验证：

$TGPUSA < TCPUS$

因为TGPUS ≤ TCPUS。如果GPU计算任务在两个模型中需要大约相同的时间：

$\sum (A + B + C)SA \approx \sum (A + B + C)S$

那么，我们得到一个简单条件：

$\sum (S + W)SA < Tidle S$

这意味着如果通信时间总和（发送TS和等待TW）小于同步模型中的GPU Tidle时间（GPU等待CPU工作的空闲时间），则SA模型更快。直观地说，SA模型的有效性依赖于SA模型中CUDA流同步通信相对于S模型中CPU发起通信加上GPU同步的相对大小。

C3条件：碎片化计算。子任务R、X、Y和Z的数量越大，执行越异步（由于C1）。(C3)：R > 0, Y ≥ 0, max(X, Z) > 0。在第6节中，我们将这些条件应用于几个MPI + CUDA应用。

Kernel-Initiated模型（KI）。负责执行CUDA内核的流式多处理器（SMs）可以直接向发送消息或等待接收完成的通信原语发布。HCA门铃寄存器和CQ映射在GPU中，CUDA线程可以使用代码内的简单值赋值和比较来分别 ringing 门铃或轮询CQ。在KI模型中，我们使用了内核融合技术，其中在单个CUDA内核中融合通信（发送或等待）和计算任务（类型A、B或C）。这种方法如果与GPUDirect RDMA结合，可能导致GPU内存一致性问题[19, NVIDIA, GPUDirect RDMA design considerations, http://docs.nvidia.com/cuda/gpudirect-rdma/#design-considerations]。为了避免这个问题，在我们的基准（第6节）中，我们使用主机内存作为通信缓冲区。Kernel-Initiated模型的典型时序如图10所示。

图10 Kernel-Initiated 通信模式时序
图10 Kernel-Initiated 通信模式时序

如SA模型中，CPU准备通信描述符，后来这些通信直接由CUDA内核KI中的线程触发（使用描述符），而不是如SA模型中由CUDA流触发。复杂性移动到CUDA内核KI，这至少需要N + M + 1个块，其中N是发送操作前计算类型A任务所需的块数，M是处理接收数据的类型C任务所需的块数加上1个块，用于轮询CQ，如图11所示。如SA模型中，任务B表示与其他通信无关的其他（可能）CUDA任务，可以由任何种类的块执行（图中所示），取决于特定算法。

图11 KI内核，CUDA块任务
图11 KI内核，CUDA块任务

在KI模型中，内核融合技术与动态调度器结合，使用原子操作来挑选每个线程块并根据以下规则将其分配到正确的任务：为了避免死锁，接收操作不得阻止发送操作启动或进展：必须始终至少有一个块等待接收和另一个块执行发送。接收在我们经验中是时间关键的，因此第一个接收器块用于等待传入消息。特别是每个线程轮询与每个远程节点关联的接收CQ。从第二个到N+1-th的块分配到A组操作，而剩余的M个块分配到C组操作，等待接收器块信号所有传入数据已被接收。使用inter-block barrier方案[36, Xiao and Feng, Inter-Block GPU Communication via Fast Barrier Synchronization]来同步接收器和taskC块，如图12所示，其中每个taskC块的线程0等待接收器块的线程0将全局内存变量设置为1，而剩余线程移动到__syncthreads() barrier。为了防止浪费CUDA块等待接收器，TaskA、TaskB和send应始终在TaskC之前执行。当发生这种情况（接收完成后），taskC块中的所有线程0将到达匹配的__syncthreads() barrier，然后开始解包接收数据。对于第4.1和4.2节引入的相同应用类，KI模型执行时间可以估计为：

图12 Inter-block barrier
图12 Inter-block barrier

$TKI = TGPUKI = \sum (LA + (Ai + Si) + Bj + (Wk + Ck) - Overlap)$

其中Ai + Si是发送者块执行任务A加上发送所花费的时间，Wk + Ck是等待数据并执行任务C所花费的时间；Bj表示与其他通信无关的其他（可能）任务（如处理内部结构），可以由任何类型块（A块、C块或其他块）执行。经验上，我们测量到TCPUKI总是可以忽略，因此我们可以认为TKI = TGPUKI。当在GPU上运行时，CUDA内核的多个块可以并发执行，因此任务A + S、任务W + C和任务B的执行可以重叠。为了表示这个大小，我们定义了Overlap，这是一个非平凡函数，表示所有任务的重叠时间，考虑几个输入参数，如GPU SM的数量、任务调度算法、特定通信模式、任务A、B和C的计算时间等。KI模型相对于SA模型的增益可以描述为：

$Gain = \sum Overlap$

最佳场景是当所有任务适合GPU上可用的逻辑CUDA块数且任务在执行的大部分时间重叠：

$Overlap \approx min(\sum (Ai + Si), \sum (Wk + Ck), \sum Bj)$

相反，最坏场景是当每个任务需要高数量的块，减少时间重叠的重要性：

$Overlap \approx 0$

这意味着A、B和C任务将几乎顺序执行，如SA模型的情况，重叠在公式(3)中不代表相对于公式(2)的真正改进。在第6节中，我们解释何时更方便使用KI模型或SA模型。

方法细节中的引用汇总。方法细节中引用的参考文献包括：[1, Agostini et al., Offloading communication control logic in GPU accelerated applications, 2017, IEEE/ACM International Symposium on Cluster, Cloud and Grid Computing]（在相关工作比较中，描述为之前工作引入改进）；[2, Ammendola et al., GPU peer-to-peer techniques applied to a cluster interconnect, 2013, IEEE International Parallel & Distributed Processing Symposium, IPDPS]（在相关工作概述中，描述为优化GPU和NIC数据路径的方法）；[3, Ammendola et al., GPU peer-to-peer techniques applied to a cluster interconnect, 2013, CASS 2013 workshop at IEEE International Parallel & Distributed Processing Symposium, IPDPS]（在动机中，描述为注意到GPUDirect RDMA性能依赖于PCIe结构）；[4, Ammendola et al., NaNet: a flexible and configurable low-latency NIC for real-time trigger systems based on GPUs, 2014, Journal of Instrumentation]（在相关工作概述中，描述为自定义FPGA-based NIC）；[12, Daoud et al., GPUrdma: GPU-side library for high performance networking from GPU kernels, 2016, International Workshop on Runtime and Operating Systems for Supercomputers]（在相关工作和LibMP中，描述为GPU侧GPI库实现，并讨论挑战）；[15, GDRcopy, https://github.com/NVIDIA/gdrcopy]（在系统要求中，描述为所需库）；[17, GPUDirect LibGDSync, http://github.com/gpudirect/libgdsync]（在系统要求中，描述为所需库）；[19, NVIDIA, GPUDirect RDMA design considerations, http://docs.nvidia.com/cuda/gpudirect-rdma/#design-considerations]（在相关工作和KI模型中，描述为内存一致性问题）；[22, Mellanox, InfiniBand Standard, http://www.mellanox.com/pdf/whitepapers/IB_Intro_WP_190.pdf]（在实施中，描述为HCA支持）；[23, Kim et al., GPUnet: Networking Abstractions for GPU Programs, 2014, USENIX Symposium on Operating Systems Design and Implementation]（在相关工作中，描述为GPU网络层）；[27, Oden and Froning, GGAS: Global GPU address spaces for efficient communication in heterogeneous clusters, 2013, IEEE International Conference on Cluster Computing, CLUSTER]（在相关工作概述中，描述为EXTOLL互连实验）；[28, Oden et al., Infiniband-Verbs on GPU: A case study of controlling an Infiniband network device from the GPU, 2014, IEEE International Parallel & Distributed Processing Symposium Workshops, IPDPSW]（在相关工作和LibMP中，描述为GPU侧IB Verbs实现和挑战）；[32, Rossetti et al., State of GPUDirect technologies, 2016, GPU Technology Conference]（在动机中，描述为GPU架构影响）；[34, Venkatesh et al., MPI-GDS: High performance MPI designs with GPUDirect-aSync for CPU–GPU control flow decoupling, 2017, International Conference on Parallel Processing]（在相关工作和动机中，描述为MPI-GDS早期结果）；[36, Xiao and Feng, Inter-Block GPU Communication via Fast Barrier Synchronization]（在KI模型中，描述为inter-block barrier方案）。这些引用出现在相应段落中，用于支持技术描述、比较方法或解释要求。

实验环境

实验使用的数据集包括：微基准测试中的消息大小从0到大消息（独立参数）；HPGMG-FV CUDA中使用8个盒子，log2(size)=4到更大尺寸的弱缩放；CoMD-CUDA中使用分布式2,048,000个原子，EAM势；BFS中使用大规模图的邻接矩阵，2D分解。模型架构关键参数包括：CUDA内核如Pack、Interior、Unpack（HPGMG-FV），力交换和原子交换内核（CoMD-CUDA），BFS中的CUDA内核和CUB函数。硬件配置为：微基准使用两个Supermicro服务器，每个有单个Tesla K40m GPU（提升时钟875 MHz）和Mellanox Connect-IB卡（56 Gbps带宽），通过FDR IB交换机连接；应用基准使用Wilkes集群（128 Dell T620服务器、256 NVIDIA Tesla K20c GPU、256 Mellanox Connect-IB卡）。软件配置为：CUDA 8.0工具包，OpenMPI 1.10.7/1.10.3，MLNX OFED 3.4/4.0（带补丁），LibGDSync，LibMP，NVIDIA显示驱动程序384或更高，nvidia_peer_memory模块，GDRcopy库；基准在Linux OS上实现，使用C/CUDA。

实验结果

微基准测试。所有微基准运行1000次预热迭代并取1000个样本。乒乓延迟基准基于点对点（发送-接收）通信，使用主机内存数据，可选执行恒定时间CUDA内核模拟GPU计算。结果显示：无内核时，MPI基线为微秒级半往返延迟（图中为全往返）；SA模型（红圈）延迟更不规则，有分段常数期和突发峰值，由于GPU前端单元轮询模式；KI模型（绿三角）较平滑；传统CPU驱动通信对小消息更快，GPU路径有约2.5 µs恒定每操作开销（图13）。通知等待在SA模型中，等待反应时间（WRT）为锯齿波函数，周期12 µs，每额外活跃流加1 µs（图15，Tesla K40m）；Pascal架构（如P100）WRT更好（图17）。带有~5 µs内核的乒乓延迟显示KI模型最佳，SA模型片段常数行为消失（图18）；MPI时间线显示接收后启动内核、同步并发送（图19）；SA显示所有任务在GPU流上约3 ms后执行，CPU/GPU重叠（图20）。

应用基准。HPGMG-FV CUDA（几何多网格求解器）使用2D模板模式：Pack内核（A任务）、同步/发送、Interior内核（B任务）、接收、Unpack内核（C任务）（图21、22）。SA模型修改移除同步，使用mp_iput_on_stream确保接收准备好，获得24%最大增益（log2(size)=4，2进程），弱缩放时增益随大小增加而减少（图25）；KI模型融合到单个内核，块如图26，获得26%最大增益，优于SA（图27）；时间减少见表2。CoMD-CUDA（分子动力学）使用EAM势，力/原子交换重复3次（x,y,z）：2 A内核、同步/发送、2 B内核、接收、2 C内核。SA模型获得25-35%通信增益（图28），总时间增益~5%（通信占7%）；KI模型类似SA，由于高块数（~1700）导致最坏场景。BFS（大规模图搜索）使用2D分解，主循环迭代未知，运行时参数计算，通信大小动态；SA实现无改进，由于同步需求和计算开销增加（CUB 3x时间，内核14%增加）。

负面测试案例。BFS中，SA模型不提供优势，因为迭代数未知需要同步，参数运行时计算要求固定网格/过大CUB，通信大小动态限制Async使用，导致TGPUSA = TGPUS，无C2增益。

结论

GPUDirect Async是一种启用GPU和第三方设备直接控制路径的技术，已在CUDA 8.0中初步发布。InfiniBand网络支持以新实验OFED verbs（Mellanox OFED 4.0）、中层抽象库LibGDSync和样本消息传递库LibMP形式出现。本文呈现的应用在GitHub可用：HPGMG-FV CUDA Async[21]和CoMD-CUDA Async[8]。总结：GPUDirect Async允许多GPU加速应用的新通信模型；GPUDirect Async不一定更快，例如当GPU空闲时间大于通信时间时。CPU可以发布越多连续异步通信期，潜在性能增益越大。主要缺点：通信参数必须在GPU流发布前已知；如果GPU过载或空闲时间小于通信时间，性能可能实际下降。GPUDirect Async仍在开发中。我们正在评估其在其他领域的效能，例如与CUDA多进程服务（MPS）结合，当多个MPI进程共享单个GPU时。正在探索新优化，如在GPU内存分配InfiniBand元素、卸载集体。我们计划分析GPUDirect Async与GPUDirect RDMA的交互，解释如何克服内存一致性问题[19]。到目前为止，我们只使用了Mellanox硬件，因此将探索新类型连接。而且我们希望使用更多节点测试异步应用，以更好地评估其可扩展性。鉴于第3.2节提到的设计选择，我们预计最终会面临可扩展性问题，例如GPU必须轮询O(N) CQ。在某种意义上，我们已经在SA模型中面临轮询的限制。我们目前正在研究改进设计以放松当前约束。我们还计划探索将GPUDirect Async与现代NIC的高级硬件功能结合，如硬件标签匹配和MPI协议卸载。