Use the RCCL device API and GIN

Use the RCCL device API and GIN#

The experimental RCCL device API lets GPU kernels communicate through a device communicator (ncclDevComm). GPU-initiated networking (GIN) extends that API with one-sided puts, signals, counters, and barriers across nodes. RCCL 2.31 incorporates Device API and GIN enhancements from upstream NCCL 2.31.2. This page describes those APIs and the limits of the AMD host-proxy backend.

To supply a GIN backend rather than consume one, see Using the RCCL GIN plugin API.

NCCL_GIN_TYPE values for AMD backends are not compatible with 2.30.7: rocSHMEM GDA moved from 5 to 6 and Anvil SDMA from 6 to 7 because NCCL 2.31 placed EFA GDA at 5. The IB proxy remains NCCL_GIN_TYPE=2.

Requirements#

GIN requires a supported network backend, symmetric memory, and collective communicator and window setup. The validated AMD path is GIN_IB_PROXY:

export NCCL_GIN_TYPE=2
export NCCL_CUMEM_ENABLE=1
export NCCL_DMABUF_ENABLE=1
export NCCL_IB_MERGE_NICS=0

The ROCm runtime and NIC driver must support exporting and registering VMM allocations. NCCL_DMABUF_ENABLE=1 is the validated registration path; peer-memory registration can also be available on supported systems. Full GIN connections require cross-rail reachability, so do not force NCCL_CROSS_NIC=0 for world-team operations. Leave ginQueueDepth at its default value of zero; the AMD host-proxy backend does not support a custom QP depth.

The following snippets abbreviate error handling with NCCLCHECK. Replace it with the application’s normal ncclResult_t error handling.

Before using GIN, query communicator properties and confirm that the selected backend and Device API are available:

ncclCommProperties properties = NCCL_COMM_PROPERTIES_INITIALIZER;
NCCLCHECK(ncclCommQueryProperties(comm, &properties));
if (!properties.deviceApiSupport ||
    properties.ginType != NCCL_GIN_TYPE_PROXY) {
  // Select a non-GIN path or report that this configuration is unsupported.
}

Allocate and register every communication buffer collectively:

void* buffer = nullptr;
ncclWindow_t window = nullptr;
NCCLCHECK(ncclMemAlloc(&buffer, bytes));
NCCLCHECK(ncclCommWindowRegister(
    comm, buffer, bytes, &window, NCCL_WIN_COLL_SYMMETRIC));

Create a device communicator#

Always initialize requirements with NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER. The initializer records the header version used to compile the application.

ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL;
reqs.ginContextCount = 4;       // Hint; inspect devComm.ginContextCount.
reqs.ginSignalCount = 1;
reqs.ginCounterCount = 1;
reqs.barrierCount = 1;          // Generic ncclBarrierSession.
reqs.worldGinBarrierCount = 1;

ncclDevComm devComm{};
NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));

GIN contexts and their signal, counter, and queue resources are allocated per device communicator. Two device communicators backed by the same host communicator therefore don’t alias those resources. ginContextCount is a request; use the value returned in devComm.ginContextCount when assigning work to contexts.

ginTrafficClass overrides the host communicator traffic class for this device communicator. The service level of the GIN connections comes from NCCL_IB_SL when it is set, otherwise from ginTrafficClass or the host communicator. Their RoCE traffic class comes from NCCL_GIN_IB_TC when it is set, then from NCCL_IB_TC, and only then from ginTrafficClass or the host communicator. Setting NCCL_GIN_IB_TC moves GIN traffic to its own RoCE traffic class while the collective and point-to-point connections keep NCCL_IB_TC. InfiniBand link layers don’t carry the traffic class, so NCCL_GIN_IB_TC has no effect there; GIN traffic can be separated only through the service level, by setting ginTrafficClass and leaving NCCL_IB_SL unset. Set these values only when the fabric administrator provides appropriate values.

Device code creates an ncclGin object for one returned context:

ncclGin gin{devComm, contextIndex, NCCL_GIN_RESOURCE_SHARING_GPU};

For example, a CTA can issue a put and wait for local queue completion:

__global__ void putKernel(
    ncclDevComm devComm, ncclWindow_t source, ncclWindow_t destination,
    size_t bytes, int peer) {
  ncclGin gin{devComm, /*contextIndex=*/0};
  if (threadIdx.x == 0) {
    gin.put(ncclTeamWorld(devComm), peer,
            destination, /*destinationOffset=*/0,
            source, /*sourceOffset=*/0, bytes,
            ncclGin_WeakSignalInc{/*signal=*/0});
  }
  gin.flush(ncclCoopCta());
}

flush makes the local source buffer reusable. It does not by itself prove remote visibility; the peer must wait for the associated signal before using the destination bytes.

NCCL_GIN_RESOURCE_SHARING_GPU permits sharing across the GPU; NCCL_GIN_RESOURCE_SHARING_CTA limits sharing to a CTA. These modes select resource-sharing behavior on direct device backends. The AMD GIN host-proxy backend uses the same proxy queue behavior for both modes.

Choose strong or weak signal semantics#

Explicit signal types state which earlier puts become visible when a peer observes the signal:

  • A weak signal guarantees visibility of only the put carrying that signal. Use it for independent per-operation completion, as in the example above. It makes no guarantee about earlier puts, although a backend can order them.

  • A strong signal also guarantees visibility of every preceding put issued to the same peer through the same GIN context. Use it as a publication point for a sequence of puts.

For example, one strong signal can publish two buffers without attaching a signal to the first put:

if (threadIdx.x == 0) {
  auto team = ncclTeamWorld(devComm);
  gin.put(team, peer, destination, 0, source, 0, firstBytes,
          ncclGin_None{});
  gin.put(team, peer, destination, secondOffset,
          source, secondOffset, secondBytes,
          ncclGin_StrongSignalInc{/*signal=*/0});
}
gin.flush(ncclCoopCta());

After the peer’s waitSignal returns, both puts are visible. This ordering is limited to the same peer and GIN context. A strong signal does not publish puts issued through another context or to another peer.

The explicit indexed actions are ncclGin_StrongSignalInc, ncclGin_StrongSignalAdd, ncclGin_WeakSignalInc, and ncclGin_WeakSignalAdd. Equivalent StrongVASignal and WeakVASignal actions target a window and offset; their window must be registered with NCCL_WIN_STRICT_ORDERING.

ginStrongSignalsRequired defaults to true in NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER. The runtime rejects device communicator creation if the selected backend cannot provide strong signals. A kernel that uses only weak signals can set the field to false to permit such a backend:

ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL;
reqs.ginSignalCount = 1;
reqs.ginStrongSignalsRequired = false; // Kernel issues only weak signals.

Do not issue a strong action from a device communicator created with ginStrongSignalsRequired=false; that usage is undefined. The older ncclGin_SignalInc, ncclGin_SignalAdd, and VA equivalents are deprecated. Their strength is selected globally by ginStrongSignalsRequired; new code should use an explicit type.

Use world-team barriers and timeouts#

Reserve worldGinBarrierCount slots to construct a world-team GIN barrier without manually allocating a barrier handle:

ncclGinBarrierSession<ncclCoopCta> barrier{
    ncclCoopCta(), gin, ncclTeamTagWorld{}, barrierIndex};

ncclResult_t result = barrier.sync(
    ncclCoopCta(), cuda::memory_order_acq_rel,
    ncclGinFenceLevel::None, timeoutCycles);

The timeout overload returns ncclTimeout if all team members don’t arrive within timeoutCycles. ncclGinFenceLevel is a bitmask: None is arrival only, Put makes inbound (and self) puts visible, Get drains local gets, and omitting the fence argument is Put | Get. Relaxed is a deprecated alias for None. Pass ncclGinAllContexts(devComm) instead of a single ncclGin when puts or gets were issued on more than one GIN context.

Barrier resources for ncclGinBarrierSession, ncclLsaBarrierSession, and ncclBarrierSession are separate. Reserve barrierCount for generic ncclBarrierSession objects; use lsaBarrierCount, railGinBarrierCount, or worldGinBarrierCount for the corresponding specialized session. barrierIndex must be smaller than the selected count. Concurrent CTAs must use distinct indices, typically blockIdx.x.

Destroy resources#

Destroy the device communicator before deregistering its application windows, then free the backing allocations. These calls are collective where their corresponding create or registration operation is collective.

NCCLCHECK(ncclDevCommDestroy(comm, &devComm));
NCCLCHECK(ncclCommWindowDeregister(comm, window));
NCCLCHECK(ncclMemFree(buffer));

Version and backend notes#

ncclDevComm is versioned. The upstream NCCL 2.30.3 and 2.30.4 release notes require applications using GIN APIs to be rebuilt with the matching release. The upstream NCCL 2.30.7 release notes describe explicit Strong and Weak signals and add ncclGinFenceLevel semantics for GIN barriers (None, Put, Get, default Put | Get). The explicit signal API is marked available since NCCL 2.30.5. RCCL accepts compatible layouts within the 2.30 family, but applications using pre-2.30 GIN device code must be rebuilt with compatible RCCL headers. The runtime rejects pre-2.30 requirements that request indexed GIN resources.

NCCL_GIN_IB_TC, added in upstream NCCL 2.30.7, applies to the connections of the IB proxy backend (NCCL_GIN_TYPE=2). The one-sided host RMA operations run on the same backend and use the same traffic class. The rocSHMEM GDA backend (NCCL_GIN_TYPE=6) doesn’t read NCCL_GIN_IB_TC, NCCL_IB_TC, or ginTrafficClass; its queue pairs take their traffic class from ROCSHMEM_GDA_TRAFFIC_CLASS. Anvil SDMA (NCCL_GIN_TYPE=7) moves data within a node without a NIC, so traffic class doesn’t apply to it.

The 128-byte, versioned GIN proxy descriptor and per-context proxy progress are internal implementation details and require no application configuration. The max_rd_atomic and max_dest_rd_atomic changes and the doca-gpunetio update in the NCCL 2.30 release apply to NVIDIA’s GDAKI backend, not the AMD GIN host-proxy backend.

The source includes experimental proxy get and nonblocking-flush paths, but they are not yet validated for production use on the AMD host-proxy backend. GIN access to elastic and multi-segment windows is also unavailable on that backend until a complete multi-segment VMM range can be exported for DMA-BUF registration. Use the validated put/signal path with single-segment symmetric windows.

Contents