1 of 4

Parallelism, collectives, and topology

Splitting a model across GPUs means picking which axis to cut along. Tensor parallelism cuts inside a layer: Megatron-LM partitions the weight matrix of a transformer's MLP block by splitting the first GEMM’s weight matrix along its rows and the input along its columns, so each GPU computes a partial result that has to be combined before the next operation can proceed. That combination is a synchronization point after every parallelized block, which is why tensor parallelism wants the fastest possible link between GPUs, typically NVLink inside a node, and degrades quickly across a slower network. Pipeline parallelism cuts the other way: whole layers are assigned to different GPUs, and activations flow between stages instead of being combined mid-layer. Megatron-LM describes its own tensor parallel approach as “orthogonal and complimentary to pipeline model parallelism,” meaning production systems compose both at once rather than choosing one.

Neither form of parallelism does anything without a way for GPUs to exchange the partial results. That is the job of a collective: an operation like all-reduce, all-gather, or reduce-scatter that every participating GPU executes together. NCCL, NVIDIA’s library for this, implements exactly these primitives: all-reduce, all-gather, reduce, broadcast, reduce-scatter, and arbitrary send/receive patterns, and is “optimized to achieve high bandwidth on platforms using PCIe, NVLink, NVswitch, as well as networking using InfiniBand Verbs or TCP/IP sockets,” supporting any number of GPUs in a single node or spread across many. Tensor parallelism’s per-layer combine step is an all-reduce; pipeline parallelism’s stage-to-stage handoff is a point-to-point send/receive. Which collective a parallelism strategy needs, and how often it needs it, is what determines whether that strategy can tolerate a slow interconnect or requires the fastest one available.

In code, a collective is less exotic than the name suggests. The first complete example in NCCL’s documentation has a single process drive four GPUs: it creates one communicator per device with ncclCommInitAll, then issues an all-reduce on every device inside a group call, which NCCL requires when one thread manages multiple GPUs. Each call names a send buffer, a receive buffer, an element count, a datatype, and a reduction operator; the library sums the four send buffers and leaves the identical result in every receive buffer. Nothing is finished when ncclGroupEnd returns: the operations are queued on CUDA streams, so the host synchronizes each stream before trusting the result.

ncclAllReduce across four GPUs, from the NCCL documentation

ncclComm_t comms[4];
int nDev = 4;
int size = 32*1024*1024;
int devs[4] = { 0, 1, 2, 3 };

//initializing NCCL
NCCLCHECK(ncclCommInitAll(comms, nDev, devs));

//calling NCCL communication API. Group API is required when using
//multiple devices per thread
NCCLCHECK(ncclGroupStart());
for (int i = 0; i < nDev; ++i)
  NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i],
      size, ncclFloat, ncclSum, comms[i], s[i]));
NCCLCHECK(ncclGroupEnd());

//synchronizing on CUDA streams to wait for completion of NCCL operation
for (int i = 0; i < nDev; ++i) {
  CUDACHECK(cudaSetDevice(i));
  CUDACHECK(cudaStreamSynchronize(s[i]));
}

//finalizing NCCL
for (int i = 0; i < nDev; ++i)
  ncclCommDestroy(comms[i]);

This is why topology is not a footnote to parallelism, it is a constraint on which parallelism strategies are viable at all. A group of GPUs connected by NVLink can absorb tensor parallelism’s frequent all-reduces; GPUs connected only by a network fabric across nodes generally cannot, which is why tensor parallelism is typically confined inside a node and pipeline or data parallelism is used to scale beyond it, where the coarser, less frequent communication tolerates the added latency.

Tensor vs. pipeline parallelism. The same four-layer model split two ways. Tensor parallelism cuts every layer in half across both GPUs, so activations must be all-reduced at every layer boundary. Pipeline parallelism gives each GPU whole layers and sends activations across the boundary once.Illustrative numbers

The size of that fastest tier is a moving target. NVIDIA’s multi-node tuning guide notes that before the GB200 NVL72, an NVLink domain topped out at eight GPUs on an HGX H200 baseboard at 900 GB/s of communication per GPU; the NVL72 rack design extends one domain to 72 Blackwell GPUs at 1.8 TB/s each, with the NVLink switch providing 130 TB/s of aggregate GPU bandwidth inside the domain, and describes that 72-GPU domain as acting as one massive GPU. For the parallelism strategies above, that is a ninefold change in where the scale-up boundary sits: a split that would have crossed the network fabric between eight-GPU nodes can now stay on NVLink across a whole rack.

NVLink is one vendor’s fabric, and open standards now exist for both directions of scaling. UALink is the scale-up one: its 200G 1.0 specification “defines a low-latency, high-bandwidth interconnect for communication between accelerators and switches in AI computing pods” and “enables 200G per lane scale-up connection for up to 1,024 accelerators within an AI computing pod.” Ultra Ethernet is the scale-out counterpart, a Linux Foundation project whose stated mission is an “Ethernet based open, interoperable, high performance, full-communications stack architecture” for AI and HPC at scale; the problems it names, multi-pathing, fast reaction to congestion, and flows “where tail latency is the figure of merit,” are this chapter’s constraints restated as networking requirements. Whichever fabric wins a given deployment, the collective at the top of the stack stays the same; what the standards change is who can build the hardware underneath it.