How DOCA GPUNetIO Unifies GPU-Initiated Networking Across the NVIDIA Software Stack

NVIDIA's DOCA GPUNetIO has evolved from a standalone GPU-centric packet-processing framework into a unified GDA-KI foundation now integrated across NVSHMEM, NCCL, Aerial 5G SDK, UCX/NIXL, NVQLink with the Holoscan Sensor Bridge operator, Holoscan Advanced Network Operator, and DeepEP/HybridEP. The SDK lets CUDA kernels directly drive Ethernet, RDMA, Verbs, and DMA operations via GPUDirect RDMA, GPUDirect Async Kernel-Initiated (GDA-KI), and GDRCopy, keeping the CPU off the application critical path, and NVIDIA is also releasing an open-source version alongside the full DOCA SDK superset. Consolidating previously separate per-library GDA-KI RDMA implementations into one shared codebase reduces duplicated engineering and speeds feature, optimization, and bug-fix propagation across the stack.

GPU applications increasingly need networking and data movement to behave like first-class GPU-controlled operations rather than host-driven services. When the CPU sits in the middle of every network transaction, it becomes a bottleneck on the critical path, adding latency and limiting how efficiently distributed applications can respond in real time. NVIDIA DOCA GPUNetIO, a GPU-centric networking SDK layer for real-time packet processing and data movement, addresses this directly. It brings together technologies such as GPUDirect RDMA, GPUDirect Async Kernel-Initiated GDA-KI and GDRCopy so CUDA kernels can directly drive Ethernet, RDMA, Verbs, and DMA operations while keeping the CPU out of the application critical path. Specifically, DOCA GPUNetIO provides both: - CPU functions on the control path to export on GPU memory transport objects created with DOCA Ethernet, DOCA Verbs, DOCA DMA, DOCA CommChannel - GPU CUDA functions on the data path to allow the creation of CUDA kernels that can manipulate the transport objects exported Since NVIDIA first introduced GPU-centric packet processing with DOCA GPUNetIO https://developer.nvidia.com/blog/inline-gpu-packet-processing-with-nvidia-doca-gpunetio/ , the framework has matured significantly, evolving from a powerful way to remove the CPU from the critical path into a unified GDA-KI foundation integrated across NVSHMEM https://github.com/NVIDIA/nvshmem , NCCL https://github.com/NVIDIA/nccl , Aerial 5G SDK https://developer.nvidia.com/aerial-sdk , UCX https://github.com/openucx/ucx / NIXL https://github.com/ai-dynamo/nixl/tree/main , NVQLink https://developer.nvidia.com/blog/nvidia-nvqlink-architecture-integrates-accelerated-computing-with-quantum-processors/ with Holoscan Sensor Bridge operator https://github.com/nvidia-holoscan/holoscan-sensor-bridge/tree/main/src/hololink/operators/gpu roce transceiver , Holoscan Advanced Network Operator https://github.com/nvidia-holoscan/holohub/tree/main/operators/advanced network , DeepEP/HybridEP https://github.com/deepseek-ai/DeepEP/tree/hybrid-ep , and others. This post covers what’s new: the open-source library, the unified architecture, and how these integrations work in practice. Why a unified GPUNetIO umbrella matters Before GPUNetIO became the shared foundation, every communication library had built its own separate implementation of GDA-KI-style, GPU-initiated RDMA. Each one worked independently, with its own assumptions, its own code, and its own maintenance burden. None of them shared anything with the others. Bringing the different GDA-KI-style implementations under the GPUNetIO umbrella provides a great advantage: instead of each library building and maintaining its own GDA-KI-based RDMA communication path, they can converge on a single GPUNetIO GDA-KI implementation that everyone can use, improve, and extend in one place. That creates a shared investment point for features, optimizations, and bug fixes, reduces duplicated engineering across SDKs, and allows innovations developed for one framework to become available to others much more quickly. In other words, GPUNetIO becomes the common GDA-KI foundation that multiple communication libraries can build on, rather than a collection of parallel, partially overlapping implementations. This unified approach is especially valuable for higher-level libraries such as NIXL/UCX, NCCL, and NVSHMEM. They are collaborating to enrich GPUNetIO, all taking and benefit from improvements made across the stack. From an ecosystem perspective, that means less duplicated code, less fragmentation in behavior, and a much stronger path for long-term evolution. GPUNetIO: Open Source vs SDK NVIDIA ships GPUNetIO in two closely related forms because the ecosystem needs both breadth and openness. The full DOCA SDK version is the superset: it is the GPUNetIO implementation documented in the DOCA Programming Guide https://networking-docs.nvidia.com/doca/sdk/doca-gpunetio and it spans the broader DOCA stack, including but not limited to Verbs, Ethernet, DMA, and Comm Channel integration, while preserving the core GPUNetIO model of moving the control path closer to the GPU and removing the CPU from the application critical path. In parallel, NVIDIA also publishes an open-source GPUNetIO project https://github.com/NVIDIA-DOCA/gpunetio as a lighter-weight, RDMA-Verbs-focused implementation designed for frameworks that want to remain fully open in how they integrate networking transports. Today, the open version is the smaller Verbs-oriented subset, while the DOCA SDK version is the superset that adds broader capabilities such as richer RDMA support, Ethernet, and DMA. That split is not about creating two divergent software stacks. It is about providing a common open-source foundation for GPU-initiated RDMA communications while still leaving a path to more advanced functionality when the DOCA SDK is present. Indeed, the open-source implementation can detect the presence of DOCA SDK https://github.com/NVIDIA-DOCA/gpunetio enable-sdk-mode at runtime and, when available, call selected closed-source DOCA SDK functions through dlopen ; if the SDK is absent, it continues to operate using the open-source implementation. This is the key architectural advantage: instead of every communication library building and maintaining its own private version of GDAKI-style RDMA plumbing, GPUNetIO becomes the shared implementation point that multiple frameworks and libraries can adopt, harden, and improve together. On the CUDA device side, the gap is intentionally small for the Verbs path: both the open and SDK variants expose a device-facing API, which keeps the GPU programming model largely aligned even as the host-side implementation scales from a lightweight open-source path to the broader DOCA SDK feature set. Programming model This section covers the key concepts and programming model elements that apply to any GPUNetIO application. CPU control path In applications that use DOCA GPUNetIO, an initial host-side CPU configuration phase is required to execute the control path, followed by a second phase in which the data path runs on the GPU through CUDA kernels. The control-path workflow generally follows this sequence: - The GPU and network devices are initialized and configured, and the required memory is allocated. - Network transport objects are created. For example, DOCA Verbs or DOCA Ethernet creates a network queue object via mlx5dv on the CPU. - A GPUNetIO CPU function exports the relevant elements of the network queue object into a descriptor stored in GPU memory and provides a GPU address for it. - The application launches a CUDA kernel, passing the GPU descriptor for the network queue as one of the input parameters. To simplify some of these operations for DOCA Verbs, a set of GPUNetIO high-level functions for example, doca gpu verbs create qp hl has been introduced to condense the steps required to create and connect RDMA QPs. These functions are present in both SDK and open-source versions. Once the control path is complete, the data path can start on the GPU, where one or more CUDA kernels use GPUNetIO CUDA functions to operate on the transport objects exported into GPU memory to send or receive traffic. GPU data path The generic structure of a CUDA kernel for GPU communication typically consists of the following steps: - Post WQEs: one or more CUDA threads post Work Queue Entries WQEs , such as RDMA Write, RDMA Read, or Ethernet Send/Recv, to the network queue object. - Ring the doorbell: one or more CUDA threads notify the network card that new WQEs are ready by writing to its registers that is, “ringing the doorbell” . - Poll CQEs optional : one or more CUDA threads wait for Completion Queue Entries CQEs to confirm that the WQEs have completed successfully. API: High-level vs low-level For Ethernet and RDMA Verbs transports, DOCA GPUNetIO offers two different levels of API: high-level and low-level. The high-level API offers implementation of complex and composite operations. As an example, on the Verbs side, the doca gpu dev verbs put signal offers a pre-implemented thread-safe combination of combining an RDMA Write Work Queue Entry WQE with RDMA Atomic Fetch & Add WQE executed per-thread scope or per-warp scope all threads in the warp cooperate to the operation posting WARP SIZE RDMA Write and just a single RDMA Atomic at the end . The function takes care about concurrent submissions of WQEs from other CUDA threads on the same network queue and rings the network card doorbell to avoid race conditions. Similarly, on the Ethernet side, another example is the doca gpu dev eth txq send , offering a thread-safe combination of posting an Ethernet Send WQE on the network queue and then ringing the doorbell at thread, warp or block scope. The same applies to the receiving side. The low-level API, by contrast, offers basic building blocks an application can use to create its own customized composite operations like posting WQEs with different opcodes, ringing the network card doorbell, poll the Completion Queue waiting for Completion Queue Entries CQE . These APIs are not thread safe, so it’s the application’s responsibility to properly synchronize the concurrent access if present to the same network queue object. Ring the doorbell Once WQEs have been posted to the network queue, the network card must be notified so it can execute them. This step is known as “ringing the doorbell.” In GPUNetIO, it refers to the CUDA kernel notifying the network card that new work is ready for execution. DOCA GPUNetIO offers several ways to ring the doorbell: - Regular doorbell: This is the canonical mode, in which the NIC registers are MMIO-mapped into the CUDA memory space, allowing CUDA threads to write to them directly. - BlueFlame: This follows the same general model as the regular doorbell, but in this case the entire WQE, rather than just a notification, is written into the NIC registers. BlueFlame is typically used in latency-sensitive applications with a small number of network queues. - CPU-assisted doorbell: In this mode, the NIC register is mapped into CPU memory rather than GPU memory. The GPU writes a notification to a shared memory area that is polled by a CPU thread. Once the CPU thread detects the GPU update, it rings the NIC doorbell. This mode is mainly useful to enable GDA-KI in systems without a direct GPU-to-NIC connection, such as DGX Spark. Examples To facilitate both API exploration and performance benchmarking, GPUNetIO offers various reference implementations. Developers can access the open-source examples https://github.com/NVIDIA-DOCA/gpunetio/tree/main/examples directly within the project repository, while the more comprehensive SDK samples https://networking-docs.nvidia.com/doca/sdk/gpunetio-sample-guide and a dedicated Ethernet-based application https://networking-docs.nvidia.com/doca/sdk/doca-gpu-packet-processing-application-guide are distributed as part of the full DOCA SDK. Furthermore, these samples have been integrated into GitHub https://github.com/NVIDIA-DOCA/doca-samples for easier accessibility in recent releases. The following sections highlight specific code snippets to illustrate how high-level and low-level primitives can be leveraged to implement similar logic using different programming semantics. Ethernet example A GPU-initiated Ethernet packet generator can be implemented using high-level API like in the code snippet below. Specifically, the high-level function doca gpu dev eth txq send facilitates the concurrent submission of WQEs across multiple CUDA blocks to a shared Send Queue, handling necessary synchronization internally to ensure thread-safe operation without requiring custom application-level locks. Also, it abstracts away the granular management of queue structures, such as tracking the specific indices of WQEs and CQEs. global void send packets struct doca gpu eth txq txq, uint8 t addr, const uint32 t mkey, const size t size, uint32 t exit cond { enum doca gpu eth send flags flags = DOCA GPUNETIO ETH SEND FLAG NONE; doca gpu dev eth ticket t out ticket; uint32 t num completed = 0; / Only the last thread in the block requests a CQE. / if threadIdx.x == blockDim.x - 1 flags = DOCA GPUNETIO ETH SEND FLAG NOTIFY; while DOCA GPUNETIO VOLATILE exit cond == 0 { doca gpu dev eth txq send<DOCA GPUNETIO ETH RESOURCE SHARING MODE GPU, DOCA GPUNETIO ETH SYNC SCOPE GPU, DOCA GPUNETIO ETH NIC HANDLER AUTO, DOCA GPUNETIO ETH EXEC SCOPE BLOCK txq, addr, mkey, size, flags, &out ticket ; / syncthreads already present in send with BLOCK scope / if threadIdx.x == 0 doca gpu dev eth txq poll completion<DOCA GPUNETIO ETH CQ POLL LAST txq, 1, DOCA GPUNETIO ETH WAIT FLAG B, &num completed ; syncthreads ; } } Applications can also utilize low-level primitives to achieve precise control over their implementation. For instance, if a packet generator allocates exactly one Send Queue per CUDA block, the internal synchronization overhead of the high-level doca gpu dev eth txq send becomes unnecessary. In such scenarios, the logic can be streamlined using the following pattern: global void send packets struct doca gpu eth txq txq, uint8 t start addr, const uint32 t mkey, const size t size, uint32 t exit cond { uint64 t wqe idx = threadIdx.x, cqe idx = 0; enum doca gpu eth send flags flags = DOCA GPUNETIO ETH SEND FLAG NONE; struct doca gpu dev eth txq wqe wqe ptr; / For simplicity, every thread always sends the same buffer / uint64 t addr = uint64 t start addr + uint64 t size threadIdx.x ; if threadIdx.x == blockDim.x - 1 flags = DOCA GPUNETIO ETH SEND FLAG NOTIFY; while DOCA GPUNETIO VOLATILE exit cond == 0 { wqe ptr = doca gpu dev eth txq get wqe ptr txq, wqe idx ; doca gpu dev eth txq wqe prepare send txq, wqe ptr, wqe idx, addr, mkey, size, flags ; syncthreads ; if threadIdx.x == blockDim.x - 1 { / Ring network card doorbell / doca gpu dev eth txq submit txq, wqe idx + 1 ; / Poll for last send completion / doca gpu dev eth txq poll completion at<DOCA GPUNETIO ETH RESOURCE SHARING MODE GPU, DOCA GPUNETIO ETH SYNC SCOPE CTA txq, cqe idx, DOCA GPUNETIO ETH WAIT FLAG B ; cqe idx++; } syncthreads ; wqe idx += blockDim.x; } Verbs example The GPUNetIO Verbs API follows a similar paradigm. For instance, developers can leverage high-level primitives to implement a performance-oriented application that replicates the logic of ib write bw using GPU-driven communication through the following put functionality: template <enum doca gpu dev verbs exec scope scope global void put bw struct doca gpu dev verbs qp qp, uint32 t num iters, uint32 t data size, uint8 t src buf, uint32 t src buf mkey, uint8 t dst buf, uint32 t dst buf mkey { doca gpu dev verbs ticket t out ticket; uint32 t lane idx = doca gpu dev verbs get lane id ; uint32 t tidx = threadIdx.x + blockIdx.x blockDim.x ; for uint32 t idx = blockIdx.x blockDim.x + threadIdx.x; idx < num iters; idx += blockDim.x gridDim.x { doca gpu dev verbs put<DOCA GPUNETIO VERBS RESOURCE SHARING MODE GPU, DOCA GPUNETIO VERBS NIC HANDLER AUTO, scope qp, doca gpu dev verbs addr{.addr = uint64 t dst buf + data size tidx , .key = uint32 t dst buf mkey}, doca gpu dev verbs addr{.addr = uint64 t src buf + data size tidx , .key = uint32 t src buf mkey}, data size, &out ticket ; syncthreads ; } } Applications can also utilize low-level primitives to achieve precise control over their implementation, such as creating a CUDA block-specific bandwidth test that manages every individual operation manually: global void write bw struct doca gpu dev verbs qp qp, uint32 t num iters, uint32 t size, uint8 t src buf, uint32 t src buf mkey, uint8 t dst buf, uint32 t dst buf mkey { uint64 t wqe idx; struct doca gpu dev verbs wqe wqe ptr; for uint32 t idx = threadIdx.x; idx < num iters; idx += blockDim.x { wqe idx = doca gpu dev verbs reserve wq slots qp, 1 ; wqe ptr = doca gpu dev verbs get wqe ptr qp, wqe idx ; doca gpu dev verbs wqe prepare write qp, wqe ptr, wqe idx, MLX5 OPCODE RDMA WRITE, DOCA GPUNETIO MLX5 WQE CTRL CQ UPDATE, 0, uint64 t dst buf + size threadIdx.x , dst buf mkey, uint64 t src buf + size threadIdx.x , src buf mkey, size ; syncthreads ; if threadIdx.x == blockDim.x - 1 doca gpu dev verbs submit<DOCA GPUNETIO VERBS RESOURCE SHARING MODE EXCLUSIVE qp, wqe idx + 1 ; syncthreads ; wqe idx += blockDim.x; } if threadIdx.x == blockDim.x - 1 doca gpu dev verbs poll cq at doca gpu dev verbs qp get cq sq qp , wqe idx - blockDim.x ; syncthreads ; } NCCL Device API for GIN NCCL GIN GPU-Initiated Networking delivers the device-side communication abstraction layer for CUDA kernels, exposing primitives such as put, get, signal, wait, and flush. Since NCCL 2.27, GIN has integrated GDA-KI as a backend leveraging the open-source GPUNetIO Verbs path. This architecture enables NCCL algorithms and device-API applications to drive RDMA operations directly from the GPU while abstracting away the low-level queue management and work-submission complexities beneath the GIN interface. The workflow begins with the CPU executing the control path during communicator initialization, where it handles RDMA queue pair creation, memory registration, and the export of transport descriptors to GPU memory. Once configured, the data path shifts entirely to the GPU. GIN operations are mapped into GPUNetIO Verbs calls, which manage WQE preparation, doorbell ringing, and completion polling. This model successfully removes the CPU from the application critical path for every network transaction. By decoupling communication semantics from transport implementation, NCCL focuses on collective algorithms while GPUNetIO provides the hardened GDA-KI plumbing. NCCL utilizes GIN-level controls to manage thread, CTA, or GPU-wide resource sharing, enabling request aggregation and optimized queue management. These capabilities align with GPUNetIO’s strengths, facilitating low-latency and high-throughput paths without fragmenting the RDMA implementation. Ultimately, GPUNetIO serves as the shared foundation where performance enhancements and NIC support can be developed once and utilized across the entire GPU communication ecosystem. To explore functional implementations of the NCCL GIN interface, developers can access the NCCL tests GitHub repository https://github.com/NVIDIA/nccl/tree/master/docs/examples/06 device api for reference code. NVSHMEM NVSHMEM is a PGAS-based programming model that provides scalable point-to-point and collective communication primitives, such as put, get, atomics, barriers, and reductions, in GPU clusters. It enables multi-GPU applications to implement fine-grained GPU-to-GPU data movement and synchronization within a CUDA kernel, providing significant gains in strong scaling of application workloads. For network communication, NVSHMEM offers several transports depending on the available network capabilities and application requirements. NVSHMEM provides an IBGDA transport, which implements GDA-KI using MLX5 direct verbs. With the release of the open-source GPUNetIO library, NVSHMEM 3.7 introduced a new GPUNetIO-based transport. By building on this common implementation layer, NVSHMEM benefits from GPUNetIO optimizations and advanced NIC features available through the DOCA SDK. The new GPUNetIO transport supports both GDA-KI communication and host-initiated communication a.k.a CPU communications via RDMA in this second scenario, GPUNetIO is only used on the control path to configure network elements while the data path is executed on the CPU via NVSHMEM native functions . Users can select the GPUNetIO transport by setting NVSHMEM REMOTE TRANSPORT=gpunetio and enable GDA-KI with NVSHMEM GPUNETIO ENABLE GDAKI=1 . The GDA-KI part of the GPUNetIO transport in NVSHMEM follows the structure of the IBGDA transport but replaces larger low-level code portions with calls to the GPUNetIO library. For example, as described in the “Ring the doorbell” section above, the doorbell construction and ringing code was entirely externalized into the GPUNetIO library. While achieving the same functionality and performance as IBGDA, the GPUNetIO transport reduces the complexity of the transport implementation significantly, showcasing the value of consolidating common functionality of the GDA-KI implementations. Depending on the communication pattern, enabling GDA-KI may significantly enhance the achievable bandwidth for small message sizes. The following experiment uses the shmem put bw benchmark of the NVSHMEM performance test suite using nvshmem double put nbi with different message sizes. The experiment uses 1 thread and 1 QP per CTA cooperative thread array with an increasing number of CTAs. The two plots below show the achieved bandwidth of the GPUNetIO transport with the first one showing the results with GDA-KI disabled, i.e., using CPU communications, and the second one showing the results using GDA-KI. As shown in Figure 5, above, when NVSHMEM uses CPU communications control CPU path is configured with GPUNetIO functions while CPU data path is NVSHMEM native , bandwidth is capped for small message sizes when scaling to more CTAs and QPs due to the CPU proxy bottleneck. In the Figure 6 instead, the same benchmark is executed with GPUNetIO backend enabling GDA-KI feature data path executed by GPU with GPUNetIO CUDA functions . The experiment confirms that the removed proxy bottleneck when using GDA-KI leads to a better CTA and QP scaling behavior. The GDA-KI variant can reach higher peak bandwidth for smaller message sizes and achieves peak bandwidth earlier in terms of message size than the non-GDA-KI variant. Previous experiments https://developer.nvidia.com/blog/improving-network-performance-of-hpc-systems-using-nvidia-magnum-io-nvshmem-and-gpudirect-async/ comparing the bandwidth performance of IBRC and IBGDA have shown similar results. Overall, GPUNetIO preserves the functionality and performance of IBGDA and IBDEVX while substantially reducing NVSHMEM’s implementation complexity and maintenance costs. NVQLink NVIDIA NVQLink https://www.nvidia.com/en-us/solutions/quantum-computing/nvqlink/ is NVIDIA’s low-latency reference architecture for connecting quantum processors control systems to GPU computing, enabling real-time quantum-classical workflows such as Quantum Error Correction QEC , feed-forward, and adaptive calibration through tight integration with NVIDIA CUDA-Q https://developer.nvidia.com/cuda-q . Holoscan Sensor Bridge HSB https://www.nvidia.com/en-us/technologies/holoscan-sensor-bridge provides the networking foundation for this approach by using RoCE and lightweight FPGA IP plus host-side control software, while NVQLink pushes that HSB-based model into a much lower-latency regime suitable for live QPU control rather than conventional sensor-ingest pipelines. The HSB network operator used by NVQLink is called GPU RoCE Transceiver. Built on the DOCA GPUNetIO and DOCA Verbs libraries, it is designed to minimize traffic-forwarding latency as much as possible. In standalone mode, the operator provides a traffic-forwarding execution path for network sanity checks and performance analysis. The operator launches a persistent CUDA kernel that uses DOCA GPUNetIO CUDA functions to continuously receive packets from a remote peer and send them back without any CPU intervention. In this setup, the remote peer is an FPGA that measures the elapsed time between when a packet is transmitted and when it is received back. To measure the network-forwarding latency of the HSB 2.7.0 https://github.com/nvidia-holoscan/holoscan-sensor-bridge/tree/release-2.7.0 GPU RoCE Transceiver operator, we used an NVIDIA IGX Thor platform with a Blackwell GPU and a ConnectX-7 network card. On the other side, connected back-to-back over an OSFP cable, we used an FPGA on a Xilinx RFSoC 4×2 board to send, receive, and measure round-trip latency. The minimum reported latency, measured as the time between packet transmission and reception at the FPGA, is approximately 2.6 microseconds while median latency is 2.7 microseconds. Get started with DOCA GPUNetIO - DOCA SDK GPUNetIO programming guide https://networking-docs.nvidia.com/doca/sdk/doca-gpunetio — full SDK documentation - Open-source GPUNetIO repository https://github.com/NVIDIA-DOCA/gpunetio — examples and open-source implementation - DOCA SDK samples on GitHub https://github.com/NVIDIA-DOCA/doca-samples — reference implementations - NCCL tests — GIN device API https://github.com/NVIDIA/nccl-tests/tree/master/src/device api/gin — GIN reference code - Holoscan Sensor Bridge GPU RoCE Transceiver https://github.com/nvidia-holoscan/holoscan-sensor-bridge/tree/main/src/hololink/operators/gpu roce transceiver — NVQLink operator