Types

The following types are used by the NCCL library.

ncclComm_t

type ncclComm_t

NCCL communicator. Points to an opaque structure inside NCCL.

ncclResult_t

type ncclResult_t

Return values for all NCCL functions. Possible values are:

ncclSuccess

(0) Function succeeded.

ncclUnhandledCudaError

(1) A call to a CUDA function failed.

ncclSystemError

(2) A call to the system failed.

ncclInternalError

(3) An internal check failed. This is due to either a bug in NCCL or a memory corruption.

ncclInvalidArgument

(4) An argument has an invalid value.

ncclInvalidUsage

(5) The call to NCCL is incorrect. This is usually reflecting a programming error.

ncclRemoteError

(6) A call failed possibly due to a network error or a remote process exiting prematurely.

ncclInProgress

(7) A NCCL operation on the communicator is being enqueued and is being progressed in the background.

Whenever a function returns an error (neither ncclSuccess nor ncclInProgress), NCCL should print a more detailed message when the environment variable NCCL_DEBUG is set to “WARN”.

ncclDataType_t

type ncclDataType_t

NCCL defines the following integral and floating data-types.

ncclInt8

Signed 8-bits integer

ncclChar

Signed 8-bits integer

ncclUint8

Unsigned 8-bits integer

ncclInt32

Signed 32-bits integer

ncclInt

Signed 32-bits integer

ncclUint32

Unsigned 32-bits integer

ncclInt64

Signed 64-bits integer

ncclUint64

Unsigned 64-bits integer

ncclFloat16

16-bits floating point number (half precision)

ncclHalf

16-bits floating point number (half precision)

ncclFloat32

32-bits floating point number (single precision)

ncclFloat

32-bits floating point number (single precision)

ncclFloat64

64-bits floating point number (double precision)

ncclDouble

64-bits floating point number (double precision)

ncclBfloat16

16-bits floating point number (truncated precision in bfloat16 format, CUDA 11 or later)

ncclFloat8e4m3

8-bits floating point number, 4 exponent bits, 3 mantissa bits (CUDA >= 11.8 and SM >= 90)

ncclFloat8e5m2

8-bits floating point number, 5 exponent bits, 2 mantissa bits (CUDA >= 11.8 and SM >= 90)

ncclRedOp_t

type ncclRedOp_t

Defines the reduction operation.

ncclSum

Perform a sum (+) operation

ncclProd

Perform a product (*) operation

ncclMin

Perform a min operation

ncclMax

Perform a max operation

ncclAvg

Perform an average operation, i.e. a sum across all ranks, divided by the number of ranks.

ncclScalarResidence_t

type ncclScalarResidence_t

Indicates where (memory space) scalar arguments reside and when they can be dereferenced.

ncclScalarHostImmediate

The scalar resides in host memory and should be dereferenced in the most immediate way.

ncclScalarDevice

The scalar resides on device visible memory and should be dereferenced once needed.

ncclConfig_t

type ncclHostCftMode_t

Values for the hostCftMode communicator configuration.

ncclHostCftDefault

Use the version-specific default.

ncclHostCftEnable

Enable host-side CFT support.

ncclHostCftDisable

Disable host-side CFT support.

ncclHostCftFallback

Try to create CFT logical endpoints. In case of an error, host-side CFT will be disabled.

type ncclConfig_t

A structure-based configuration users can set to initialize a communicator; a newly created configuration must be initialized by NCCL_CONFIG_INITIALIZER.

NCCL_CONFIG_INITIALIZER

A configuration macro initializer which must be assigned to a newly created configuration.

blocking

This attribute can be set as integer 0 or 1 to indicate nonblocking or blocking communicator behavior correspondingly. Blocking is the default behavior.

cgaClusterSize

Set Cooperative Group Array (CGA) size of kernels launched by NCCL. This attribute can be set between 0 and 8, and the default value is 4 since sm90 architecture and 0 for older architectures.

minCTAs

Set the minimal number of CTAs NCCL should use for each kernel. Set to a positive integer value, up to 32. The default value is 1.

maxCTAs

Set the maximal number of CTAs NCCL should use for each kernel. Set to a positive integer value, up to 32. The default value is 32.

netName

Specify the network module name NCCL should use for network communication. The value of netName must match exactly the name of the network module (case-insensitive). NCCL internal network module names are “IB” (generic IB verbs) and “Socket” (TCP/IP sockets). External network plugins define their own names. The default value is undefined, and NCCL will choose the network module automatically.

splitShare

Specify whether to share resources with child communicator during communicator split. Set the value of splitShare to 0 or 1. The default value is 0. When the parent communicator is created with splitShare=1 during ncclCommInitRankConfig, the child communicator can share internal resources of the parent during communicator split. Split communicators are in the same family. When resources are shared, aborting any communicator can result in other communicators in the same family becoming unusable. Irrespective of whether sharing resources or not, users should always abort/destroy all no longer needed communicators to free up resources. Note: when the parent communicator has been revoked, resource sharing during split is disabled regardless of this flag.

shrinkShare

Specify whether to share resources with child communicator during communicator shrink. Set the value of shrinkShare to 0 or 1. The default value is 0. Note: when shrink is used with NCCL_SHRINK_ABORT, the value of shrinkShare is ignored and no resources are shared. When the parent communicator has been revoked, resource sharing is also disabled. The behavior of this flag is similar to splitShare, see above.

trafficClass

Set the traffic class (TC) to use for network operations on the communicator. The meaning of TC is specific to the network plugin in use by the communicator (e.g. IB networks use service level, RoCE networks use type of service). Assigning different TCs to each communicator can benefit workloads which overlap communication. TCs are defined by the system configuration and should be greater than or equal to 0. Note that environment variables, such as NCCL_IB_SL and NCCL_IB_TC, take precedence over user-specified TC values. To utilize user-defined TCs, ensure that these environment variables are unset.

collnetEnable

Set 1/0 to enable/disable IB SHARP on the communicator. The default value is 0 (disabled).

CTAPolicy

Set the policy for the communicator. The full list of supported policies can be found in NCCL Communicator CTA Policy Flags. The default value is NCCL_CTA_POLICY_DEFAULT.

nvlsCTAs

Set the total number of CTAs NCCL should use for NVLS kernels. Set to a positive integer value. By default, NCCL will automatically determine the best number of CTAs based on the system configuration.

commName

Specify the user defined name for the communicator. The communicator name can be used by NCCL to enrich logging and profiling.

nChannelsPerNetPeer

Set the number of network channels to be used for pairwise communication. The value must be a positive integer and will be round up to the next power of 2. The default value is optimized for the AlltoAll communication pattern. Consider increasing the value to increase the bandwidth for send/recv communication.

graphUsageMode

Set the graph usage mode for the communicator. It support three possible values: 0 (no graphs), 1 (one graph) and 2 (either multiple graphs or mix of graph and non-graph). The default value is 2. If NCCL_GRAPH_STREAM_ORDERING or graphStreamOrdering disables capture-time stream ordering (0), graph mixing must be off—use graphUsageMode 0 or 1 only; graphUsageMode=2 must not be combined with ordering 0 (see NCCL_GRAPH_STREAM_ORDERING).

graphStreamOrdering

(since 2.30)

Per-communicator setting, honored unless NCCL_GRAPH_STREAM_ORDERING is set to 0 or 1. 1 keeps NCCL’s default capture-time serialization of communication kernels. 0 disables it for this communicator—kernels are placed on the capture stream and the application must guarantee correct ordering (see NCCL_GRAPH_STREAM_ORDERING).

Defaults to NCCL_CONFIG_UNDEF_INT (inherits NCCL_GRAPH_STREAM_ORDERING). Setting that environment variable to 0 or 1 overrides this field.

graphStreamOrdering=0 requires graphUsageMode 0 or 1 (mixing off). If it is combined with graphUsageMode=2, NCCL emits a warning and forces graphStreamOrdering to 1; communicator creation still succeeds. See NCCL_GRAPH_STREAM_ORDERING. NCCL enforces this compatibility fallback within a single communicator only. Communicators created by ncclCommSplit with splitShare share one internal serialization event with their parent, so ordering 0 on one and graphUsageMode=2 on another that shares those resources is equally unsupported and is not diagnosed.

Mixed values on one GPU: A communicator set to 1 still receives NCCL’s internal serialization for its own kernels, but NCCL does not insert cross-communicator ordering with a peer set to 0—its kernels may overlap in situations NCCL would have serialized. Use 0 only when the application guarantees ordering of all NCCL communication kernels that may run concurrently on the GPU.

launchOrderImplicit

(since 2.31)

Per-communicator request for NCCL_LAUNCH_ORDER_IMPLICIT. 1 enables implicit launch ordering for this communicator; 0 disables it. NCCL_CONFIG_UNDEF_INT is the default and has the same effective behavior as 0.

Communicators with different effective values can coexist. Overlap safety is about communication operations that may run concurrently on the same GPU:

• Operations on disabled/default communicators retain the existing multiple-communicator ordering guarantees.

• Operations on enabled communicators may overlap with operations on other enabled communicators if the application follows the host-side ordering requirements described for NCCL_LAUNCH_ORDER_IMPLICIT.

• Operations on enabled communicators must not overlap with operations on disabled/default communicators. The application must order or synchronize those operations so they do not overlap, or configure the communicators consistently.

NCCL logs an INFO message if a CUDA context has used both enabled and disabled/default effective values, but it still initializes the communicator.

If NCCL_LAUNCH_ORDER_IMPLICIT is set in the environment, it overrides this field before initialization.

maxP2pPeers

(since 2.30)

Set the maximum number of peers any rank will concurrently communicate with using P2P communication. Where it applies, RCCL divides the P2P channel pool among this many peers rather than among all ranks, so a smaller value gives each peer more channels, influencing send/recv and the send/recv-based collectives (all-to-all, scatter, gather). It does not restrict which peers a rank is allowed to communicate with.

The value is read in two places only: the per-peer channel tiling enabled by RCCL_SATURATE_P2P_NCHANNELS (on by default for gfx1250 only), and the multi-node per-peer reduction, which requires more than one node and nChannelsPerNetPeer left unset. A single-node job on another architecture with default settings is unaffected by this field.

Defaults to NCCL_CONFIG_UNDEF_INT, which resolves to the number of ranks in the communicator. Values greater than the number of ranks are capped to the communicator size, with a message logged at NCCL_DEBUG=INFO. Any other value less than one is rejected: ncclCommInitRankConfig returns ncclInvalidArgument. The field is ignored when config.version is below 2.30.

Setting the NCCL_P2P_MAX_PEERS environment variable overrides this field. Values <= 0 in that variable are ignored rather than rejected, and the field keeps its value.

numRmaCtx

(since 2.31)

Number of one-sided RMA communication contexts to provision on the communicator. The ctx argument of ncclPutSignal(), ncclSignal(), and ncclWaitSignal() must lie in [0, numRmaCtx). The default value is 1.

numRmaSig

(since 2.31)

Set the number of one-sided RMA signal indexes available per context. The default value is 1. Host one-sided RMA operations such as ncclPutSignal(), ncclSignal(), and ncclWaitSignal() use sigIdx values in the range [0, numRmaSig).

rmaEagerInit

(since 2.31)

Controls when the collective one-sided RMA signal setup is initialized. With 0 (default), it is initialized during the first window registration (ncclCommWindowRegister()), a collective point. Use 1 to initialize it at communicator-init time instead; this is required if a communicator issues ncclSignal() or ncclWaitSignal() without first registering a window, which otherwise returns ncclInvalidUsage.

If NCCL_RMA_EAGER_INIT is set in the environment, it overrides this field before initialization.

hostCftMode

(since 2.31)

Controls support for host-side Compute Fabric Transport (CFT) queries. ncclHostCftEnable creates the communicator’s unicast and multicast logical endpoints during the first ncclCommWindowRegister() call on a CFT-capable communicator. ncclHostCftDisable disables support, and ncclHostCftFallback tries to create logical endpoints and disables host-side CFT in case of error. ncclHostCftDefault selects the library-defined default behavior.

nvlsHostMode

(since 2.31)

Controls host-side use of NVLink SHARP (NVLS) for a single communicator. The value is a bitmask that can independently disable the NVLS collective transport and multimem use by NCCL’s internal symmetric kernels and Copy Engine (CE) collectives. Explicit Device API multimem requests remain available. Selectively disabling host NVLS components on small or infrequently used communicators can preserve limited hardware multicast resources for other communicators that can utilize them more effectively.

Supported values (ncclNvlsHostMode_t):

• ncclNvlsHostModeDefault — Library-defined default. Currently equivalent to ncclNvlsHostModeEnable.

• ncclNvlsHostModeEnable — Enable all host NVLS components.

• ncclNvlsHostModeDisableTransport — Disable the NVLS collective transport and its registered-buffer optimization. The NVLS, NVLS_TREE, and multi-RPN PAT algorithms will be disabled.

• ncclNvlsHostModeDisableSymmetricMultimem — Disable multimem use by NCCL’s internal symmetric kernels and CE collectives. UC symmetric kernels and CE fallbacks remain available.

• ncclNvlsHostModeDisable — Disable all present and future host NVLS components. This value sets all non-reserved bits for forward compatibility.

The component flags can be combined with the bitwise OR operator. These flags do not affect the communicator’s reported Device API multimem capability and do not prevent explicit Device API calls from allocating multimem resources.

ncclCollConfig_t

type ncclCollConfig_t

Per-collective configuration passed through an nccl*Config collective API. Every ncclCollConfig_t must be initialized with NCCL_COLLCONFIG_INITIALIZER before any member is set. Passing NULL to a config-taking collective is equivalent to calling the corresponding collective without a configuration.

The application owns the configuration and any storage it references, and must keep them valid for the duration of each call that uses them. Set the same configuration on every rank unless a member is documented otherwise.

For options that also have environment-variable and communicator-level forms, NCCL applies values in the following order: environment variable, per-collective configuration, communicator configuration, and library default. Leave resource and CTA-policy members at NCCL_CONFIG_UNDEF_INT to fall through to the next source in this list.

NCCL_COLLCONFIG_INITIALIZER

(since 2.31)

Initializes a new ncclCollConfig_t with default values.

ncclConfigExt_t *ext

(since 2.31)

Head of a linked list of vendor-specific ncclConfigExt_t options, or NULL. The official NCCL library does not read this member; it exists for vendor libraries layered on top of NCCL to forward their own vendor-defined options through the same call.

int minCTAs

(since 2.31)

Per-collective override of minCTAs in ncclConfig_t. Accepts the same value range as NCCL_MIN_CTAS; an out-of-range value is ignored and NCCL falls through to the next priority source. If the resolved value exceeds the resolved maxCTAs, NCCL resets minCTAs to 1.

int maxCTAs

(since 2.31)

Per-collective override of maxCTAs in ncclConfig_t. Accepts the same value range as NCCL_MAX_CTAS; an out-of-range value is ignored and NCCL falls through to the next priority source. The resolved value is further clamped to the communicator’s maxCTAs and applies only to this collective, including when calls are grouped.

int nvlsCTAs

(since 2.31)

Per-collective override of nvlsCTAs in ncclConfig_t. This setting applies only to NVLS algorithms.

int cgaClusterSize

(since 2.31)

Per-collective override of cgaClusterSize in ncclConfig_t. Accepts the same value range as NCCL_CGA_CLUSTER_SIZE; an out-of-range value is ignored and NCCL falls through to the next priority source. All collectives in a group must use the same value; inconsistent values result in undefined behavior.

const char *algSelection

(since 2.31)

A case-insensitive expression that filters the algorithms and kernels NCCL may use for this call. NULL or an empty string requests automatic selection.

Each name is matched as a prefix from the start of a known name. For example, ring selects RING_LL, RING_LL128, and RING_SIMPLE, but does not select SYMK_RailRing_LsaSTMC; similarly, nvls selects both NVLS_SIMPLE and NVLSTREE_SIMPLE. The supported operators are + (AND), , (OR), and ^ (NOT); parentheses are not supported. For example, tree,ring permits either family, ring + ^RING_LL128 permits Ring except LL128, and ^symk excludes symmetric kernels.

The selectable built-in names and the collectives for which they are valid are:

• AllReduce: TREE_LL, TREE_LL128, TREE_SIMPLE, RING_LL, RING_LL128, RING_SIMPLE, COLLNET_DIRECT_SIMPLE, COLLNET_CHAIN_SIMPLE, NVLS_SIMPLE, NVLSTREE_SIMPLE, SYMK_AGxLL_R, SYMK_AGxLLMC_R, SYMK_RSxTmaLD_AGxTmaST, SYMK_RSxLD_AGxST, and SYMK_RSxLDMC_AGxSTMC.

• AllGather: RING_LL, RING_LL128, RING_SIMPLE, COLLNET_DIRECT_SIMPLE, NVLS_SIMPLE, PAT_SIMPLE, SYMK_LL, SYMK_LLMC, SYMK_TmaST, SYMK_ST, SYMK_TmaSTMC, SYMK_STMC, and SYMK_RailRing_LsaSTMC.

• ReduceScatter: RING_LL, RING_LL128, RING_SIMPLE, COLLNET_DIRECT_SIMPLE, NVLS_SIMPLE, PAT_SIMPLE, SYMK_LL, SYMK_TmaLD, SYMK_LD, SYMK_LDMC, SYMK_RailA2A_LsaLD, and SYMK_RailA2A_LsaLDMC.

• Broadcast and Reduce: RING_LL, RING_LL128, and RING_SIMPLE.

The SYMK_* names refer to symmetric-memory kernels built on the NCCL Device API. LL is a low-latency protocol that tags data with an inline flag instead of a separate synchronization step; LD/ST are ordinary load/store data movement and TmaLD/TmaST use the Tensor Memory Accelerator for bulk asynchronous copies; the MC suffix (as in LLMC, LDMC, STMC) means the load, store, or reduction is done through NVLink SHARP multimem (multicast) instead of per-peer accesses; and Rail/Lsa names are multi-node kernels that combine a local NVLink-domain (Lsa) step with a network exchange between same-rail ranks across nodes.

Algorithm names are implementation details and may evolve. AllToAll, Gather, and Scatter lower to point-to-point sends and receives and expose no selectable names; Copy Engine paths are likewise not selectable here. Request a Copy Engine path with CTAPolicy instead.

int forceAlgSelection

(since 2.31)

Controls what happens when algSelection is invalid or cannot be honored. With the default value, 1, NCCL returns an error. Set it to 0 to fall back to automatic selection instead.

int CTAPolicy

(since 2.31)

Per-collective override of the communicator’s CTA policy. A set value replaces the communicator value; see NCCL Communicator CTA Policy Flags. If both NCCL_CTA_POLICY_ZERO and NCCL_CTA_POLICY_EFFICIENCY are set, Zero-CTA takes precedence. Copy Engine paths are selected through this policy, not algSelection.

uint64_t userProfilerTag

(since 2.31)

An opaque value delivered with the call’s events to profiler plugins that support profiler interface v7. Zero denotes an untagged call; older profiler interfaces ignore the tag.

The tag does not affect execution. NCCL reserves values with the most-significant bit set for possible future use; applications should use values with that bit clear.

cudaEvent_t launchCompletionEvent

(since 2.32)

A caller-owned CUDA event that NCCL records at the collective kernel-launch completion boundary. Create the event with cudaEventCreateWithFlags and cudaEventDisableTiming. It must not be an interprocess or interop event.

A NULL event requests no update. Event presence must match across ranks: every rank in the corresponding collective must provide a rank-local non-NULL event, or every rank must provide NULL. A group supports at most one non-NULL event per communicator.

Keep the event alive through all queued waits and graph executions that use it. For a grouped collective, enqueue a dependent cudaStreamWaitEvent only after ncclGroupEnd returns ncclSuccess. With a nonblocking communicator, first poll ncclCommGetAsyncError until it returns ncclSuccess; do not enqueue the wait after an error.

NCCL records the event only for kernel launches; Copy Engine and RMA work do not record it.

CUDA 12.3 or newer provides launch-completion semantics. On older CUDA platforms, NCCL emits a warning and records the event before the kernel launch; this fallback does not provide a launch-completion guarantee. CUDA launch completion is best effort and does not guarantee that dependent work overlaps the collective. Graph capture is supported with a capture-compatible event; sharing a recorded event across graphs is not supported.

type ncclConfigExt_t

A vendor-specific key-value option that can be linked through its next member and attached to ncclCollConfig_t.ext. Each option has a key consisting of key.vendorId and key.optionId and one value in the val.i, val.s, or val.raw union member. The list is unordered and must not contain duplicate keys.

Vendors should choose a non-zero vendorId under 2^24 that is unlikely to collide with other vendor libraries, for example by picking a random number; optionId then distinguishes options within that vendor’s own namespace.

ncclSimInfo_t

type ncclSimInfo_t

This struct will be used by ncclGroupSimulateEnd() to return information about the calls.

NCCL_SIM_INFO_INITIALIZER

NCCL_SIM_INFO_INITIALIZER is a configuration macro initializer which must be assigned to a newly created ncclSimInfo_t struct.

estimatedTime

Estimated time for the operation(s) in the group call will be returned in this attribute.

ncclCommMemStat_t

type ncclCommMemStat_t

Memory statistic selectors for ncclCommMemStats().

ncclStatGpuMemSuspend

Communicator allocated GPU memory that can be released via suspend (bytes).

ncclStatGpuMemSuspended

Whether communicator allocated GPU memory is currently suspended (0 = active, 1 = suspended).

ncclStatGpuMemPersist

Communicator allocated GPU memory that cannot be suspended (bytes).

ncclStatGpuMemTotal

Total communicator allocated GPU memory that is tracked by NCCL (bytes).

ncclWindow_t

type ncclWindow_t

NCCL window object for window registration and deregistration.