Markdown

NCCL collectives and algorithm selection¶

Scope: how NCCL implements all-reduce, all-gather, reduce-scatter, and broadcast, how it selects an algorithm (Ring/Tree/CollNet/NVLS) and protocol (Simple/LL/LL128) by message size and topology, the key tuning env vars (NCCL_ALGO, NCCL_PROTO, NCCL_NTHREADS, NCCL_BUFFSIZE), communicator lifecycle cost and reuse, and how to validate with nccl-tests bus bandwidth.

What it is¶

NCCL (NVIDIA Collective Communications Library) is a many-to-many communication library providing optimized collectives (all-reduce, all-gather, broadcast, reduce-scatter) used by groups of GPUs to share data.¹ It underpins most multi-GPU training in the NVIDIA ecosystem: each GPU computes gradients on its data shard, then NCCL all-reduces those gradients so every GPU updates weights with the averaged result.¹

The four collectives in scope:

all-reduce: sum (or reduce) a tensor across all ranks; every rank ends with the full reduced result. The gradient-sync primitive for data-parallel training. Ring-based algorithms decompose it as a reduce-scatter followed by an all-gather; tree-based algorithms instead decompose it as a reduce up to a root followed by a broadcast down. All-gather and broadcast are distinct collectives (see below), not interchangeable names for the same phase.¹
all-gather: every rank contributes a shard; every rank ends with the concatenation of all shards. Used in FSDP to re-materialize sharded parameters.
reduce-scatter: reduce across ranks, but each rank keeps only its slice of the result. The reduction half of FSDP gradient handling.
broadcast: one root rank sends identical data to all ranks (e.g. updated weights). NCCL can use NVSwitch hardware multicast for one-hop broadcast inside an NVLink domain.¹

NCCL runs over PCIe, NVLink, NVSwitch, InfiniBand, and TCP sockets, and automatically chooses the fastest path between any two GPUs.¹ At communicator init it inspects message size, interconnect topology, and GPU generation to pick the algorithm+protocol combination per collective.¹

Why use it¶

Communication, not compute, is the scaling wall. The same all-reduce can run at tens of GB/s or hundreds of GB/s depending purely on whether NCCL routes over NVLink versus PCIe.¹ A topology-unaware single ring across four GPUs split over two PCIe switches forces a stage over the slow inter-switch link; a hierarchical approach keeps the bulk on NVLink. In the book's profiled example this is the difference between 60% SM utilization at 100 ms/iter and 90% SM utilization at 70 ms/iter.¹

Algorithm and protocol choice is message-size-dependent. Small messages are latency-dominated (startup cost dominates); large messages are bandwidth-dominated (byte-movement dominates).¹ Picking the wrong one (or letting a misconfiguration force a fallback) silently costs an order of magnitude without crashing.

When to use it (and when not)¶

Needed whenever multiple GPUs must agree on a tensor: DDP/FSDP gradient sync, tensor-parallel partial-sum reduction, weight broadcast at init. NCCL is the correct backend for any NVIDIA multi-GPU collective; it is PyTorch's default.¹

Tuning the selection (overriding NCCL_ALGO/NCCL_PROTO) is rarely needed. NCCL's automatic selection is good; manual override is for troubleshooting, research experiments, or a profiled, confirmed pathology such as unexpectedly high cross-node latency.¹ Set explicit values to pin behavior across NCCL upgrades; defaults can change between versions and are hard to debug when they do.¹

Not the right tool for point-to-point inference transfers (KV-cache movement): NCCL send()/recv() exist but are less optimized than NIXL for one-to-one tail latency.¹ See Disaggregated Inference. Never use the CPU-bound Gloo backend for GPU training; it staged through host memory over TCP and runs an order of magnitude slower.¹

Architecture¶

flowchart TB
    C["Collective and message size"] --> A{"Algorithm"}
    A -->|"bandwidth-bound"| R["Ring"]
    A -->|"latency-bound"| T["Tree or NVLSTree"]
    A -->|"hierarchical fabric"| N["CollNet or NVLS"]
    A -->|"AllGather or ReduceScatter"| P["PAT"]
    R --> Q{"Protocol"}
    T --> Q
    N --> Q
    P --> Q
    Q --> L["LL or LL128"]
    Q --> S["Simple"]

How to use it¶

Algorithm selection (topology- and size-driven)¶

NCCL's primary collective algorithms:¹

Ring: GPUs form a logical ring; the all-reduce decomposes into a reduce-scatter phase and an all-gather phase, each n-1 steps moving a data_size/n slice per link. Total per-link traffic across both phases is 2*(n-1) * data_size/n bytes, matching the AllReduce busbw correction factor documented below. Perfectly balances load and is bandwidth-optimal, but latency scales with hop count O(N). Best for large messages (bandwidth-dominated).¹
Tree / NVLSTree: reduces up to a root over a spanning tree, then broadcasts the result back down, completing all-reduce in O(log N) dependency steps. This is a hierarchical reduce-then-broadcast decomposition, not the ring's fixed data_size / n reduce-scatter and all-gather schedule. NCCL still chunks and pipelines buffers across channels, so do not infer one full-buffer transfer per tree edge. Trees favor small messages; they may not saturate all links on large ones. NVLSTree enables NVLink SHARP (NVLS) offload.¹⁴
CollNet / CollTree: two-level hierarchical collectives: a high-throughput local algorithm inside each fast domain (node / NVSwitch island), then one leader per group joins a second-level tree across groups over RDMA, pipelined. Low internode latency plus full intranode bandwidth; can offload to InfiniBand SHARP when the NCCL-SHARP plugin is enabled.¹
PAT (Parallel Aggregated Trees): a Bruck-derived algorithm introduced in NCCL 2.23 for AllGather and ReduceScatter, not a standalone all-reduce algorithm. It gives a logarithmic number of network steps for small sizes, progressively increasing transfers as size grows to keep buffering minimal, and works for any rank count (no power-of-two requirement, unlike recursive doubling). Early versions require one GPU per node.⁹

Rule of thumb: small messages (tens of MB) favor trees (fewer steps); large messages favor ring (bandwidth).¹ Keep as much traffic as possible on the fastest interconnect (NVLink/NVSwitch intranode); minimize PCIe and inter-NUMA hops.¹

The valid NCCL_ALGO values per the official NCCL docs are Ring, Tree, CollnetChain, CollnetDirect, NVLS, NVLSTree, PAT; a ^ prefix excludes rather than includes. Unset (the default) lets NCCL choose from node topology and architecture.² (The book uses illustrative spellings such as NVLSTree,PAT; prefer the official token list on disagreement.¹)

# Override only for troubleshooting / A-B testing. Set BEFORE ncclCommInitRank.
export NCCL_ALGO=Tree          # force tree (latency-dominated small messages)
export NCCL_ALGO=^Ring         # exclude ring, let NCCL pick among the rest

Protocol selection (Simple / LL / LL128)¶

Independent of algorithm, NCCL picks a wire protocol that trades latency against bandwidth. Valid NCCL_PROTO values per the official docs are LL, LL128, Simple (^ excludes):²

LL ("low latency"): lowest latency, lowest peak bandwidth; for the smallest messages.
LL128: low latency tuned for 128-byte granularity; high bandwidth on NVLink-class links. Only available on platforms that support it.
Simple: highest peak bandwidth, higher fixed latency; for large messages.

Default is unset, which enables all supported protocols: LL,LL128,Simple where LL128 is supported, LL,Simple otherwise.²

export NCCL_PROTO=Simple        # force the bandwidth-optimal protocol
export NCCL_PROTO=^LL128         # exclude LL128 (e.g. suspected LL128 corruption)

Key tuning env vars¶

Variable	Purpose	Default (official)
`NCCL_ALGO`	Allowed collective algorithm(s)	unset → auto by topology²
`NCCL_PROTO`	Allowed protocol(s)	unset → all supported²
`NCCL_NTHREADS`	CUDA threads per block; one block per channel	512 on recent GPUs, 256 on some older ones²
`NCCL_BUFFSIZE`	Per-GPU-pair communication buffer, bytes	4194304 (4 MiB)²

NCCL_NTHREADS valid values are 64, 128, 256, 512.² NCCL_BUFFSIZE takes integer bytes (powers of 2 recommended).² The book cautions that raising NCCL_BUFFSIZE can improve large-all-reduce bandwidth but must be sized carefully: too high causes GPU memory pressure; start at 4 MiB and increase stepwise while monitoring memory.¹

export NCCL_NTHREADS=512
export NCCL_BUFFSIZE=4194304     # 4 MiB; raise stepwise, watch GPU memory

Related channel tunables (leave at default unless profiled): NCCL_MIN_CTAS / NCCL_MAX_CTAS control how many subrings/channels (CTAs, Cooperative Thread Arrays) NCCL uses; each channel is one CUDA block, so more channels cost more GPU resources, though too few can leave throughput on the table. These replace the deprecated (since NCCL 2.17, still functional) NCCL_MIN_NCHANNELS / NCCL_MAX_NCHANNELS names. Both accept integers up to 64 (32 prior to NCCL 2.25) and are platform-dependent by default.³ On NVSwitch systems NCCL auto-tunes channel count by topology and message size.¹

How to develop with it¶

Confirm the path before trusting throughput¶

NCCL falls back silently. Always confirm the intended algorithm/protocol/transport is actually active:¹

export NCCL_DEBUG=INFO                 # log NET/IB paths, algo/proto, fallbacks
export NCCL_TOPO_DUMP_FILE=/tmp/nccl_topo.xml   # dump detected topology
export NCCL_SOCKET_IFNAME=ib0          # bootstrap handshake over the IB HCA

A red flag for a silent fallback: during all-reduce, GPU utilization drops and CPU utilization spikes: the CPU is copying data instead of GPUDirect RDMA.¹ Do not leave debug-only kills (NCCL_P2P_DISABLE=1, NCCL_SHM_DISABLE=1) set in production; they force host-staged copies and collapse intranode bandwidth from hundreds of GB/s to tens.¹ See RDMA and RoCE Performance Tuning and NCCL Hang / Collective Stall.

How to maintain it¶

Validate with nccl-tests bus bandwidth¶

nccl-tests reports two bandwidths. Algorithm bandwidth (algbw) = input size / time. Bus bandwidth (busbw) corrects algbw for the number of ranks so the result is comparable to hardware peak independent of rank count.⁴ The official correction factors are:⁴

Collective	busbw = algbw x
AllReduce	`2*(n-1)/n`
AllGather	`(n-1)/n`
ReduceScatter	`(n-1)/n`
Broadcast	`1`

where n is the number of ranks. Compare measured busbw against the link's hardware peak; a large gap means a degraded path or a suboptimal algorithm/protocol.

This standard-library model checks the official nccl-tests correction factors, including invalid rank counts and collective names:

def bus_bandwidth(collective, algorithm_bandwidth, ranks):
    if ranks < 1:
        raise ValueError("ranks must be positive")
    if algorithm_bandwidth < 0:
        raise ValueError("bandwidth cannot be negative")

    factors = {
        "all_reduce": 2 * (ranks - 1) / ranks,
        "all_gather": (ranks - 1) / ranks,
        "reduce_scatter": (ranks - 1) / ranks,
        "broadcast": 1.0,
    }
    if collective not in factors:
        raise ValueError(f"unknown collective: {collective}")
    return algorithm_bandwidth * factors[collective]


assert bus_bandwidth("all_reduce", 100.0, 8) == 175.0
assert bus_bandwidth("all_gather", 100.0, 8) == 87.5
assert bus_bandwidth("reduce_scatter", 100.0, 8) == 87.5
assert bus_bandwidth("broadcast", 100.0, 8) == 100.0
assert bus_bandwidth("all_reduce", 100.0, 1) == 0.0

for args in (
    ("all_reduce", 100.0, 0),
    ("unknown", 100.0, 8),
    ("broadcast", -1.0, 8),
):
    try:
        bus_bandwidth(*args)
    except ValueError:
        pass
    else:
        raise AssertionError(f"invalid input accepted: {args}")

print("NCCL bus-bandwidth validation: all asserts passed")

Executed output:

NCCL bus-bandwidth validation: all asserts passed

Run a size sweep on one node (8 GPUs), then across nodes via MPI (binary must be built with MPI=1):⁵

# Single node, 8 GPUs: sweep 8 B -> 128 MiB, doubling each step
./build/all_reduce_perf -b 8 -e 128M -f 2 -g 8

# 64 GPUs across 8 nodes (8 GPUs/node), 1 GPU per process
mpirun -np 64 -N 8 ./build/all_reduce_perf -b 8 -e 8G -f 2 -g 1

Flags: -b minimum size, -e maximum size, -f size multiplication factor, -g GPUs per thread.⁵ The sweep exposes the size-dependent crossover where NCCL shifts from latency-oriented (tree, LL/LL128) to bandwidth-oriented (ring, Simple) behavior; watch busbw rise toward link peak as message size grows.

Treat NCCL warnings as actionable: unable to enable P2P, falling back to copy and NET/Socket: using Ethernet interface eth0 both indicate the fast path was not taken; track down the cause rather than ignoring it.¹ Re-validate busbw after every NCCL upgrade; performance usually improves but defaults can shift and require retuning.¹

How to run it in production¶

Communicator lifecycle: initialize once¶

A communicator (ncclCommInitRank, or PyTorch's torch.distributed.init_process_group) is expensive to create: every rank exchanges unique IDs and addresses, and NCCL builds the rings/trees and allocates buffers for the group. The cost grows worse than linearly with rank count because setup needs an all-to-all handshake among the participants.⁶

The anti-pattern is calling init_process_group/destroy_process_group inside the training loop instead of once at startup. On a 2-rank job the book measures 48 ms of init/destroy overhead versus a 0.5 ms all-reduce, i.e. the collective itself is over 98% dominated by avoidable setup cost; on a 32-GPU job, naive per-rank communicator creation can take 2-3 minutes instead of 2-3 seconds.⁶ PyTorch's DDP already does this correctly (one init_process_group call, reused for every step); the mistake shows up in hand-rolled multi-communicator code for model/pipeline parallelism, where a naive implementation re-derives a subgroup communicator every iteration instead of caching it.

# Anti-pattern: measured ~48 ms/iter of init+destroy overhead added to a 0.5 ms all-reduce
# Fix: call init_process_group() once outside the loop, destroy_process_group() once at exit.

For subgroups (tensor-parallel or pipeline-parallel process groups), create each with torch.distributed.new_group() once at startup and reuse the handle; never create and destroy a subgroup communicator per iteration.⁶ If you must create several communicators together at once (dynamic membership, staged initialization), NCCL's C++ API batches the handshake with ncclGroupStart() / ncclCommInitRank() (repeated per communicator) / ncclGroupEnd(), which amortizes the all-to-all setup cost across the batch instead of paying it once per communicator.⁶ As of this writing PyTorch does not expose fully dynamic communicator membership at runtime without a full teardown; every rank must call creation and destruction in lockstep or the job hangs.⁶

Persistent user buffer registration (zero-copy)¶

NCCL supports registering your own tensor buffers with a communicator via ncclCommRegister() / ncclCommDeregister(), so collectives operate directly on those buffers instead of staging through internal channel buffers. If any rank in a collective uses a registered buffer, every rank must; for some algorithms the buffer's offset from its head must also match across ranks.⁷ Registration is a one-time setup cost, so it pays off for long-lived, repeatedly-reduced buffers (persistent gradient buffers, weight tensors) rather than one-shot transfers. It is the prerequisite for the fastest SHARP/NVLS paths, both on-node and off-node; see SHARP: In-Network Reduction for the fabric-level payoff.

Profiling communicators: the NCCL profiler plugin API¶

Beyond NCCL_DEBUG=INFO logging, NCCL exposes a plugin API (NCCL_PROFILER_PLUGIN, loaded like other NCCL plugins) that lets a profiler observe the internal timeline of group, collective, and point-to-point events, plus proxy-thread activity, as GPU counts scale into the range where a flat log becomes unreadable.⁸ The plugin configures a 32-bit event-activation bitmask (one bit per event category) and implements five callbacks:⁸

init: sets up the plugin context and declares which event categories to capture.
startEvent: receives an event descriptor from NCCL, allocates an event object, and returns an opaque handle.
stopEvent: marks an event complete so its resources can be recycled.
recordEventState: updates an in-flight event as it transitions between states.
finalize: releases all plugin resources once profiling ends.

This is the integration point third-party profilers (PyTorch Kineto included) use to attribute NCCL activity precisely; when the plugin is not loaded, Kineto can still gather coarser NCCL activity through CUPTI and NVTX.⁸

Failure modes¶

A forced algorithm or protocol excludes the best path for the deployed topology and size.
NCCL selects sockets or host staging instead of NVLink or GPUDirect RDMA.
Ranks call collectives in different orders or with incompatible counts and datatypes.
A process group is created inside the iteration loop, adding setup latency and failure surface.
Registered-buffer requirements differ across ranks.
A bandwidth comparison uses algbw where link-normalized busbw is required.

References¶

Chris Fregly, AI Systems Performance Engineering (O'Reilly, 2026), Chapter 4, "Tuning Distributed Networking Communication" — NCCL collectives, topology awareness, communication algorithms (Ring/Tree/CollTree/CollNet/PAT), NVLS/SHARP, and environment-variable gotchas.
NVIDIA NCCL — Environment Variables: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html (NCCL_ALGO, NCCL_PROTO, NCCL_NTHREADS, NCCL_BUFFSIZE, NCCL_MIN_CTAS/NCCL_MAX_CTAS values and defaults).
NVIDIA nccl-tests — Performance metrics (busbw correction factors): https://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md
NVIDIA nccl-tests — Usage (all_reduce_perf flags, MPI invocation): https://github.com/NVIDIA/nccl-tests/blob/master/README.md
NVIDIA Developer Blog — New Scaling Algorithm and Initialization with NCCL 2.23 (PAT is Bruck-derived, for AllGather/ReduceScatter): https://developer.nvidia.com/blog/new-scaling-algorithm-and-initialization-with-nvidia-collective-communications-library-2-23/

Reference templates only. The commands, env vars, and bandwidth formulas here are transcribed from the book and official NVIDIA documentation; they have not been hardware-tested in this knowledge base. Validate on your own fabric before relying on any number.

Fregly, AI Systems Performance Engineering, Ch. 4. ↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩
NVIDIA NCCL Environment Variables, https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html ↩↩↩↩↩↩↩↩↩
NVIDIA NCCL Environment Variables: NCCL_MIN_CTAS/NCCL_MAX_CTAS control the minimum/maximum CTAs (channels) NCCL uses per communicator, settable up to 64 (32 prior to NCCL 2.25) and also configurable via ncclCommInitRankConfig; they replace NCCL_MIN_NCHANNELS/NCCL_MAX_NCHANNELS, deprecated since NCCL 2.17 but still functional. https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html ↩
NVIDIA nccl-tests PERFORMANCE.md, https://github.com/NVIDIA/nccl-tests/blob/master/doc/PERFORMANCE.md ↩↩↩
NVIDIA nccl-tests README.md, https://github.com/NVIDIA/nccl-tests/blob/master/README.md ↩↩
Fregly, Ch. 4, "NCCL Communicator Lifecycle and Environment Gotchas": communicator init exchanges unique IDs/addresses and builds rings/trees across all ranks, with worse-than-linear scaling by rank count; measured 48 ms init/destroy overhead vs. a 0.5 ms all-reduce on 2 ranks, and 2-3 minutes vs. 2-3 seconds on 32 GPUs if communicators are created per rank; create subgroups once via new_group()/ncclGroupStart()+ncclCommInitRank()+ncclGroupEnd() and reuse; PyTorch does not support dynamic membership changes without a full teardown, and all ranks must create/destroy in lockstep or the job hangs. ↩↩↩↩↩
Fregly, Ch. 4, "Persistent NCCL User Buffers and Zero-Copy Registration": ncclCommRegister()/ncclCommDeregister() let collectives operate directly on registered application buffers; if any rank registers, all ranks must, and buffer offsets from head must match across ranks for some algorithms; registration is essential to the best SHARP paths for both on-node (NVLS) and off-node (InfiniBand). ↩
Fregly, Ch. 4, "Profiling and Debugging NCCL": the NCCL profiler plugin API, loaded via NCCL_PROFILER_PLUGIN, exposes a 32-bit event-activation bitmask and five callbacks (init, startEvent, stopEvent, recordEventState, finalize) so third-party profilers such as PyTorch Kineto can capture a hierarchical, low-overhead timeline of group/collective/point-to-point/proxy events; without the plugin, Kineto still gathers coarser NCCL activity via CUPTI and NVTX. ↩↩↩
NVIDIA Developer Blog, "New Scaling Algorithm and Initialization with NVIDIA Collective Communications Library 2.23": "The PAT algorithm is a variation of the Bruck algorithm, which features a logarithmic number of network steps for small sizes at scale, progressively increasing the number of network transfers as sizes increase, to keep buffering needs minimal"; applies to AllGather and ReduceScatter, works for any rank count, and initial support requires one GPU per node. https://developer.nvidia.com/blog/new-scaling-algorithm-and-initialization-with-nvidia-collective-communications-library-2-23/ ↩