---
title: GPUDirect Async Kernel-Initiated (GDAKI)
url: https://www.emergentmind.com/topics/gpudirect-async-kernel-initiated-gdaki
type: topic
---

# GPUDirect Async Kernel-Initiated (GDAKI)

GPUDirect Async Kernel-Initiated (GDAKI) is a backend architecture in NVIDIA’s NCCL GPU-Initiated Networking (GIN) stack that enables CUDA device-side code to directly initiate and manage network RDMA operations, fully bypassing the CPU during both communication and completion handling. The approach achieves low-latency, fine-grained GPU-to-GPU communication integrated seamlessly into NCCL’s production-grade collective infrastructure, leveraging DOCA GPUNetIO for direct GPU–NIC interaction. GDAKI is particularly relevant to modern AI workloads with tight computation-communication coupling requirements, such as Mixture-of-Experts (MoE) models, where traditional host-initiated paradigms incur prohibitive coordination overhead [2511.15076].

## 1. Hardware-Software Architecture

GDAKI operates through a hardware-software pipeline involving the following core components:

- **GPU**: Provides streaming multiprocessors (SMs), CUDA cores, on-GPU memory for work-queue entries (WQEs), and PCIe/NVLink BAR mappings into the device address space.
- **NIC**: Typically an InfiniBand/RoCE adapter (e.g., ConnectX-6 Dx or newer), exposing doorbell registers, send/receive QPs, and CQs mapped into GPU-visible BARs.
- **NIC DMA Engines**: Fetch WQEs from GPU memory, drive one-sided RDMA transactions (READ/WRITE/ATOMICS), and post completion notifications back into GPU-visible memory.
- **DOCA GPUNetIO**: API and library providing kernel-side verbs for GPU→NIC communication with minimal overhead.

The data path follows this sequence:
1. **Host-side Initialization**: The host (via `ncclCommInitRank`) allocates collective symmetric memory windows, initializes GDAKI context using DOCA GPUNetIO to set up NIC QPs, CQs, and BAR locations, exchanges lkey/rkey handles, and makes the context device-accessible.
2. **Device-side Issuance**: CUDA kernels instantiate an `ncclGin` object referencing the GIN context. Application threads enqueue RDMA WQEs (encoding source, destination, length, rkey, and notification flags) in an on-GPU circular buffer.
3. **NIC Notification**: Device-side code writes to the NIC's doorbell register (by PCIe/NVLink BAR), which triggers the NIC DMA engines to read WQEs and issue network operations.
4. **Completion Polling**: Upon RDMA completion, the NIC posts completion entries to a GPU-visible CQ buffer. Threads poll this buffer using APIs such as `gin.readSignal` or `gin.waitSignal` to detect operation completion and ensure remote visibility.

This fully device-initiated flow provides “zero-CPU-overhead” communication in production settings, eliminating host thread interruptions and CPU context-switching.

## 2. Programming Model and Device API Integration

The device-facing interface for GDAKI is encapsulated in the `ncclGin` class of the NCCL Device API, which maps directly onto DOCA GPUNetIO device verbs. The API supports collective operations at the remote memory window granularity, including atomic updates and remote signaling for synchronization.

Example kernel-side sequence for a remote PUT with notification:

```cpp
__global__ void stageAndSend(ncclDevComm devComm,
                             ncclWindow_t localWin, size_t localOff,
                             uint64_t remoteAddr, int peer, int signalId) {
  // 1. Initialize device context
  ncclGin gin(devComm, /*contextIndex=*/0);

  // 2. Remote write with signaling
  gin.put(/*team=*/ncclTeamWorld(devComm),
          peer,
          /*dstWindow*/remoteAddr, /*dstOffset*/0,
          /*srcWindow*/localWin, /*srcOffset*/localOff,
          /*bytes*/N,
          ncclGin_SignalInc{signalId});

  // 3. Optional completion notification
  gin.waitSignal(ncclCoopCta(), signalId, /*expected=*/1);
  gin.resetSignal(signalId);
}
```

Internally, operations like `gin.put` become inlined GPUNetIO calls such as `gpi_write()` or `gpi_atomic()`, plus a BAR store to ring the NIC’s doorbell. Completion APIs (`readSignal`, `waitSignal`, `readCounter`, `waitCounter`) poll GPU-visible memory. There is no host-involved proxy or kernel launches required for progress or completion notification.

## 3. Performance Characteristics and Modeling

Latency and bandwidth in GDAKI are governed by three principal terms in a first-order model:

\[
T_{\mathrm{total}}(S) = T_{\mathrm{issue}} + T_{\mathrm{network}}(S) + T_{\mathrm{complete}}
\]

Where:
- \(T_{\mathrm{issue}}\): Time to write WQE and ring doorbell (~hundreds of ns)
- \(T_{\mathrm{network}}(S) = L_0 + \frac{S}{B}\): On-wire network latency (\(L_0\) ≈ 130 ns for InfiniBand) plus message serialization at bandwidth \(B\) (e.g., 400 Gb/s)
- \(T_{\mathrm{complete}}\): Time to poll completion queue entry (tens–hundreds of ns).

For round-trip operations, the network term is doubled and two issue-complete pairs are incurred. For steady-state k-message pipelines, bandwidth asymptotically approaches \(B\), subject to NIC DMA fetch efficiency and GPU memory performance.

## 4. Empirical Microbenchmark and Application Results

Extensive benchmarking on H100 GPUs and InfiniBand networks demonstrates:

| Mode                              | GDAKI (DOCA GPUNetIO) | NCCL Proxy | NVSHMEM IBGDA | NVSHMEM IBRC |
|------------------------------------|-----------------------|------------|---------------|--------------|
| Point-to-Point RTT (4–128 B, µs)   | 16.7                  | 18.0       | 24.3          | 16.0         |
| DeepEP Aggregate RDMA (GB/s)       | ~84 (dispatch)        | –          | ~84           | –            |
| 8-GPU E2E Latency (40.6 µs)        | 40.6                  | –          | 41.4          | –            |
| 2-Node E2E Latency (µs)            | 142.5                 | –          | 157.0         | –            |

Key findings:
- GDAKI direct GPU→NIC path outperforms the CPU-assisted proxy by ~1–2 µs and host-initiated RDMA by ~7–8 µs for small messages.
- Bandwidth matches or exceeds (within 1–2%) best-known figures at scale (~400 Gb/s per link).
- DeepEP Mixture-of-Experts kernels on GDAKI see up to 9% lower latency in multi-node low-latency mode, with high-throughput mode matched exactly.
- CPU involvement is eliminated: zero CPU threads per communicator, freeing host cores for other use.
- Multi-node “combine” performance is within 1–3% of NVSHMEM.

## 5. Hardware and Software Prerequisites

The following requirements must be met for GDAKI operation:

- **NIC**: NVIDIA ConnectX-6 Dx or newer, or BlueField-2/3 DPU with GPUNetIO support.
- **PCIe/NVLink Topology**: GPU and NIC must share the same root complex, ensuring peer-to-peer BAR mappings and low-latency path.
- **GPUs**: Volta-class (sm_70+) or newer; CUDA toolkit 12.x or higher.
- **Kernel Modules and Drivers**: `nv_peer_mem` or `dmabuf` for GPUDirect RDMA, OFED/MOFED stack matching NIC firmware (e.g., NVIDIA OFED 5.x).
- **Software Stack**: DOCA SDK with appropriate `doca-gpunetio` library version aligned to firmware.

This enables direct access and control of NIC hardware resources from CUDA devices.

## 6. Integration into NCCL and MoE Libraries

GDAKI is integrated into NCCL as a plugin backend:

- **NCCL Core**: At communicator initialization, the runtime detects GPUNetIO capability, loads the GDAKI plugin, allocates QPs/BARs, performs memory registration (for lkey/rkey and address handles), and distributes this context to all participating ranks.
- **Memory Windows**: Application code registers CUDA pointers with NCCL via `ncclCommWindowRegister`, resulting in a `ncclWindow_t` handle via DOCA GPUNetIO.
- **DeepEP Example**: Replaces NVSHMEM one-sided routines with `ncclGin` methods. Mapping pointer arithmetic to (window, offset) pairs and instantiating multiple communicators if required (e.g., 24 QPs per rank).
- **CUDA Kernel Usage**: Importing `<nccl_device_api.h>`, device code calls `ncclGin` methods for put and signaling; host-side proxies or thread launches are not needed.

Integration is transparent for higher-level frameworks including PyTorch, JAX, and TensorRT-LLM, and typically requires only minor kernel changes.

## 7. Summary of Observed Benefits

GDAKI achieves the following experimentally validated benefits:

- Sub-17 µs round-trip latency for small (4–128 B) messages.
- Bandwidth scalability up to line-rate (≈400 Gb/s per link) across all tested collective modes.
- Elimination of host-side synchronization and context switching, resulting in zero CPU thread overhead for communication progress.
- Application-level performance (MoE on DeepEP) at or above that of competing approaches, and up to 9% latency reduction in multi-node scenarios with minimal code changes for adoption.
- Fully device-side controllable collective operations, enabling tight fusion of communication and computation within CUDA kernels [2511.15076].

GDAKI, in conjunction with DOCA GPUNetIO and NCCL's collective runtime, represents a tightly integrated architecture for efficient one-sided RDMA communication in GPU-centric distributed AI workloads.

Source: https://www.emergentmind.com/topics/gpudirect-async-kernel-initiated-gdaki