LearnProject Deep DivesDeep Dive - NCCL

⚡HardFine-Tuning & Training

Deep Dive - NCCL

Read NCCL as the communication engine beneath distributed AI: collective contracts, rings and trees, topology discovery, CUDA streams, transports, profiling, and hang diagnosis.

34 min read

Learning path

Step 185 of 196 in the full curriculum

Deep Dive - DeepGEMM Deep Dive - Megatron

Personalize this lesson

Adapt explanations and teaching visuals to your background and preferred voice.

The last chapter examined one GPU's matmul. Now suppose four GPUs each spend 8 ms computing local gradients. The step then waits 6 ms to combine a 256 MiB (mebibyte) gradient buffer from each GPU. These are illustrative timings, not a hardware benchmark.

A standalone communication test may call that transfer fast, while a real training step exposes more or less time depending on whether the move overlaps other work. If one GPU enters a different communication operation, the other three can wait at a call that looks like a network failure.

NCCL (NVIDIA Collective Communications Library, pronounced "nickel") is the library that coordinates that movement on NVIDIA systems. It provides collective operations such as all-reduce, reduce-scatter, all-gather, broadcast, and all-to-all, plus point-to-point send and receive. NCCL isn't a training framework or scheduler. PyTorch, DeepSpeed, Megatron, and serving runtimes decide what to communicate; NCCL decides how to move it across available GPU links.^{[1]Reference 1NVIDIA Collective Communication Library User Guidehttps://docs.nvidia.com/deeplearning/nccl/user-guide/index.html}

Animated four-rank all-reduce of the first vector slot. Values 1, 10, 100, and 1000 move into a logical sum; four copies of 1111 then sit on ranks 0 through 3. — Follow the first slot of the running example: 1, 10, 100, and 1000 become 1111 on every rank. The center sigma is logical ownership, not a claim that NCCL ships every byte through one physical reducer.

The collective contract

A rank is one participant in a communication group. A communicator records that group and maps ranks to devices. Each rank calls the same collective in the same order, with compatible counts and data types, the same reduction operation, and the same root when the operation has one. A root is a rank index, not a GPU device number. Breaking that agreement can cause a hang, crash, or corrupted result.^{[2]Reference 2NCCL Collective Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html}

For sum all-reduce over $k$ ranks, rank $r$ contributes a vector $\mathbf{x}^{(r)}$ with $N$ elements. Every rank receives:

\mathbf{y}_i = \sum_{r=0}^{k-1} \mathbf{x}^{(r)}_i, \qquad i \in \{0, \ldots, N-1\}.

Here, $k$ is rank count, $N$ is elements per input vector, and $i$ selects one element. This example requests ncclSum, not an average. NCCL also provides ncclAvg. The caller must choose the reduction and scaling needed by its training objective, without dividing twice.^{[3]Reference 3NCCL Typeshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/types.html}

Choose the operation from desired ownership

The output placement matters as much as the arithmetic. Use all-reduce when every rank needs the whole result. Reduce-scatter fits one reduced shard per rank. Existing shards become a complete tensor through all-gather.

Operation	Input on each rank	Output ownership	Common AI use
All-reduce	Same-shaped tensor	Full reduced tensor on every rank	Replicated data-parallel gradients
Reduce-scatter	Same-shaped tensor	One reduced chunk per rank	Sharded gradients in FSDP or ZeRO
All-gather	One local chunk	Concatenated chunks on every rank	Parameter or activation reconstruction
Broadcast	Full tensor on root	Root tensor copied to all ranks	Seeds, metadata, or initial weights
All-to-all	Destination chunks	One chunk from every source	MoE token routing
Send / receive	Point-to-point buffer	Named peer only	Pipeline stage activations

Broadcast is the simple copy in that table: rank 0 keeps a buffer, and the others receive it. The animations below follow the operations whose ownership is easier to mix up, using the same four-rank numbers you'll compute by hand next.

Animated NCCL reduce-scatter on the four-rank running example. Matching chunk indices reduce to 1111, 2222, 3333, and 4444; rank r keeps only chunk r. — Color marks chunk index. Matching slots reduce to 1111, 2222, 3333, and 4444, then rank r keeps only chunk r. Nobody finishes with a full tensor.

Animated NCCL all-gather on the four-rank running example. Shards 1111, 2222, 3333, and 4444 leave their owners, concatenate in rank order, and then copies of that full tensor sit on every rank. — All-gather starts where reduce-scatter ends. The four owned shards concatenate in rank order, then every rank holds `[1111, 2222, 3333, 4444]`.

Animated NCCL all-to-all. Sixteen labeled payloads move from source-rank rows to destination-rank rows, preserving source identity while ownership changes. — Read color as destination and the label as a source-to-destination address. Rank 2 finishes with one rank-2-bound payload from every source. This is an equal-chunk exchange, not a reduction.

NCCL's all-to-all uses equal-sized chunks. MoE routing may produce unequal token counts, so a runtime needs padding, matched point-to-point transfers, or another variable-size dispatch implementation. Don't assume the equal-chunk picture alone implements a complete MoE router. The guide also documents Gather (collect at a root) and Scatter (distribute from a root).^{[2]Reference 2NCCL Collective Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html}

An all-reduce can be implemented as reduce-scatter followed by all-gather. The first half adds corresponding chunks and leaves one finished chunk on each rank. The second half circulates those finished chunks until every rank has the full result. This is an equivalence of operations, not a claim that every NCCL algorithm executes these two API calls. Floating-point summation order can also change the last bits.^{[2]Reference 2NCCL Collective Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html}

Work one all-reduce by hand

Use four ranks with four-element vectors:

Rank	Input vector
0	`[1, 2, 3, 4]`
1	`[10, 20, 30, 40]`
2	`[100, 200, 300, 400]`
3	`[1000, 2000, 3000, 4000]`

Elementwise addition gives [1111, 2222, 3333, 4444]. NCCL's reduce-scatter contract places that result by rank index: rank 0 keeps 1111, rank 1 keeps 2222, rank 2 keeps 3333, and rank 3 keeps 4444. All-gather then shares those four owned chunks so every rank reconstructs the same vector.^{[2]Reference 2NCCL Collective Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html}

In fully sharded training, such as FSDP or ZeRO stage 3, each rank can keep a reduced gradient shard, update its optimizer and parameter shards, then gather parameters for later computation. That later gather is of updated parameters, not necessarily the same gradient buffer. Earlier ZeRO stages have different ownership rules.

The sum is easy on paper. Moving those bytes without a central reducer that would serialize every link is the actual constraint. A ring is the bandwidth-friendly way to do that.

Why rings work well for large messages

A ring orders ranks so each sends to one neighbor and receives from another. Split a message into $k$ chunks. Reduce-scatter circulates and combines chunks for $k-1$ steps. All-gather circulates the completed chunks for another $k-1$ steps.

Follow only chunk x2 (the third slot: 3, 30, 300, 3000) so the hops stay readable. The partial sum starts on rank 3 and walks clockwise until rank 2, which is the owner of index 2:

Neighbor hop	Rank that adds	Running sum for `x2`
start	3	`3000`
1	0	`3000 + 3 = 3003`
2	1	`3003 + 30 = 3033`
3	2	`3033 + 300 = 3333`

After those three reduce-scatter hops, rank 2 owns 3333. Three all-gather hops then copy that finished shard to ranks 3, 0, and 1. Other chunks move on the same ring at the same time, offset by one slot, which is why a large buffer can keep every link busy.

Animated ring all-reduce following tensor chunk x2. The partial sum travels R3 to R0 to R1 to R2 during reduce-scatter and finishes as 3333 on rank 2. All-gather then circulates that completed shard from rank 2 around the ring. — The token is chunk `x2` only. Reduce-scatter ends on rank 2 with `3333`, matching NCCL's rank-index ownership rule. All-gather then starts from that owner. The other three chunks are in flight on the same ring, just not drawn.

The next snippet is a CPU, standard-library model of that schedule, not a GPU kernel. It represents each chunk by one integer, so it accepts exactly $k$ values from each of $k$ ranks. It checks ownership and reconstruction without modeling CUDA streams, network progress, or floating-point error.

ring_allreduce.py

inputs = [
    [1, 2, 3, 4],
    [10, 20, 30, 40],
    [100, 200, 300, 400],
    [1000, 2000, 3000, 4000],
]
k = len(inputs)
expected = [sum(row[i] for row in inputs) for i in range(k)]

def ring_reduce_scatter(bufs):
    k = len(bufs)
    if not k or any(len(row) != k for row in bufs):
        raise ValueError("expected k nonempty ranks with k scalar chunks each")
    state = [list(row) for row in bufs]
    for step in range(k - 1):
        send = [state[r][(r - step - 1) % k] for r in range(k)]
        for r in range(k):
            recv_idx = (r - step - 2) % k
            state[r][recv_idx] += send[(r - 1) % k]
    return [state[r][r] for r in range(k)]

def ring_all_gather(owned):
    k = len(owned)
    if not k:
        raise ValueError("expected at least one owned shard")
    result = [[None] * k for _ in range(k)]
    for r, value in enumerate(owned):
        result[r][r] = value
    for step in range(k - 1):
        snapshot = [row[:] for row in result]
        for r in range(k):
            src = (r - 1) % k
            chunk = (src - step) % k
            result[r][chunk] = snapshot[src][chunk]
    return result

owned = ring_reduce_scatter(inputs)
gathered = ring_all_gather(owned)

rank = 3
total = inputs[rank][2]
print("x2 path")
print(f"  rank {rank} holds {total}")
while rank != 2:
    rank = (rank + 1) % k
    added = inputs[rank][2]
    total += added
    owned_note = " (owns chunk 2)" if rank == 2 else ""
    print(f"  rank {rank} adds {added} -> {total}{owned_note}")

print(f"owned shards: {owned}")
print(f"rank 0 after all-gather: {gathered[0]}")
print(f"send volume factor 2*(k-1)/k = {2 * (k - 1) / k:.2f}")

assert owned == expected == [1111, 2222, 3333, 4444]
assert all(row == expected for row in gathered)

# The functions derive rank count from their arguments, not the example above.
for ranks in (1, 2, 3, 5):
    rows = [[10 * r + i for i in range(ranks)] for r in range(ranks)]
    reduced = [sum(row[i] for row in rows) for i in range(ranks)]
    assert ring_reduce_scatter(rows) == reduced
    assert ring_all_gather(reduced) == [reduced] * ranks
for invalid in ([], [[1, 2]], [[1], [2, 3]]):
    try:
        ring_reduce_scatter(invalid)
    except ValueError:
        pass
    else:
        raise AssertionError("invalid chunk layout accepted")
print("rank-count and shape checks passed")

Ring all-reduce toy

x2 path
  rank 3 holds 3000
  rank 0 adds 3 -> 3003
  rank 1 adds 30 -> 3033
  rank 2 adds 300 -> 3333 (owns chunk 2)
owned shards: [1111, 2222, 3333, 4444]
rank 0 after all-gather: [1111, 2222, 3333, 4444]
send volume factor 2*(k-1)/k = 1.50
rank-count and shape checks passed

For message size $M$ , each rank in a ring all-reduce sends about

V_{send} = 2\frac{k-1}{k}M.

Each rank receives the same volume. As $k$ grows, that factor approaches $2$ , so the traffic looks like $2M$ per rank, not $kM$ . The optimal-bandwidth result is for a point-to-point communication model, as analyzed by Patarasuk and Yuan. It isn't a universal byte lower bound for hardware with in-network reduction or multicast offload.^{[4]Reference 4Bandwidth Optimal All-Reduce Algorithms for Clusters of Workstationshttps://doi.org/10.1016/j.jpdc.2008.09.002}

The same factor reappears in the nccl-tests busbw column. For this four-rank toy, the factor is $1.5$ ; the measurement section turns that traffic factor into a concrete receipt and separates it from application time.^{[5]Reference 5NCCL Tests Performance Reportinghttps://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md}

Rings pay in latency. The algorithm needs $2(k-1)$ neighbor steps, so a tiny message on thousands of ranks spends too much time starting rounds. Pipelining many chunks hides link delay for large buffers, but it can't erase startup cost for small ones.

Why trees help small and medium messages

A binary tree reduces values toward a root and broadcasts the answer back down. Tree depth grows roughly with $\log_2 k$ , which cuts round count compared with a ring. One plain tree can overload internal ranks, so NCCL added complementary double binary trees that split data and swap leaf versus internal-node roles.^{[6]Reference 6Massively Scale Your Deep Learning Training with NCCL 2.4https://developer.nvidia.com/blog/massively-scale-deep-learning-training-nccl-2-4/}

No algorithm wins for every tensor and topology. Rings favor sustained bandwidth. Trees favor fewer rounds. The NCCL 2.31 documentation names Ring, Tree, CollnetChain, CollnetDirect, NVLS, NVLSTree, and PAT as selectable algorithm families, plus Simple, LL, and LL128 protocols.^{[7]Reference 7NCCL Environment Variableshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html}

Availability depends on the operation, hardware, network plugins, registration, and runtime tuning. Treat the names as candidates that explain a log, not as a menu that every host can run.

NCCL normally chooses algorithms and protocols from its topology and performance model. Forcing NCCL_ALGO or NCCL_PROTO can be useful for an experiment, but a fixed override can become wrong after hardware, message shapes, or NCCL versions change. Forcing LL128 on a path that doesn't support it can corrupt data, which is why the docs discourage protocol overrides except for diagnosis.^{[7]Reference 7NCCL Environment Variableshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html}

Workload shape	Likely pressure	What to inspect
Large gradient buckets	Link bandwidth	Ring or fabric-offload path, channel count, bus bandwidth
Tiny tensor-parallel reductions	Startup latency	Tree-like path, protocol, launch gaps
MoE all-to-all	Imbalance and network injection	Per-peer bytes, hot experts, NIC rails
Pipeline send/receive	Bubble and ordering	Stage timeline, peer pairing, stream waits
FSDP all-gather	Burst overlap	Bucket timing, parameter prefetch, compute coverage

A communicator is an ordered program

NCCL doesn't launch processes. A launcher such as MPI, Slurm, or torchrun starts them; application or framework bootstrap code exchanges a unique communicator ID and agrees on rank-to-device mapping. In a single process, ncclCommInitAll can handle this setup for a device list. Each rank within a communicator must use a distinct CUDA device.^{[8]Reference 8NCCL: Creating a Communicatorhttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/communicators.html}

Each process then binds a rank to a CUDA device and creates its communicator. Only after that agreement can the ranks enqueue matching work.

After communicator creation, each rank can enqueue this C++ collective with its own buffers and CUDA stream, an ordered GPU work queue. This is an integration fragment, not a standalone executable: the caller supplies valid device buffers, a communicator, a stream, and an error-checking NCCLCHECK macro. It requires a CUDA/NCCL system and wasn't executed by the CPU examples.

all_reduce_step.cu

NCCLCHECK(ncclAllReduce(
    local_gradient,
    global_gradient,
    element_count,
    ncclFloat32,
    ncclSum,
    communicator,
    communication_stream));

CUDA stream timeline: a produced bucket becomes ready, reduce-scatter overlaps independent later compute, and a shard reader waits for communication completion. — NCCL may start after a bucket-ready event while later backward kernels use independent buffers. The shard reader waits for the reduced-shard event. Updating parameters can require additional dependencies on backward computation that still reads those parameters.

For a successful call on a default blocking communicator outside a group, host return means work was enqueued, not completed. The producer must make local_gradient ready before communication reads it. A later CUDA event, stream dependency, or synchronization establishes when another kernel may consume global_gradient. Keep both buffers alive and avoid conflicting writes until the relevant work completes.^{[9]Reference 9NCCL CUDA Stream Semanticshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/streams.html}

Grouped calls can defer enqueue until ncclGroupEnd. A nonblocking communicator can return ncclInProgress; follow the documented completion/error polling before treating its stream work as enqueued. These are different boundaries from GPU completion.^{[10]Reference 10NCCL Group Callshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/groups.html}^{[3]Reference 3NCCL Typeshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/types.html}

When one thread manages several GPUs, group their matching calls with ncclGroupStart and ncclGroupEnd. A plain loop can block on its first device before the thread reaches the other ranks. Grouping doesn't relax collective order, including the documented order across communicators within a group.^{[10]Reference 10NCCL Group Callshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/groups.html}

Ordering stays global even when work is asynchronous. If one rank enters collective A then B while another enters B then A, stream scheduling can't repair the mismatch. Framework process groups and bucket schedulers exist partly to keep this distributed instruction stream identical across ranks.

💡 Key insight: An NCCL hang often begins before NCCL. One rank took a different branch, hit an out-of-memory error, loaded data slowly, or skipped a bucket. Peers then wait inside the next collective they all expected to share.

Topology turns one API into different routes

Two calls to ncclAllReduce can exercise very different hardware. Inside one server, bytes may travel through direct GPU peer access, NVLink, or NVSwitch. Across servers, the path can include PCIe, a network interface card (NIC), InfiniBand or RoCE, switches, and a remote host's reverse path.

Two separate paths: local GPU to NVSwitch to peer GPU, and cross-node GPU through PCIe, NIC A, network fabric, NIC B and PCIe to a remote GPU. NVSwitch is not a PCIe bridge. — Compare the short local route with the cross-node route. The same collective can be limited by NVLink locally, then by PCIe placement, NIC choice, rail balance, or network fabric once it leaves the host.

Read the matrix before changing a setting

On an NVIDIA host with the driver tools installed, start with host evidence before setting NCCL overrides. Run the topology query:

topology.sh

nvidia-smi topo -m

nvidia-smi topo -m prints connections among GPUs and NICs, along with CPU and memory affinities. Its labels classify path relationships; they aren't achieved-bandwidth measurements.^{[11]Reference 11nvidia-smi documentationhttps://docs.nvidia.com/deploy/nvidia-smi/index.html}

The exact matrix depends on the host. This excerpt is an illustrative shape, not a measurement:

nvidia-smi topo -m (illustrative)

        GPU0 GPU1 GPU2 GPU3 NIC0 NIC1
GPU0      X  NV4  SYS  SYS  PIX  SYS
GPU1     NV4   X  SYS  SYS  PIX  SYS
GPU2     SYS SYS   X  NV4  SYS  PIX
GPU3     SYS SYS  NV4   X   SYS  PIX
NIC0     PIX PIX  SYS SYS   X   SYS
NIC1     SYS SYS  PIX  PIX  SYS   X

Read the common cells as path hints:

Cell	Path class	First hypothesis to test
`NV4`	A bonded set of four NVLinks	GPU traffic stays on the NVLink fabric
`PIX`	At most one PCIe switch	GPU and NIC share a short PCIe route
`PXB`	Multiple PCIe switches	Extra PCIe hops may add contention
`PHB`	A PCIe host bridge	The CPU or root complex sits on the route
`NODE`	PCIe host bridges within one NUMA node	Placement may cross a host-bridge boundary
`SYS`	PCIe plus an interconnect between NUMA nodes	Remote NUMA or fabric traffic may dominate

NV4 means four bonded links, not four times a measured rate. The matrix also doesn't draw switch nodes.

Separate the local GPU fabric from NIC access. A GPU may reach another GPU through NVSwitch, while reaching a NIC through PCIe. NVSwitch isn't a PCIe-to-network bridge. NCCL's PXN path can relay through an intermediate GPU with better NIC access; verify that distinct route in platform topology and logs. A host bridge on the route doesn't by itself mean the payload was copied through CPU memory.^{[7]Reference 7NCCL Environment Variableshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html}

NCCL turns that topology hypothesis into a schedule. It discovers GPU, CPU, PCIe, NVLink, and network relationships, searches candidate graph layouts, and tunes a plan. GPUDirect remote direct memory access (RDMA) can let a NIC access GPU memory without staging the payload through ordinary CPU copies. Shared-memory and host-staged paths still matter when direct peer access isn't available.

Topology awareness can't create bandwidth that hardware lacks. A GPU behind the wrong PCIe root, a disabled link, a slow network interface, or oversubscribed switches still cap the collective. The library can select a better route only among paths it can discover and use.

Read one call through the source tree

The pinned source snapshot makes the runtime boundary concrete.^{[12]Reference 12NCCL Source Repositoryhttps://github.com/NVIDIA/nccl/tree/5067397c2676d5aed50042fc39e5c8ee96eb0027} Start with src/collectives.cc, where public functions such as ncclAllReduce package operation metadata.

The snapshot is commit 5067397c2676d5aed50042fc39e5c8ee96eb0027. The diagram is a responsibility map, not a literal stack trace: transport connections can be prepared before a device kernel runs, and not every path needs host proxy progress.

Diagram showing collectives.cc API + arguments, enqueue.cc work planning, graph/tuning.cc algo + protocol, and device/all_reduce.h CUDA choreography. — collectives.cc API + arguments, enqueue.cc work planning, graph/tuning.cc algo + protocol, and device/all_reduce.h CUDA choreography.

Use this source map instead of searching for one giant all-reduce loop:

Source area	Responsibility	Question to carry
`src/init.cc`, `bootstrap.cc`	Communicator setup and peer discovery	Which ranks and devices joined?
`src/graph/topo.cc`, `search.cc`	Physical topology and graph search	Which links can form rings or trees?
`src/graph/tuning.cc`	Algorithm, protocol, and channel model	Why did this message get this plan?
`src/enqueue.cc`, `collectives.cc`	Host API validation and work creation	What exactly was enqueued?
`src/device/*.h`, `primitives.h`	Device-side copy and reduction steps	Which rank sends, receives, or reduces next?
`src/transport/*`	P2P, shared memory, sockets, InfiniBand, plugins	Which physical route carries bytes?
`src/proxy.cc`	Host proxy progress for transports that need it	Is progress waiting on host or network work?
`src/debug.cc`, `src/ras/`	Logs and reliability signals	What evidence identifies the failed rank or link?
`src/nccl_device/`, `src/gin/`	Device-initiated collectives and GPU-side networking	Did a host enqueue start this, or did a device kernel?

The conventional host all-reduce path and the device-initiated interfaces aren't the same entry point. Reading only ncclAllReduce shows the contract, not the chosen schedule. Reading only a CUDA primitive shows local mechanics, not why the runtime picked that path.

How higher-level systems use NCCL

NCCL is shared infrastructure across training and inference. Its importance grows as model code partitions more tensors, because each partition boundary creates a communication event.

System pattern	NCCL traffic	Scaling failure when communication wins
Data-parallel training	Gradient all-reduce	More replicas yield diminishing tokens/sec
Fully sharded FSDP or ZeRO-3	Reduce-scatter plus parameter all-gather	Memory fits, but layer gathers stall compute
Tensor parallelism	Frequent activation reductions or gathers	Wide TP across slow links hurts every layer
Pipeline parallelism	Send/receive activations and gradients	Stage bubbles or peer waits grow
Expert parallelism	All-to-all token dispatch and combine	Hot experts create uneven peer traffic
Multi-GPU inference	Tensor-parallel reductions, KV or state exchange	Inter-token latency rises despite idle FLOPs

DeepSpeed and Megatron build groups, buckets, overlap schedules, and sharding policies above NCCL. PyTorch exposes NCCL through torch.distributed. Specialized expert-routing libraries can use other communication implementations; don't infer that every GPU collective is an NCCL call. None of those layers remove the need to understand ownership and physical paths.

Choose the optimization target from operation shape. If tensor parallelism emits a small all-reduce in every transformer layer, lowering its launch latency may matter more than peak bandwidth. If data parallelism emits a few huge buckets, saturating network rails and overlapping buckets with backward compute matter more.

Measure bytes, rounds, and exposed time

Suppose six rounds each add 2 microseconds of startup, while 1 MB traverses a 50 GB/s path. The rough estimate is 12 microseconds of startup plus 20 microseconds of payload time, or 32 microseconds total. These decimal units and invented inputs illustrate a model, not an NCCL prediction.

Writing that calculation generally separates startup latency from payload time:

T_{collective} \approx R\alpha + \frac{V}{B},

where $R$ is communication rounds, $\alpha$ is per-round startup cost, $V$ is bytes crossing the measured path, and $B$ is achieved bandwidth. Real NCCL schedules add channels, protocol overhead, topology contention, and synchronization, but this equation tells you whether reducing rounds or moving bytes faster is the bigger target.

Small payloads make $R\alpha$ visible. Large payloads make $V/B$ visible. Keep that split in mind when you vary message size, rank count, or topology.

Turn one `all_reduce_perf` row into evidence

nccl-tests is a separate NVIDIA repository that checks both NCCL correctness and performance. First build it against your installed CUDA and NCCL (make, with CUDA_HOME and NCCL_HOME when needed); multi-process tests also require an MPI-enabled build. On one host with four visible GPUs, pin the scan to one payload size before comparing routes. This hardware-dependent command wasn't run for the lesson's CPU checks:

all_reduce_perf

./build/all_reduce_perf -b 256M -e 256M -g 4 -n 20 -w 5 -c 1 -T 60

Here, -g 4 uses four GPUs per thread (one thread and process in this launch), -b and -e select 256 MiB, -n 20 requests 20 timed iterations, -w 5 adds five untimed warmups, and -c 1 requests one correctness-check iteration. -T 60 sets the test timeout in seconds. Check flags against your checkout: this description follows the public README consulted in September 2026.^{[13]Reference 13NCCL Testshttps://github.com/NVIDIA/nccl-tests}

The following receipt has the shape of current all_reduce_perf output. Its 6.00 ms value is chosen for arithmetic, not measured from a named GPU, driver, NCCL build, or network:

all_reduce_perf (illustrative row)

# 4 ranks · float32 sum · 5 warmups · 20 timed iterations · check=1
#                                      out-of-place                    in-place
# size (B) count (elements) type redop root  time  algbw  busbw #wrong  time  algbw  busbw #wrong
#                                               (us)  (GB/s) (GB/s)          (us)  (GB/s) (GB/s)
 268435456   67108864      float sum   -1   6000.0  44.74  67.11    0   6000.0  44.74  67.11    0

Start with units. size is 268,435,456 bytes, and count is 67,108,864 float32 elements. Since 6000 microseconds is 0.006 seconds, algorithmic bandwidth is $268{,}435{,}456 / 0.006 = 44.74$ GB/s. It answers, "How quickly did this logical payload complete?"^{[5]Reference 5NCCL Tests Performance Reportinghttps://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md}

For a four-rank all-reduce, the test's operation-based correction is $2(k-1)/k = 2(4-1)/4 = 1.5$ . Thus busbw is $44.74 \times 1.5 = 67.11$ GB/s.^{[5]Reference 5NCCL Tests Performance Reportinghttps://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md}

Read that as a normalized link-oriented score for comparing the same operation across systems. It isn't a packet counter or a promise that every physical mechanism moved exactly 1.5 payloads. The performance guide derives this factor for point-to-point collective traffic and lists different factors for reduce-scatter, all-gather, broadcast, reduce, and all-to-all.^{[5]Reference 5NCCL Tests Performance Reportinghttps://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md}

The two halves of the row also answer a buffer question. Out-of-place uses separate send and receive buffers. In-place aliases them for all-reduce, which NCCL defines as sendBuff == recvBuff.^{[14]Reference 14NCCL In-place Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/inplace.html}

#wrong reports elements that failed the requested correctness check. A zero count is necessary evidence for this row, but it doesn't cover a later stream or buffer-lifetime race in the application.^{[13]Reference 13NCCL Testshttps://github.com/NVIDIA/nccl-tests}

For four ranks, 268435456 bytes divided by 0.006 seconds gives 44.74 GB/s algorithm bandwidth. Multiplying by 2(4-1)/4, or 1.5, gives 67.11 GB/s bus bandwidth. This is normalization, not measured link traffic. — The arithmetic connects logical bytes, elapsed time, and a collective-specific normalization. Neither bandwidth number identifies the physical route. Values are illustrative; application overlap needs a separate timeline.

In checkouts supporting -I 1, it adds per-iteration CUDA-event timing summaries. That shows timing spread, not proof of which rank caused a delay. Inspect per-rank traces for attribution. With only 20 samples, the reported nearest-rank p99 can equal the maximum; don't treat it as a stable tail estimate. Keep timing separate from correctness checks.^{[13]Reference 13NCCL Testshttps://github.com/NVIDIA/nccl-tests}

Attach five pieces of evidence to a real performance claim:

Evidence	Record beside the row
Hardware and software	GPU model and count, host count, topology, CUDA, NCCL, framework, and `nccl-tests` versions
Workload	Payload bytes, element count, rank count, and any aggregation or batch shape
Precision and path	Datatype, reduction op, in-place mode, and any forced algorithm or protocol
Baseline	Exact comparison command and equivalent tuning on the alternate route or build
Correctness	Check mode, tolerance or status, and `#wrong` for each buffer mode

The receipt isolates one collective. In an idealized training schedule, hiding 4 ms of a 6 ms collective leaves 2 ms exposed. Real overlap can slow compute or communication through resource contention, so establish the actual critical path from a trace.

Both schedules can produce the same receipt, so compare it with a profiler timeline and the stream event that makes the reduced buffer readable. algbw or busbw alone isn't a tokens-per-second or step-time result.

Run tests on the same hosts, GPUs, container settings, and network rails as the workload. A clean row establishes an isolated fabric baseline; it doesn't prove that model code, data loading, bucket timing, or stream dependencies expose the same time.

🎯 Production tip: Keep a small performance envelope by message size and rank count. One headline number hides the exact region where tensor-parallel latency or giant gradient buckets fail.

A clean receipt answers whether the fabric can move isolated bytes. It doesn't tell you why a live step exposed time or stopped, so the next question is causal: did ranks disagree, did a process fail, or did the route break?

Diagnose a hang without guessing

First classify whether ranks disagree, a process failed, or the network path is broken. A timeout stack that ends in NCCL doesn't distinguish those causes.

Symptom	Likely class	First evidence
One rank exited or OOMed	Missing participant	Earliest per-rank error, not peer timeouts
Same step hangs every run	Collective order or shape mismatch	Per-rank operation sequence and tensor metadata
Multi-node only	Interface, firewall, NIC, or RDMA path	Bootstrap and `NET` logs, host connectivity
Slow after scale-out	Topology or straggler	Timeline by rank, link counters, `nccl-tests`
Container init failure	Shared memory or memlock	`/dev/shm`, limits, NCCL warning
Wrong data without hang	Buffer lifetime or stream dependency	CUDA events, producer and consumer streams

Set NCCL_DEBUG=WARN for explicit errors. Escalate briefly to NCCL_DEBUG=INFO with focused subsystems such as INIT,NET,GRAPH, and write one log per host and process when needed. NCCL documentation also covers shared-memory limits, network-interface selection, container configuration, and asynchronous error handling.^{[15]Reference 15NCCL Troubleshootinghttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting.html}

Then compare all ranks at the first divergence. Check world size, rank-to-device mapping, collective name, count, dtype, root, communicator, and sequence number. If those match, test the path outside the application with nccl-tests and inspect the topology NCCL discovered.

Finally, restore automatic tuning after the experiment. Variables such as NCCL_SOCKET_IFNAME and NCCL_IB_HCA can be stable cluster configuration. Algorithm, protocol, and channel overrides are usually diagnostic controls, not permanent cargo cult.

Strengths and limits

NCCL is a small collective API over many NVIDIA GPU and network topologies. It fuses communication and reduction on device paths, picks algorithms at runtime, works from one process or many, and exposes plugins for network, tuner, profiler, and newer device-side interfaces.

Its limits are equally important:

NVIDIA GPU focus makes NCCL the wrong backend for CPU-only or non-NVIDIA collectives.
Collective correctness is distributed. NCCL can't repair mismatched control flow above it.
Topology selection can't fix bad cabling, oversubscription, NUMA placement, or unhealthy links.
Communication kernels consume GPU resources and can interfere with compute when overlap is poorly scheduled.
Environment overrides can hide one machine issue while hurting another.
Fast transport doesn't choose a good model-parallel strategy. A framework still owns groups, tensor layouts, and bucket timing.

Call NCCL when the job is on NVIDIA GPUs and bytes have to move inside a node or across nodes. Stay on the framework's process-group API unless you're building a runtime, fused kernel, or communication library that needs NCCL directly.

NVIDIA project, public source

NVIDIA develops NCCL and publishes the source in the NVIDIA/nccl repository. The repo accepts signed-off contributions under a Developer Certificate of Origin, asks large changes to begin with an issue and design discussion, and keeps tests in the separate NVIDIA/nccl-tests project.^{[12]Reference 12NCCL Source Repositoryhttps://github.com/NVIDIA/nccl/tree/5067397c2676d5aed50042fc39e5c8ee96eb0027}

The public contributor history records individual work; it isn't a permanent ownership roster or support contract.^{[16]Reference 16NVIDIA NCCL Contributorshttps://github.com/NVIDIA/nccl/graphs/contributors}

The pinned snapshot's license file needs a careful reading. Most project source uses Apache License 2.0, some parts retain original BSD terms, and borrowed files can carry their own licenses. "NCCL is BSD" or "every NCCL file is Apache-2.0" both erase that file-level boundary.^{[17]Reference 17NCCL Source Licensehttps://github.com/NVIDIA/nccl/blob/5067397c2676d5aed50042fc39e5c8ee96eb0027/LICENSE.txt}

Research roots rather than one NCCL paper

NCCL doesn't have one canonical paper that defines the whole current implementation. Its core ideas come from collective-communication research plus years of topology, kernel, and transport engineering.

Patarasuk and Yuan derived a bandwidth lower bound for all-reduce and described a ring algorithm that reaches it for large messages.^{[4]Reference 4Bandwidth Optimal All-Reduce Algorithms for Clusters of Workstationshttps://doi.org/10.1016/j.jpdc.2008.09.002} NVIDIA's NCCL 2.4 engineering account explains why double binary trees reduce latency growth while retaining high aggregate bandwidth, and why NCCL switches between tree and ring paths.^{[6]Reference 6Massively Scale Your Deep Learning Training with NCCL 2.4https://developer.nvidia.com/blog/massively-scale-deep-learning-training-nccl-2-4/}

Later hardware paths add NVSwitch collectives, in-network reduction, registration, multiple rails, and device-initiated communication. Those features change how an algorithm maps to hardware, but not the basic questions: who owns each chunk, how many rounds run, how many bytes cross each bottleneck, and when the consumer stream may read the result.

Source-reading exercise

Pick one all-reduce from a PyTorch profiler trace. Record tensor bytes, communicator size, duration, and whether compute overlaps it. Capture nvidia-smi topo -m beside the trace, then run nccl-tests near that message size on the same ranks.

Use the pinned source map to answer four questions:

Which public function packages the operation?
Which algorithm and protocol names appear in logs?
Which graph and transport files correspond to the physical path?
Which rank first diverges in the CPU operation-log example below?

Produce a small evidence table rather than a screenshot of a green benchmark. Connect logical operation, source path, topology labels, payload bytes, algbw, busbw, overlap time, correctness status, and first failure evidence. If the standalone row looks strong but the application still stalls, the table should make that boundary visible.

Without GPUs, compare planned collective logs instead. Each tuple below is (operation, count, dtype, reduction, root) for one communicator; None means an argument doesn't apply. This teaching checker compares completed logs. It neither runs NCCL nor detects a live network failure. A missing call is exercised locally, without leaving distributed processes hanging.

collective_log.py

from itertools import zip_longest

def first_divergence(logs):
    if not logs:
        raise ValueError("at least one rank log is required")
    missing = object()
    for step, calls in enumerate(zip_longest(*logs, fillvalue=missing)):
        if any(call != calls[0] for call in calls[1:]):
            return step
    return None

reduce = ("all_reduce", 4, "float32", "sum", None)
broadcast = ("broadcast", 4, "float32", None, 0)
baseline = [reduce, broadcast]
assert first_divergence([baseline] * 4) is None
cases = {
    "missing call": [reduce],
    "reordered calls": [broadcast, reduce],
    "different count": [("all_reduce", 3, "float32", "sum", None), broadcast],
    "different root": [reduce, ("broadcast", 4, "float32", None, 1)],
}
for name, changed in cases.items():
    logs = [baseline, baseline, changed, baseline]
    step = first_divergence(logs)
    assert step == (1 if name in ("missing call", "different root") else 0)
    print(f"{name}: first disagreement at step {step}")

payload_bytes, seconds, ranks = 256 * 2**20, 0.006, 4
algbw = payload_bytes / seconds / 1e9
busbw = algbw * 2 * (ranks - 1) / ranks
assert round(algbw, 2) == 44.74 and round(busbw, 2) == 67.11
print(f"illustrative receipt: algbw={algbw:.2f}, busbw={busbw:.2f} GB/s")

Find the first disagreement

missing call: first disagreement at step 1
reordered calls: first disagreement at step 0
different count: first disagreement at step 0
different root: first disagreement at step 1
illustrative receipt: algbw=44.74, busbw=67.11 GB/s

Agreement doesn't establish correctness: all ranks could agree on the wrong operation, and a correct log says nothing about buffer lifetime. Compare these contracts with the intended model computation and actual stream dependencies.

Takeaways

A collective is an ordered agreement across ranks, not a remote function call.
All-reduce decomposes into reduce-scatter plus all-gather, which explains sharded-training traffic.
Rings achieve efficient point-to-point byte counts; trees cut rounds. Hardware offload and actual topology can change the best choice.
nvidia-smi topo -m gives path classes and affinities; it doesn't give achieved bandwidth.
For point-to-point all-reduce accounting, nccl-tests busbw is algbw times $2(k-1)/k$ , the same factor as ring send volume.
algbw describes an isolated logical payload. It isn't a substitute for an application's exposed step time after overlap.
CUDA streams make collectives asynchronous, but explicit dependencies still protect buffer correctness.
Debug the earliest rank divergence, then validate the fabric separately.

Mastery quiz

Next Step

Continue to Deep Dive - Megatron

NCCL gives you the communication substrate. Megatron shows how tensor, pipeline, data, context, and expert parallel groups turn those collectives into one large-model training program.

PreviousDeep Dive - DeepGEMM

Share this article