A collective is a communication operation in which every GPU in a group participates: each one contributes data, receives data, or both. The workhorse of distributed deep learning is one such operation. After every training step, the gradients computed independently on each GPU must be summed and the sum handed back to all of them, so that every GPU applies the same update to its copy of the model. That pattern is an all-reduce, and on a large training run it can consume as much wall-clock time as the arithmetic it synchronizes.
The standard collectives are as follows:
Broadcast sends one GPU’s buffer to every other GPU.
Reduce combines a buffer from every GPU with an operator (sum, maximum, and so on) and leaves the result on one GPU; All-Reduce leaves the result on all of them.
All-Gather concatenates each GPU’s buffer so that every GPU ends with the whole.
Reduce-Scatter reduces as All-Reduce does but partitions the result, leaving each GPU one slice.
An all-reduce is exactly a reduce-scatter followed by an all-gather, and the fastest implementations exploit that decomposition to move each byte the minimum number of times.
NCCL (the NVIDIA Collective Communications Library,
pronounced “nickel”) implements these operations for NVIDIA GPUs. Its
purpose is to hide the topology of Section 9.7 behind a small interface:
a program creates a communicator holding one rank per GPU,
calls ncclAllReduce(), ncclBroadcast(), and the rest, and NCCL selects
an algorithm matched to how the GPUs are physically connected. Within a
node, it moves data over NVLink; across nodes, it uses GPUDirect RDMA
over InfiniBand or RoCE; where no faster link is available, it falls
back to PCIe. It arranges the ranks into rings and trees so that each
link runs near its bandwidth and the collective’s cost grows slowly as
GPUs are added.
A NCCL collective executes as a CUDA kernel enqueued into a stream,
so it composes with the rest of a program’s work. A gradient all-reduce
can be issued into its own stream and overlapped with the backward-pass
kernels that are still producing later gradients, hiding the
communication behind computation. Several collectives bracketed by
ncclGroupStart() and ncclGroupEnd() are fused into a single launch,
which matters when many small tensors must be exchanged at once.
SHARP takes the reduction off the endpoints altogether. On an InfiniBand fabric built from NVIDIA Quantum switches, the Scalable Hierarchical Aggregation and Reduction Protocol performs the summation inside the switch as the data flows through it, rather than on the GPUs. The switch ASIC adds the contributions arriving on its ports and forwards one reduced result, so an all-reduce moves roughly half the data an endpoint-based ring would: each node sends its data in once and receives the answer, instead of pushing the full volume around a ring. NCCL drives SHARP through a plugin, so an application collects the benefit without changing a single collective call.
Multi-node CUDA programs still need a way to launch a process on each
node and give each one an identity, and that layer is almost always MPI.
A job started with mpirun (or a scheduler’s
srun) brings up one rank per GPU across the cluster, and
the ranks find one another through MPI’s communicator. A NCCL program
bootstraps on exactly this: one rank calls ncclGetUniqueId() and
broadcasts the result over MPI, and every rank passes that identifier to
ncclCommInitRank() to join the communicator.
MPI also moves data in its own right. A CUDA-aware MPI
accepts device pointers directly in MPI_Send(), MPI_Recv(), and the
collective calls, and carries the bytes over GPUDirect RDMA (Section
9.7) with no staging through host memory – the same transport NCCL rides
on. The usual division of labor is that MPI handles process launch,
ranks, and irregular point-to-point messaging, while NCCL handles the
dense, regular collectives that dominate deep-learning training; a large
training job runs both, MPI underneath and NCCL for the gradient
exchange.