NVIDIA DOCA GPUNetIO Programming Guide

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

A1 主要贡献

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

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

CPU中心方法：

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模块输出示例：

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服务同步