Fault tolerance in RCCL

Fault tolerance in RCCL#

Large-scale jobs running across many AMD GPUs and nodes must survive failures such as a network link going down, an ECC error, a node crash, or a process that exits unexpectedly. RCCL provides APIs to detect these conditions, release the affected resources, and keep running without restarting the whole job.

RCCL inherits the NCCL error-handling model and adds AMD-specific extensions on top of it. The APIs are declared in rccl.h and use the standard NCCL signatures.

Note

Fault tolerance relies on communicators created in non-blocking mode. With config.blocking = 0, every RCCL call (except ncclCommDestroy() and ncclCommAbort()) returns immediately, so the application can react to a hang or a failure instead of being stuck inside a collective.

Error handling and communicator abort#

Every RCCL call returns an ncclResult_t code. Set NCCL_DEBUG=WARN to print a human-readable message on error, or NCCL_DEBUG=INFO to also print the call stack.

Error	Description	Handling
`ncclSuccess`	No error.	None.
`ncclUnhandledCudaError`	Error during a HIP (ROCm runtime) call.	Fatal. Abort the communicator and re-create it.
`ncclSystemError`	Error during a system call, for example a network failure.	Fatal. Abort the communicator and re-create it.
`ncclInternalError`	A bug inside RCCL.	Fatal. Abort the communicator and report the issue to AMD.
`ncclInvalidArgument`	An argument is invalid, for example a `NULL` pointer.	Non-fatal. The call had no effect; the communicator is still usable.
`ncclInvalidUsage`	The sequence of RCCL calls is invalid.	Fatal. Abort the communicator and re-create it.
`ncclRemoteError`	A remote process exited or a network error occurred.	Fatal. Abort the communicator, then shrink out the failed ranks or re-create it.
`ncclInProgress`	The call is still running (non-blocking mode).	Poll with `ncclCommGetAsyncError()`.
`ncclTimeout`	An operation exceeded its time limit, for example an unresponsive peer.	Fatal. Abort the communicator, then shrink out the failed ranks or re-create it.

A fatal error applies to all communicators in the same group. To recover, call ncclCommAbort() on every affected communicator and re-create it.

Asynchronous errors#

Network failures are not reported by the originating call; they surface later through ncclCommGetAsyncError(). In non-blocking mode, poll it until the communicator leaves the ncclInProgress state, yielding the CPU between checks.

#include <sched.h>

// Wait for a non-blocking communicator to reach a terminal state.
ncclResult_t waitForAsyncResult(ncclComm_t comm) {
    ncclResult_t state = ncclInProgress;
    while (state == ncclInProgress) {
        ncclResult_t err = ncclCommGetAsyncError(comm, &state);
        if (err != ncclSuccess) return err;   // failed to query the state
        if (state == ncclInProgress) sched_yield();
    }
    return state;
}

When waiting for a collective to finish, query the stream and poll for asynchronous errors at the same time, instead of blocking in hipStreamSynchronize, which can hang forever if a peer has failed.

int streamSyncWithAbort(hipStream_t stream, ncclComm_t comm) {
    while (true) {
        hipError_t hipErr = hipStreamQuery(stream);
        if (hipErr == hipSuccess) return 0;                 // collective done
        if (hipErr != hipErrorNotReady) return 1;           // HIP failure

        ncclResult_t asyncErr = ncclSuccess;
        ncclCommGetAsyncError(comm, &asyncErr);
        if (asyncErr != ncclSuccess) {                      // peer/network died
            ncclCommAbort(comm);
            return 2;
        }
        sched_yield();
    }
}

Aborting symmetric-memory collectives#

Collectives that run on symmetric memory are served by the built-in symmetric kernels, which exchange data directly between peer GPUs instead of going through a proxy thread. A rank that never reaches the collective therefore leaves its peers parked inside a device-side spin loop, where the host cannot see them.

Because there is no proxy thread and no device-side timeout on this path, neither ncclCommGetAsyncError() nor a communicator timeout will ever report the missing peer. Deciding when to give up is therefore the caller’s responsibility: drive the abort from your own deadline or an out-of-band watchdog, not from a loop waiting for an async error that never arrives.

ncclCommAbort() releases the parked kernels and the stream drains, with a short delay because the abort flag is only sampled periodically. The same rules as in Recovering from a failure apply: the communicator must be non-blocking and no thread may be inside an RCCL call when ncclCommAbort() is invoked.

Note

Symmetric memory requires NCCL_CUMEM_ENABLE=1 and a host platform that supports virtual memory management. When symmetric memory is unavailable, collectives fall back to the general kernels, whose abort path is handled by the proxy thread instead.

Recovering from a failure#

RCCL offers several recovery strategies, depending on how much of the job you want to preserve:

Abort and re-create the whole communicator – a coarse, job-wide restart.
Shrink to drop the failed ranks and keep running with the survivors.
Grow to bring replacement ranks back in and return to full size.
Revoke (RCCL-specific) to abort in-flight work without destroying the communicator, so it can be reused as a parent for a later shrink or grow.

Abort requires non-blocking communicators, and no thread may be inside an RCCL call when ncclCommAbort() is invoked.

Shrinking a communicator#

ncclCommShrink() creates a new communicator by removing ranks from an existing one. Only the ranks that remain call it; the excluded ranks must not. Remaining ranks are renumbered contiguously.

// Drop one failed rank and continue on the survivors.
int        excludeList[] = { failedRank };
int        excludeCount  = 1;
ncclComm_t newComm       = nullptr;

if (myRank != failedRank) {
    ncclResult_t res = ncclCommShrink(comm, excludeList, excludeCount,
                                      &newComm, nullptr, NCCL_SHRINK_ABORT);
    if (res != ncclSuccess) {
        ncclCommAbort(comm);
        return res;
    }

    // newComm has (nRanks - excludeCount) ranks, renumbered 0..N-1.
    int newRank = -1, newSize = 0;
    ncclCommUserRank(newComm, &newRank);
    ncclCommCount(newComm, &newSize);

    // ... run collectives on newComm ...

    ncclCommDestroy(newComm);
}
ncclCommDestroy(comm);   // parent no longer needed

Shrink flags:

NCCL_SHRINK_DEFAULT – the parent has no in-flight work, or it was already aborted by ncclCommRevoke().
NCCL_SHRINK_ABORT – abort in-flight work on the parent first, then shrink. Use this when shrinking directly after a failure, without a preceding revoke.

Growing a communicator#

ncclCommGrow() creates a new communicator by adding ranks. A coordinator generates a unique ID with ncclCommGetUniqueId() and distributes it to the new ranks out of band (here using MPI_Bcast).

Existing root (coordinator): comm set, uniqueId = &growId, rank = -1.
Existing non-root ranks: comm set, uniqueId = NULL, rank = -1.
New ranks: comm = NULL, uniqueId = &growId, rank = the assigned rank (a non-NULL uniqueId is required here).

// Grow a communicator of `existing` ranks up to `newTotal` ranks.
ncclComm_t newComm = nullptr;

// 1. Coordinator generates the grow ID; the new ranks need it out of band.
ncclUniqueId growId{};
if (myRank == 0) {
    ncclCommGetUniqueId(comm, &growId);
}
MPI_Bcast(&growId, sizeof(growId), MPI_BYTE, 0, MPI_COMM_WORLD);

// 2. Each rank calls grow with the arguments for its role.
if (myRank == 0) {
    // Existing root: passes the ID.
    ncclCommGrow(comm, newTotal, &growId, -1, &newComm, nullptr);
} else if (myRank < existing) {
    // Existing non-root: pass NULL (grow contract in rccl.h).
    ncclCommGrow(comm, newTotal, nullptr, -1, &newComm, nullptr);
} else {
    // New rank: comm = NULL, passes the ID and its assigned rank.
    hipSetDevice(localDevice);
    ncclCommGrow(nullptr, newTotal, &growId, myRank, &newComm, nullptr);
}

For a non-blocking grow, pass a config with blocking = 0 and poll for completion before using the new communicator.

ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
config.blocking = 0;

if (myRank == 0) {
    ncclCommGrow(comm, newTotal, &growId, -1, &newComm, &config);
} else if (myRank < existing) {
    ncclCommGrow(comm, newTotal, nullptr, -1, &newComm, &config);
} else {
    ncclCommGrow(nullptr, newTotal, &growId, myRank, &newComm, &config);
}

ncclResult_t asyncErr = ncclInProgress;
while (asyncErr == ncclInProgress) {
    ncclCommGetAsyncError(newComm, &asyncErr);
}
// asyncErr == ncclSuccess once the grow has completed.

ncclCommDestroy(comm);   // parent can be released after a successful grow

Notes:

Existing ranks keep their rank numbers; new ranks are numbered from the parent size upward.
The grow ID is single-use.
The parent must have no outstanding operations when grow is called.

Revoking a communicator#

ncclCommRevoke() aborts in-flight collectives without destroying the communicator, so it can recover from a peer that failed mid-collective. Output buffers of an aborted collective contain undefined data.

The typical flow is collective, revoke, shrink, continue. Because revoke already aborts in-flight work, the shrink uses NCCL_SHRINK_DEFAULT.

// A collective is in flight on the parent when a peer fails.
ncclAllReduce(sendBuf, recvBuf, count, ncclFloat, ncclSum, parent, stream);

// 1. Revoke aborts the in-flight collective but keeps `parent` usable.
ncclCommRevoke(parent, NCCL_REVOKE_DEFAULT);
MPI_Barrier(MPI_COMM_WORLD);   // all healthy ranks agree to recover

// 2. Build a smaller communicator from the survivors (DEFAULT, not ABORT).
ncclComm_t child = NCCL_COMM_NULL;
if (myRank != failedRank) {
    ncclCommShrink(parent, excludeList, excludeCount,
                   &child, nullptr, NCCL_SHRINK_DEFAULT);
}
MPI_Barrier(MPI_COMM_WORLD);

// 3. Continue on the child communicator.
if (child != NCCL_COMM_NULL) {
    ncclAllReduce(sendBuf, recvBuf, count, ncclFloat, ncclSum, child, stream);
    hipStreamSynchronize(stream);
    ncclCommDestroy(child);
}
ncclCommDestroy(parent);

Behavior:

New collectives on a revoked communicator return ncclInvalidUsage.
A revoked communicator stays valid as a parent for ncclCommSplit() and ncclCommShrink(), and can be torn down with ncclCommDestroy().
Revoking twice returns ncclInvalidArgument.
ncclCommFinalize() on a revoked communicator returns ncclInvalidUsage – use ncclCommDestroy() instead.
revokeFlags must be NCCL_REVOKE_DEFAULT.

Finalizing and destroying#

For a clean shutdown, call ncclCommFinalize() to drain outstanding operations, then ncclCommDestroy() to free resources.

// Clean shutdown: drain, then free.
ncclCommFinalize(comm);                 // moves comm to ncclInProgress

ncclResult_t state = ncclInProgress;
while (state == ncclInProgress) {       // wait until globally quiescent
    ncclCommGetAsyncError(comm, &state);
}

ncclCommDestroy(comm);                  // non-blocking once state is ncclSuccess

Do not access a communicator after it is destroyed. Use ncclCommAbort() instead when the communicator is in a bad state and outstanding operations cannot be drained.