NVIDIA DOCA GPUNetIO Unifies GPU-Initiated Networking Across Software Ecosystem

NVIDIA DOCA GPUNetIO Unifies GPU-Initiated Networking Across Software Ecosystem NVIDIA DOCA GPUNetIO provides a unified GPU-centric networking foundation that enables CUDA kernels to directly execute network operations, eliminating CPU bottlenecks on the critical data path. DOCA GPUNetIO combines GPUDirect RDMA, GDA-KI, and GDRCopy to let CUDA kernels directly control Ethernet, RDMA Verbs, and DMA transactions. The SDK separates control path (CPU setup and descriptor export) from data path (CUDA kernel operations) and is available as a full SDK or lightweight open-source project. Integrated across libraries such as NCCL and NVSHMEM, GPUNetIO replaces fragmented per-library implementations with a consolidated GDA-KI infrastructure. NVQLink and Holoscan Sensor Bridge leverage GPUNetIO to achieve microsecond-scale round-trip latencies for real-time quantum and sensor workloads. Architecture and Core Execution Model By shifting network transactions away from CPU mediation, DOCA GPUNetIO removes latency bottlenecks on the critical path. The framework combines technologies including GPUDirect RDMA, GPUDirect Async Kernel-Initiated (GDA-KI), and GDRCopy to enable direct GPU control over Ethernet, RDMA Verbs, and DMA. Its architecture splits operations into a CPU-driven control path—which initializes devices, provisions transport queues, and exports descriptors into GPU memory—and a GPU-driven data path where CUDA kernels execute packet operations directly. The library provides both a comprehensive DOCA SDK implementation and a lightweight open-source version focused on RDMA Verbs that can dynamically bind to full SDK routines at runtime. Device APIs and Work Submission Modes CUDA device kernels interact with network queues via high-level or low-level APIs. High-level functions provide pre-built, thread-safe mechanisms that aggregate operations—such as combining RDMA writes with atomics or submitting Ethernet packets—across thread, warp, or block scopes. Low-level APIs expose raw building blocks for posting Work Queue Entries (WQEs), manual queue index handling, and completion polling. To dispatch operations, GPUNetIO supports three doorbell mechanisms: Regular MMIO register mapping into GPU memory, BlueFlame inline writing for low latency with few queues, and a CPU-assisted doorbell fallback for systems lacking direct GPU-NIC links. Integration in NCCL and NVSHMEM Historically, distinct communication frameworks developed isolated GDA-KI implementations. DOCA GPUNetIO unifies this layer across the ecosystem. In NCCL 2.27+, the GPU-Initiated Networking (GIN) abstraction integrates GPUNetIO to let CUDA kernels execute RDMA calls natively, yielding speedups over legacy host-proxied communication in collectives like Reduce-Scatter. Similarly, NVSHMEM 3.7 introduced a GPUNetIO transport matching the performance of its MLX5-based IBGDA backend. In multi-node benchmarks on B300 and ConnectX-8 platforms, GPUNetIO eliminates the bandwidth plateaus caused by CPU proxies when scaling small message transfers. Real-Time Quantum Workflows via NVQLink GPUNetIO also supports ultra-low-latency real-time control through NVIDIA NVQLink and the Holoscan Sensor Bridge (HSB). HSB uses a GPU RoCE Transceiver operator driven by GPUNetIO and DOCA Verbs to maintain persistent CUDA kernels that process and echo traffic without CPU intervention. Connected over RoCE between an IGX Thor system and an FPGA platform, the transceiver demonstrates round-trip latencies as low as 2.6 to 2.7 microseconds, fulfilling the stringent response requirements of quantum error correction and classical-quantum control loops.

NVIDIA DOCA GPUNetIO：统一 GPU 驱动网络通信框架及生态整合 NVIDIA 推出了以 GPU 为中心的网络 SDK 层 DOCA GPUNetIO，通过整合 GPUDirect RDMA、GPUDirect Async Kernel-Initiated (GDA-KI) 与 GDRCopy，使 CUDA 核函数能够直接执行以太网、RDMA、Verbs 和 DMA 操作，消除 CPU 处于关键路径上的性能瓶颈并统一底层通信实现。 DOCA GPUNetIO 在控制路径利用 CPU 导出传输对象描述符，在数据路径完全由 CUDA 核函数直接驱动网络传输。 除完整的闭源 DOCA SDK 外，NVIDIA 提供了轻量级开源版本，两者在设备端对齐 API，开源版本支持运行时通过 dlopen 调用闭源功能。 支持高层与底层两类 API，提供常规 MMIO、BlueFlame 及 CPU 辅助三种 Doorbell 触发模式以适配不同硬件拓扑。 NCCL 2.27+ 引入 GIN 并结合 GPUNetIO 实现 GDA-KI 后端，显著减少集合通信延迟。 NVSHMEM 3.7+ 增加基于 GPUNetIO 的传输层，在保持与 IBGDA 性能一致的同时大幅降低代码维护复杂度。 架构与开发模式 DOCA GPUNetIO 解决了各通信库重复实现 GPU 端网络驱动的碎片化问题。CPU 仅用于初始化设备、配置队列与导出传输描述符到 GPU 内存；随后 CUDA 核函数接管数据路径，直接提交 WQE、敲击 Doorbell 并轮询 CQE。SDK 提供高层抽象（线程/Warp/Block 作用域且线程安全）与低层基础接口（手动管理索引与同步），并支持常规 MMIO、BlueFlame 以及针对无直接 GPU-NIC 连接拓扑（如 DGX Spark）的 CPU 辅助模式。 NCCL 与 NVSHMEM 生态整合 自 NCCL 2.27 起，GIN 层集成基于 GPUNetIO Verbs 的 GDA-KI 后端，消除控制和数据路径间的 CPU 参与，在 Reduce-Scatter 等集合操作中相比传统 CPU 代理通信带来显著加速。NVSHMEM 3.7 引入 GPUNetIO 传输实现（通过环境变量控制开启），将底层的 Doorbell 构造与提交例程抽离，消除了 CPU 代理带来的小包瓶颈，在 B300 与 ConnectX-8 平台上实现了更早达到峰值带宽的强扩展性。 实时量子控制系统（NVQLink）应用 在面向量子处理器低延迟控制的 NVQLink 架构中，Holoscan Sensor Bridge (HSB) 2.7.0 构建了基于 DOCA GPUNetIO 的 GPU RoCE 收发算子。通过运行常驻 CUDA 核函数实现数据包的零 CPU 干预回传，在 Blackwell 架构 IGX Thor 平台与 FPGA 的测试中，实现了约 2.6 微秒的最小往返时延与 2.7 微秒的中位数时延。

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 , DeepEP/HybridEP and others. This post covers what’s new: the open-source library, the unified architecture, and how these integrations work in practice. Figure 1. DOCA GPUNetIO enables CPU-side control path management for memory and device exported objects alongside GPU-side data path CUDA functions, building upon the foundational capabilities of additional DOCA SDK libraries 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. Figure 2. GPUNetIO open-source and DOCA SDK share similar CUDA Device API. When the DOCA SDK is available, the open-source host API may call SDK functions via dlopen if the documented environment variable is set with the correct DOCA SDK path; otherwise it runs standalone 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. Figure 3. In the GPUDirect Async Kernel-Initiated workflow, CUDA threads post WQEs to the network queue, synchronize, and ring the NIC doorbell. The NIC fetches and executes the WQEs, then writes completions back to the CQ for GPU threads to poll 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. Figure 4. GIN GDA-KI backend speedup over the legacy CPU-communication path in NCCL 2.29 Reduce-Scatter, across message sizes from 4K to 16GB. Each line represents a different scale: N = nodes, G = total GPUs 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. Figure 5: The bandwidth of the GPUNetIO transport with GDA-KI disabled is capped for small message sizes when scaling to more CTAs and QPs due to the CPU proxy bottleneck. Benchmarks executed on two nodes with B300+CX8 using Spectrum-X 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). Figure 6. The bandwidth of the GPUNetIO transport can scale to a higher bandwidth for smaller message sizes when the number of CTAs is increasing. Benchmarks were executed on two nodes with B300+CX8 using Spectrum-X 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. Figure 7. HSB GPU RoCE Transceiver operator forwarding traffic to a FPGA via network connection 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. Figure 8. Network-forwarding round-trip latency analysis of the HSB 2.7.0 GPU RoCE Transceiver operator 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