NCCL 2.29 – One-Sided Host APIs

[gab9talavera]

One-Sided Host APIs

One-Sided Host APIs introduce host-initiated one-sided Remote Memory Access (RMA) APIs in NCCL 2.29 that allow a rank to directly write data into remote memory without requiring a matching operation from the target rank. These APIs support both scale-up (NVLink/PCIe) and scale-out (network) communication. The implementation avoids Streaming Multiprocessor (SM) involvement, enabling SM resources to remain available for computation. For scale-up transfers, data movement is handled by GPU Copy Engines over NVLink or PCIe, while scale-out transfers are coordinated by a CPU proxy over the network.

A typical usage pattern starts with ncclPutSignal, which writes to a peer's registered window and updates a remote signal. The target rank can subsequently synchronize on this signal using ncclWaitSignal. Memory exposed for one-sided access must first be registered with the communicator via ncclCommWindowRegister.

// Allocate symmetric buffers and register windows.
void *sendbuff, *recvbuff;
NCCLCHECK(ncclMemAlloc(&sendbuff, bytes));
NCCLCHECK(ncclMemAlloc(&recvbuff, bytes));

ncclWindow_t sendWindow, recvWindow;
NCCLCHECK(ncclCommWindowRegister(comm, sendbuff, bytes, &sendWindow, NCCL_WIN_COLL_SYMMETRIC));
NCCLCHECK(ncclCommWindowRegister(comm, recvbuff, bytes, &recvWindow, NCCL_WIN_COLL_SYMMETRIC));

int peer = (rank == 0) ? 1 : 0;
int ctx = 0;     // Context index
int sigIdx = 0;  // Signal index
unsigned int flags = 0; // Reserved

ncclWaitSignalDesc_t waitDesc = {.opCnt = 1, .peer = peer, .sigIdx = sigIdx, .ctx = ctx};

if (rank == 0) {
  // Rank 0: put then wait.
  NCCLCHECK(ncclPutSignal(sendbuff, count, datatype, peer, recvWindow, 0,
                          sigIdx, ctx, flags, comm, stream));
  NCCLCHECK(ncclWaitSignal(1, &waitDesc, comm, stream));
} else {
  // Rank 1: wait then put.
  NCCLCHECK(ncclWaitSignal(1, &waitDesc, comm, stream));
  NCCLCHECK(ncclPutSignal(sendbuff, count, datatype, peer, recvWindow, 0,
                          sigIdx, ctx, flags, comm, stream));
}

For more complex communication patterns involving multiple peers, ncclPutSignal calls can be batched using ncclGroupStart / ncclGroupEnd. Grouping allows multiple one-sided operations to be launched concurrently, reducing host overhead and improving launch efficiency. Within a group, put operations targeting different peers or different contexts may execute out of order and overlap, and no ordering should be inferred across peers. NCCL guarantees ordering per peer and per context: operations issued to the same peer and context complete in program order. A signal update also provides completion semantics, guaranteeing that all preceding puts to the same peer and context have been delivered before the signal is observed by the receiver.

For barrier-style or fan-in synchronization across many ranks, grouped ncclPutSignal operations should be paired with a corresponding set of ncclWaitSignal descriptors on the receiving side.

One-sided APIs require CUDA 12.5 or higher. For additional details and usage guidelines, see the NCCL documentation on one-sided communication.

—

Authored by Zhenhao He (@zhenhaohe)