NVIDIA DOCA GPUNetIO Programming Guide

发表时间: 2023-10 · NVIDIA Programming Guide v2.2.0

速读

一句话结论 本文介绍了 NVIDIA 的 DOCA GPUNetIO 库，通过构建以 GPU 为中心的网络数据包处理架构，将 CPU 从关键路径中移除，从而实现了低延迟、零丢包的实时 GPU 网络数据包收发。

要解决什么问题 传统的网络数据包处理采用以 CPU 为中心的方法，CPU 始终处于关键路径上。具体而言，CPU 需要负责协调网卡（NIC）将数据包通过 GPUDirect RDMA 技术接收到 GPU 内存中，随后还要负责通知在 GPU 上等待新数据包的 CUDA 核心。这种机制的致命卡点在于，当系统处于低功耗平台，或者网络中的客户端数量大幅增加时，CPU 的调度和通知开销会迅速增加，极易成为整个系统的性能瓶颈。一旦 CPU 发生阻塞，GPU 的强大算力就无法得到充分发挥，导致系统无法实现以最低延迟达到零丢包吞吐量的目标。因此，如何将 CPU 从实时数据包处理的关键路径中彻底移除，是提升 GPU 网络应用性能的核心问题。

怎么做的 核心思路是实现以 GPU 为中心的解决方案，让 CUDA 内核直接控制网卡进行数据包的收发，完全绕开 CPU 的干预。整个应用被划分为两个阶段：CPU 仅负责初始化的设置阶段（包括设备配置、内存分配、导出对象的 GPU 句柄并启动 CUDA 内核），而主数据路径则完全由 GPU 和网卡直接交互。该方法由几个关键部件构成。首先是 GPUDirect 异步内核启动网络通信机制，它允许 CUDA 内核直接调用设备函数与网卡交互。其次是 GPUDirect RDMA，负责将数据包直接写入连续的 GPU 内存区域。为了管理这些内存，系统设计了智能内存分配机制，支持分配四种不同可见性的内存类型（如仅 GPU 可见的 DOCA_GPU_MEM_GPU 或 CPU 与 GPU 共享的 DOCA_GPU_MEM_CPU_GPU）。在数据接收端，GPU 端的接收函数会将数据包存入一个循环的缓冲区数组中。若函数返回接收数量 $N$ 和起始索引 $X$，则表示数组中范围 $[X, X + N - 1]$ 内的缓冲区已被连续填满。在数据发送端，提供了强模式和弱模式的发送函数，并支持精确发送调度，允许在发送队列中设置基于时间戳的等待屏障。为了在不同的 CUDA 内核之间（或 GPU 与 CPU 之间）同步状态，方法引入了标准化的信号量机制。信号量采用轮询与更新机制，更新方填充数据包信息并将状态标志设为就绪，轮询方等待就绪信号、处理数据后将状态设为完成。在典型的双核流水线模型中，一个 CUDA 内核专门负责从网卡接收数据包，并通过信号量将数据包的内存索引和数量分派给第二个 CUDA 内核，第二个内核读取信号量后直接在 GPU 显存中完成实时处理，全程无需 CPU 参与。

效果如何 实验在支持 GPUDirect RDMA 的硬件拓扑上搭建，要求 GPU 和网卡（如 ConnectX-6 Dx 或 DPU 融合卡）之间具有专用的 PCIe 连接，并在 Ubuntu 系统下基于 CUDA 12.1 和 DOCA 库运行。由于本文档为编程指南，实验部分主要用于验证核心功能而非进行吞吐量跑分，因此未提供与其他路线的基线对比。在量化功能测试中，示例程序生成了大小为 1kB 的原始以太网数据包，重点验证了精确发送调度功能。程序成功实现了基于时间戳的批量发送，共发送 8 个批次、每批 32 个数据包，并能通过参数精确控制每批之间的发送时间间隔。该方案也存在明确的局限性和使用代价：首先，它对运行环境要求苛刻，必须禁用 IOMMU，且不支持 KVM 等虚拟化环境，仅建议在裸金属或 Docker 中运行；其次，硬件拓扑必须保证 GPU 和网卡间的直接通信，必要时需禁用 PCIe 访问控制服务，且网卡固件有严格的版本门槛；最后，在软件工程层面，框架不提供自动的锁机制，开发者必须自行管理循环缓冲区的大小和消费速度，以防止未处理的数据包被新流入的数据覆盖。

A1 主要贡献

本文介绍了一种用于实时GPU处理网络数据包的技术，该技术适用于信号处理、网络安全、信息收集和输入重建等应用领域。传统的CPU中心（CPU-centric）方法中，CPU处于关键路径上，负责协调网卡（NIC）将数据包接收到GPU内存（通过GPUDirect RDMA），并通知在GPU上等待新数据包的CUDA核心。在低功耗平台或客户端数量增加时，CPU很容易成为瓶颈，这限制了GPU的性能发挥，导致无法实现以最低延迟达到零丢包吞吐量的目标。

为解决此问题，本文档介绍了新的DOCA GPUNetIO库，它旨在通过实现GPU中心（GPU-centric）的解决方案来优化性能，将CPU从关键路径中移除。

CPU中心方法：

CPU中心方法
CPU中心方法

GPU中心方法：

GPU中心方法
GPU中心方法

DOCA GPUNetIO通过以下特性实现GPU中心的解决方案：

  • GPUDirect异步内核启动网络（GDAKIN）通信：允许CUDA内核调用GPUNetIO设备函数来直接与NIC交互进行收发，无需CPU干预关键路径。
  • GPUDirect RDMA：将数据包直接接收到连续的GPU内存区域。
  • 信号量（Semaphores）：提供接收实体和进行实时数据包处理的CUDA内核之间的标准化I/O通信协议。
  • 智能内存分配：分配对齐的GPU内存缓冲区，并使其可被CPU直接访问。
  • 库的结合：集成了CUDA和DPDK的gpudev库（该库需要GDRCopy库），这些已内嵌在DOCA发布的DPDK中。
  • GPU上的以太网协议管理。

NVIDIA的Morpheus和Aerial 5G SDK等应用已在积极使用DOCA GPUNetIO。更深入的技术和动机介绍可参考NVIDIA博客文章《Inline GPU Packet Processing with NVIDIA DOCA GPUNetIO》。

A3 背景知识与环境配置

先决条件

软件与安装：DOCA GPUNetIO目前仅包含在适用于Ubuntu 20.04和Ubuntu 22.04的DOCA for Host软件包中，且要求CUDA 12.1或更高版本。安装所有GPUNetIO组件的命令为：apt install -y doca-gpu doca-gpu-dev。

系统配置：必须禁用IOMMU。可以通过修改grub命令行GRUB_CMDLINE_LINUX_DEFAULT="iommu=off"来显式禁用。此外，该库已在裸金属和Docker环境中测试，不建议在虚拟化环境（如KVM）中使用。

2.1. 硬件拓扑

GPUDirect-RDMA友好拓扑：为了最大化GPU和NIC之间的内部吞吐量，系统内部硬件拓扑应支持GPUDirect RDMA。假设应用程序在主机的CPU核心上运行，GPU和NIC之间必须有专用的PCIe连接。这可以通过两种方式实现：
1. 将一个额外的PCIe交换机连接到主机系统总线的一个PCIe插槽。
2. 使用一个集成的DPU卡，该卡向主机暴露GPU和NIC。

硬件拓扑选项
硬件拓扑选项

可以使用lspci -tvvv或nvidia-smi topo -m命令检查系统拓扑。在某些主机系统上，必须禁用PCIe访问控制服务（ACS）以确保NIC和GPU之间的直接通信。

2.1.1. 选项1：以太网模式的网卡

固件要求：NVIDIA® ConnectX®固件版本必须为22.36.1010或更高。DOCA GPUNetIO允许CUDA内核在处理以太网协议时控制网卡，因此系统上的ConnectX NIC必须设置为以太网模式。

# 启动MST mst start mst status -v

MST模块输出示例：
MST模块

2.1.2. 选项2：DPU融合卡

固件要求：DPU固件版本必须为24.35.2000或更高。为了在主机上运行的应用程序能够使用DPU融合卡上的GPU和NIC，需要将DPU配置为在NIC模式下运行。

# 启用MST
sudo mst start
sudo mst status

# MST设备:
# ------------
# /dev/mst/mt41686_pciconf0 - PCI configuration cycles access.
#   domain:bus:dev.fn=0000:b8:00.0
#   addr.reg=88
#   data.reg=92
#   cr_bar.gw_offset=-1
#   Chip revision is: 01

# 将DPU融合卡上的GPU暴露给主机。对于BF2，偏移量为4；对于BF3，偏移量为8
sudo mlxconfig -d /dev/mst/mt41686_pciconf0 --yes s PCI_DOWNSTREAM_PORT_OWNER[4]=0x0

# 将BlueField-2端口设置为以太网模式（非InfiniBand）
sudo mlxconfig -d /dev/mst/mt41686_pciconf0 --yes set LINK_TYPE_P1=2 LINK_TYPE_P2=2

# 将BlueField-2设置为在DPU（嵌入式CPU）模式下运行
sudo mlxconfig -d /dev/mst/mt41686_pciconf0 --yes set INTERNAL_CPU_MODEL=1 INTERNAL_CPU_PAGE_SUPPLIER=1 INTERNAL_CPU_ESWITCH_MANAGER=1 INTERNAL_CPU_IB_VPORT0=1 INTERNAL_CPU_OFFLOAD_ENGINE=DISABLED

# 精确调度相关设置
sudo mlxconfig -d /dev/mst/mt41686_pciconf0 --yes set ACCURATE_TX_SCHEDULER=1 REAL_TIME_CLOCK_ENABLE=1

# 冷重启
sudo ipmitool power cycle

# 验证DPU固件更改是否已应用
sudo mlxconfig -d /dev/mst/mt41686_pciconf0 q LINK_TYPE_P1 LINK_TYPE_P2 INTERNAL_CPU_MODEL INTERNAL_CPU_PAGE_SUPPLIER INTERNAL_CPU_ESWITCH_MANAGER INTERNAL_CPU_IB_VPORT0 INTERNAL_CPU_OFFLOAD_ENGINE ACCURATE_TX_SCHEDULER REAL_TIME_CLOCK_ENABLE

验证输出示例：
DPU配置验证

2.2. PCIe配置

ACS禁用：在某些主机系统上，必须禁用PCIe访问控制服务（ACS）以确保NIC和GPU之间的直接通信。更多信息请参考相关文档。

2.3. GPU配置

CUDA与驱动：主机上必须安装CUDA Toolkit 12.1或更高版本。建议启用持久模式以减少应用程序初始延迟：nvidia-smi -pm 1。

内核模块：为了允许NIC使用GPU内存收发数据包，需要加载NVIDIA内核模块 nvidia-peermem（使用modprobe nvidia-peermem）。为了让CPU能够直接访问GPU内存（无需CUDA API），DPDK和DOCA需要系统上安装GDRCopy内核模块。

# 运行nvidia-peermem内核模块
sudo modprobe nvidia-peermem

# 安装GDRCopy
sudo apt install -y check kmod
git clone https://github.com/NVIDIA/gdrcopy.git /opt/mellanox/gdrcopy
cd /opt/mellanox/gdrcopy
make

# 运行gdrdrv内核模块
./insmod.sh

# 再次检查nvidia-peermem和gdrdrv模块是否正在运行
$ lsmod | egrep gdrdrv
gdrdrv                 24576  0
nvidia              55726080  4 nvidia_uvm,nvidia_peermem,gdrdrv,nvidia_modeset

# 导出库路径
export LD_LIBRARY_PATH=${LD_LIBRARY_PATH}:/opt/mellanox/gdrcopy/src

# 确保CUDA库路径在环境变量中
export PATH="/usr/local/cuda-12/bin:${PATH}"
export LD_LIBRARY_PATH="/usr/local/cuda-12/lib:/usr/local/cuda-12/lib64:${LD_LIBRARY_PATH}"
export CPATH="$(echo /usr/local/cuda-12/targets/{x86_64,sbsa}-linux/include | sed 's/ /:/'):${CPATH}"

BAR1映射空间：在GPU网络应用中，一个好的实践是通过RSS将传入流量分散到不同的接收队列中，以提高处理传入数据包的并行度。因此，重要的是要检查BAR1映射空间是否足够大，以容纳多个接收队列。可以使用nvidia-smi来验证GPU的BAR1映射空间。
检查GPU BAR1映射空间

3. 架构

应用阶段划分：一个GPU数据包处理网络应用可以分为两个基本阶段：
1. CPU上的设置：包括设备配置、内存分配、启动CUDA内核等。
2. 主数据路径：GPU和NIC交互以执行其功能。

GPUNetIO构建块：DOCA GPUNetIO提供不同的构建块，其中一些与DOCA以太网库结合使用，以创建一个完全在GPU上运行的完整流水线。在CPU的设置阶段，应用程序必须：
1. 在CPU上准备所有对象。
2. 为它们导出一个GPU句柄。
3. 启动一个CUDA内核，并将对象的GPU句柄传递给它，以便在数据路径期间使用该对象。

库组成：因此，DOCA GPUNetIO由两个库组成：
* libdoca_gpunetio：包含由CPU调用的函数，用于准备GPU、分配内存和对象。
* libdoca_gpunetio_device：包含由GPU在CUDA内核中于数据路径期间调用的函数。

链接注意事项：DOCA GPUNetIO共享库的pkgconfig文件是doca-gpu.pc。然而，DOCA GPUNetIO CUDA设备静态库/opt/mellanox/doca/lib/x86_64-linux-gnu/libdoca_gpunetio_device.a没有pkgconfig文件，因此如果需要使用DOCA GPUNetIO CUDA设备函数，必须将其显式链接到CUDA应用程序。

典型流程图：下图展示了典型的应用流程。
DOCA GPUNetIO典型流程

有关使用DOCA GPUNetIO发送和接收以太网数据包的示例，请参阅《NVIDIA DOCA GPU Packet Processing Application Guide》。

A2 方法细节

4. API

本节详细介绍与CPU和GPU上的主要DOCA GPUNetIO API相关的特定结构和操作。GPUNetIO的头文件包括：
* doca_gpunetio.h – CPU函数
* doca_gpunetio_dev_buf.cuh – 用于管理DOCA缓冲区数组的GPU函数
* doca_gpunetio_dev_eth_rxq.cuh – 用于管理DOCA以太网接收队列的GPU函数
* doca_gpunetio_dev_eth_txq.cuh – 用于管理DOCA以太网发送队列的GPU函数
* doca_gpunetio_dev_sem.cuh – 用于管理DOCA GPUNetIO信号量的GPU函数

对象导出为GPU句柄：所有与GPUNetIO一起使用的DOCA核心和以太网对象都有一个GPU导出函数，以获取该对象的GPU句柄。以下是一些示例：
* doca_buf_array 导出为 doca_gpu_buf_arr:

struct doca_mmap *mmap;
struct doca_buf_arr *buf_arr_cpu;
struct doca_gpu_buf_arr *buf_arr_gpu;
doca_mmap_create(NULL, &(mmap));
/* 填充并启动mmap */
doca_buf_arr_create(mmap, &buf_arr_cpu);
/* 填充并启动buf arr属性。在GPU上设置数据路径 */
/* 将buf array CPU句柄导出为buf array GPU句柄 */
doca_buf_arr_get_gpu_handle(buf_arr_cpu, &(buf_arr_gpu));
/* 要使用GPU句柄，请将其作为CUDA内核的参数传递 */
cuda_kernel<<<...>>>(buf_arr_gpu, ...);
  • doca_eth_rxq 导出为 doca_gpu_eth_rxq:
struct doca_mmap *mmap;
struct doca_eth_rxq *eth_rxq_cpu;
struct doca_gpu_eth_rxq *eth_rxq_gpu;
doca_eth_rxq_create(&eth_rxq_cpu);
/* 填充并启动以太网接收队列属性。在GPU上设置数据路径 */
/* 将以太网接收队列CPU句柄导出为以太网接收队列GPU句柄 */
doca_eth_rxq_get_gpu_handle(eth_rxq_cpu, &(eth_rxq_gpu));
/* 要使用GPU句柄，请将其作为CUDA内核的参数传递 */
cuda_kernel<<<...>>>(eth_rxq_gpu, ...);

4.1 doca_gpu_mem_type

内存类型枚举：此枚举列出了所有可以用GPUNetIO分配的可能内存类型。语法上，DOCA_GPU_MEM_前缀后的字符串表示<内存所在位置>_<谁有访问权限>。

enum doca_gpu_mem_type {
    DOCA_GPU_MEM_GPU = 0,
    DOCA_GPU_MEM_GPU_CPU = 1,
    DOCA_GPU_MEM_CPU = 2,
    DOCA_GPU_MEM_CPU_GPU = 3,
};
  • DOCA_GPU_MEM_GPU：内存位于GPU上，仅能从GPU访问。
  • DOCA_GPU_MEM_GPU_CPU：内存位于GPU上，但CPU也可以访问。
  • DOCA_GPU_MEM_CPU：内存位于CPU上，仅能从CPU访问。
  • DOCA_GPU_MEM_CPU_GPU：内存位于CPU上，但GPU也可以访问。
    DOCA_GPU_MEM_GPU_CPU内存类型的典型用途是从CPU向GPU发送通知（例如，一个CUDA内核周期性地检查CPU设置的退出条件是否满足）。

4.2. doca_gpu_create

创建GPUNetIO句柄：这是GPUNetIO应用程序必须调用的第一个函数，用于在GPU设备上创建一个句柄。该函数初始化一个指向struct doca_gpu *类型内存结构的指针。
doca_error_t doca_gpu_create(const char *gpu_bus_id, struct doca_gpu **gpu_dev);
* gpu_bus_id：要使用的GPU设备的PCIe地址，格式为<PCIe-bus>:<device>.<function>。
* gpu_dev [out]：指向该GPU设备的GPUNetIO句柄。
用户可以使用lspci或nvidia-smi命令获取PCIe地址。

4.3. doca_gpu_mem_alloc

内存分配函数：这个CPU函数用于分配不同类型的内存。
doca_error_t doca_gpu_mem_alloc(struct doca_gpu *gpu_dev, size_t size, size_t alignment, enum doca_gpu_mem_type mtype, void **memptr_gpu, void **memptr_cpu)
* gpu_dev：GPUNetIO设备句柄。
* size：要分配的内存区域大小（字节）。
* alignment：内存地址对齐方式。如果为0，则使用默认值。
* mtype：要分配的内存类型。
* memptr_gpu [out]：如果内存在GPU上或对GPU可见，则用于从GPU修改该内存的GPU指针。
* memptr_cpu [out]：如果内存在CPU上或对CPU可见，则用于从CPU修改该内存的CPU指针。如果内存是仅GPU的，可以为NULL。
警告：请确保在正确的设备上使用正确的指针！如果应用程序尝试从CPU使用memptr_gpu地址访问内存，将导致段错误。

4.4. doca_gpu_semaphore_create

创建信号量：创建一个新的DOCA GPUNetIO信号量实例。信号量由一个项目列表组成，每个项目默认包含一个状态标志、数据包数量以及doca_gpu_buf_arr中doca_gpu_buf的索引。例如，信号量可用于一个CUDA内核接收数据包到与以太网接收队列对象doca_gpu_eth_rxq关联的doca_gpu_buf_arr中，并将数据包信息分派给第二个进行处理的CUDA内核。另一种用法是在不同实体（如两个CUDA内核或一个CUDA内核和一个CPU线程）之间交换数据。这允许CUDA内核将处理结果提供给CPU，以便CPU编译统计报告。为此，可以为信号量中的每个项目关联一个自定义的应用定义结构，使其成为消息传递对象。
信号量通信模型

通信逻辑：通过信号量通信的实体必须采用轮询/更新机制：
* 更新方:
1. 填充信号量的下一个项目（数据包信息和/或自定义应用信息）。
2. 将状态标志设置为READY。

  • 轮询方:
    1. 等待下一个项目的状态标志变为READY。
    2. 读取并处理信息。
    3. 将状态标志设置为DONE。

doca_error_t doca_gpu_semaphore_create(struct doca_gpu *gpu_dev, struct doca_gpu_semaphore **semaphore)

  • gpu_dev：GPUNetIO句柄。
  • semaphore [out]：与GPU设备关联的GPUNetIO信号量句柄。

4.5. doca_gpu_semaphore_set_memory_type

设置信号量内存类型：此函数定义信号量分配的内存类型。如果应用程序仅在CUDA内核间共享数据包信息，建议使用DOCA_GPU_MEM_GPU。如果需要从CUDA内核向CPU共享信息，建议使用DOCA_GPU_MEM_CPU_GPU。
doca_error_t doca_gpu_semaphore_set_memory_type(struct doca_gpu_semaphore *semaphore, enum doca_gpu_mem_type mtype)

4.6. doca_gpu_semaphore_set_items_num

设置信号量项目数：此函数定义信号量中的项目数量。
doca_error_t doca_gpu_semaphore_set_items_num(struct doca_gpu_semaphore *semaphore, uint32_t num_items)

4.7. doca_gpu_semaphore_set_custom_info

关联自定义信息：此函数将应用程序特定的结构与信号量项目关联。内存类型的选择同上。
doca_error_t doca_gpu_semaphore_set_custom_info(struct doca_gpu_semaphore *semaphore, uint32_t nbytes, enum doca_gpu_mem_type mtype)

4.8. doca_gpu_semaphore_get_status

从CPU获取状态：从CPU查询信号量项目的状态。如果信号量分配为DOCA_GPU_MEM_GPU，此函数将导致段错误。

4.9. doca_gpu_semaphore_get_custom_info_addr

从CPU获取自定义信息地址：从CPU检索与信号量项目关联的自定义信息结构的地址。如果信号量或自定义信息分配为DOCA_GPU_MEM_GPU，此函数将导致段错误。

4.10. doca_gpu_dev_eth_rxq_receive_*

GPU端接收函数：为了在CUDA内核中获取数据包，DOCA GPUNetIO提供了不同范围的接收函数：_block（CUDA块）、_warp（CUDA warp）和_thread（CUDA线程）。

__device__ doca_error_t doca_gpu_dev_eth_rxq_receive_block(struct doca_gpu_eth_rxq *eth_rxq, uint32_t max_rx_pkts, uint64_t timeout_ns, uint32_t *num_rx_pkts, uint64_t *doca_gpu_buf_idx);
__device__ doca_error_t doca_gpu_dev_eth_rxq_receive_warp(struct doca_gpu_eth_rxq *eth_rxq, uint32_t max_rx_pkts, uint64_t timeout_ns, uint32_t *num_rx_pkts, uint64_t *doca_gpu_buf_idx);
__device__ doca_error_t doca_gpu_dev_eth_rxq_receive_thread(struct doca_gpu_eth_rxq *eth_rxq, uint32_t max_rx_pkts, uint64_t timeout_ns, uint32_t *num_rx_pkts, uint64_t *doca_gpu_buf_idx);

参数说明：
* eth_rxq：以太网接收队列的GPU句柄。
* max_rx_pkts：允许接收的最大数据包数。
* timeout_ns：返回前等待数据包的纳秒数。
* num_rx_pkts [out]：实际接收到的数据包数。
* doca_gpu_buf_idx [out]：此函数接收的第一个数据包的DOCA缓冲区索引。

使用要点：同一范围（线程、warp或块）内的CUDA线程必须在同一个接收队列上调用该函数。输出参数num_rx_pkts和doca_gpu_buf_idx必须对范围内的所有线程可见（例如，对于warp和块使用CUDA共享内存）。对于_block范围，调用接收函数的块必须至少有32个CUDA线程（即一个warp）。函数在timeout_ns到达或接收到最大数据包数时退出。

数据存储：接收到的每个数据包都进入内部创建并与以太网队列关联的doca_gpu_buf_arr。输出参数指示实际接收了多少数据包（num_rx_pkts）以及第一个接收到的数据包在doca_gpu_buf_array中的索引（doca_gpu_buf_idx）。数据包连续存储，如果函数返回num_rx_pkts=N和doca_gpu_buf_idx=X，则表示doca_gpu_buf_arr中范围[X, .. ,X + (N-1)]内的所有doca_gpu_buf都已填充数据包。
接收数据包到缓冲区数组

循环缓冲区：DOCA缓冲区数组以循环方式处理，一旦最后一个DOCA缓冲区被数据包填充，队列将回到第一个DOCA缓冲区。应用程序无需锁定或释放doca_gpu_buf_arr缓冲区。应用程序有责任在数据包被覆盖之前消费它们，需要适当调整DOCA缓冲区数组的大小并在多个接收队列之间进行扩展。

4.11. doca_gpu_dev_eth_txq_send_*

GPU端发送函数：为了从CUDA内核发送数据包，DOCA GPUNetIO为将数据包入队到以太网txq提供了强模式和弱模式。两种模式的范围都是单个CUDA线程，每个线程从doca_gpu_buf_arr中填充并入队一个不同的doca_gpu_buf。通常建议使用强模式。

强模式：
强模式发送
__device__ doca_error_t doca_gpu_dev_eth_txq_send_enqueue_strong(struct doca_gpu_eth_txq *eth_txq, const struct doca_gpu_buf *buf_ptr, const uint32_t nbytes)

弱模式：在弱模式下，开发者必须指定一个队列描述符编号，用于将数据包入队，并确保队列中没有描述符为空，该编号会以16位掩码回绕。
弱模式发送
__device__ doca_error_t doca_gpu_dev_eth_txq_send_enqueue_weak(const struct doca_gpu_eth_txq *eth_txq, const struct doca_gpu_buf *buf_ptr, const uint32_t nbytes, const uint32_t ndescr)

4.12. doca_gpu_dev_eth_txq_wait_*

精确发送调度：要启用精确发送调度，必须在发送队列中设置“按时等待”屏障（基于时间戳），然后再入队更多数据包。与send_*函数类似，wait_*函数也有强模式和弱模式。通常建议使用强模式。

强模式：
__device__ doca_error_t doca_gpu_dev_eth_txq_wait_time_enqueue_strong(struct doca_gpu_eth_txq *eth_txq, const uint64_t wait_on_time_value)

弱模式：
__device__ doca_error_t doca_gpu_dev_eth_txq_wait_time_enqueue_weak(struct doca_gpu_eth_txq *eth_txq, const uint64_t wait_on_time_value, const uint32_t ndescr)

4.13. doca_gpu_dev_eth_txq_commit_*

提交发送队列：在将所有要发送的数据包和时间屏障入队后，必须在txq队列上调用一个提交函数。必须根据send_*和wait_*中使用的入队模式（强或弱）来使用正确的提交函数。

强模式：
__device__ doca_error_t doca_gpu_dev_eth_txq_commit_strong(struct doca_gpu_eth_txq *eth_txq)

弱模式：
__device__ doca_error_t doca_gpu_dev_eth_txq_commit_weak(struct doca_gpu_eth_txq *eth_txq, const uint32_t descr_num)

使用流程：在一个范围（CUDA块或CUDA warp）内，只有一个CUDA线程可以在多次入队操作后调用此函数。典型流程是：1. 范围内的所有线程将数据包入队到发送队列。2. 同步点。3. 范围内只有一个线程执行发送队列提交。

4.14. doca_gpu_dev_eth_txq_push

推送至网卡：提交后，发送队列中的项目必须被实际推送到网卡。
__device__ doca_error_t doca_gpu_dev_eth_txq_push(struct doca_gpu_eth_txq *eth_txq)

使用流程：在一个范围（CUDA块或CUDA warp）内，只有一个CUDA线程可以在多次入队或提交操作后调用此函数。典型流程是：1. 范围内的所有线程将数据包入队到发送队列。2. 同步点。3. 范围内只有一个线程执行发送队列提交。4. 范围内只有一个线程执行发送队列推送。

5. 构建块

本节解释了创建DOCA GPUNetIO应用程序时使用的基本构建块背后的一般概念。

5.1. 初始化GPU和NIC

设备设置：当DOCA GPUNetIO与NIC结合使用以发送或接收以太网流量时，必须执行以下操作来正确设置应用程序和设备：

uint16_t dpdk_port_id;
struct doca_dev *ddev;
struct doca_gpu *gdev;
char *eal_param[3] = {"", "-a", "00:00.0"};
/* 使用空设备初始化DPDK。DOCA设备稍后将热插拔网卡。 */
rte_eal_init(3, eal_param);
/* 在特定网卡上创建DOCA设备 */
doca_dpdk_port_probe(&ddev);
get_dpdk_port_id_doca_dev(&ddev, &dpdk_port_id);
/* 在特定GPU上创建GPUNetIO句柄 */
doca_gpu_create(gpu_pcie_address, &gdev);

5.2. 以太网接收队列

创建与配置：如果DOCA应用程序必须接收以太网数据包，则必须创建接收队列。接收队列以循环方式工作：在创建时，每个接收队列都与应用程序在GPU上分配的DOCA缓冲区数组相关联。缓冲区数组的每个DOCA缓冲区都有一个固定的最大尺寸。

/* 启动DPDK设备 */
rte_eth_dev_start(dpdk_port_id);
/* 初始化DOCA Flow */
struct doca_flow_port_cfg port_cfg;
port_cfg.port_id = port_id;
doca_flow_init(port_cfg);
doca_flow_port_start();
struct doca_eth_rxq *eth_rxq_cpu;
struct doca_gpu_eth_rxq *eth_rxq_gpu;
struct doca_mmap *mmap;
void *gpu_buffer;
/* 创建DOCA以太网接收队列 */
doca_eth_rxq_create(&eth_rxq_cpu);
/* 设置以太网接收队列属性 */
/* ... */
/* 在GPU内存中创建DOCA mmap，用于与此以太网队列关联的DOCA缓冲区数组 */
doca_mmap_create(&mmap);
doca_gpu_mem_alloc(gdev, buffer_size, alignment, DOCA_GPU_MEM_GPU, (void **)&gpu_buffer, NULL);
doca_mmap_start(mmap);
doca_eth_rxq_set_pkt_buffer(eth_rxq_cpu, mmap, 0, buffer_size);
/* 启动以太网队列对象 */
/* 导出接收队列的GPU句柄 */
doca_eth_rxq_get_gpu_handle(eth_rxq_cpu, &eth_rxq_gpu);

强制要求：必须将DOCA Flow管道与接收队列关联，否则应用程序无法接收任何数据包。

5.3. 以太网发送队列

创建与配置：如果DOCA应用程序必须发送以太网数据包，则必须创建发送队列，并结合doca_gpu_buf_arr从GPU内存准备和发送数据包。

struct doca_eth_txq *eth_txq_cpu;
struct doca_gpu_eth_txq *eth_txq_gpu;
/* 创建DOCA以太网发送队列 */
doca_eth_txq_create(&eth_txq_cpu);
/* 设置发送队列属性 */
/* 导出发送队列的GPU句柄 */
doca_eth_rxq_get_gpu_handle(eth_txq_cpu, &eth_txq_gpu);
/* 创建DOCA mmap以定义DOCA buf array的内存布局和类型 */
struct doca_mmap *mmap;
doca_mmap_create(&mmap);
/* 设置DOCA mmap属性 */
/* 创建DOCA buf arr并将其导出到GPU */
struct doca_buf_arr *buf_arr;
struct doca_gpu_buf_arr *buf_arr_gpu;
doca_buf_arr_create(mmap, &buf_arr);
/* 设置DOCA buf array属性 */
/* 导出buf arr的GPU句柄 */
doca_buf_arr_get_gpu_handle(buf_arr, &buf_arr_gpu);

5.4. 信号量

用途与创建：如果DOCA应用程序必须在CUDA内核之间或从CUDA内核与CPU线程之间分派数据包信息，则必须创建信号量。信号量是一个项目列表，根据用例分配在GPU或CPU上，但对GPU和CPU都可见。默认情况下，每个信号量项目可以保存其状态（FREE, READY, HOLD, DONE, ERROR）、接收到的数据包数量以及doca_gpu_buf_arr中doca_gpu_buf的索引。如果信号量必须用于与CPU交换数据，首选的内存布局是DOCA_GPU_MEM_CPU_GPU。如果信号量仅在CUDA内核之间需要，DOCA_GPU_MEM_GPU是最佳的内存布局。还可以附加自定义结构以传递更多特定于应用程序的信息。

#define SEMAPHORE_ITEMS 1024
/* 应用定义的自定义结构，用于通过信号量项目传递信息 */
struct custom_info {
    int a;
    uint64_t b;
};

/* 用于从GPU向CPU共享信息的信号量 */
struct doca_gpu_semaphore *sem_to_cpu;
struct doca_gpu_semaphore_gpu *sem_to_cpu_gpu;
doca_gpu_semaphore_create(gdev, &sem_to_cpu);
doca_gpu_semaphore_set_memory_type(sem_to_cpu, DOCA_GPU_MEM_CPU_GPU);
doca_gpu_semaphore_set_items_num(sem_to_cpu, SEMAPHORE_ITEMS);
/* 这是可选的 */
doca_gpu_semaphore_set_custom_info(sem_to_cpu, sizeof(struct custom_info), DOCA_GPU_MEM_CPU_GPU);
doca_gpu_semaphore_start(sem_to_cpu);
doca_gpu_semaphore_get_gpu_handle(sem_to_cpu, &sem_to_cpu_gpu);

/* 用于在GPU CUDA内核间共享信息，无CPU参与的信号量 */
struct doca_gpu_semaphore *sem_to_gpu;
struct doca_gpu_semaphore_gpu *sem_to_gpu_gpu;
doca_gpu_semaphore_create(gdev, &sem_to_gpu);
doca_gpu_semaphore_set_memory_type(sem_to_gpu, DOCA_GPU_MEM_GPU);
doca_gpu_semaphore_set_items_num(sem_to_gpu, SEMAPHORE_ITEMS);
/* 这是可选的 */
doca_gpu_semaphore_set_custom_info(sem_to_gpu, sizeof(struct custom_info), DOCA_GPU_MEM_GPU);
doca_gpu_semaphore_start(sem_to_gpu);
doca_gpu_semaphore_get_gpu_handle(sem_to_gpu, &sem_to_gpu_gpu);

5.5. GPU上的数据路径

此时，应用程序已创建并初始化了GPU执行数据路径所需的所有对象，以使用GPUNetIO发送或接收数据包。

5.5.1. 接收与处理

双核流水线模型：在此示例中，应用程序必须使用一个接收器CUDA内核从不同队列接收数据包，并将数据包信息分派给负责数据包处理的第二个CUDA内核。
* CPU端代码：CPU启动CUDA内核并在信号量上等待输出。

#define CUDA_THREADS 512
#define CUDA_BLOCKS 1
int semaphore_index = 0;
enum doca_gpu_semaphore_status status;
struct custom_info *gpu_info;

/* 在CPU上 */
cuda_kernel_receive_dispatch<<<CUDA_THREADS, CUDA_BLOCKS, ..., stream_0>>>(eth_rxq_gpu, sem_to_gpu_gpu);
cuda_kernel_process<<<CUDA_THREADS, CUDA_BLOCKS, ..., stream_1>>>(eth_rxq_gpu, sem_to_cpu_gpu, sem_to_gpu_gpu);

while(/* condition */) {
    doca_gpu_semaphore_get_status(sem_to_cpu, semaphore_index, &status);
    if (status == DOCA_GPU_SEMAPHORE_STATUS_READY) {
        doca_gpu_semaphore_get_custom_info_addr(sem_to_cpu, semaphore_index, (void **)&(gpu_info));
        report_info(gpu_info);
        doca_gpu_semaphore_set_status(sem_to_cpu, semaphore_index, DOCA_GPU_SEMAPHORE_STATUS_FREE);
        semaphore_index = (semaphore_index+1) % SEMAPHORE_ITEMS;
    }
}
  • GPU端代码：两个CUDA内核在不同的流上运行。
// 接收与分派内核
cuda_kernel_receive_dispatch(eth_rxq_gpu, sem_to_gpu_gpu) {
    __shared__ uint32_t rx_pkt_num;
    __shared__ uint64_t rx_buf_idx;
    int semaphore_index = 0;
    doca_gpu_dev_eth_rxq_receive_block(eth_rxq_gpu, MAX_NUM_RECEIVE_PACKETS, TIMEOUT_RECEIVE_NS, &rx_pkt_num, &rx_buf_idx);
    if (threadIdx.x == 0 && rx_pkt_num > 0) {
        doca_gpu_dev_sem_set_packet_info(sem_to_gpu_gpu, semaphore_index, DOCA_GPU_SEMAPHORE_STATUS_READY, rx_pkt_num, rx_buf_idx);
        semaphore_index = (semaphore_index+1) % SEMAPHORE_ITEMS;
    }
}

// 处理内核
cuda_kernel_process(eth_rxq_gpu, sem_to_cpu_gpu, sem_to_gpu_gpu) {
    __shared__ uint32_t rx_pkt_num;
    __shared__ uint64_t rx_buf_idx;
    int semaphore_index = 0;
    int thread_buf_idx = 0;
    struct doca_gpu_buf *buf_ptr;
    uintptr_t buf_addr;
    struct custom_info *gpu_info;

    while (/* exit condition */) {
        if (threadIdx.x == 0) {
            do {
                result = doca_gpu_dev_sem_get_packet_info_status(sem_to_gpu_gpu, semaphore_index, DOCA_GPU_SEMAPHORE_STATUS_READY, &rx_pkt_num, &rx_buf_idx);
            } while(result != DOCA_ERROR_NOT_FOUND /* && other exit condition */);
        }
        __syncthreads();
        thread_buf_idx = threadIdx.x;
        while (thread_buf_idx < rx_pkt_num) {
            /* 从接收队列的GPU缓冲区中获取DOCA GPU缓冲区 */
            doca_gpu_dev_eth_rxq_get_buf(eth_rxq_gpu, rx_buf_idx + thread_buf_idx, &buf_ptr);
            /* 获取DOCA GPU缓冲区的内存地址 */
            doca_gpu_dev_buf_get_addr(buf_ptr, &buf_addr);
            
            /* 这里的原子操作是因为整个CUDA块访问同一个到CPU的信号量。更智能的实现可以在warp级别使用多个信号量等来避免此原子操作 */
            int semaphore_index_tmp = atomicAdd_block(&semaphore_index, 1);
            semaphore_index_tmp = semaphore_index_tmp % SEMAPHORE_ITEMS;
            doca_gpu_dev_sem_get_custom_info_addr(sem_to_cpu_gpu, semaphore_index_tmp, (void **)&gpu_info);
            populate_custom_info(buf_addr, gpu_info);
            doca_gpu_dev_sem_set_status(sem_to_cpu_gpu, semaphore_index_tmp, DOCA_GPU_SEMAPHORE_STATUS_READY);
        }
    }
}

此代码可以用下图表示（当使用多个队列和/或信号量时）：
多队列/信号量接收处理模型

单核模型：接收和分派到另一个CUDA内核不是必需的。一个更简单的场景可以只有一个CUDA内核接收和处理数据包。这种方法的缺点是两次接收之间的时间取决于CUDA内核处理接收到的数据包所花费的时间。
单核接收处理模型

5.5.2. 生产与发送

GPU生产数据并发送：在此示例中，GPU生成一些数据，将其存储到数据包中，然后通过网络发送出去。
* CPU端代码：CPU启动CUDA内核并继续执行其他工作。

#define CUDA_THREADS 512
#define CUDA_BLOCKS 1
/* 在CPU上 */
cuda_kernel_produce_send<<<CUDA_THREADS, CUDA_BLOCKS, ..., stream_0>>>(eth_txq_gpu, buf_arr_gpu);
/* 做其他事情 */
  • GPU端代码：CUDA内核用有意义的数据填充数据包并发送它们。在以下示例中，范围是CUDA块，因此每个块使用一个不同的DOCA以太网发送队列。
cuda_kernel_produce_send(eth_txq_gpu, buf_arr_gpu) {
    uint64_t doca_gpu_buf_idx = threadIdx.x;
    struct doca_gpu_buf *buf;
    uintptr_t buf_addr;
    uint32_t packet_len;
    while (/* exit condition */) {
        /* 每个CUDA线程从doca_gpu_buf_arr中检索doca_gpu_buf */
        doca_gpu_dev_buf_get_buf(buf_arr_gpu, doca_gpu_buf_idx, &buf);
        /* 获取doca_gpu_buf中数据包的内存地址 */
        doca_gpu_dev_buf_get_addr(buf, &buf_addr);
        /* 应用程序生成数据并在doca_gpu_buf中制作数据包 */
        populate_packet(buf_addr, &packet_len);
        /* 将数据包入队到发送队列 */
        doca_gpu_dev_eth_txq_send_enqueue_strong(eth_txq_gpu, buf, packet_len);
        /* 同步点 */
        __syncthreads();
        /* 块中只有一个CUDA线程必须提交和推送发送队列 */
        if (threadIdx.x == 0) {
            doca_gpu_dev_eth_txq_commit_strong(eth_txq_gpu);
            doca_gpu_dev_eth_txq_push(eth_txq_gpu);
        }
        /* 同步点 */
        __syncthreads();
        /* 假设块中的所有线程都在发送队列中推送了一个数据包 */
        doca_gpu_buf_idx += CUDA_THREADS;
    }
}

A4 实验环境与结果

实验环境

  • 数据集/模型: 示例程序不使用特定数据集或模型。它生成并发送大小为1kB的原始以太网数据包到虚拟以太网地址10:11:12:13:14:15。

  • 硬件配置:

    • GPU: 通过PCIe地址指定（例如ca:00.0），需要与NIC构成GPUDirect-RDMA友好拓扑。
    • NIC: NVIDIA ConnectX-6 Dx或更新版本（支持精确发送调度），通过PCIe地址指定（例如17:00.0）。ConnectX固件需22.36.1010或更新，或DPU固件需24.35.2000或更新。
    • CPU: 主机CPU。
    • 连接: GPU和NIC之间有专用的PCIe连接。

  • 软件配置:

    • 操作系统: Ubuntu 20.04 或 22.04。
    • 核心库: DOCA for Host，CUDA Toolkit 12.1+，DPDK（随DOCA提供）。
    • 内核模块: 需加载nvidia-peermem和gdrdrv。
    • 系统配置: 禁用IOMMU。使用phc2sys服务进行CPU与NIC的时钟同步。
    • 代码实现: 示例代码使用meson和ninja进行构建。

实验结果

实验内容: 该示例演示了GPUNetIO应用中的精确发送调度（Accurate Send Scheduling）功能。该功能允许NVIDIA NIC根据应用程序提供的时间戳在未来的特定时间发送数据包。示例程序会发送8个批次的数据包，每批包含32个原始以太网数据包，并通过命令行选项-t设置每批之间的发送时间间隔。

时钟同步: 在运行示例之前，必须使用phc2sys服务同步