Deep Dive - NCCL

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. Use these assumed timings to work through the communication cost.

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}

Final frame5.2s loop

View larger

The calls below follow one pinned source map. As of October 10, 2026, the newest public NCCL library release is 2.32.3, published September 17, and the live guide documents that version. Initial Rubin support still lacks a tuned performance model; benchmark the target topology before estimating speed.^{[2]Reference 2NCCL v2.32.3-1 releasehttps://github.com/NVIDIA/nccl/releases/tag/v2.32.3-1}

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.^{[3]Reference 3NCCL 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:

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.^{[4]Reference 4NCCL 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.

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.

Final frame5.2s loop

View larger

Final frame5.2s loop

View larger

Final frame5.2s loop

View larger

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).^{[3]Reference 3NCCL 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 describes an operation-level equivalence; NCCL can produce the same result with other algorithms. Floating-point summation order can also change the last bits.^{[3]Reference 3NCCL 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:

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.^{[3]Reference 3NCCL 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:

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.

Final frame8.5s loop

View larger

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.

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")

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

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.^{[5]Reference 5Bandwidth 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.^{[6]Reference 6NCCL 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.^{[7]Reference 7Massively 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.32.3 documentation names Ring, Tree, CollnetChain, CollnetDirect, NVLS, NVLSTree, and PAT as selectable algorithm families, plus Simple, LL, and LL128 protocols.^{[8]Reference 8NCCL 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.^{[8]Reference 8NCCL Environment Variableshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html}

Those rows still assume every rank entered the same operation. What makes that agreement real when each process only sees its own call?

A communicator is an ordered program

Nothing in the algorithm name enforces the agreement those rows assume. NCCL doesn't launch processes. A launcher such as MPI, Slurm, or torchrun starts them. Application or framework bootstrap code then 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.^{[9]Reference 9NCCL: 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. A runnable program also needs valid device buffers, a communicator, a stream, an error-checking NCCLCHECK macro, and a CUDA/NCCL system. The CPU examples do not execute this integration fragment.

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

Final frame6.7s loop

View larger

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.^{[10]Reference 10NCCL 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.^{[11]Reference 11NCCL Group Callshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/groups.html[4]Reference 4NCCL Typeshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/types.html}

Grouping can also change the overlap you thought you had. When a group mixes streams, NCCL joins their dependencies before its kernel and blocks those streams until the kernel completes. Two stream names inside one group don't provide independent progress. Keep unrelated work outside that synchronization point when its dependencies permit it.^{[10]Reference 10NCCL CUDA Stream Semanticshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/streams.html}

NCCL 2.32.3 adds ncclCollConfig_t::launchCompletionEvent. The event reports the launch boundary. A shard reader must wait for completion of the collective's writes on the communication stream before reading the reduced buffer. Its header requires a caller-owned, rank-local event created with cudaEventDisableTiming, consistent NULL/non-NULL participation across ranks, and at most one non-NULL event per communicator in a group. CUDA versions before 12.3 record it before launch. Keep the launch event alive through queued waits and captured-graph executions.^{[12]Reference 12NCCL 2.32.3 public header: launch-completion event contracthttps://github.com/NVIDIA/nccl/blob/v2.32.3-1/src/nccl.h.in#L215}

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.^{[11]Reference 11NCCL 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.

View larger

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:

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. Measure achieved bandwidth with a workload that matches the target traffic.^{[13]Reference 13nvidia-smi documentationhttps://docs.nvidia.com/deploy/nvidia-smi/index.html}

The exact matrix depends on the host. This illustrative excerpt shows how to read a host-dependent matrix:

        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:

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.^{[8]Reference 8NCCL 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.^{[14]Reference 14NCCL 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.

View larger

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

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.

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:

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:

./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 checked October 10, 2026.^{[15]Reference 15NCCL Testshttps://github.com/NVIDIA/nccl-tests}

This receipt has the shape of current all_reduce_perf output. The supplied 6.00 ms value lets us check the arithmetic; replace it with a measurement from your GPU, driver, NCCL build, and network:

# 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?"^{[6]Reference 6NCCL 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.^{[6]Reference 6NCCL 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 estimates traffic under the guide's normalization; inspect hardware counters to measure physical transfers. 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.^{[6]Reference 6NCCL 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.^{[16]Reference 16NCCL In-place Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/inplace.html}

That equality doesn't describe in-place reduce-scatter or all-gather. Their local shard sits at the rank's offset inside the full buffer. With a typed float* data, rank 2, and one element per shard, reduce-scatter receives at data + 2; all-gather sends from data + 2. Pointer arithmetic advances by elements, so this is an eight-byte offset for float32. Passing the base address as all-gather's local input would read rank 0's slot instead.^{[16]Reference 16NCCL In-place Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/inplace.html}

The CPU example uses an actual shared-memory view to show that layout. It models the completed reduce-scatter write, then compares the correct and incorrect all-gather input slices. Verify CUDA buffer readiness with the pinned NCCL implementation and an on-device test.

from array import array

rank, count = 2, 1
data = array("f", [100.0, 200.0, 300.0, 400.0])
assert data.itemsize == 4
full = memoryview(data)
start = rank * count
local = full[start:start + count]
local[0] = 3333.0  # Model rank 2's completed reduce-scatter output.

assert local.obj is data and data[2] == 3333.0
correct_input = list(local)
wrong_input = list(full[:count])
assert correct_input == [3333.0] and wrong_input == [100.0]
print(f"rank {rank} shard offset: {start * data.itemsize} bytes")
print(f"all-gather local input: {correct_input}")
print(f"using the base address reads: {wrong_input}")

rank 2 shard offset: 8 bytes
all-gather local input: [3333.0]
using the base address reads: [100.0]

#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.^{[15]Reference 15NCCL Testshttps://github.com/NVIDIA/nccl-tests}

View larger

In checkouts supporting -I 1, per-iteration CUDA-event summaries show timing across the measured iterations. Use per-rank traces to attribute a delay. With only 20 samples, nearest-rank p99 can equal the maximum, so collect more iterations for a stable tail estimate. Keep timing separate from correctness checks.^{[15]Reference 15NCCL Testshttps://github.com/NVIDIA/nccl-tests}

Attach five pieces of evidence to a real performance claim:

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. Use a clean row as an isolated fabric baseline; measure model code, data loading, bucket timing, and stream dependencies in the full workload.

🎯 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.

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.^{[17]Reference 17NCCL Troubleshootinghttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting.html}

In 2.32.3, ATTN messages expose important nonfatal fallbacks. Don't set NCCL_DEBUG=ATTN across a mixed-version deployment: versions before 2.32 don't recognize that level and fall back to no debug logging. Keep NCCL_DEBUG=WARN and add NCCL_DEBUG_LEVELS=ATTN; newer versions add the notices, while older versions retain their warning baseline.^{[18]Reference 18NCCL 2.32.3 Logginghttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting/logging.html}

RAS adds GPU progress counters and a watchdog DMA mirror to investigate stalled kernels, alongside expanded link and device diagnostics. Use this evidence to distinguish a launched kernel that stopped progressing from ranks that never agreed on the operation.^{[2]Reference 2NCCL v2.32.3-1 releasehttps://github.com/NVIDIA/nccl/releases/tag/v2.32.3-1}

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.

The 2.32.3 release lists communicator support for multiple MIG instances. The guide dates its experimental path to 2.31: distinct instances remain distinct CUDA devices, require a CUDA 13.0 or newer driver, and don't support NVLink P2P or NVLS. MNNVL must be disabled. Don't transfer a full-GPU topology's performance expectations to this setup.^{[2]Reference 2NCCL v2.32.3-1 releasehttps://github.com/NVIDIA/nccl/releases/tag/v2.32.3-1[9]Reference 9NCCL: Creating a Communicatorhttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/communicators.html}

Optional TLS in 2.32.3 covers NCCL-owned socket traffic when built with OpenSSL 3 and configured through ncclSetEncryption. Configure encryption separately for external transports and plugin traffic.^{[2]Reference 2NCCL v2.32.3-1 releasehttps://github.com/NVIDIA/nccl/releases/tag/v2.32.3-1}

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.^{[14]Reference 14NCCL 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.^{[19]Reference 19NVIDIA 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.^{[20]Reference 20NCCL 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 under their communication model.^{[5]Reference 5Bandwidth Optimal All-Reduce Algorithms for Clusters of Workstationshttps://doi.org/10.1016/j.jpdc.2008.09.002} The result assumes associative, commutative reduction and omits startup and reduction-compute costs. Analyze latency optimality and floating-point reproducibility separately. 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.^{[7]Reference 7Massively 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:

1. Which public function packages the operation?
2. Which algorithm and protocol names appear in logs?
3. Which graph and transport files correspond to the physical path?
4. 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.

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")

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

Verify correctness by comparing the collective with the intended model computation and checking actual stream dependencies and buffer lifetime; all ranks can agree on the same incorrect operation.

Takeaways

Follow one slot from four local vectors to the shared sum. Then separate a ring's send volume from the bus-bandwidth score, and a missing rank from a slow route.

• 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