NCCL's Device API
August 04, 2026
Prerequisites
- You know GPU chip organization well and have comfort with CUDA programming and concepts like threadblocks.
- You know what collective communication is and the core primitive patterns.
- You understand that NCCL is a stream-aware instantiation of GPU-centric MPI.
Hardware Background
Motivation
The problem is pretty straightforward: there’s some computational workload (e.g. running inference on a large language model), and it’s takes too much time or memory to run on a single GPU. Multi-GPU programming is all about getting data from the memory of one GPU to another: everything boils down to this one single task!
Chips already have an internal memory hierarchy that needs to be exploited cleverly for the best performance. You can think about multi-GPU programming as expanding this hierarchy even further beyond HBM: there’s the intranode domain, internode domain, and (increasingly) inter-data-center domain (in industry people like to call this scaling up, scaling out, and scaling “across” respectively).
Like the chip memory hierarchy, your available bandwidth goes down substantially the farther away your GPUs are from each other.
Intranode communication (LSA-Accessible + Multimem)
Things are easy when all your GPUs are in the same NVLink domain (i.e. they occupy the same NVIDIA-specified rack architecture). Here, we can leverage an amazing property called unified virtual addressing. Essentially, you can pretend that all your GPU memory addresses are lined up one after the other, and then directly use the pointers to peer GPUs to load and store data just like normal.
If your NVIDIA system has an NVSwitch that is “SHARP-enabled”, you can also use a special multimem capability that accelerates broadcasts and reductions in the switch itself, allowing your GPU SMs to stay busy with other important work.
Normally this all requires IPC and some nasty cuMemMapping yourself, but NCCL now has a simplified way to do all this in a much more straightforward way.
Internode communication (GIN)
If you want to communicate between two different GPU racks, you have to do so through dedicated network interface cards (NICs). These NICs are connected through a specific kind of networking “fabric” (e.g. Infiniband or Ethernet), and they leverage a special technology called RDMA which allows NICs to directly access and transmit data in GPU HBM without the involvement of the OS kernel, which stands in contrast to traditional computer networking.
This has fascinating implications for the internode programming model: instead of “sockets” you talk to other NICs using the concept of “queue pairs” (QPs), where there is a continuous “queue” of send messages, receive messages, and completion messages that the CPU can poll.
The typical data movement workflow is:
- Rank A allocates a QP, rank B allocates a QP, and they connect to each other.
- Under the hood, a QP is just some pinned memory buffers with bookkeeping that the NIC continuously polls and monitors.
- This QP is point to point: so one QP needs to be allocated per connection.
- Rank A put WQEs (“wookies”) into a send queue.
- Rank A rings the doorbell on the NIC” (i.e. writes to a memory mapped “doorbell” register on the device), and this triggers the NIC to read the WQE and execute the network transaction.
Now this description glosses over something rather important: Rank A is ringing the doorbell, but who is doing it - the CPU or GPU?
Traditionally, this work is done by the CPU using the concept of a proxy thread: this is a dedicated CPU thread whose sole job is to initiate network operations (i.e. read and write data to and from and the NIC).
Thanks to some great work by NVIDIA, there’s new technologies like GDAKI and DOCA, which allow the GPU device itself to work with the NIC buffer. This way, the host is not in the critical path at all! This is what GIN builds on to do device-side comms.
Important Note: This was incredibly confusing when I was first learning GIN, but note that the word “context” can mean two very different things:
-
a “NCCL context” or “GIN context” is essentially an object that wraps around 1+ QPs (along with some other connection-specific info), with GIN handling most of the WQE enqueuing and doorbell operations. So intuitively, a (NCCL) “context” can reach EVERY rank, but a specific QP can only reach ONE rank. A NCCL context must therefore contain only one QP per peer: you need multiple contexts for more independent QPs to the same set of peers.
-
a “CTX” (used by NCCL internally) essentially refers to a QP itself. For the rest of this blog post, a “context” refers to a NCCL context, and “CTX” refers to an individual QP connection.
NCCL Device API Programming Model
Host Setup
Before you can even think about device-initiated comms, you have to do some important setup on the host itself. You have to establish this idea of a NCCL communication window, which is essentially a lookup table (for each Rank i, we want to know (ptr_i, rkey_i) where ptr_i is the memory address of rank i, and rkey_i is rank i’s associated NIC “permission token”).
This way, you can allocate memory however you please (via cudaMalloc or PyTorch tensors). If you want to this buffer to later be globally accessible, you can register it as a communication window!
Why this separation? NVIDIA already has another communication library (NVSHMEM) where allocation and registration are the exact same thing: nvshmem_malloc both allocates GPU memory and places it into a symmetric heap, with a pointer that’s valid across all PEs.
Here’s a canonical example of what host-side setup looks like:
int main(int argc, char** argv) {
// MPI init
// standard NCCL bootstrap using ncclUniqueId
// standard NCCL initialization:
ncclComm_t comm;
ncclCommInitRank(&comm, nRanks, id, myRank);
// GIN-specific setup of communicators
// This is a device communicator: it's a separate subset of the host communicator specifically used for device API operations.
ncclDevComm devComm;
ncclDevCommCreate(&devComm, comm, /* enable GIN */ nullptr);
// Rank i allocates GPU buffers like normal
float *sendBuf, *recvBuf;
cudaMalloc(&sendBuf, N * sizeof(float));
cudaMalloc(&recvBuf, N * sizeof(float));
// ... Fill the sendBuf with something valuable
// COLLECTIVE registration: ranks call this together and exchange (pts + rkeys)
// We take the previous memory buffers we allocated them, and register them as "communication windows" - symmetric memory!
ncclWindow_t sendWin, recvWin;
ncclCommWindowRegister(comm, sendBuf, N * sizeof(float), &sendWin);
ncclCommWindowRegister(comm, recvBuf, N * sizeof(float), &recvWin);
// When launching the kernel, pass in the ncclWindow_t objects!
ringExchange<<<1, 32>>>(devComm, sendWin, recvWin, myRank, nRanks);
cudaDeviceSynchronize();
}
LSA/Multimem Kernel
NCCL LSA and multimem is incredibly simple and straightforward: NCCL basically just gives you a wrapper around exisitng UVA pointers or multimem pointers respectively. You can see examples here.
Once you have the pointer you want, it becomes trivial to do intranode programming by just dereferencing these device memory addresses.
GIN Kernel
Here’s an example of a kernel using GIN for inter-node communication via NICs. Doing this kind of thing used to require a lot more code!
__global__ void ringExchange(ncclDevComm devComm,
ncclWindow_t sendWin,
ncclWindow_t recvWin,
int myRank, int nRanks)
{
ncclGin gin(devComm, 0); // every gpu thread calls this constructor.
// This is a lightweight object acts as a handle through which to issue network operations.
int peer = (myRank + 1) % nRanks;
gin.put(
ncclTeamWorld(devComm), // communicator group
peer, // the peer I want to put into
recvWin, myRank * N * sizeof(float), // DST + OFFSET: I'm putting it into THEIR recvWin buffer, at offset *function of myRank*
sendWin, 0, // SRC + OFFSET: I'm getting this data from MY sendWin at offset 0
N * sizeof(float), // amount of data to put
ncclGin_SignalInc{0} // optional: lets you signal at the very end
);
}
GIN data movement operations
To do data movement, we only care about 3 core functions: put, get, and signal. Everything else (e.g. putValue) is for convenience.
void put(
ncclTeam team, // (1) which team/group of peer indices I'm looking at
int peer, // (2) which rank within `team` to write into
ncclWindow_t dstWnd, // (3) the PEER's registered window (where to write)
size_t dstOffset, // (4) byte offset within that peer's window
ncclWindow_t srcWnd, // (5) MY registered window (where to read from)
size_t srcOffset, // (6) byte offset within my own window
size_t bytes, // (7) how much data to move
// ---- everything below is OPTIONAL, has defaults ----
RemoteAction remoteAction, // (8) what to do on the PEER once data lands
// e.g. ncclGin_StrongSignalInc{sigId}
// default: ncclGin_None (no action)
LocalAction localAction, // (9) what to do on MY side once the send
// is locally consumed
// e.g. ncclGin_CounterInc{counterId}
// default: ncclGin_None
Coop coop, // (10) which threads are cooperating on this
// ONE call (default: ncclCoopThread —
// a single thread issues it; this is
// the recommended/normal case)
DescriptorSmem descriptor, // (11) scratch/staging shared memory the
// implementation may use internally
// to build the descriptor faster
// (an optimization hook, not something
// you need to reason about content-wise)
cuda::thread_scope alreadyReleased, // (12) a hint: "I've already done a memory
// fence up to this scope, don't redo it"
// — lets you skip redundant fencing if
// you know your data is already released
// at this scope
cuda::thread_scope expected_scope, // (13) the scope at which YOU need the source
// data to be visible before this put reads
// it (drives automatic fence insertion)
SegmentType bufType // (14) what kind of memory the src/dst
// buffers physically are
// default: ncclGin_SegmentDevice
// (plain GPU memory — the normal case)
);
void get(
ncclTeam team, // (1) same as put
int peer, // (2) same as put
ncclWindow_t remoteWnd, // (3) the PEER's window — SOURCE this time
size_t remoteOffset, // (4) offset in their window to read from
ncclWindow_t localWnd, // (5) MY window — DESTINATION this time
size_t localOffset, // (6) offset in my window to write into
size_t bytes, // (7) how much to pull
// ---- optional, with defaults ----
Coop coop = ncclCoopThread{}, // (8) same meaning as put
DescriptorSmem descriptor = ncclGin_None{}, // (9) same meaning as put
uint32_t optFlags = ncclGinOptFlagsDefault, // (10) generic flags — doc doesn't
// enumerate what's available;
// default is presumably fine
// for normal use
SegmentType bufType = ncclGin_SegmentDevice{} // (11) same meaning as put
);
// for signal, we ONLY tell you the (team + peer) to signal. This avoids having to put(payload_size=0).
void ncclGin::signal(
ncclTeam team, // (1) same as put/get
int peer, // (2) same as put/get
RemoteAction remoteAction, // (3) what signal to bump on the peer
// e.g. ncclGin_StrongSignalInc{0}
// — this is the WHOLE POINT of the call
Coop coop, // (4) same meaning as put
DescriptorSmem descriptor, // (5) same meaning as put
cuda::thread_scope alreadyReleased, // (6) same meaning as put
cuda::thread_scope expected_scope // (7) same meaning as put
);
Example of two ranks using the signal API:
__global__ void signalDemo(ncclDevComm devComm, int myRank)
{
// Bind to context 0 — every rank does this the same way
ncclGin gin(devComm, 0);
// Signal ID 0 — an arbitrary pre-allocated "mailbox slot" both
// ranks agree to use. Doesn't need to be "registered" like a
// window — it's just an index into a fixed-size signal array
// that GIN sets up per-context.
constexpr uint32_t SIGNAL_ID = 0;
if (myRank == 0) {
// ---- SENDER ----
// No data transfer at all — just tell rank 1 "bump signal 0"
gin.signal(
ncclTeamWorld(devComm), // team: everyone in the communicator
/* peer = */ 1, // target rank
ncclGin_StrongSignalInc{SIGNAL_ID} // action: increment signal 0 on rank 1
);
printf("Rank 0: signal sent\n");
}
if (myRank == 1) {
// ---- RECEIVER ----
// Block until signal 0's value reaches (at least) 1
gin.waitSignal(ncclCoopThread{}, SIGNAL_ID, /* least = */ 1);
// At this point, the increment has happened —
// safe to assume rank 0's signal() has "landed"
printf("Rank 1: got the signal! doing my thing now\n");
// Reset it so it can be reused next round
gin.resetSignal(SIGNAL_ID);
}
}
GIN flush operations
These functions will block execution until the operations occuring on my local side are fully complete. Note I have no way of knowing what operations or data are completed on my peer without some kind of signal/notification!
- flush / flushAsync / wait: asks if all pending ops on this context are done.
- readCounter / waitCounter / resetCounter: asks if this specific op/pipeline stage is done. Very fine-grained, this is done per-operation.
void flush(
Coop coop, // (1) which threads are jointly waiting
// here (must match how you issued
// the ops you're flushing — see below)
cuda::memory_order ord = cuda::memory_order_acquire
// (2) the memory-ordering strength applied
// once flush returns — default is
// "acquire," meaning: after this call
// returns, any subsequent reads by this
// thread are guaranteed to see the
// settled state (standard C++/CUDA
// acquire semantics)
);
// Example: pull data from a peer, then flush to know it's safe to read
ncclGin gin(devComm, 0);
gin.get(team, peer, remoteWin, 0, myWin, 0, bytes);
gin.flush(ncclCoopThread{}); // blocks until the pulled data is actually in myWin
// beyond this flush call, all source PUT buffers are safe to reuse!
// we can't say much about completion on remote peers: only that my send buffer is free to overwrite. For GET it's whether the data is actually in my local buffer.
float val = *(float*)myBufPtr; // NOW safe to read
flushAsync is very similar except it’s nonblocking: to block you just have call an associated .wait().
void flushAsync(
ncclTeam team, // (1) team defining valid peer indices
uint32_t peer, // (2) flush only ops targeting THIS peer
// (unlike plain flush, which covers
// everything the coop issued — this
// is scoped to one specific peer)
ncclGinRequest_t *request, // (3) OUT parameter: caller-provided,
// initialized by this call — this is
// the handle you'll later pass to wait()
Coop coop = ncclCoopThread{}, // (4) same meaning as flush
uint32_t optFlags = ncclGinOptFlagsDefault, // (5) generic flags, default fine
DescriptorSmem descriptor = ncclGin_None{} // (6) same scratch-space hook as put/get
);
ncclGinRequest_t req;
gin.flushAsync(team, peer, &req);
// ... do other useful work here while the flush is in-flight ...
gin.wait(req); // now actually block until it's done
GIN notificaiton operations
You can use the following functions to notify a remote peer about completion:
// Sender:
gin.signal(team, peer, someAction); // "hey peer, bump your signal"
// Receiver, pick one:
gin.waitSignal(coop, signalId, 1); // block until it happens
uint64_t v = gin.readSignal(signalId); // or just check right now, no blocking
// Receiver, cleanup:
gin.resetSignal(signalId); // reset to 0 for next round
Note that there’s some complexity in what’s passed in as the signal:
-
Signal itself can be (1) a small int ID, or (2) (window, offset) tuple. Use the latter if the signal’s location needs to be computed dynamically at a specific address: this is usually kind of niche and most apps don’t need it.
- “Weak signal” only guarantees this specific signal’s data is visible (what you expect)
- “Strong signal” guarantees this signal, as well as everything I sent to said peer earlier, has arrived.
GIN barrier operations
Barriers are a simple idea: no rank leaves the barrier until every rank has arrived at it. In GIN, you construct a ncclGinBarrierSession and then call .sync on that object.
__global__ void barrierDemo(ncclDevComm devComm, int myRank)
{
ncclGin gin(devComm, 0);
// ---- BARRIER ----
// No rank gets past this line until EVERY rank has reached it.
ncclGinBarrierSession bar(
ncclCoopThread{}, // just this one thread participating
gin,
ncclTeamWorld(devComm), // sync across ALL ranks
devComm.railGinBarrier, // where the barrier's internal state lives
/* index = */ 0 // barrier slot id
);
bar.sync(ncclCoopThread{}, cuda::memory_order_acquire);
// ---- Everyone is guaranteed to be HERE now ----
printf("Rank %d: phase 2 starting, everyone is ready\n", myRank);
}