GPU-centric Communication on NVIDIA GPU Clusters with InfiniBand: A Case Study with OpenSHMEM

发表时间: 2017-12 · HiPC'17 (IEEE 24th International Conference on High Performance Computing)

作者/机构: S. Potluri, A. Goswami, D. Rossetti and C. J. Newburn (NVIDIA Corporation), M. Gorentla Venkata and N. Imam (Oak Ridge National Laboratory)

A1 主要贡献

本文旨在解决在大型GPU集群中，传统MPI+CUDA编程模型下CPU和GPU之间交替进行计算和通信所带来的开销和资源利用率不足的问题。传统的模型中，计算和通信阶段是分离的，这会导致重复的内核启动开销，并且在通信期间GPU资源闲置，在计算期间网络资源闲置。尽管可以通过代码重构来重叠计算与通信，但这会增加代码复杂性，且随着每个GPU处理的问题规模变小，其收益也会减弱。

为了克服这些限制，本文的研究目标是探索在GPU内部直接发起通信的能力，从而实现长时间运行的内核，提高GPU利用率，并避免启动和同步开销。具体来说，本文基于先前为PCIe或NVLink互连GPU开发的NVSHMEM库（一个OpenSHMEM的实现），将其扩展到支持通过InfiniBand（IB）进行设备端通信。

本文的主要贡献如下：

  • 设计了GPU端的IB verbs：在CUDA中实现了IB verbs，并深入探讨了多种设计方案。设计中特别考虑了NVIDIA GPU的特性，如松散内存模型、自动内存访问合并以及线程层次结构，并解决了由此产生的正确性问题。
  • 基于GPU端IB verbs设计了OpenSHMEM：将所提出的GPU端IB verbs用于实现OpenSHMEM运行时，并讨论了如何透明地或通过API扩展利用批量化（batching）和模板化（templating）等优化。
  • 详尽的性能评估：通过微基准测试和2D-stencil小程序，对设计进行了详细评估。微基准测试表明，单个NVIDIA Pascal P100 GPU仅用其56个流式多处理器（SM）中的一到两个就能饱和网络带宽，单GPU单IB EDR适配器可实现约9000万条消息/秒的吞吐量。
  • 应用级性能验证：将基于NVSHMEM的2D-stencil应用与使用GPUDirect RDMA的CUDA-aware MPI实现进行比较，在2到4个节点上，对于足以填满P100 GPU占用率的输入规模，观察到了23%到42%的性能提升。这表明，当所有计算都在GPU上运行时，从通信路径中移除CPU可以带来显著的性能增益。

A3 背景知识

A. InfiniBand Verbs

IB verbs是用于对InfiniBand（IB）适配器进行编程的用户模式API。本文主要关注用于发起IB数据移动操作和轮询其完成的API。作者在CUDA中实现了这两种API的原型，以便在NVIDIA GPU上运行的内核可以直接向IB适配器发出数据传输操作并轮询其完成。IB的用户级协议是围绕队列定义的。在InfiniBand上发起通信操作，首先需要在网络接口卡（NIC）的DMA引擎可访问的内存中准备一个工作队列元素（Work Queue Element, WQE）。然后，CPU向DMA引擎发送一个消息以表明此意图。其中有一个中间步骤，即将指向命令的指针写入一个称为“门铃记录”（doorbell record）的缓冲区中。如果由于拥塞或冲突导致“敲门铃”（ringing of the doorbell）操作被错过，适配器会使用这个记录来访问命令。图1以图形方式展示了这些步骤。这意味着，在具有松散内存模型的系统上，必须在这三个步骤之间放置适当的内存屏障以确保顺序。通信操作的完成是通过在完成队列（Completion Queue, CQ）上轮询来检测的。完成一个操作后，网络适配器会生成一个完成队列条目（Completion Queue Entry, CQE）来提供WQE的状态。本文为CUDA内核提供了IB verbs的实现。

图1. 通过IB verbs向网络适配器发布命令的步骤
图1. 通过IB verbs向网络适配器发布命令的步骤

B. OpenSHMEM

OpenSHMEM是一个分区全局地址空间（PGAS）库接口规范。它提供了一个PGAS视角的执行上下文和内存模型。OpenSHMEM中的执行上下文是一个由称为处理单元（Processing Element, PE）的整数标识的操作系统进程。一个OpenSHMEM程序每个PE拥有私有地址空间，并在所有PE之间拥有共享地址空间。OpenSHMEM中的共享地址空间表现为对称对象，可被OpenSHMEM程序中的所有PE访问。对称对象通过OpenSHMEM中的一个集体分配操作进行分配。OpenSHMEM提供了与其他PE进行通信和同步的例程。OpenSHMEM中的通信主要是一边倒（one-sided）的。它提供了Put、Get和原子操作的多种变体，用于访问和修改位于远程PE上的对称数据对象。它还提供了排序和集体通信API。

A2 方法细节

III. 使用OPENSHMEM进行GPU发起的通信

NVSHMEM提出了一种利用GPU发起的通信的编程方法，其动机源于GPU面向吞吐量的架构。GPU被设计为支持成千上万的线程以实现最大的并行吞吐量。这些线程极其轻量级，成千上万的线程（以称为warp的组）排队等待工作。当一个warp等待内存访问时，另一个warp会取而代之开始执行。由于为所有活动线程分配了独立的寄存器集，因此无需交换寄存器或状态。因此，这种执行模型具有内在的延迟隐藏能力和最小的调度开销。随着并行度的增加，GPU架构可以拥有足够的状态来隐藏不仅是到本地GPU设备内存的延迟，还包括通过网络到远程GPU内存的延迟。GPU发起的通信可以利用GPU硬件的这种内在能力。此外，这也提高了可编程性，因为开发人员不必依赖混合模型来协调和重叠应用程序的不同阶段。从CUDA内核内部发起并同步的通信不仅减少了对CPU的依赖，还避免了限制强扩展性的现有同步开销。使用像OpenSHMEM这样的PGAS编程模型更适合GPU发起的通信。OpenSHMEM的一边倒通信模型仅要求接口的调用方（源端）处于活动状态，这与GPU上的大规模并行和动态调度模型相匹配。

IV. 设计与实现

A. GPU端IB verbs：

本文在先前研究的基础上，深入探讨了GPU特性对IB verbs性能的影响。其他研究人员如Oden等人【6, 2014 IEEE International Parallel & Distributed Processing Symposium Workshops, Phoenix, AZ, USA, May 19-23, 2014. IEEE Computer Society, 2014】和Daoud等人【7, Gpurdma: Gpu-side library for high performance networking from gpu kernels, 2016, ROSS ’16】已对GPU的IB verbs进行了原型实现。本文的工作更进一步，讨论了GPU的warp级合并和松散GPU内存模型等特性如何影响其性能。同时，本文还讨论了如何使用GPU端verbs来实现OpenSHMEM通信模型，并解释了如何在实现类似OpenSHMEM的模型时透明地使用模板化（templating）和批处理（batching）等性能优化。此外，本文考虑了先前工作忽略的接收端一致性问题，并在SHMEM的背景下解决了它。最后，本文介绍了使用GPU端IB verbs实现2D-stencil小程序的过程，该程序是典型迭代式HPC应用的代表。关于IB verbs的背景知识和向网络适配器发出IB命令的操作序列在第二节A部分有详细解释。

Warp级合并。在CUDA中，线程被分组为warp，这不仅与计算相关，也与内存访问相关。当一个warp内的线程写入内存中的连续地址时，这些访问会自动合并为更大的事务（通常等于缓存行大小）。无论是访问GPU全局内存、CPU/主机内存还是对等PCIe设备的内存，都会发生这种情况。在IB中发出操作涉及准备一个命令并将其写入命令队列（图1中的步骤1）。一种优化是将数据与命令内联，并作为上述操作的一部分将其写入命令队列。可内联的数据大小存在网络适配器特定的限制。另一种可能的优化是在“敲门铃”（图1中的步骤3）时。这涉及将完整的命令写入门铃寄存器（而不仅仅是写入包含命令位置的前8个字节）。这可以省去适配器从内存中读取命令的步骤。将命令写入队列的操作将涉及对主机或GPU内存的写入，具体取决于命令队列的创建位置。写入门铃则涉及通过PCIe向网络适配器上的寄存器进行写入。在这两种情况下，同一warp内的线程可以被用来集体发出写操作，从而生成单个到内存或通过PCIe的事务。

为利用GPU并行性与合并能力而扩展的SHMEM API。当前的OpenSHMEM标准只提供可由单个执行线程调用的API。为了利用GPU上的并行性并从合并能力中受益，作者通过向SHMEM API添加一个线程组参数来扩展它们。例如，修改后的shmem_put API如下所示：

void shmemx_put(TYPE *dest, const TYPE *source, size_t nelems, int pe, int thread_group);

线程组可以是特定于线程模型的。在CUDA的上下文中，它们可以是THREAD、WARP、CTA或GRID。作者期望设备端通信API在WARP或更高粒度上被调用。这将允许在SHMEM中利用如上所述的合并能力。在文中的示例中，作者使用了WARP级的SHMEM操作。

GPU内存模型的影响。NVIDIA GPU具有高度松散的内存模型，这意味着由CUDA线程发出的读或写操作可能会被重排序，除非它们是针对同一地址，或者存在适当的数据/控制依赖，或者存在适当的内存屏障。PTX（NVIDIA GPU的并行线程执行ISA）中的内存屏障有一个作用域（sys, gl, cta）限定符，指定了由屏障分隔的操作将按顺序对哪些实体可见。例如，membar.sys保证排序对GPU中的所有其他线程以及所有客户端（包括通过PCIe与GPU交互的客户端，如本文中的IB适配器）都可见。发布一个IB操作涉及三次写操作（图1中的步骤1、2和3），这些操作需要按顺序对网络适配器可见才能正确执行。如第二节A部分所述，所有这些写操作都发生在不同的内存位置：命令队列、门铃记录和门铃寄存器。CUDA线程必须插入内存屏障来保证它们的顺序。图2中的表格显示了根据命令队列和门铃记录所在的内存位置所需的内存屏障。当命令队列和门铃记录都在主机内存中时，需要在更新这些位置之间以及更新门铃之前使用membar.sys来保证顺序。对于下一种组合，当命令队列在全局内存中时，一个membar.gl就足以替代membar1。这是基于membar.gl保证了全局内存更新相对于GPU上的PCIe端点的可见性顺序。任何来自NIC到全局内存的访问都由这个PCIe端点服务，因此会观察到相同的顺序。然而，这并非ISA规范所保证，并可能在未来的架构中改变。当门铃记录在主机内存中时，更新之后需要一个membar.sys。在这种情况下，在“敲门铃”之后需要一个membar.sys，以防止它与来自下一个事务的对门铃记录的写入发生重排序。这可能导致下一个命令在完全写入之前就被提取。当命令队列和门铃记录都在GPU内存中时，membar.gl足以替代membar1和membar2。

图2. 根据命令队列和门铃记录的位置所需的内存屏障
图2. 根据命令队列和门铃记录的位置所需的内存屏障

命令队列和门铃记录的理想位置取决于通信的发起方。作者期望OpenSHMEM运行时可以为GPU使用维护独立的端点资源。近期已有提案建议在OpenSHMEM中设置多个通信上下文，每个上下文具有不同的使用特性（例如，单线程与多线程）。一个给定的上下文将用于主机还是设备，可以成为一个新的参数，允许运行时适当地分配IB资源。

完成轮询。虽然Daoud等人在其早期工作中【7, Gpurdma: Gpu-side library for high performance networking from gpu kernels, 2016, ROSS ’16】实验了将数据结构放置在主机或设备内存上，但他们没有比较使用CPU线程轮询完成与让GPU线程在CUDA内核内轮询完成的情况。作者实现了这些替代方案，并在后续章节中比较了它们的性能。这些设计选择在编程模型方面对用户是透明的。即使使用CPU轮询IB完成，GPU上的线程仍然可以从CUDA内核内部查询SHMEM通信API的完成情况（例如，通过调用shmem_quiet）。在这种情况下，CPU将充当代理，轮询完成并在完成后通知GPU线程。

图3. 有批处理和无批处理的操作序列
图3. 有批处理和无批处理的操作序列

批量提交（Batched Posts）。尽管将IB数据结构放置在设备内存中可以降低内存屏障的强度，但它们仍然占据了在GPU上发起IB通信操作时间的相当一部分。这一点在下一节的性能数据中有所说明。为了减少内存屏障对吞吐量的影响，作者实现了一种在向IB适配器提交之前批量处理多个操作的设计。图3显示了不使用批处理和使用批处理发出一系列操作的区别。它减少了所需的内存屏障数量。假设操作发生在同一个可靠的端点上，用最后一个操作的值更新门铃记录并敲响门铃，也会触发之前未发出的操作。这也减少了对门铃记录和门铃寄存器的写入次数。减少写入带来的好处在Daoud等人的早期工作中已经提出【7, Gpurdma: Gpu-side library for high performance networking from gpu kernels, 2016, ROSS ’16】。具体来说，第五节的表I和表II中的“QP on GPU memory + no inline + thread”配置紧密代表了Daoud等人的工作。在这项工作中，作者还考虑了其他可能的情况以及给定情况下对内存屏障的影响。

在OpenSHMEM中透明应用批量优化。OpenSHMEM应用程序通常会在阻塞等待完成（使用shmem_quiet）之前发出多个操作。这使得上述优化可以在OpenSHMEM运行时中透明地应用。当调用nvshmem_quiet时，所有待处理的操作都必须被发出并轮询。

模板化操作（Templated Operations）。由于GPU核心的串行性能较低，对于小消息而言，发出操作的成本占总通信时间的相当一部分。如果可以提前准备命令队列，这个成本可以被降低。然而，这需要知道本地和远程内存地址、数据传输大小等信息。在实现像屏障（barriers）这样的同步操作时，这种优化是可能的，因为它们可以使用内部内存并重复执行相同的操作。作者在2D-stencil小程序的实现中应用了这一点，稍后会解释。虽然这不能应用于当前标准中的OpenSHMEM通信API，但允许持久性通信句柄（persistent communication handles）的扩展可以从中受益。持久性句柄在固定网格的迭代应用中很有用，其中每个进程在每次迭代中都从相同的内存位置向相同的进程集发送数据。

目标端一致性（Target-side Consistency）。对于一个已经在运行的GPU内核中的线程，不能保证它们能以特定的顺序和在特定的时间间隔内看到由CPU或其他设备对内存所做的更新。CUDA没有提供强制执行这种一致性的方法，作者称之为目标端一致性。这个问题在Daoud等人的早期工作中被强调【7, Gpurdma: Gpu-side library for high performance networking from gpu kernels, 2016, ROSS ’16】，但作者没有提供解决方案。本文通过确保通信的完成和接收数据的使用跨越一个内核边界来处理这个一致性问题。例如，在2D-stencil的实现中，每次迭代结束时执行一个shmem_barrier，它保证了在返回之前发出的所有通信操作都已完成。这也标志着一个CUDA内核的结束。接收到的数据在下一次迭代的内核启动中被访问。这次启动提供了所需的一致性保证。每次迭代启动一个内核并不一定意味着每次迭代结束时都必须与CPU进行同步。一系列迭代（内核启动）可以排队到一个CUDA流上，相对于CPU异步执行。CUDA内核启动有相关的软件和硬件开销。将多个内核排队到一个流上可以重叠这些阶段，从而提供延迟隐藏。需要进行收敛检查的应用程序将不得不推测性地启动一个或多个迭代，以便从排队中受益。这些好处将在下一节中更详细地解释。

B. 2D stencil小程序

本文实现了两个版本的2D stencil小程序，一个使用MPI，另一个使用基于GPU Verbs的SHMEM。在MPI版本中，所有通信都由CPU通过MPI库发起。而在SHMEM版本中，通信直接由GPU使用GPU Verbs发起。这避免了对于内核执行时间短的小输入尺寸至关重要的内核启动开销。此外，也避免了为确保数据一致性而在CPU上的MPI库与由CUDA内核在GPU内存中产生的数据之间进行的必要同步。

图4. 1D进程网格布局
图4. 1D进程网格布局
图5. 2D进程网格布局
图5. 2D进程网格布局

1) 数据分区：每个rank或PE获得一个数据方块，在所有rank或PE之间平均分配。进程以所有可能的一维和二维布局排列。图4显示了可能的一维网格布局（未显示转置网格布局）。在这种情况下，当有3个或更多进程时，任何rank最多与两个邻居通信。二维网格布局仅在2个或更多节点上才可能，如图5所示（未显示8个节点的转置布局）。在这种情况下，位于角落的进程与另外两个邻居通信，位于边界的进程与三个邻居通信，而内部的进程与四个邻居通信。因此，每个进程的通信量不会扩展。

2) 数据传输的粒度：每个CTA（Cooperative Thread Array）被分配一部分边界区域进行处理。一旦边界计算完成，它需要将数据发送到远程PE。边界交换的粒度可以根据应用变化。已有提案允许用户为SHMEM远程内存访问API指定同步粒度。例如，一个CTA可以选择每次在边界上滑动到新位置时执行一次put操作。如果多个warp参与数据传输，这需要CTA内部的同步，可以通过CUDA中的__syncthreads()原语轻松实现。然而，从GPU发布许多小尺寸消息的开销可能会损害性能。或者，CTA可以使用批处理来避免多次小写入的开销，并将其分配的整个边界区域作为单个消息发送。最后，一个CTA可以等到所有其他CTA完成各自边界段的计算，然后将所有发往特定邻居的消息合并成一个由其中一个CTA发出的单一消息。如果有多个邻居需要发送数据，可以选择一些CTA并分配它们向特定邻居传输的任务。

A4 实验环境

  • 微基准测试平台:

    • 硬件: 两台配备14核Intel Xeon Broadwell (E5-2680v4) CPU的机器，搭载NVIDIA Tesla P100 GPU，通过Mellanox ConnectX-4 MT27700 EDR InfiniBand (IB)适配器连接。NIC和GPU通过一个PLX 9797 PCIe交换机连接，绕过PCIe主桥，使得GPU和NIC之间可用的数据传输带宽达到12 GB/s。
    • 软件: Ubuntu 14.04，CUDA驱动版本375.66，OFED版本3.4。

  • 2D-stencil小程序平台:

    • 硬件: 一个八节点的集群，每台机器配备16核Intel Xeon Haswell (E5-2698v3) CPU，搭载NVIDIA Tesla P100 GPU，通过Mellanox Connect-IB MT27600 FDR InfiniBand (IB)适配器连接。
    • 软件: CentOS 7.3，CUDA驱动版本375.39，OFED版本4.0。
    • 基线对比: 使用支持GPUDirect RDMA的CUDA-aware MPI实现MVAPICH2-GDR 2.2。

A4 实验结果

评估分为两部分：微基准测试和2D stencil小程序应用级性能测试。

A. 微基准测试

  1. 延迟（Latency）:

    • 提交延迟 (Post Latency): 测量提交一个RDMA写操作的延迟（对应shmem_put_nbi）。如表I所示，最低延迟为2.26 µs，在（1）QP（队列对）位于GPU内存中和（2）使用一个warp（32个线程）执行提交时获得。Warp版本因并行写入64字节的WQE而比单线程版本快。QP在GPU内存中时，避免了较慢的PCIe访问，延迟更低。
    • 乒乓延迟 (Ping-pong Latency): 测量一次完整的往返通信延迟的一半。如表II所示，最低延迟为3.16 µs，在（1）QP位于GPU内存，（2）数据和门铃都内联（inlined），以及（3）使用一个warp提交写入时观察到。内联数据和门铃避免了NIC从内存中获取数据和命令的延迟。
    • Ping延迟: 测量提交RDMA写并轮询其完成的总时间（对应shmem_put后跟shmem_quiet）。如图6所示，最低延迟为4.97 µs（8字节消息），在（1）QP和CQ（完成队列）都位于GPU内存，（2）由GPU线程轮询完成，（3）数据和门铃都内联，以及（4）使用一个warp提交时获得。
    • 延迟分解: 表III显示，通过放宽运行时安全检查（如流控制和错误处理），可以进一步降低延迟。表IV显示，将WQE的准备（prepare）和提交（post）分离可以避免一个内存屏障，从而优化延迟。

    表I. GPU发起的RDMA写操作的提交延迟
    表II. GPU发起的RDMA写操作的乒乓延迟
    图6. 使用不同QP/CQ位置和轮询代理的Ping延迟
    表III. 放宽运行时安全检查时的Ping延迟分解
    表IV. 准备、提交和轮询的Ping延迟分解

  2. 吞吐量和带宽 (Throughput and Bandwidth):

    • 多Warp/线程的影响:

      • 带宽: 如图7所示，单个warp需要128KB或更大的消息才能饱和12 GB/s的带宽。相比之下，使用一个SM（流式多处理器）上的所有warp，仅用4KB大小的消息就能饱和带宽，与CPU线程相当。这表明仅需P100 GPU（共56个SM）的一小部分计算资源即可饱和网络。
      • 吞吐量: 如图8所示，单线程/warp的峰值吞吐量远低于CPU。通过增加warp数量可以提高吞吐量。当使用P100上每个SM的16个warp时，小消息（<=128字节）的最高消息率达到约1600万条/秒。

    • 批量提交的影响:

      • 吞吐量: 如图8所示，批量提交显著提高了吞吐量，峰值消息率达到约9000万条/秒。消息率在8字节后下降是由于实现中的内联数据大小限制，可以修复。

    图7. 使用不同数量QP从GPU获得的带宽
    图8. 不同配置下的吞吐量（消息数/秒）

B. 2D-stencil小程序

通过改变每个GPU上的矩阵大小来模拟强扩展性实验，并将基于GPU Verbs的SHMEM版本与基于MVAPICH2-GDR的MPI版本进行比较。

  • 实现差异: MPI版本使用5个不同的内核来交错CPU通信和GPU计算/打包，而GPU Verbs版本使用一个融合了计算和通信的内核。由于资源消耗增加，GPU Verbs版本的内核占用率（224个CTA）低于MPI版本（448个CTA）。
  • 性能结果: 如图9所示，GPU Verbs版本的性能表现出三个阶段：
    1. 小问题规模（低于占用率）: 当每个GPU的矩阵大小不足以填满占用率时（例如，<= 512x512），GPU Verbs版本比MPI版本快24%到42%。这是因为避免了多次内核启动的开销，这在内核执行时间短的情况下至关重要。
    2. 中等问题规模（接近占用率）: 在1Kx1K到2Kx2K的矩阵大小范围内，GPU Verbs版本的性能比MPI版本低约10%。这是由于其较低的内核占用率，导致并发运行的CTA数量少于MPI版本。
    3. 大问题规模（计算密集型）: 从4Kx4K开始，计算时间占主导地位，两个版本的性能趋于一致。GPU Verbs版本因其融合内核带来更好的数据重用而略微领先约7%，并一直保持到P100 GPU内存所能容纳的最大输入尺寸32Kx32K。这表明即使GPU Verbs主要为小尺寸优化，在处理大尺寸问题时，为支持核内通信而进行的内核重写也未引入性能开销。
图9. 使用GPU Verbs与使用GDR MVAPICH的2D stencil加速比
图9. 使用GPU Verbs与使用GDR MVAPICH的2D stencil加速比

A7 补充细节

VI. 相关工作

先前工作主要集中在减少CPU参与的开销，但CPU仍是通信的发起者。NVIDIA的GPUDirect技术【8, Gpudirect, 2015】【9, Gpudirect rdma, 2015】通过消除数据到CPU内存的拷贝来减少通信开销，但CPU仍然负责发起通信和同步。GPUDirect Async【10, Gpudirect: Integrating the gpu with a network interface, 2015】允许GPU按顺序触发和同步网络操作与计算内核，但通信操作仅在内核边界调度，导致在强扩展应用中开销仍然显著。本文则探索在CUDA内核内部进行通信作为替代方法。

已有多项工作致力于在MPI和PGAS模型中高效使用GPU。Wang等人【11, Mvapich2-gpu: Optimized gpu to gpu communication for infiniband clusters, 2011】和Potluri等人【12, Efficient inter-node mpi communication using gpudirect rdma for infiniband clusters with nvidia gpus, 2013】在MVAPICH2【13, MVAPICH】中提出了CUDA-aware MPI，允许使用标准MPI接口移动GPU数据。Aji等人【14, MPI-ACC: An Integrated and Extensible Approach to Data Movement in Accelerator-Based Systems, 2012】提出使用数据类型属性来识别GPU内存。Potluri等人【15, Extending openshmem for gpu computing, 2013】为OpenSHMEM提出了CUDA-aware方法。这些工作都处理CPU发起的、涉及GPU设备内存的通信。Cunningham等人【16, Gpu programming in a high level language: Compiling x10 to cuda, 2011】提出了使用X10语言的APGAS模型。Miyoshi等人【17, Flat: A gpu programming framework to provide embedded mpi, 2012】的FLAT GPU框架允许在GPU内核中嵌入MPI调用。本文的方法则让GPU用户在熟悉的CUDA模型中，使用OpenSHMEM来表达GPU间的数据移动。

已有研究人员实现了GPU的IB verbs原型。Oden等人【6, 2014 IEEE International Parallel & Distributed Processing Symposium Workshops, Phoenix, AZ, USA, May 19-23, 2014. IEEE Computer Society, 2014】观察到IB verbs不适合GPU发起的通信。Daoud等人【7, Gpurdma: Gpu-side library for high performance networking from gpu kernels, 2016, ROSS ’16】展示了如何规避开销并用GPU发起的IB verbs实现更好性能。本文的工作通过考虑GPU的特性和功能（如内存模型和合并）来推动现有技术的发展。作者讨论了如何将这些设计用于OpenSHMEM通信模型，并移植了一个对延迟敏感且代表一类常见HPC应用的2D-Stencil小程序。

与dCUDA的对比。Gysi等人【18, dCUDA: Hardware Supported Overlap of Computation and Communication, 2016】提出了dCUDA，一个通过CPU卸载（CPU on-loading）提供GPU端通信的框架。与此不同，本文专注于IB verbs的GPU原生实现及其性能。

A5 结论

本文基于先前在PCIe或NVLink上实现GPU端OpenSHMEM的工作，将这一能力扩展到了使用Mellanox InfiniBand适配器的NVIDIA GPU集群。作者在CUDA中实现了IB verbs，并评估了多种设计方案，重点考虑了GPU的松散内存模型、自动内存访问合并和线程层次结构，并解决了由此产生的正确性问题。这些设计被透明地或通过API扩展应用到NVSHMEM中。

通过微基准测试和2D-stencil小程序的详细评估，本文展示了显著的性能成果。在单个P100 GPU和单个IB EDR适配器上，实现了约9000万条消息/秒的吞吐量。在2D-stencil小程序上，与基于CUDA-aware MPI的实现相比，获得了23%到42%的性能加速。这些结果表明，通过在GPU内部发起通信来消除CPU在通信路径中的作用，可以为完全在GPU上运行的计算带来显著的性能增益。