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.

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.