Device API - CFT

Compute Fabric Transport

Compute Fabric Transport (CFT) support is available starting with NCCL 2.31. CFT exposes CUDA fabric logical endpoints to device kernels so they can move data across a CFT-capable fabric using fabric instructions. NCCL provides the team, logical endpoint, transfer, reduction, and barrier helpers needed to use those endpoints from the Device API.

CFT is intended for systems where all participating ranks can create compatible CUDA fabric logical endpoints. Kernels that use the CFT helpers require CUDA Toolkit support for fabric PTX instructions and must be compiled for architectures that support those instructions.

CFT Communicator Requirements

Request CFT resources when creating the device communicator with ncclDevCommCreate().

cftCaps

Bitmask of NCCL_CFT and NCCL_CFT_MULTIMEM values. If CFT resources are requested on a communicator where not all ranks support CFT, ncclDevCommCreate() fails.

cftBarrierCount

Number of CFT barriers to allocate for ncclCftBarrierSession. This should match the number of independently addressed barrier slots used by the kernel, commonly one per CTA with blockIdx.x as the barrier index.

ncclDevComm devComm;
ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;

reqs.cftCaps = NCCL_CFT | NCCL_CFT_MULTIMEM;
reqs.cftBarrierCount = nCTAs;

NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));

Use NCCL_CFT to request unicast CFT logical endpoints. Use NCCL_CFT_MULTIMEM when the kernel also uses multicast CFT operations or multimem CFT barriers.

NCCL_CFT_NONE

No CFT capability is requested.

NCCL_CFT

Request unicast CFT logical endpoint support.

NCCL_CFT_MULTIMEM

Request multicast CFT logical endpoint support.

CFT Teams

CFT operations may use ncclTeam_t and ncclTeam values to address peers.

ncclTeam_t ncclTeamCft(ncclComm_t comm, ncclCftTeamMode_t mode)

Returns the host-side unicast CFT team containing ranks that are part of the CFT unicast group. comm is the host-side communicator and mode is a selector used to filter the subset of CFT unicast group ranks that are part of the team (useful to build hierarchical communication patterns).

ncclTeam_t ncclTeamCftMultimem(ncclComm_t comm)

Returns the host-side multicast CFT team containing all the ranks that are part of the CFT multicast group.

ncclTeam ncclTeamCft(ncclDevComm const &comm, ncclCftTeamMode_t mode)

Returns the device-side unicast CFT team containing ranks that are part of the CFT unicast group. comm is the device-side communicator and mode is a selector used to filter the subset of CFT unicast group ranks that are part of the team (useful to build hierarchical communication patterns).

ncclTeam ncclTeamCftMultimem(ncclDevComm const &comm)

Returns the device-side multicast CFT team containing all the ranks that are part of the CFT multicast group.

type ncclCftTeamMode_t

Selects the CFT team layout returned by ncclTeamCft().

NCCL_CFT_TEAM_FLAT

Flat CFT team includes all the ranks that are part of the CFT unicast group.

NCCL_CFT_TEAM_HIER_MULTIMEM

Hierarchical CFT team includes ranks that are part of the CFT unicast group and share the same index across different multicast CFT groups.

NCCL_CFT_TEAM_HIER_LSA

Hierarchical CFT team includes ranks that are part of the CFT unicast group and share the same index across different LSA groups.

Logical Endpoint Queries

CFT data movement routines operate on a logical endpoint ID and an offset within that endpoint. NCCL provides host-side and device-side helpers for translating registered windows and peer ranks into those values.

ncclResult_t ncclGetCftDeviceLeInfo(ncclWindow_t window, size_t offset, int peerCft, ncclTeam_t cftTeam, ncclCftLeId *leId, size_t *leOffset)

Host-side query for the logical endpoint leId and leOffset corresponding to window, byte offset, and peerCft within cftTeam.

ncclResult_t ncclGetPeerDeviceLeInfo(ncclWindow_t window, size_t offset, int peerWorld, ncclCftLeId *leId, size_t *leOffset)

Host-side query for the logical endpoint leId and leOffset corresponding to window, byte offset and peerWorld within the world team.

ncclResult_t ncclGetMultimemDeviceLeInfo(ncclWindow_t window, size_t offset, ncclCftLeId *leId, size_t *leOffset)

Host-side query for the multicast logical endpoint leId and leOffset corresponding to window and byte offset.

void ncclGetCftLeInfo(ncclWindow_t w, size_t offset, int peerCft, ncclTeam cftTeam, ncclDevComm const &comm, ncclCftLeId *leId, size_t *leOffset)

Device-side query for the logical endpoint leId and leOffset corresponding to window, byte offset, peerCft within cftTeam and device comm.

void ncclGetPeerLeInfo(ncclWindow_t w, size_t offset, int peerWorld, ncclDevComm const &comm, ncclCftLeId *leId, size_t *leOffset)

Device-side query for the logical endpoint leId and leOffset corresponding to window, byte offset, peerWorld within the world team and device comm.

void ncclGetMultimemLeInfo(ncclWindow_t w, size_t offset, ncclDevComm const &comm, ncclCftLeId *leId, size_t *leOffset)

Device-side query for the multicast logical endpoint leId and leOffset corresponding to window, byte offset and device comm.

Host-side CFT queries require the communicator to create CFT logical endpoints before the query is issued. Set hostCftMode in ncclConfig_t when creating the communicator if the application needs to query CFT endpoint information from host code before creating a CFT-enabled device communicator.

CFT Operations

struct ncclCftSmem

A user-managed shared memory state object used by ncclCft to track progress of outstanding CFT operations. The object can track a specific amount of in-flight data, limiting the number of outstanding CFT operations. Specifically, the object can track up to 16MB of non-fetching CFT operations (i.e., put, red and their multicast variants) and up to 1MB of fetching CFT operations (i.e., get, pullred).

template<typename Coop>
class ncclCft

A class encompassing major elements of CFT support.

ncclCft(Coop coop, ncclCftSmem &cftSmem)

Initializes a new ncclCft object. cftSmem is a user-managed shared memory state used by threads in coop to make progress on fabric operations and to check for their completion.

void put(OpCoop coop, ncclCftLeId leId, size_t leOffset, void *smemSource, uint32_t bytes)

Initiates an asynchronous transfer of bytes bytes from the shared memory buffer referenced by smemSource to the unicast logical endpoint referenced by leId at offset leOffset.

coop can include a subset of the threads in the cooperative group used to create the CFT object. Cooperative groups initiating fabric operations independently under the same CFT object share its state.

void putCpMask(OpCoop coop, ncclCftLeId leId, size_t leOffset, void *smemSource, uint32_t bytes, uint16_t cpMask)

Variant of ncclCft::put() with an explicit copy mask. The copy mask is a 16-bit integer where bits correspond to the bytes of every 16-byte word being copied. A bit is set to 1 if the corresponding byte ought to be copied to the destination logical endpoint, or it is set to 0 otherwise.

void putMultimem(OpCoop coop, ncclCftLeId leId, size_t leOffset, void *smemSource, uint32_t bytes)

Multicast variant of ncclCft::put() where leId references a logical endpoint for which the the NCCL_CFT_MULTIMEM bit was set in the cftCaps bitmask at device communicator creation time.

void putMultimemCpMask(OpCoop coop, ncclCftLeId leId, size_t leOffset, void *smemSource, uint32_t bytes, uint16_t cpMask)

Multicast variant of ncclCft::put() with an explicit copy mask.

void get(OpCoop coop, ncclCftLeId leId, size_t leOffset, void *smemDestination, uint32_t bytes)

Initiates an asynchronous transfer of bytes bytes from the unicast logical endpoint referenced by leId at offset leOffset to the shared memory buffer referenced by smemDestination.

coop can include a subset of the threads in the cooperative group used to create the CFT object. Cooperative groups initiating fabric operations independently under the same CFT object share its state.

void red(OpCoop coop, ncclCftLeId leId, size_t leOffset, RedOp const &red, void *smemSource, uint32_t bytes)

Initiates an asynchronous reduction operation (see Reduction Operations for reduction tags and data types) for bytes bytes from the shared memory buffer referenced by smemSource to the unicast logical endpoint referenced by leId at offset leOffset.

coop can include a subset of the threads in the cooperative group used to create the CFT object. Cooperative groups initiating fabric operations independently under the same CFT object share its state.

void redMultimem(OpCoop coop, ncclCftLeId leId, size_t leOffset, RedOp const &red, void *smemSource, uint32_t bytes)

Multicast reduction variant of ncclCft::red() where leId references a logical endpoint for which the NCCL_CFT_MULTIMEM bit was set in the cftCaps bitmask at device communicator creation time.

void pullRed(OpCoop coop, ncclCftLeId leId, size_t leOffset, RedOp const &red, void *smemDestination, uint32_t bytes)

Initiates an asynchronous reduction operation (see Reduction Operations for reduction tags and data types) for bytes bytes from the multicast logical endpoint referenced by leId at offset leOffset to the shared memory buffer referenced by smemDestination. leId references a logical endpoint for which the NCCL_CFT_MULTIMEM bit was set in the cftCaps bitmask at device communicator creation time.

coop must be ncclCoopWarp.

void submit(OpCoop coop)

Submits previously initiated CFT operations for all the threads in coop, allowing them to make forward progress. Every cooperative group (OpCoop) that initiated fabric operations must call submit before ncclCft::flushSmem() and ncclCft::flush().

void flushSmem(OpCoop coop)

Makes all the threads in coop wait for CFT operations that consume shared memory data. All the shared memory buffers that source data to outstanding fabric operations can be reused after the function returns. Users can call this function to wait for shared memory buffers to be consumed by previously initiated CFT operations before initiating new ones, and call ncclCft::flush() only when they need to wait for the completion of all the outstanding CFT operations, e.g., because they reached the number of in-flight bytes that the CFT shared memory state object can track.

void flush(OpCoop coop, bool *hasReport = nullptr, uint32_t *report = nullptr)

Flushes outstanding CFT operations for all the threads in coop and waits for their completion. When flush returns shared memory buffers are safe to reuse. coop must include the threads of all cooperative groups that initiated fabric operations under the same CFT object. No other cooperative group can initiate fabric operations until flush has completed. If the operation succeeds hasReport is set to false and report is not modified. If the operation fails hasReport is set to true and report is set to an error code that can be decoded via the cudaFabricOpErrorStatusGet/Count.

smemSource and smemDestination pointers must be 16-byte aligned, and bytes must be a multiple of 16. CFT operates on shared memory; a typical kernel stages data through shared memory, issues one or more operations, then calls ncclCft::submit() and ncclCft::flush() or ncclCft::flushSmem() as required by the reuse and ordering rules of the kernel.

Reduction Operations

Reduction fabric operations use operation tag types:

struct ncclCftOpSum

struct ncclCftOpAnd

struct ncclCftOpXor

struct ncclCftOpOr

struct ncclCftOpMin

struct ncclCftOpMax

Refer to CFT PTX documentation for supported reductions and data types.

CFT Barriers

Applications request CFT barrier resources with cftBarrierCount in ncclDevCommRequirements, then use ncclCftBarrierSession from device code. The barrier count should cover every barrier index used by the kernel.

template<typename Coop>
class ncclCftBarrierSession

A CFT-backed barrier session.

ncclCftBarrierSession(Coop coop, ncclDevComm const &comm, uint32_t index, bool multimem = false)

Initializes a new CFT barrier session. coop represents a cooperative group (see Thread Groups) of threads that synchronize throught the barrier. comm is the device communicator for which barrier resources were requested. index selects the barrier slot from the device communicator. multimem can be set to true to select multicast CFT barrier resources (requires CFT barriers to be created with NCCL_CFT_MULTIMEM capability).

void arrive(Coop coop, cuda::memory_order order, ncclMemProxyType producer, ncclMemProxyType consumer)

Signals arrival at the barrier. order dictates the semantics used to order memory accesses from different proxies around the barrier. producer indicates the type of proxy that stored the data into global memory. consumer indicates the type of proxy that will be loading the producer data from global memory.

void wait(Coop coop, cuda::memory_order order, ncclMemProxyType producer, ncclMemProxyType consumer)

Waits for the barrier to complete. It can be used to establish a happens-before relationship between producer and consumer when order complements the memory ordering semantics of the matching ncclCftBarrierSession::arrive().

void sync(Coop coop, cuda::memory_order order, ncclMemProxyType producer, ncclMemProxyType consumer)

Combines ncclCftBarrierSession::arrive() and ncclCftBarrierSession::wait().

ncclResult_t wait(Coop coop, cuda::memory_order order, ncclMemProxyType producer, ncclMemProxyType consumer, uint64_t timeoutCycles)

ncclResult_t sync(Coop coop, cuda::memory_order order, ncclMemProxyType producer, ncclMemProxyType consumer, uint64_t timeoutCycles)

Timeout-returning variants of ncclCftBarrierSession::wait() and ncclCftBarrierSession::sync().

void ncclMemFence(Coop coop, cuda::memory_order order, ncclMemProxyType producer, ncclMemProxyType consumer, ncclMemFenceScope scope)

Issues a memory fence between producer and consumer proxy domains for the participating threads. coop represents a cooperative group of threads which memory accesses are ordered around the fence. order dictates the semantics used to order memory accesses from different proxies around the fence. producer indicates the type of proxy that stored the data into memory. consumer indicates the type of proxy that will be loading the producer data from memory. scope indicates the point of consistency for the memory accesses around the fence.

enum class ncclMemProxyType

Describes which memory proxy produced or consumes memory around a CFT barrier or a fence.

enumerator Generic

Indicates that the method of memory access is the SM load/store unit.

enumerator Fabric

Indicates that the method of memory access is the fabric.

enum class ncclMemFenceScope

Describes the scope of a CFT memory fence.

enumerator Cta

Indicates that the point of consistency for memory accesses is the SM shared memory.

enumerator Sys

Indicates that the point of consistency for memory accesses is the L2 cache.