NVIDIA DOCA GPUNetIO是 NVIDIA DOCA SDK 中的一个库，专门为实时内联 GPU 数据包处理而设计。它结合了GPUDirect RDMA和GPUDirect Async等技术，能够创建以 GPU 为中心的应用程序，其中 CUDA 内核可以直接与网络接口卡（NIC）通信，用于发送和接收数据包，绕过 CPU 并将其排除在关键路径之外。
DOCA GPUNetIO 的核心原理和用途已在前几篇文章《Inline GPU Packet Processing with NVIDIA DOCA GPUNetIO》和《Realizing the Power of Real-Time Network Processing with NVIDIA DOCA GPUNetIO》以及DOCA GPUNetIO 编程指南中进行了讨论。
此前，DOCA GPUNetIO与DOCA Ethernet和DOCA Flow一起，仅限于处理以太网传输层上的数据包传输。随着 DOCA 2.7 的推出，现在有一组扩展的 API使 DOCA GPUNetIO 能够直接从 GPU CUDA 内核使用 RoCE 或 InfiniBand 传输层支持 RDMA 通信。
这篇文章探讨了由 GPU CUDA 内核和 DOCA GPUNetIO 控制的新的远程直接内存访问（RDMA）功能，并对其与性能测试（perftest）微基准的性能进行了比较。
注意，RDMA 首字母缩写描述了一种协议，该协议允许从一台计算机的存储器到另一台计算机存储器的远程直接存储器访问，而不涉及任何一台计算机中的操作系统。操作示例包括 RDMA 写入和 RDMA 读取。不能将其与GPUDirect RDMA混淆，后者与 RDMA 协议无关。GPUDirect RDMA 是 NVIDIA 在 GPUDirect 技术家族中启用的技术之一，使网卡能够绕过 CPU 内存副本和操作系统例程，直接发送或接收访问 GPU 内存的数据。GPUDirect RDMA 可以由任何使用以太网、InfiniBand 或 RoCE 的网络框架启用。
具有 GPUNetIO 的 RDMA GPU 数据路径
RDMA 提供了在两个主机的主内存之间的直接访问，而不涉及操作系统、缓存或存储。这使得数据传输具有高吞吐量、低延迟和低 CPU 利用率。这是通过注册并共享本地内存区域，以便远程主机知道如何访问它。
两个对等方需要通过 RDMA 交换数据的应用程序通常遵循三个基本步骤：
- 步骤 1–本地配置：每个对等端在本地创建 RDMA 队列和内存缓冲区，以便与其他对等端共享这些资源。
- 步骤 2–交换信息： 使用带外（OOB）机制（例如，Linux 套接字），对等端交换有关 RDMA 队列和要远程访问的内存缓冲区的信息。
- 步骤 3–数据路径：两个对等方使用远程内存地址执行 RDMA 读、写、发送和接收，以交换数据。
DOCA RDMA 库按照上面列出的三个步骤通过 InfiniBand 或 RoCE 实现 RDMA 通信，所有这些步骤都是用 CPU 执行的。随着新GPUNetIO RDMA功能的引入，应用程序可以在 GPU 上执行步骤 3，使用 CUDA 内核管理 RDMA 应用程序的数据路径，而步骤 1 和 2 保持不变，因为它们与 GPU 数据路径无关。
将 RDMA 数据路径移动到 GPU 上的好处与以太网用例中的好处相同。在数据处理发生在 GPU 上的网络应用程序中，将网络通信从 CPU 卸载到 GPU，使其能够成为应用程序的主控制器，消除与 CPU 交互所需的额外延迟，知道数据何时准备就绪以及数据位于何处，这也释放了 CPU 周期。此外，GPU 可以同时并行管理多个 RDMA 队列，例如，每个 CUDA 块可以在不同的 RDMA 队列上发布 RDMA 操作。
IB Verbs 和 DOCA GPUNetIO 性能测试
在 DOCA 2.7 中，引入了一个新的 DOCA GPUNetIO RDMA 客户机-服务器代码示例，以显示新 API 的使用情况并评估其正确性。这篇文章分析了 GPUNetIO RDMA 函数与 IB Verbs RDMA 函数之间的性能比较，重现了众所周知的 perftest 套件中的一个微基准。
简而言之，perftest 是一组微基准点，用于使用基本的 RDMA 操作测量 RDMA 带宽（BW）和两个对等点（服务器和客户端）之间的延迟尽管网络控制部分发生在 CPU 中，但可以通过启用 GPUDirect RDMA 并指定--use_cuda标志来指定数据是否驻留在 GPU 内存中。
一般来说，RDMA 写单向 BW 基准测试（即 ib_write_bw）在每个 RDMA 队列上发布一个针对相同大小消息的写请求列表，用于固定迭代次数，并命令 NIC 执行发布的写操作，这就是所谓的“按门铃”程序。为了确保所有写入都已发出，在进入下一次迭代之前，它轮询完成队列，等待每个写入都已正确执行的确认。然后，对于每个消息大小，可以检索发布和轮询所花费的总时间，并以 MB/s 为单位计算 BW。
图 1 显示了 IB 谓词ib_write_bw性能测试主循环。在每次迭代中，CPU 发布一个 RDMA 写入请求列表，命令 NIC 执行这些请求（按门铃），然后等待完成后移动到下一次迭代。启用 CUDA 标志后，要写入的数据包将从 GPU 内存本地获取，而不是从 CPU 内存。
ib_write_bw性能测试主回路实验是用 DOCA 库复制ib_write_bw微基准标记，使用 DOCA RDMA 作为 CPU 上的控制路径以建立客户端-服务器连接，并使用 DOCA GPUNetIO RDMA 作为数据路径，在 CUDA 内核内发布写入。这种比较并不完全一致，因为 perftest 使用 GPUDirect RDMA 来传输数据，但网络通信由 CPU 控制，而 DOCA GPUNetIO 同时使用 GPUDirect RDMA 和 GPUDirect Async 来控制网络通信和来自 GPU 的数据传输。目标是证明 DOCA GPUNetIO RDMA 性能与 IB Verbs 性能测试相当，后者被视为基线。
为了重现ib_write_bw数据路径并测量针对每个消息大小发布 RDMA 写入操作所花费的时间，CPU 记录一个 CUDA 事件，启动rdma_write_bw CUDA 内核，然后记录第二个 CUDA 事件。这应该可以很好地近似 CUDA 内核使用 DOCA GPUNetIO 函数发布 RDMA 写入所用的时间（以毫秒为单位），如下面的代码段 1 所示。
Int msg_sizes[MAX_MSG] = {....};for (int msg_idx = 0; msg_idx < MAX_MSG; msg_idx++) { do_warmup(); cuEventRecord(start_event, stream); rdma_write_bw<<<num_queue, msg_per_size, 0, stream>>>(msg_sizes[msg_idx], …); cuEventRecord(end_event, stream); cuEventSynchronize(end_event); cuEventElapsedTime(&total_ms, start_event, end_event); calculate_result(total_ms, msg_sizes[msg_idx], …)} |
在下面的代码段 2 中，CUDA 内核rdma_write_bw，按照弱模式使用 DOCA GPUNetIO 设备函数，对于给定数量的迭代，并行发布一系列 RDMA 写入，每个 CUDA 块中的 CUDA 线程发布一个写操作，按照弱模式。
__global__ void rdma_write_bw(struct doca_gpu_dev_rdma *rdma_gpu, const int num_iter, const size_t msg_size, const struct doca_gpu_buf_arr *server_local_buf_arr, const struct doca_gpu_buf_arr *server_remote_buf_arr){ struct doca_gpu_buf *remote_buf; struct doca_gpu_buf *local_buf; uint32_t curr_position; uint32_t mask_max_position; doca_gpu_dev_buf_get_buf(server_local_buf_arr, threadIdx.x, &local_buf); doca_gpu_dev_buf_get_buf(server_remote_buf_arr, threadIdx.x, &remote_buf); for (int iter_idx = 0; iter_idx < num_iter; iter_idx++) { doca_gpu_dev_rdma_get_info(rdma_gpu, &curr_position, &mask_max_position); doca_gpu_dev_rdma_write_weak(rdma_gpu, remote_buf, 0, local_buf, 0, msg_size, 0, DOCA_GPU_RDMA_WRITE_FLAG_NONE, (curr_position + threadIdx.x) & mask_max_position); /* Wait all CUDA threads to post their RDMA Write */ __syncthreads(); if (threadIdx.x == 0) { /* Only 1 CUDA thread can commit the writes in the queue to execute them */ doca_gpu_dev_rdma_commit_weak(rdma_gpu, blockDim.x); /* Only 1 CUDA thread can flush the RDMA queue waiting for the actual execution of the writes */ doca_gpu_dev_rdma_flush(rdma_gpu); } __syncthreads(); } return;} |
图 2 描述了代码段 2。在每次迭代时，GPU CUDA 内核并行发布一个 RDMA 写入请求列表，每个 CUDA 块中的 CUDA 线程一个。在同步所有 CUDA 线程后，只有线程 0 命令 NIC 执行写入并等待完成，然后刷新队列，最后再进行下一次迭代。
为了比较性能，为 IB Verbs perftest 和 DOCA GPUNetIO perftest 设置了相同的参数：1 个 RDMA 队列，2048 次迭代，每次迭代执行 512 次 RDMA 写入，并测试消息大小从 64 字节到 4096 字节。
已在 Dell R750 机器上执行基准测试，该机器配备NVIDIA H100 GPU和ConnectX-7网卡（RoCE 模式），通过系统 PCIe 总线连接（无专用 PCIe 交换机）。如下图所示，perftest 的两种实现所花费的总时间是完全可比较的（图 3），以及以 MB/s 为单位报告的峰值 BW（图 4）。由于代码中不同逻辑的性质，时间和 BW 是用不同的方法来测量的，IB Verbs perftest 使用系统时钟，而 DOCA GPUNetIO perftest 则依赖于 CUDA 事件，后者可能具有不同的内部时间测量开销。
请注意，像 perftest 这样的应用程序并不是显示 GPU 利用率优势的最佳工具，因为可实现的并行化量非常低。DOCA GPUNetIO 性能测试 RDMA 写入是以并行方式发布在队列中的（512 个写入，每个写入由不同的 CUDA 线程执行），发布所花费的时间微不足道，约 4 微秒。大部分性能测试时间是由 NIC 实际执行 RDMA 写入、通过网络发送数据和返回正反馈所花费。
这个实验可以被认为是成功的，因为它证明了使用 DOCA GPUNetIO RDMA API 与使用常规 IB Verbs 相比不会引入任何相关开销，并且在运行相同类型的工作负载和工作流时可以达到性能目标。ISV 开发人员和最终用户可以使用 DOCA GPUNetIO RDMA，获得 GPUDirect 异步技术将通信控制卸载到 GPU 的好处。
这种架构选择提供了以下好处：
- 一个可扩展性更强的应用程序，能够同时并行管理多个 RDMA 队列（通常每个 CUDA 块一个队列）。
- 利用 GPU 提供的高度并行性的能力，该 GPU 拥有多个 CUDA 线程在不同数据上并行工作，并能以尽可能低的延迟在同一队列上发布 RDMA 操作。
- 更少的内部总线事务（例如 PCIe），因为不需要同步 GPU 上的工作与 CPU 活动。CPU 不再负责发送或接收 GPU 必须处理的数据。
准备好深入了解 DOCA GPUNetIO 了吗？查看官方 DOCA GPUNetIO 程序员指南 以获取全面的见解和技术细节。
要了解 DOCA GPUNetIO 的更多功能，请参阅使用 NVIDIA DOCA GPUNetIO 的内联 GPU 数据包处理。要了解 DOCA GPUNetIO 如何彻底改变 GPU 通信，请参阅使用 NVIDIA DOCA GPUNetIO 实现实时网络处理的强大功能。
