Read NCCL as the communication engine beneath distributed AI: collective contracts, rings and trees, topology discovery, CUDA streams, transports, profiling, and hang diagnosis.
Personalize this lesson
Adapt explanations and teaching visuals to your background and preferred voice.
Four GPUs finish computing four gradient vectors. None can update the model yet. Every GPU needs the same global sum, and moving those bytes can take longer than producing them.
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]
A rank is one participant in a communication group. A communicator records the group of ranks and their mapping to devices. Each rank calls the same collective in the same order, with matching element counts and data types. Breaking that agreement can cause a hang, crash, or corrupted result.[2]
For sum all-reduce over ranks, rank contributes a vector with elements. Every rank receives:
Here, is rank count, is elements per input vector, and selects one element. The operation combines values and puts the same answer on every rank. A training framework may divide the sum by before applying a mean gradient, but that scaling convention sits above NCCL.
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 |
An all-reduce can be expressed 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.[2]
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]. During reduce-scatter, rank 0 can finish and own 1111, rank 1 can own 2222, and so on. All-gather then shares those four owned chunks so all ranks reconstruct the same vector.
That decomposition explains why sharded training prefers reduce-scatter plus later all-gather. Fully Sharded Data Parallel and ZeRO don't need to keep the full reduced gradient on every rank. They keep one gradient shard, update one optimizer shard, and gather parameters only when computation needs them.
A ring orders ranks so each sends to one neighbor and receives from another. Split a message into chunks. Reduce-scatter circulates and combines chunks for steps; all-gather circulates the completed chunks for another steps.
For message size , the payload sent by each rank in a ring all-reduce is approximately:
Each rank receives the same payload volume. As grows, the sent volume approaches , not . This bandwidth property is why ring all-reduce remains strong for large buffers when links can be kept busy.[3]
Rings pay in latency. The algorithm needs 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.
A binary tree reduces values toward a root and broadcasts the answer back down. Tree depth grows roughly with , 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.[4]
No algorithm wins for every tensor and topology. Rings favor sustained bandwidth. Trees favor fewer rounds. Current NCCL source also names CollNet, NVLS, NVLS Tree, and PAT algorithm families, plus Simple, LL, and LL128 protocols. Availability depends on operation, hardware, network plugins, registration, and runtime tuning.[5]
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.
| 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 |
NCCL doesn't launch ranks or distribute a unique communicator ID by itself. A launcher such as MPI, Slurm, or torchrun starts processes and shares bootstrap information. Each process then binds a rank to a CUDA device and creates its communicator.
After communicator creation, each rank enqueues this C++ collective with its own buffers and CUDA stream.
1NCCLCHECK(ncclAllReduce(
2 local_gradient,
3 global_gradient,
4 element_count,
5 ncclFloat32,
6 ncclSum,
7 communicator,
8 communication_stream));The host call returning doesn't mean the bytes have arrived. It means NCCL enqueued work on communication_stream. A later CUDA event, stream dependency, or synchronization establishes when another kernel may consume global_gradient.[6]
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: A 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.
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.
NCCL discovers GPU, CPU, PCIe, NVLink, and network relationships, searches candidate graph layouts, and tunes a schedule. GPUDirect 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.
The pinned source snapshot makes the runtime boundary concrete.[5] Start with src/collectives.cc, where public functions such as ncclAllReduce package operation metadata. From there, enqueue and planner code select work, graph code maps topology, device headers define collective kernels, and transport code moves chunks.
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? |
One conceptual call crosses all those layers. 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.
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 add little tokens/sec |
| FSDP or ZeRO | 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. DeepEP uses newer device-side communication interfaces for expert routing. PyTorch exposes NCCL through torch.distributed. None of those layers remove the need to understand the communicator and physical path.
A useful optimization starts 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.
A simple communication estimate separates startup latency from payload time:
where is communication rounds, is per-round startup cost, is bytes crossing the measured path, and 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.
nccl-tests reports algorithmic bandwidth (algbw) and bus bandwidth (busbw). Algorithmic bandwidth uses logical payload size divided by time. Bus bandwidth applies an operation-specific correction so results better reflect traffic on links and can be compared with hardware limits.[7]
Run tests on the same hosts, GPUs, container settings, and network rails as the workload. A strong standalone all-reduce result doesn't prove training overlap is good. It does establish whether the fabric can carry the bytes before model code, data loading, and bucket timing complicate the trace.
🎯 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.
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.[8]
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.
NCCL gives framework authors a small collective API over many NVIDIA GPU and network topologies. It fuses communication and reduction work on device paths, chooses algorithms dynamically, supports single-process and multi-process applications, and exposes plugins for network, tuner, profiler, and newer device-side interfaces.
Its limits are equally important:
Use NCCL when the system runs on NVIDIA GPUs and needs high-performance intra-node or multi-node communication. Use the framework's process-group API unless you're building a runtime, fused kernel, or communication library that needs NCCL directly.
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.[5]
Public history shows long-running work from Sylvain Jeaugey (sjeaugey) and current contributions from NVIDIA and community accounts including xiakun-lu, nv-lschneider, kwen2501, kgioioso, and others. Treat that list as a source-snapshot view, not a permanent team roster.[9]
The current 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.[10]
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.[3] 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.[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.
Pick one all-reduce from a PyTorch profiler trace. Record tensor bytes, communicator size, duration, and whether compute overlaps it. Then run nccl-tests near that message size on the same ranks.
Use the pinned source map to answer four questions:
The artifact isn't a screenshot of a green benchmark. Produce a small table that connects logical operation, source path, topology path, expected bytes, measured time, and first failure evidence.
Answer every question, then check your score. Score 75% or higher to mark this lesson complete.
8 questions remaining.
NVIDIA Collective Communication Library User Guide
NVIDIA · 2026
NCCL Collective Operations
NVIDIA · 2026
Bandwidth Optimal All-Reduce Algorithms for Clusters of Workstations
Patarasuk, P., & Yuan, X. · 2009 · Journal of Parallel and Distributed Computing
Massively Scale Your Deep Learning Training with NCCL 2.4
Jeaugey, S. · 2019
NCCL Source Repository
NVIDIA and NCCL Contributors · 2026
NCCL CUDA Stream Semantics
NVIDIA · 2026
NCCL Tests
NVIDIA and NCCL Contributors · 2026
NCCL Troubleshooting
NVIDIA · 2026
NVIDIA NCCL Contributors
NVIDIA and NCCL Contributors · 2026
NCCL Source License
NVIDIA · 2026
Questions and insights from fellow learners.