GPUNetIO and UROM

1What GPUNetIO changes in the data path

DOCA GPUNetIO, Beta in 3.5.0, lets “CUDA kernels … control network communications to send or receive data” — the model NVIDIA calls GPUDirect Async Kernel-Initiated networking.^⁠5⁠1 It stands on three pieces: GPUDirect RDMA so the NIC can write and read GPU memory directly, GDRCopy so the CPU can touch GPU memory when needed, and the BlueField DMA engine for the GPU-to-CPU copy sample.^⁠1⁠3 The CPU keeps the control plane: it creates the device objects, the queues and the DOCA Flow rules that steer traffic to GPU-owned queues, then launches a kernel that runs for the life of the application.^⁠14

flowchart LR; W[Wire] --> NIC["NIC queues on BlueField-3 (DPU mode) or ConnectX-7/8"]; NIC -- "GPUDirect RDMA via dmabuf (nvidia-open) or nvidia-peermem" --> GMEM["GPU memory: doca_gpu_mem_alloc(DOCA_GPU_MEM_TYPE_GPU)"]; GMEM --> K["CUDA kernel: doca_gpu_dev_eth_rxq_receive_block()"]; K -- "doca_gpu_dev_eth_txq_send_enqueue_strong / commit / push" --> NIC; CPU["CPU control plane: doca_gpu_create, doca_eth_rxq_create, doca_eth_rxq_get_gpu_handle, doca_flow pipes"] -.-> NIC; CPU -.-> K

The gpunetio_simple_receive host program shows the control plane in order: doca_flow_init and doca_flow_port_start first, then doca_gpu_create on the GPU PCIe address, doca_eth_rxq_create and doca_eth_rxq_set_type, a GPU-resident packet buffer from doca_gpu_mem_alloc sized with doca_eth_rxq_estimate_packet_buf_size and attached with doca_eth_rxq_set_pkt_buf, doca_ctx_set_datapath_on_gpu and doca_eth_rxq_get_gpu_handle, and finally the pipes (create_udp_pipe, then the control pipe with doca_flow_pipe_control_add_entry and doca_flow_entries_process) that steer UDP to that queue; a doca_gpu_dmabuf_fd call appears where dmabuf mapping is used.^⁠14 The kernel side, gpunetio_simple_receive_kernel.cu, is one __global__ function that receives with the block, warp or thread execution scope (DOCA_GPUNETIO_ETH_EXEC_SCOPE_BLOCK/WARP/THREAD), reads packets with doca_gpu_dev_eth_rxq_get_pkt_addr, and uses DOCA_GPUNETIO_VOLATILE for the exit flag.^⁠13

2Prerequisites: drivers, topology, BAR1 and the fallback

The guide’s hardware and system setup section is a checklist. Drivers: “Install the nvidia-open drivers to utilize dmabuf. If the environment relies on the closed-source NVIDIA drivers (cuda-drivers), the system does not support dmabuf; use nvidia-peermem instead.”^⁠1 Topology: “ensure the GPU and NIC share a PIX or PXB topology (connected via a single PCIe bridge). Avoid NODE or SYS topologies, as routing data across NUMA nodes or SMP interconnects reduces efficiency”, read from nvidia-smi topo -m.^⁠1 BAR1: “If an application fails to map memory buffers to the NIC, verify the GPU’s BAR1 size using nvidia-smi -q. Enable Resizable BAR in the system BIOS”.^⁠1 Without GPUDirect RDMA at all (the guide names DGX Spark), “use the DOCA_GPU_MEM_TYPE_CPU_GPU memory type and enable CPU proxy mode”.^⁠1

Programming prerequisites follow: an active CUDA context on the target device before any DOCA GPUNetIO call, obtained cheaply with cudaFree(0); DPDK’s gpudev and hugepages for the Ethernet path; Ethernet samples run with sudo, whereas “Verbs, RDMA, and DMA operations can run without root access” if the NVIDIA driver is loaded with NVreg_RegistryDwords="PeerMappingOverride=1;".^⁠1 The supported matrix in the guide is GPUs H100, A100 and L40, and NICs BlueField-3 in DPU mode, ConnectX-7 and ConnectX-8.^⁠1 The samples README asks you to set the GPU architecture in meson.build (sm_80 Ampere, sm_89 L40, sm_90 H100) and buildtype = 'release', and to export PATH with /usr/local/cuda/bin, CPATH with the CUDA targets include, PKG_CONFIG_PATH including /opt/mellanox/dpdk/lib/…/pkgconfig and /opt/mellanox/doca/lib/…/pkgconfig, and LD_LIBRARY_PATH including /opt/mellanox/gdrcopy/src.^⁠2⁠1

3API shape, samples, and honest numbers

Host APIs: doca_gpu_create, doca_gpu_mem_alloc with DOCA_GPU_MEM_TYPE_GPU, GPU_CPU or CPU_GPU, and for Ethernet doca_ctx_set_datapath_on_gpu() followed by doca_eth_rxq_get_gpu_handle().^⁠1 Device APIs: doca_gpu_dev_eth_rxq_receive_block/warp/thread(), doca_gpu_dev_eth_txq_send_enqueue_strong(), doca_gpu_dev_eth_txq_commit/push(), semaphores doca_gpu_dev_sem_*(), doca_gpu_dev_rdma_write(), the doca_gpu_dev_verbs_*() family and doca_gpu_dev_verbs_submit_bf() for BlueFlame doorbells (3.4.0 and later).^⁠1 RDMA and Verbs join through the same handoff you saw in M6.1: doca_ctx_set_datapath_on_gpu() with doca_rdma_get_gpu_handle(), and doca_gpu_verbs_export_qp() on the Verbs side.^⁠7⁠8 The 3.4.0 additions listed in the guide are DOCA_GPUNETIO_VERBS_GPU_CODE_OPT_BF_UNRELIABLE, doca_gpu_dev_verbs_wqe_prepare_atomic_ext, doca_gpu_dev_verbs_fence_acquire_nvidia_nic and Gbps reporting in the bandwidth samples.^⁠1

The 3.5.0 sample tree has gpunetio_simple_receive, gpunetio_simple_send, gpunetio_send_wait_time (Accurate Send Scheduling, ConnectX-6 Dx or later, needs phc2sys; run as sudo ./build/doca_gpunetio_send_wait_time -n 17:00.0 -g ca:00.0 -t 5000000), gpunetio_dma_memcpy (requires a DPU because it uses the DMA engine), gpunetio_rdma_client_server_write (RDMA by device name, e.g. -d mlx5_0, with an optional -cm mode), and the Verbs micro-benchmarks gpunetio_verbs_{get_bw,put_bw,put_counter_bw,put_counter_lat,put_signal_bw,put_signal_lat,twosided_bw,write_bw,write_lat}; the reference application is gpu_packet_processing, enabled with the meson option enable_gpu_packet_processing.^{⁠12⁠2⁠11} DOCA Bench can drive doca_gpunetio::eth_tx|eth_rx|eth_bidir with --gpu and the doca_gpunetio.threads-per-block attribute.^⁠9

The numbers you will be asked about come from an NVIDIA blog first published 2022-12-14 (updated 2023-05-08): an IP-checksum kernel on a single queue receiving “3B packets of 1 KB size at ~100 Gbps (~11.97 Mpps)” with zero loss, 512 B packets at “86 Gbps (~20.9 Mpps)” limited by the T-Rex generator, and forwarding at “~90 Gbps” with one RX and one TX queue — measured on an A100X with ConnectX-6 Dx/BlueField-2X, Ubuntu 20.04, MOFED 5.8, CUDA 11.8 and a PCIe Gen3 sender.^⁠3 Label them exactly that way: older generation, NVIDIA-measured, not DOCA 3.5.0 on BlueField-3 or H100.^⁠3

4UROM: offloading pieces of HPC and AI frameworks

DOCA UROM is “a framework for offloading a portion of parallel computing tasks, such as those related to HPC or AI workloads and frameworks, from the host to the NVIDIA DPUs”.^⁠4 Three parts: the UROM daemon (a DOCA service running on the DPU that spawns workers), the host library, and plugins — “A UROM offload plugin is where developers of AI/HPC offloads implement their own offloading logic while using DOCA UROM as the transport layer and resource manager”.^⁠4 Reference plugins are UCC (collectives), RDMO and Graph.^⁠4 In 3.5.0 the library is Alpha in the libraries index while the UROM service is GA in the services index — quote both when a customer asks “is it supported”.^⁠5⁠6

The host API has three object families. Service: doca_urom_service_create, doca_urom_service_set_dev, doca_urom_service_get_plugins_list, doca_urom_service_get_workers_by_gid_task_allocate_init. Worker: doca_urom_worker_create, _set_service, _set_plugins, _set_gid, doca_urom_worker_cmd_task_allocate_init. Domain: doca_urom_domain_create, _set_workers, _add_buffer, _set_oob.^⁠4 A plugin exports urom_plugin_get_version() and urom_plugin_get_iface(), whose interface provides open(), addr(), worker_cmd(), progress() and notif_pack().^⁠4 Prerequisite: “UCX is required for the communication channel between the host and DPU parts of DOCA UROM based on TCP socket transport”, with UCX_NET_DEVICES selecting the devices.^⁠4

The samples make the two halves visible: plugin samples (worker_sandbox, which offloads the UCX tagged send/receive API, and worker_graph) build into a host component linked to the program and a DPU component compiled as worker_<name>.so and “loaded at runtime by the DOCA UROM service (daemon, workers)”; the programs are urom_multi_workers_bootstrap and urom_ping_pong, the latter launched with Open MPI, both taking -d/--device <IB device name>.^⁠10 The reference application is urom_rdmo.^⁠12 The service’s container YAML and image tag were not captured on the fetched pages; check the NGC listing before writing a deployment step.^⁠4