How DOCA GPUNetIO Unifies GPU-Initiated Networking Across the NVIDIA Software 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, 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, NCCL, Aerial 5G SDK, UCX/NIXL, NVQLink with Holoscan Sensor Bridge operator , Holoscan Advanced Network Operator , DeepEP/HybridEP, 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 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 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 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 directly within the project repository, while the more comprehensive SDK samples and a dedicated Ethernet-based application are distributed as part of the full DOCA SDK. Furthermore, these samples have been integrated into GitHub 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 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 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 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. Holoscan Sensor Bridge (HSB) 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 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 — full SDK documentation
• Open-source GPUNetIO repository — examples and open-source implementation
• DOCA SDK samples on GitHub — reference implementations
• NCCL tests — GIN device API — GIN reference code
• Holoscan Sensor Bridge GPU RoCE Transceiver — NVQLink operator