Compute Fabric Transport

Starting with NCCL 2.31, the Device API includes Compute Fabric Transport (CFT) helpers for kernels that communicate through CUDA fabric logical endpoints. CFT is useful when an application launches custom device-side communication kernels on CFT-capable systems and needs direct fabric put, get, reduction, or barrier operations.

Requirements

CFT kernels require hardware, driver, and CUDA Toolkit support for fabric logical endpoints and fabric PTX instructions. In this release, the example CFT barrier program documents the practical requirements as CUDA Toolkit 13.3 and SM_100 architectures.

Before using CFT in an application:

Allocate and register symmetric memory windows with NCCL.
Create a device communicator with CFT capabilities in ncclDevCommRequirements.
Use CFT team helpers to address peers by CFT team rank.
Query logical endpoint IDs and offsets before issuing CFT operations.
Match the number of requested CFT barriers to the number of barrier slots used by the kernel.

Device Communicator Setup

The host code requests CFT capability when creating the device communicator.

ncclDevComm devComm;
ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;

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

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

Request NCCL_CFT_MULTIMEM in addition to NCCL_CFT when a kernel uses multicast CFT operations or multimem CFT barriers:

reqs.cftCaps = NCCL_CFT | NCCL_CFT_MULTIMEM;

If a communicator cannot provide the requested CFT capability on all ranks, ncclDevCommCreate() fails.

Teams and Endpoints

CFT operations may address peers through CFT teams rather than directly through world ranks. Use ncclTeamCft() or ncclTeamCftMultimem() on the host side, and the corresponding device overloads in kernels, to obtain the team layout.

Use logical endpoint query helpers to translate a registered window, byte offset, and peer into the values consumed by ncclCft operations. For device-side kernels, ncclGetCftLeInfo() accepts a CFT-team peer rank and ncclGetPeerLeInfo() accepts a world-rank peer.

ncclTeam cftTeam = ncclTeamCft(devComm);
ncclCftLeId leId;
size_t leOffset;

ncclGetCftLeInfo(win, byteOffset, peerCft, cftTeam, devComm,
                 &leId, &leOffset);

Host code can use ncclGetCftDeviceLeInfo(), ncclGetPeerDeviceLeInfo(), and ncclGetMultimemDeviceLeInfo(). If host-side endpoint queries are needed before creating a CFT-enabled device communicator, configure hostCftMode in ncclConfig_t during communicator initialization.

CFT Operations

A CFT kernel stages data through shared memory, creates a ncclCft object, issues operations, and then submits and flushes the work.

__global__ void cftPutKernel(ncclDevComm devComm, ncclWindow_t win) {
  __shared__ ncclCftSmem cftSmem;
  __shared__ alignas(16) char smem[128];

  ncclCoopCta coop;
  ncclCft<ncclCoopCta> cft{coop, cftSmem};

  ncclTeam cftTeam = ncclTeamCft(devComm);
  int peer = (cftTeam.rank + 1) % cftTeam.nRanks;

  ncclCftLeId leId;
  size_t leOffset;
  ncclGetCftLeInfo(win, 0, peer, cftTeam, devComm, &leId, &leOffset);

  cft.put(coop, leId, leOffset, smem, sizeof(smem));
  cft.submit(coop);
  cft.flush(coop);
}

CFT source and destination shared memory pointers must be 16-byte aligned, and the byte count must be a multiple of 16. See Device API - CFT for the full list of put, get, multicast, reduction, and pull-reduce helpers.

CFT Barriers

CFT barriers synchronize device threads across the CFT team. The common setup is to request one barrier per CTA and use blockIdx.x as the barrier index.

__global__ void cftBarrierKernel(ncclDevComm devComm) {
  ncclCoopCta coop;
  ncclCftBarrierSession<ncclCoopCta> bar{coop, devComm, blockIdx.x};

  bar.sync(coop, cuda::memory_order_acq_rel,
           ncclMemProxyType::Generic, ncclMemProxyType::Fabric);
}

For multicast CFT barriers, create the device communicator with NCCL_CFT_MULTIMEM and pass multimem=true to the barrier session constructor.

CFT Counted Operations

Counted operations (see CFT counted operations) enable latency optimized communication by removing the ‘data + memory fence + flag update’ synchronization pattern used in CFT put/red. In CFT counted operations, the flag is replaced by a user-managed counter in symmetric memory, appropriately registered with a CFT counted window (i.e., a window created with the NCCL_WIN_CFT_COUNTED flag). The counter update is ordered after the data, so that observing the counter update guarantees that the data can be accessed. Memory fencing between counter and data access at the destination is needed.

Like other CFT operations, shared and global memory operands must be 16-byte aligned and the size of the counted operation must be a multiple of 16 bytes. The user-managed counter in global memory should be 256-byte aligned, for best performance. The number of concurrently active counters per GPU should not exceed 32, for best performance.

The following example code shows how the user-managed counter should be allocated and registered with the symmetric memory window, and how it is used in a counted put operation. It is important to note that the initiator of the put computes the counter offset from the leOffset, while the recepient uses the counter’s VA pointer to poll for updates using waitCounted. It should also be noticed that waitCounted is a convenience API providing cooperative thread synchronization, memory ordering between counter updates and data accesses, and aborts/timeouts handling. Users can implement their own checks at the recipient rank, provided that they handle all the aforementioned requirements.

// Host Code
size_t dataCount = 4;
size_t dataSize = ALIGN_UP(dataCount * sizeof(float), /*alignment*/size_t(16u));
size_t counterSize = ALIGN_UP(/*counter*/sizeof(uint64_t), /*alignment*/size_t(256u));
size_t memSize = dataSize + counterSize;
void* buffer;
ncclMemAlloc(&buffer, memSize);

ncclWindow_t win;
ncclCommWindowRegister(comm, buffer, memSize, &win, NCCL_WIN_COLL_SYMMETRIC | NCCL_WIN_CFT_COUNTED);

ncclDevComm devComm;
ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
reqs.cftCaps = NCCL_CFT;
cftCountedKernel<<</*ctas*/1, /*threads*/512, /*smemBytes*/dataSize>>>(devComm, win, dataCount);

// Device Code
__global__ void cftCountedKernel(ncclDevComm devComm, ncclWindow_t win, size_t count) {
  ncclCoopCta coop;
  ncclTeam team = ncclTeamCft(devComm);

  // Get local counter pointer
  size_t dataSize = dataCount * sizeof(float);
  float* buf = (float*)ncclGetLocalPointer(win, /*offset*/0);
  uintptr_t base = reinterpret_cast<uintptr_t>(buf);
  uintptr_t end = ALIGN_UP(base, uintptr_t(16u)) + uintptr_t(dataSize);
  uintptr_t counterAddr = ALIGN_UP(end, uintptr_t(256u));
  uint64_t* counterPtr = reinterpret_cast<uint64_t*>(counterAddr);

  extern __shared__ alignas(16) unsigned char smemScratch[];
  float* smem = reinterpret_cast<float*>(smemScratch);
  smem[0] = smem[1] = smem[2] = smem[3] = 1.0f;
  __shared__ ncclCftSmem cftSmem;
  ncclCft<ncclCoopCta> cft { coop, cftSmem };

  if (team.rank == 0) {
    ncclCftLeId leId;
    size_t leOffset;
    ncclGetCftLeInfo(win, /*offset*/0, /*peer*/1, team, devComm, &leId, &leOffset);
    size_t dataOffset = ALIGN_UP(leOffset, size_t(16u));
    size_t counterOffset = ALIGN_UP(dataOffset + dataSize, size_t(256u));
    cft.putCounted(coop, leId, dataOffset, counterOffset, smem, /*bytes*/dataSize);
    cft.submit(coop);
    cft.flush(coop);
  } else /*if (team.rank == 1)*/ {
    cft.waitCounted(coop, cuda::memory_order_acquire, ncclMemProxyType::Generic, counterPtr, /*bytes*/dataSize, nullptr);
  }
}

Examples

See the CFT barrier example under docs/examples/06_device_api/04_cft_barrier for a complete runnable CFT setup and kernel.

Cross-proxy Fences

Cross-proxy fences enforce memory ordering between operations issued by different producer and consumer proxies. The common use case for fence is ordering of data and flag updates to establish a happens-before relationship between producer stores and consumer loads to the same memory region. In the following example, the producer is the Fabric proxy, while the consumer (ommitted) is the Generic proxy. The consumer polls on the flag using relaxed loads and, after observing the flag update, issues a fence with memory_order_acquire, Fabric producer, Generic consumer and memory scope Sys.

__global__ void cftPutKernel(ncclDevComm devComm, ncclWindow_t win) {
  __shared__ ncclCftSmem cftSmem;
  __shared__ alignas(16) char smem[128];
  __shared__ alignas(16) int flag[4];

  ncclCoopCta coop;
  ncclCft<ncclCoopCta> cft{coop, cftSmem};

  ncclTeam cftTeam = ncclTeamCft(devComm);
  int peer = (cftTeam.rank + 1) % cftTeam.nRanks;

  ncclCftLeId leId;
  size_t leOffset;
  ncclGetCftLeInfo(win, 0, peer, cftTeam, devComm, &leId, &leOffset);

  cft.put(coop, leId, leOffset, smem, sizeof(smem));
  cft.submit(coop);
  cft.flush(coop);

  ncclMemFence(coop, cuda::memory_order_release, ncclMemProxyType::Fabric, ncclMemProxyType::Generic, ncclMemFenceScope::Sys);

  cft.put(coop, leId, leOffset + flagOffset, flag, 16);
  cft.submit(coop);
  cft.flush(coop);
}

Another use case for fence is ordering accesses to shared memory between the Generic proxy and the Fabric proxy. In the following example, the Generic proxy updates shared memory and makes the updates visible to the Fabric proxy using a fence.

__shared__ncclCftSmem cftSmem;
ncclCoopCta coop;
ncclCft<ncclCoopCta> cft{coop, cftSmem};

__shared__ alignas(16) payload[4];

payload[0] = payload[1] = payload[2] = payload[3] = 1;

ncclMemFence(coop, cuda::memory_order_release, ncclMemProxyType::Generic, ncclMemProxyType::Fabric, ncclMemFenceScope::Cta);

cft.put(coop, leId, leOffset, payload, 16);
cft.submit(coop);
cft.flush(coop);