RCCL library specification

Contents

RCCL library specification#

This document provides details of the API library.

Communicator functions#

ncclResult_t ncclGetUniqueId(ncclUniqueId *uniqueId)

Generates an ID for ncclCommInitRank.

Generates an ID to be used in ncclCommInitRank. ncclGetUniqueId should be called once by a single rank and the ID should be distributed to all ranks in the communicator before using it as a parameter for ncclCommInitRank.

Parameters:

uniqueId – [out] Pointer to where uniqueId will be stored

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommInitRank(ncclComm_t *comm, int nranks, ncclUniqueId commId, int rank)

Creates a new communicator (multi thread/process version).

Rank must be between 0 and nranks-1 and unique within a communicator clique. Each rank is associated to a CUDA device, which has to be set before calling ncclCommInitRank. ncclCommInitRank implicitly syncronizes with other ranks, so it must be called by different threads/processes or use ncclGroupStart/ncclGroupEnd.

Parameters:

  • comm – [out] Pointer to created communicator

  • nranks – [in] Total number of ranks participating in this communicator

  • commId – [in] UniqueId required for initialization

  • rank – [in] Current rank to create communicator for

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommInitAll(ncclComm_t *comm, int ndev, const int *devlist)

Creates a clique of communicators (single process version).

This is a convenience function to create a single-process communicator clique. Returns an array of ndev newly initialized communicators in comm. comm should be pre-allocated with size at least ndev*sizeof(ncclComm_t). If devlist is NULL, the first ndev HIP devices are used. Order of devlist defines user-order of processors within the communicator.

Parameters:

  • comm – [out] Pointer to array of created communicators

  • ndev – [in] Total number of ranks participating in this communicator

  • devlist – [in] Array of GPU device indices to create for

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommDestroy(ncclComm_t comm)

Frees local resources associated with communicator object.

Destroy all local resources associated with the passed in communicator object

Parameters:

comm – [in] Communicator to destroy

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommAbort(ncclComm_t comm)

Abort any in-progress calls and destroy the communicator object.

Frees resources associated with communicator object and aborts any operations that might still be running on the device.

Parameters:

comm – [in] Communicator to abort and destroy

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommCount(const ncclComm_t comm, int *count)

Gets the number of ranks in the communicator clique.

Returns the number of ranks in the communicator clique (as set during initialization)

Parameters:

  • comm – [in] Communicator to query

  • count – [out] Pointer to where number of ranks will be stored

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommCuDevice(const ncclComm_t comm, int *device)

Get the ROCm device index associated with a communicator.

Returns the ROCm device number associated with the provided communicator.

Parameters:

  • comm – [in] Communicator to query

  • device – [out] Pointer to where the associated ROCm device index will be stored

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommUserRank(const ncclComm_t comm, int *rank)

Get the rank associated with a communicator.

Returns the user-ordered “rank” associated with the provided communicator.

Parameters:

  • comm – [in] Communicator to query

  • rank – [out] Pointer to where the associated rank will be stored

Returns:

Result code. See Result Codes for more details.

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.

ncclResult_t ncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError)

Checks whether the comm has encountered any asynchronous errors.

Query whether the provided communicator has encountered any asynchronous errors

Parameters:

  • comm – [in] Communicator to query

  • asyncError – [out] Pointer to where result code will be stored

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommFinalize(ncclComm_t comm)

Finalize a communicator.

ncclCommFinalize flushes all issued communications and marks communicator state as ncclInProgress. The state will change to ncclSuccess when the communicator is globally quiescent and related resources are freed; then, calling ncclCommDestroy can locally free the rest of the resources (e.g. communicator itself) without blocking.

Parameters:

comm – [in] Communicator to finalize

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommSplit(ncclComm_t comm, int color, int key, ncclComm_t *newcomm, ncclConfig_t *config)

Create one or more communicators from an existing one.

Creates one or more communicators from an existing one. Ranks with the same color will end up in the same communicator. Within the new communicator, key will be used to order ranks. NCCL_SPLIT_NOCOLOR as color will indicate the rank will not be part of any group and will therefore return a NULL communicator. If config is NULL, the new communicator will inherit the original communicator’s configuration

Parameters:

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

  • color – [in] Color to assign this rank

  • key – [in] Key used to order ranks within the same new communicator

  • newcomm – [out] Pointer to new communicator

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

Returns:

Result code. See Result Codes for more details.

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.

ncclResult_t ncclCommRevoke(ncclComm_t comm, int revokeFlags)

Revoke a communicator without destroying it.

Aborts in-flight collectives by raising the comm’s abort flag, stops the proxy service, and rejects subsequently enqueued collectives with ncclInvalidUsage. The abort flag is cleared once the asynchronous revoke job completes, so the communicator remains valid as a parent for ncclCommSplit / ncclCommShrink / ncclCommGrow, and may be torn down via ncclCommDestroy / ncclCommAbort. Because in-flight collectives are aborted (not drained), revoke can recover from peers that have failed mid-collective; output buffers of an aborted collective contain undefined data. ncclCommFinalize on a revoked communicator is invalid and returns ncclInvalidUsage. Pass NCCL_REVOKE_DEFAULT for revokeFlags; other values are rejected with ncclInvalidArgument.

Parameters:

  • comm – [in] Communicator to revoke

  • revokeFlags – [in] Reserved; must be NCCL_REVOKE_DEFAULT

Returns:

Result code. See Result Codes for more details.

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.

ncclResult_t ncclCommSuspend(ncclComm_t comm, int flags)

Suspend communicator operations to free resources.

Releases the resources selected by flags. The communicator cannot be used until ncclCommResume is called.

Parameters:

  • comm – [in] Communicator to suspend

  • flags – [in] Bitmask of NCCL_SUSPEND_* flags (e.g. NCCL_SUSPEND_MEM)

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclCommResume(ncclComm_t comm)

Resume a previously suspended communicator.

Reacquires every resource that was released by the matching ncclCommSuspend call.

Parameters:

comm – [in] Communicator to resume

Returns:

Result code. See Result Codes for more details.

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)

ncclResult_t ncclCommMemStats(ncclComm_t comm, ncclCommMemStat_t stat, uint64_t *value)

Query communicator memory statistics.

Parameters:

  • comm – [in] Communicator to query

  • stat – [in] One of ncclCommMemStat_t values

  • value – [out] Pointer to receive the memory statistic value

Returns:

Result code. See Result Codes for more details.

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).

ncclResult_t ncclReduce(const void *sendbuff, void *recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, hipStream_t stream)

Reduce.

Reduces data arrays of length count in sendbuff into recvbuff using op operation. recvbuff* may be NULL on all calls except for root device. root* is the rank (not the HIP device) where data will reside after the operation is complete. In-place operation will happen if sendbuff == recvbuff.

Parameters:

  • sendbuff – [in] Local device data buffer to be reduced

  • recvbuff – [inout] Data buffer where result is stored (only for root rank). May be null for other ranks.

  • count – [in] Number of elements in every send buffer

  • datatype – [in] Data buffer element datatype

  • op – [in] Reduction operator type

  • root – [in] Rank where result data array will be stored

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclBcast(void *buff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream)

(Deprecated) Broadcast (in-place)

Copies count values from root to all other devices. root is the rank (not the CUDA device) where data resides before the operation is started. This operation is implicitly in-place.

Parameters:

  • buff – [inout] Input array on root to be copied to other ranks. Output array for all ranks.

  • count – [in] Number of elements in data buffer

  • datatype – [in] Data buffer element datatype

  • root – [in] Rank owning buffer to be copied to others

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclBroadcast(const void *sendbuff, void *recvbuff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream)

Broadcast.

Copies count values from sendbuff on root to recvbuff on all devices. root* is the rank (not the HIP device) where data resides before the operation is started. sendbuff* may be NULL on ranks other than root. In-place operation will happen if sendbuff == recvbuff.

Parameters:

  • sendbuff – [in] Data array to copy (if root). May be NULL for other ranks

  • recvbuff – [inout] Data array to store received array

  • count – [in] Number of elements in data buffer

  • datatype – [in] Data buffer element datatype

  • root – [in] Rank of broadcast root

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclAllReduce(const void *sendbuff, void *recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream)

All-Reduce.

Reduces data arrays of length count in sendbuff using op operation, and leaves identical copies of result on each recvbuff. In-place operation will happen if sendbuff == recvbuff.

Parameters:

  • sendbuff – [in] Input data array to reduce

  • recvbuff – [inout] Data array to store reduced result array

  • count – [in] Number of elements in data buffer

  • datatype – [in] Data buffer element datatype

  • op – [in] Reduction operator

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclReduceScatter(const void *sendbuff, void *recvbuff, size_t recvcount, ncclDataType_t datatype, ncclRedOp_t op, ncclComm_t comm, hipStream_t stream)

Reduce-Scatter.

Reduces data in sendbuff using op operation and leaves reduced result scattered over the devices so that recvbuff on rank i will contain the i-th block of the result. Assumes sendcount is equal to nranks*recvcount, which means that sendbuff should have a size of at least nranks*recvcount elements. In-place operations will happen if recvbuff == sendbuff + rank * recvcount.

Parameters:

  • sendbuff – [in] Input data array to reduce

  • recvbuff – [inout] Data array to store reduced result subarray

  • recvcount – [in] Number of elements each rank receives

  • datatype – [in] Data buffer element datatype

  • op – [in] Reduction operator

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclAllGather(const void *sendbuff, void *recvbuff, size_t sendcount, ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream)

All-Gather.

Each device gathers sendcount values from other GPUs into recvbuff, receiving data from rank i at offset i*sendcount. Assumes recvcount is equal to nranks*sendcount, which means that recvbuff should have a size of at least nranks*sendcount elements. In-place operations will happen if sendbuff == recvbuff + rank * sendcount.

Parameters:

  • sendbuff – [in] Input data array to send

  • recvbuff – [inout] Data array to store the gathered result

  • sendcount – [in] Number of elements each rank sends

  • datatype – [in] Data buffer element datatype

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclSend(const void *sendbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, hipStream_t stream)

Send.

Send data from sendbuff to rank peer. Rank peer needs to call ncclRecv with the same datatype and the same count as this rank. This operation is blocking for the GPU. If multiple ncclSend and ncclRecv operations need to progress concurrently to complete, they must be fused within a ncclGroupStart / ncclGroupEnd section.

Parameters:

  • sendbuff – [in] Data array to send

  • count – [in] Number of elements to send

  • datatype – [in] Data buffer element datatype

  • peer – [in] Peer rank to send to

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclRecv(void *recvbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, hipStream_t stream)

Receive.

Receive data from rank peer into recvbuff. Rank peer needs to call ncclSend with the same datatype and the same count as this rank. This operation is blocking for the GPU. If multiple ncclSend and ncclRecv operations need to progress concurrently to complete, they must be fused within a ncclGroupStart/ ncclGroupEnd section.

Parameters:

  • recvbuff – [inout] Data array to receive

  • count – [in] Number of elements to receive

  • datatype – [in] Data buffer element datatype

  • peer – [in] Peer rank to send to

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclGather(const void *sendbuff, void *recvbuff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream)

Gather.

Each rank sends count elements from sendbuff to the root rank. On the root rank, data from rank i is placed at recvbuff + i*count. On non-root ranks, recvbuff is not used. root is the rank where data will be gathered.

In-place operations will happen if sendbuff == recvbuff + root * count.

Parameters:

  • sendbuff – [in] Data array to send

  • recvbuff – [inout] Data array to recv

  • count – [in] Number of elements

  • datatype – [in] Data buffer element datatype

  • root – [in] Rank of gather root

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclScatter(const void *sendbuff, void *recvbuff, size_t count, ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream)

Scatter.

On the root rank, count elements from sendbuff+i*count are sent to rank i. On non-root ranks, sendbuff is not used. Each rank receives count elements into recvbuff. root is the rank that will distribute the data.

In-place operations will happen if recvbuff == sendbuff + root * count.

Parameters:

  • sendbuff – [in] Data array to send

  • recvbuff – [inout] Data array to recv

  • count – [in] Number of elements

  • datatype – [in] Data buffer element datatype

  • root – [in] Rank of scatter root

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclAllToAll(const void *sendbuff, void *recvbuff, size_t count, ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream) __attribute__((deprecated("ncclAllToAll is replaced with ncclAlltoAll and will be removed in the future")))

All-To-All.

Device (i) send (j)th block of data to device (j) and be placed as (i)th block. Each block for sending/receiving has count elements, which means that recvbuff and sendbuff should have a size of nranks*count elements. In-place operation is NOT supported. It is the user’s responsibility to ensure that sendbuff and recvbuff are distinct.

Deprecated:

ncclAllToAll is replaced with ncclAlltoAll and will be removed in the future.

Parameters:

  • sendbuff – [in] Data array to send (contains blocks for each other rank)

  • recvbuff – [inout] Data array to receive (contains blocks from each other rank)

  • count – [in] Number of elements to send between each pair of ranks

  • datatype – [in] Data buffer element datatype

  • comm – [in] Communicator group object to execute on

  • stream – [in] HIP stream to execute collective on

Returns:

Result code. See Result Codes for more details.

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.

ncclResult_t ncclGroupStart()

Group Start.

Start a group call. All calls to RCCL until ncclGroupEnd will be fused into a single RCCL operation. Nothing will be started on the HIP stream until ncclGroupEnd.

Returns:

Result code. See Result Codes for more details.

ncclResult_t ncclGroupEnd()

Group End.

End a group call. Start a fused RCCL operation consisting of all calls since ncclGroupStart. Operations on the HIP stream depending on the RCCL operations need to be called after ncclGroupEnd.

Returns:

Result code. See Result Codes for more details.

Library functions#

ncclResult_t ncclGetVersion(int *version)

Return the RCCL_VERSION_CODE of RCCL in the supplied integer.

This integer is coded with the MAJOR, MINOR and PATCH level of RCCL.

Parameters:

version – [out] Pointer to where version will be stored

Returns:

Result code. See Result Codes for more details.

const char *ncclGetErrorString(ncclResult_t result)

Returns a string for each result code.

Returns a human-readable string describing the given result code.

Parameters:

result – [in] Result code to get description for

Returns:

String containing description of result code.

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
Contents