使用 DPDK 和 GPUdev 在 GPUs上增强内联数据包处理

原创

ssbandjl

修改于 2024-05-03 11:01:41

1320

修改于 2024-05-03 11:01:41

文章被收录于专栏：DPUDPU

2022 年 4 月 28 日, 原作者 Elena Agostini

使用 GPU 进行网络数据包内联处理是一种数据包分析技术，可用于许多不同的应用领域：信号处理、网络安全、信息收集、输入重建等。

这些应用程序类型的主要要求是尽快将接收到的数据包移动到 GPU 内存中，以触发负责对它们执行并行处理的 CUDA 内核。

总体思路是创建一个连续的异步管道，能够将来自网卡的数据包直接接收到 GPU 内存中。您还可以使用 CUDA 内核来处理传入数据包，而无需同步 GPU 和 CPU。

有效的应用程序工作流程涉及使用无锁通信机制在以下组件之间创建协调的连续异步管道：

网络控制器(网卡)，用于向 GPU 内存提供接收到的网络数据包
CPU用于查询网络控制器以获取有关接收到的数据包的信息
GPU接收数据包信息并将其直接处理到 GPU 内存中

图 1 显示了使用 NVIDIA GPU 和 ConnectX 网卡的加速内联数据包处理应用程序的典型数据包工作流程场景。

图 1. 典型的内联数据包处理工作流程场景

在这种情况下，避免延迟至关重要。不同组件之间的通信优化得越多，系统的响应能力就越强，吞吐量也越高。一旦所需的资源可用，每个步骤都必须内联进行，而不会阻塞任何其他等待的组件。

您可以清楚地识别两种不同的流程：

数据流：优化网卡和GPU之间通过PCIe总线的数据（网络数据包）交换。
控制流程：CPU协调GPU和网卡。

数据流

关键是优化网络控制器和 GPU 之间的数据移动（发送或接收数据包）。它可以通过 GPUDirect RDMA(GDR) 技术实现，该技术使用 PCI Express 总线接口的标准功能，在 NVIDIA GPU 和第三方对等设备（例如网卡）之间建立直接数据路径。

GPUDirect RDMA 依赖于 NVIDIA GPU 在 PCI Express 基址寄存器 (BAR) 区域上公开部分设备内存的能力。有关更多信息，请参阅CUDA 工具包文档中的使用 GPUDirect RDMA 开发 Linux 内核模块。现代服务器平台上的 GPUDirect RDMA 基准测试一文对使用不同系统拓扑的标准 IB 动词执行网络操作（发送和接收）时的 GPUDirect RDMA 带宽和延迟进行了更深入的分析。

图 2. NVIDIA GPUDirect RDMA 使用 PCI Express 标准功能在 GPU 和第三方对等设备之间提供直接数据交换路径

要在 Linux 系统上启用 GPUDirect RDMA，需要nvidia-peermem模块（在 CUDA 11.4 及更高版本中提供）。图 3 显示了最大化 GPUDirect RDMA 内部吞吐量的理想系统拓扑：GPU 和 NIC 之间的专用 PCIe 交换机，而不是通过与其他组件共享的系统 PCIe 连接。

图 3. 最大化网络控制器和 GPU 之间的内部数据吞吐量的理想拓扑

控制流

CPU 是协调和同步网络控制器和 GPU 之间活动的主要参与者，用于唤醒 NIC 将数据包接收到 GPU 内存中，并通知 CUDA 工作负载有新数据包可供处理。

在处理 GPU 时，强调 CPU 和 GPU 之间的异步性非常重要。例如，考虑一个简单的应用程序在主循环中执行以下三个步骤：

接收数据包。
处理数据包。
发回修改后的数据包。

在这篇文章中，我介绍了在此类应用程序中实现控制流的四种不同方法，包括优点和缺点。

方法一

图 4 显示了最简单但效率最低的方法：单个 CPU 线程负责接收数据包，启动 CUDA 内核来处理它们，等待 CUDA 内核完成，并将修改后的数据包发送回网络控制器。

图 4. 单 CPU 将数据包传递到 CUDA 内核并等待完成以执行下一步的工作流程

如果数据包处理不是那么密集，则此方法的性能可能比仅使用 CPU 处理数据包而不涉及 GPU 更差(该方案适合密集型数据包)。例如，您可能具有高度并行性来解决数据包上的困难且耗时的算法。

方法2

在这种方法中，应用程序将 CPU 工作负载拆分为两个 CPU 线程：一个用于接收数据包并启动 GPU 处理，另一个用于等待 GPU 处理完成并通过网络传输修改后的数据包（图 5）。

图 5. 拆分 CPU 线程以通过 GPU 处理数据包

这种方法的一个缺点是为每个突发的累积数据包启动一个新的 CUDA 内核。 CPU 必须为每次迭代的 CUDA 内核启动延迟付出代价。如果GPU不堪重负，数据包处理可能无法立即执行，从而导致延迟。(需要协调CPU核GPU之间的处理流程)

方法三

图 6 显示了第三种方法，该方法涉及使用 CUDA 持久内核。

图 6. 使用持久 CUDA 内核的内联数据包处理。

CUDA 持久内核是一个预启动的内核，它正忙于等待来自 CPU 的通知：新数据包已到达并准备好进行处理。当数据包准备好时，内核通知第二个 CPU 线程它可以继续发送它们。

实现此通知系统的最简单方法是使用繁忙等待标志更新机制在 CPU 和 GPU 之间共享一些内存。虽然 GPUDirect RDMA 旨在从第三方设备直接访问 GPU 内存，但您可以使用这些相同的 API 来创建 GPU 内存的完全有效的 CPU 映射。 CPU 驱动的复制的优点是开销较小。现在可以通过GDRCopy库启用此功能。

直接映射 GPU 内存以进行信号传输使得 CPU 可以修改内存，并且轮询期间 GPU 的延迟成本更低。您还可以将该标志放在从 GPU 可见的 CPU 固定内存中，但 CUDA 内核轮询 CPU 内存标志会消耗更多 PCIe 带宽并增加整体延迟。

这种快速解决方案的问题在于它存在风险并且不受 CUDA 编程模型的支持。 GPU 内核无法被抢占。如果编写不正确，持久内核可能会永远循环。此外，长时间运行的持久内核可能会失去与其他 CUDA 内核、CPU 活动、内存分配状态等的同步。

它还拥有 GPU 资源（例如，流式多处理器），这可能不是最佳选择，以防 GPU 确实忙于其他任务。如果您使用 CUDA 持久内核，您确实必须很好地处理您的应用程序。

方法4

最后一种方法是前一种方法的混合解决方案：使用CUDA 流内存操作来等待或更新通知标志，并在 CUDA 流上预启动每组接收到的数据包一个 CUDA 内核。

图 7. 使用模型组合进行内联数据包处理的混合方法

这种方法的不同之处在于，GPU 硬件轮询（使用cuStreamWaitValue）内存标志，而不是阻塞 GPU 流式多处理器，并且仅当数据包准备就绪时才会触发数据包的处理内核。

同样，当处理内核结束时，cuStreamWriteValue通知负责发送的CPU线程数据包已经处理完毕。

这种方法的缺点是，应用程序必须不时地用新的cuStreamWaitValue+ CUDA kernel +cuStreamWriteValue按序填充GPU，以免因未准备好处理更多数据包的空流而浪费执行时间。这里的 CUDA Graph 可能是在流上重新发布的好方法。

不同的方法适合不同的应用模式。

DPDK 和 GPUdev

数据平面开发套件( DPDK) 是一组库，可帮助加速在各种 CPU 架构和不同设备上运行的数据包处理工作负载。

在 DPDK 21.11 中，NVIDIA 引入了一个名为 GPUdev 的新库，在 DPDK 的背景下引入 GPU 的概念，并增强 CPU、网卡和 GPU 之间的对话。 GPUdev 在 DPDK 22.03 中扩展了更多功能。

GPUdev库的目标如下：

介绍由 DPDK 通用库管理的 GPU 设备的概念。
实现基本的 GPU 内存交互，隐藏 GPU 特定的实现细节。
缩小网卡、GPU 设备和 CPU 之间的差距，增强通信。
简化 DPDK 与 GPU 应用程序的集成。
通过通用层公开 GPU 驱动程序特定的功能。

对于 NVIDIA 特定的 GPU，GPUdev 库功能是通过CUDA 驱动程序 DPDK 库在 DPDK 驱动程序级别实现的。要启用NVIDIA GPU 上所有gpudev可用功能，DPDK 必须构建在具有 CUDA 库和 GDRCopy 的系统上。

借助这个新库提供的功能，您可以使用 GPU 轻松实现内联数据包处理，同时处理数据流和控制流。

DPDK 在内存池（一块连续的内存块）中接收数据包。通过以下指令序列，您可以启用 GPUDirect RDMA 在 GPU 内存中分配内存池，并将其注册到设备网络中。

struct rte_pktmbuf_extmem gpu_mem; 
?
gpu_mem.buf_ptr = rte_gpu_mem_alloc(gpu_id, gpu_mem.buf_len, alignment)); 
?
/* Make the GPU memory visible to DPDK, DPDK从GPU中分配内存 */ 
rte_extmem_register(gpu_mem.buf_ptr, gpu_mem.buf_len, 
                            NULL, gpu_mem.buf_iova, NV_GPU_PAGE_SIZE); 
?
/* Create DMA mappings on the NIC */ 
rte_dev_dma_map(rte_eth_devices[PORT_ID].device, gpu_mem.buf_ptr, 
                                   gpu_mem.buf_iova, gpu_mem.buf_len)); 
?
/* Create the actual mempool */ 
struct rte_mempool *mpool = rte_pktmbuf_pool_create_extbuf(... , &gpu_mem, ...);

图 8 显示了 mempool 的结构：

图 8. 用于内联数据包处理的内存池结构

对于控制流，要启用CPU和GPU之间的通知机制，可以使用gpudev通信列表：CPU内存和CUDA内核之间的共享内存结构。列表中的每个项目都可以保存接收到的数据包的地址 ( mbufs) 和一个用于更新处理该项目的状态的标志（准备好数据包、完成处理等）。

struct rte_gpu_comm_list {
    /** DPDK GPU ID that will use the communication list. */
    uint16_t dev_id;
    
    /** List of mbufs populated by the CPU with a set of mbufs. */
    struct rte_mbuf **mbufs;
    
    /** List of packets populated by the CPU with a set of mbufs info. */
    struct rte_gpu_comm_pkt *pkt_list;
    
    /** Number of packets in the list. */
    uint32_t num_pkts;
    
    /** Status of the packets’ list. CPU pointer. */
    enum rte_gpu_comm_list_status *status_h;
    
    /** Status of the packets’ list. GPU pointer. */
    enum rte_gpu_comm_list_status *status_d;
};

伪代码示例：

struct rte_mbuf * rx_mbufs[MAX_MBUFS]; 
?
int item_index = 0; 
?
struct rte_gpu_comm_list *comm_list = rte_gpu_comm_create_list(gpu_id, NUM_ITEMS); 
?
 while(exit_condition) { 
    ... 
    // Receive and accumulate enough packets, 接收并积累足够量的数据包
    nb_rx += rte_eth_rx_burst(port_id, queue_id, &(rx_mbufs[0]), rx_pkts); 
?
    // Populate next item in the communication list, 填充通信列表中的下一项 
    rte_gpu_comm_populate_list_pkts(&(p_v->comm_list[index]), rx_mbufs, nb_rx); 
    ... 
    index++; 
}

为简单起见，假设应用程序遵循 CUDA 持久内核场景( CUDA persistent kernel scenario)，CUDA 内核上的轮询端将类似于以下代码示例：

/* CUDA 持久内核 */
__global__ void cuda_persistent_kernel(struct rte_gpu_comm_list *comm_list, int comm_list_entries) 
{ 
?
    int item_index = 0; 
    uint32_t wait_status; 
?
    /* GPU kernel keeps checking exit condition as it can’t be preempted, GPU 内核不断检查退出条件，因为它无法被抢占 */ 
    while (!exit_condition()) { 
        wait_status = RTE_GPU_VOLATILE(comm_list[item_index].status_d[0]); 
        if (wait_status != RTE_GPU_COMM_LIST_READY) 
            continue; 
?
         if (threadIdx.x < comm_list[item_index]->num_pkts) { 
            /* Each CUDA thread processes a different packet, 每个 CUDA 线程处理不同的数据包 */ 
            packet_processing(comm_list[item_index]->addr, comm_list[item_index]->size, ..); 
        }
?
        __syncthreads(); 
?
         /* Notify packets in the items have been processed, 通知项目中的数据包已被处理*/ 
        if (threadIdx.x == 0) { 
            RTE_GPU_VOLATILE(comm_list[item_index].status_d[0]) = RTE_GPU_COMM_LIST_DONE; 
            __threadfence_system(); 
        } 
?
         /* Wait for new packets on the next communication list entry, 等待下一个通信列表条目上的新数据包 */ 
        item_index = (item_index+1) % comm_list_entries; 
    } 
}

图 9. 持久内核中轮询端的伪代码示例工作流程

NVIDIA 使用 DPDKgpudev库进行内联数据包处理的具体用例是在Aerial 应用程序框架中，用于构建高性能、软件定义的 5G 应用程序。在这种情况下，数据包必须在 GPU 内存中接收并根据 5G 特定的数据包标头重新排序，从而可以在重新排序的有效负载上开始信号处理。

图 10. Aerial 5G 软件中使用 DPDK gpudev进行内联数据包处理用例

l2fwd-nv 应用程序

为了提供如何实现内联数据包处理和使用 DPDK 库的实际示例gpudev，l2fwd-nv示例代码已发布在/NVIDIA/l2fwd-nv GitHub 存储库上。这是普通 DPDK 示例的扩展，l2fwd增强了 GPU 功能。l2fwd-nv程序的作用是接收数据包、交换每个数据包的 MAC 地址（源和目标）并传输修改后的数据包。

L2fwd-nv提供了本文中讨论的所有方法的实现示例以进行比较：

仅CPU
每组数据包的 CUDA 内核
CUDA持久内核
CUDA 图形

作为示例，图 11 显示了具有 DPDKgpudev对象的 CUDA 持久内核的时间线。

图 11. 使用 DPDK gpudev对象的**CUDA 持久内核的示例时间线

为了测量l2fwd-nvDPDKtestpmd数据包生成器的性能，图 12 中使用了两台背对背连接的千兆字节服务器和 CPU：Intel Xeon Gold 6240R、PCIe gen3 专用交换机、Ubuntu 20.04、MOFED 5.4 和 CUDA 11.4。

图 12. 用于测试 l2fwd-nv 性能的两个千兆字节服务器配置

图 13 显示，当对数据包使用 CPU 或 GPU 内存时，峰值 I/O 吞吐量是相同的，因此使用其中一种内存并没有固有的损失。这里的数据包不做任何修改就转发。

图 13. 峰值 I/O 吞吐量相同

为了突出不同 GPU 数据包处理方法之间的差异，图 14 显示了方法 2（每组数据包的 CUDA 内核）和方法 3（CUDA 持久内核）之间的吞吐量比较。两种方法都将数据包大小保持为 1024 字节，在触发 GPU 工作交换数据包的 MAC 地址之前改变累积数据包的数量。

图 14. GPU 数据包处理方法之间的差异

对于这两种方法，每次迭代 16 个数据包会导致控制平面中的交互过多，并且无法实现峰值吞吐量。每次迭代有 32 个数据包，持久内核可以跟上峰值吞吐量，而每次迭代的单独启动仍然有太多的控制平面开销。对于每次迭代 64 和 128 个数据包，两种方法都能够达到峰值 I/O 吞吐量。这里的吞吐量测量不是零丢失数据包。