RCCL library specification

RCCL library specification#

This document provides details of the API library.

Communicator functions#

Warning

doxygenfunction: Cannot find function “ncclGetUniqueId” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommInitRank” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommInitAll” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommDestroy” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommAbort” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommCount” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommCuDevice” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommUserRank” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Fault tolerance and communicator management#

These functions let applications detect failures and recover by resizing or rebuilding communicators. See Fault tolerance in RCCL for usage guidance.

Warning

doxygenfunction: Cannot find function “ncclCommGetAsyncError” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommFinalize” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommSplit” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

ncclResult_t ncclCommShrink(ncclComm_t comm, int *excludeRanksList, int excludeRanksCount, ncclComm_t *newcomm, ncclConfig_t *config, int shrinkFlags)

Shrink existing communicator.

Ranks in excludeRanksList will be removed form the existing communicator. Within the new communicator, ranks will be re-ordered to fill the gap of removed ones. If config is NULL, the new communicator will inherit the original communicator’s configuration. The flag enables NCCL to adapt to various states of the parent communicator, see NCCL_SHRINK flags.

Parameters:

  • comm – [in] Original communicator object for this rank

  • excludeRanksList – [in] List of ranks to be exluded

  • excludeRanksCount – [in] Number of ranks to be excluded

  • newcomm – [out] Pointer to new communicator

  • config – [in] Config file for new communicator. May be NULL to inherit from comm

  • shrinkFlags – [in] Flag to adapt to various states of the parent communicator (see NCCL_SHRINK flags)

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommGrow(ncclComm_t comm, int nRanks, const ncclUniqueId *uniqueId, int rank, ncclComm_t *newcomm, ncclConfig_t *config)

Grow a communicator by adding new ranks.

Creates a larger communicator from an existing one plus newly joining ranks. The unique ID obtained from ncclCommGetUniqueId must be distributed to the new ranks. Parameter usage:

  • Existing non-root ranks: comm set, uniqueId = NULL, rank = -1

  • Existing root rank: comm set, uniqueId = &id, rank = -1

  • New ranks: comm = NULL, uniqueId = &id, rank = assigned The unique ID is consumed upon a successful grow and cannot be reused. If config is NULL, the new communicator inherits the original communicator’s configuration.

Parameters:

  • comm – [in] Existing communicator, or NULL for newly joining ranks

  • nRanks – [in] Total number of ranks in the new communicator

  • uniqueId – [in] Unique ID from ncclCommGetUniqueId; NULL on existing non-root ranks

  • rank – [in] Rank in the new communicator for joining ranks; -1 for existing ranks

  • newcomm – [out] Pointer to the new communicator

  • config – [in] Config for the new communicator. May be NULL to inherit from comm

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommGetUniqueId(ncclComm_t comm, ncclUniqueId *uniqueId)

Generate a per-communicator unique ID for growing a communicator.

Generates a unique ID on an existing communicator. The ID must be distributed to the ranks joining through ncclCommGrow. Constraints:

  • A new ID cannot be generated while a previous ID is unconsumed.

  • Each ID can only be used once (no reuse after consumption).

  • The grow operation must complete before calling this again.

Parameters:

  • comm – [in] Existing communicator that coordinates the grow

  • uniqueId – [out] Pointer to the generated unique ID

Returns:

Result code. See Result Codes for more details.

Warning

doxygenfunction: Cannot find function “ncclCommRevoke” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Communicator suspend and resume#

These functions release and reacquire the resources held by a communicator that is temporarily idle, and report per-communicator memory statistics. They are useful when an application wants to free GPU memory held by an inactive communicator, and then later resume collective operations on the same communicator.

Releasing the physical backing of a suspended communicator requires virtual memory management (VMM) support enabled through NCCL_CUMEM_ENABLE. Without it, ncclCommSuspend and ncclCommResume succeed but don’t release GPU memory. See Suspending and resuming a communicator for the full list of prerequisites.

Use ncclGroupStart and ncclGroupEnd when suspending or resuming multiple communicators from one thread. Requests run at ncclGroupEnd after all ranks of each affected communicator synchronize.

NCCL_SUSPEND_MEM

Communicator suspend flag: release dynamic GPU memory allocations.

Warning

doxygenfunction: Cannot find function “ncclCommSuspend” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclCommResume” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

enum ncclCommMemStat_t

Communicator memory statistic selector.

Identifier passed to ncclCommMemStats to choose which memory counter to read.

Values:

enumerator ncclStatGpuMemSuspend

Allocated GPU memory that can be suspended (bytes)

enumerator ncclStatGpuMemSuspended

GPU memory suspended? (0=active, 1=suspended)

enumerator ncclStatGpuMemPersist

Allocated GPU memory that cannot be suspended (bytes)

enumerator ncclStatGpuMemTotal

Total allocated GPU memory tracked by NCCL (bytes)

Warning

doxygenfunction: Cannot find function “ncclCommMemStats” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Collective communication operations#

Collective communication operations must be called separately for each communicator in a communicator clique.

They return when operations have been enqueued on the hipstream.

Since they may perform inter-CPU synchronization, each call has to be done from a different thread or process, or need to use Group Semantics (see below).

Warning

doxygenfunction: Cannot find function “ncclReduce” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclBcast” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclBroadcast” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclAllReduce” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclReduceScatter” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclAllGather” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclSend” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclRecv” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclGather” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclScatter” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclAllToAll” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Group semantics#

When managing multiple GPUs from a single thread, and since NCCL collective calls may perform inter-CPU synchronization, we need to “group” calls for different ranks/devices into a single call.

Grouping NCCL calls as being part of the same collective operation is done using ncclGroupStart and ncclGroupEnd. ncclGroupStart will enqueue all collective calls until the ncclGroupEnd call, which will wait for all calls to be complete. Note that for collective communication, ncclGroupEnd only guarantees that the operations are enqueued on the streams, not that the operation is effectively done.

Both collective communication and ncclCommInitRank can be used in conjunction of ncclGroupStart/ncclGroupEnd.

Warning

doxygenfunction: Cannot find function “ncclGroupStart” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclGroupEnd” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Library functions#

Warning

doxygenfunction: Cannot find function “ncclGetVersion” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Warning

doxygenfunction: Cannot find function “ncclGetErrorString” in doxygen xml output for project “RCCL 2.30.4 Documentation” from directory: /build/rccl-kH3VPW/rccl-10.0/docs/doxygen/xml

Types#

There are few data structures that are internal to the library. The pointer types to these structures are given below. The user would need to use these types to create handles and pass them between different library functions.

typedef struct ncclComm *ncclComm_t

Opaque handle to communicator.

A communicator contains information required to facilitate collective communications calls

struct ncclUniqueId

Opaque unique id used to initialize communicators.

The ncclUniqueId must be passed to all participating ranks

Enumerations#

This section provides all the enumerations used.

enum ncclResult_t

Result type.

Return codes aside from ncclSuccess indicate that a call has failed

Values:

enumerator ncclSuccess

No error

enumerator ncclUnhandledCudaError

Unhandled HIP error

enumerator ncclSystemError

Unhandled system error

enumerator ncclInternalError

Internal Error - Please report to RCCL developers

enumerator ncclInvalidArgument

Invalid argument

enumerator ncclInvalidUsage

Invalid usage

enumerator ncclRemoteError

Remote process exited or there was a network error

enumerator ncclInProgress

RCCL operation in progress

enumerator ncclTimeout

Operation timed out

enumerator ncclNumResults

Number of result types

enum ncclRedOp_t

Reduction operation selector.

Enumeration used to specify the various reduction operations ncclNumOps is the number of built-in ncclRedOp_t values and serves as the least possible value for dynamic ncclRedOp_t values constructed by ncclRedOpCreate functions.

ncclMaxRedOp is the largest valid value for ncclRedOp_t and is defined to be the largest signed value (since compilers are permitted to use signed enums) that won’t grow sizeof(ncclRedOp_t) when compared to previous RCCL versions to maintain ABI compatibility.

Values:

enumerator ncclSum

Sum

enumerator ncclProd

Product

enumerator ncclMax

Max

enumerator ncclMin

Min

enumerator ncclAvg

Average

enumerator ncclNumOps

Number of built-in reduction ops

enumerator ncclMaxRedOp

Largest value for ncclRedOp_t

enum ncclDataType_t

Data types.

Enumeration of the various supported datatype

Values:

enumerator ncclInt8
enumerator ncclChar
enumerator ncclUint8
enumerator ncclInt32
enumerator ncclInt
enumerator ncclUint32
enumerator ncclInt64
enumerator ncclUint64
enumerator ncclFloat16
enumerator ncclHalf
enumerator ncclFloat32
enumerator ncclFloat
enumerator ncclFloat64
enumerator ncclDouble
enumerator ncclBfloat16
enumerator ncclFloat8e4m3
enumerator ncclFloat8e5m2
enumerator ncclNumTypes