LeetLLM
My PlanLearnGlossaryTracksPracticeBlog
LeetLLM

Your go-to resource for mastering AI & LLM systems.

Product

  • Learn
  • Glossary
  • Tracks
  • Practice
  • Blog
  • RSS

Legal

  • Terms of Service
  • Privacy Policy

© 2026 LeetLLM. All rights reserved.

All Topics
Your Progress
0%

0 of 177 articles completed

🛠️Computing Foundations0/9
Git, Shell, Linux for AIDocker for Reproducible AIPython for AI EngineeringNumPy and Tensor ShapesCUDA for ML TrainingMPS & Metal for ML on MacData Structures for AISQL and Data ModelingAlgorithms for ML Engineers
📊Math & Statistics0/8
Gradients and BackpropVectors, Matrices & TensorsLinear Algebra for MLAdam, Momentum, SchedulersProbability for Machine LearningStatistics and UncertaintyDistributions and SamplingHypothesis Tests, Intervals, and pass@k
📚Preparation & Prerequisites0/13
Neural Networks from ScratchCNNs from ScratchTraining & BackpropagationSoftmax, Cross-Entropy & OptimizationRNNs, LSTMs, GRUs, and Sequence ModelingAutoencoders and VAEsThe Transformer Architecture End-to-EndLanguage Modeling & Next TokensFrom GPT to Modern LLMsPrompt Engineering FundamentalsCalling LLM APIs in ProductionFirst AI App End-to-EndThe LLM Lifecycle
🧮ML Algorithms & Evaluation0/11
Linear Regression from ScratchLogistic Regression and MetricsDecision Trees, Forests, and BoostingReinforcement Learning BasicsValidation and LeakageClustering and PCACore Retrieval AlgorithmsDecoding AlgorithmsExperiment Design and A/B TestingPyTorch Training LoopsDataset Pipelines and Data Quality
📦Production ML Systems0/6
Feature Engineering for Production MLBatch and Streaming Feature PipelinesGradient Boosted Trees in ProductionRanking and Recommendation SystemsForecasting and Anomaly DetectionMonitoring Predictive Models
🧪Core LLM Foundations0/8
The Bitter Lesson & ComputeBPE, WordPiece, and SentencePieceStatic to Contextual EmbeddingsPerplexity & Model EvaluationFile Ingestion for AIChunking StrategiesLLM Benchmarks & LimitationsInstruction Tuning & Chat Templates
🧰Applied LLM Engineering0/24
Dimensionality Reduction for EmbeddingsCoT, ToT & Self-Consistency PromptingFunction Calling & Tool UseMCP & Tool Protocol StandardsContext EngineeringPrompt Injection DefenseResponsible AI GovernanceData Labeling and Human FeedbackEvaluating AI AgentsProduction RAG PipelinesHybrid Search: Dense + SparseReranking and Cross-Encoders for RAGRAG Evaluation for Reliable AnswersLLM-as-a-Judge EvaluationBias & Fairness in LLMsHallucination Detection & MitigationLLM Observability & MonitoringExperiment Tracking with MLflow and W&BPrompt Optimization with DSPyModel Versioning & DeploymentSemantic Caching & Cost OptimizationLLM Cost Engineering & Token EconomicsModel Gateways, Routing, and FallbacksDesign an Automated Support Agent
🎓Portfolio Capstones0/9
Capstone: Delivery ETA PredictionCapstone: Product RankingCapstone: Demand ForecastingCapstone: Image Damage ClassifierCapstone: Production ML PipelineCapstone: Document QACapstone: Eval DashboardCapstone: Fine-Tuned ClassifierCapstone: Reproducible ML Study
🧠Transformer Deep Dives0/8
Sentence Embeddings & Contrastive LossEmbedding Similarity & QuantizationScaled Dot-Product AttentionVision Transformers and Image EncodersPositional Encoding: RoPE & ALiBiLayer Normalization: Pre-LN vs Post-LNMechanistic InterpretabilityDecoding Strategies: Greedy to Nucleus
🧬Advanced Training & Adaptation0/16
Scaling Laws & Compute-Optimal TrainingPre-training Data at ScaleBuild GPT from Scratch LabJAX for PyTorch ResearchersContinued Pretraining for Domain ShiftSynthetic Data PipelinesSupervised Fine-Tuning PipelineMixed Precision TrainingDistributed Training: FSDP & ZeROLoRA & Parameter-Efficient TuningReward Modeling from Preference DataRLHF & DPO AlignmentConstitutional AI & Red TeamingRLVR & Verifiable RewardsKnowledge Distillation for LLMsModel Merging and Weight Interpolation
🤖Advanced Agents & Retrieval0/16
Vector DB Internals: HNSW & IVFAdvanced RAG: HyDE & Self-RAGGraphRAG & Knowledge GraphsRAG Security & Access ControlStructured Output GenerationReAct & Plan-and-ExecuteGuardrails & Safety FiltersCode Generation & SandboxingComputer-Use / GUI / Browser AgentsHuman-in-the-Loop Agent ArchitectureAI Coding Workflow with AgentsAgent Memory & PersistenceAgent Failure & RecoveryRecursive Language Models (RLM)Multi-Agent OrchestrationCapstone: Production Agent
⚡Inference & Production Scale0/19
Inference: TTFT, TPS & KV CacheMulti-Query & Grouped-Query AttentionKV Cache & PagedAttentionPrefix Caching and Prompt CachingFlashAttention & Memory EfficiencyContinuous Batching & SchedulingScaling LLM InferenceModel Parallelism for LLM InferenceModel Quantization: GPTQ, AWQ & GGUFLocal LLM DeploymentSLM Specialization & Edge DeploymentSpeculative DecodingLong Context Window ManagementMixture of Experts ArchitectureMamba & State Space ModelsReasoning & Test-Time ComputeAdvanced MLOps & DevOps for AIGPU Serving & AutoscalingA/B Testing for LLMs
🏗️System Design Capstones0/9
Content Moderation SystemCode Completion SystemMulti-Tenant LLM PlatformLLM-Powered Search EngineVision-Language Models & CLIPMultimodal LLM ArchitectureDiffusion Models: Images & TextReal-Time Voice AI AgentReasoning & Test-Time Compute
🎤AI Lab Interviewing0/4
AI Lab Coding Interview: Python SystemsAI Lab System Design InterviewAI Lab Behavioral InterviewAI Lab Technical Presentation
🔬Project Deep Dives0/17
Deep Dive - vLLMDeep Dive - SkyRLDeep Dive - FlashAttentionDeep Dive - FlashInferDeep Dive - DeepGEMMDeep Dive - NCCLDeep Dive - MegatronDeep Dive - DeepSpeedDeep Dive - RayDeep Dive - MLflowDeep Dive - PyTorchDeep Dive - TransformersDeep Dive - SGLangDeep Dive - slimeDeep Dive - DeepEPDeep Dive - TinkerDeep Dive - Light-PEFT
Back to Topics
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.

19 min read
Learning path
Step 166 of 177 in the full curriculum
Deep Dive - DeepGEMMDeep Dive - Megatron

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]Reference 1NVIDIA Collective Communication Library User Guidehttps://docs.nvidia.com/deeplearning/nccl/user-guide/index.html

Four-rank all-reduce shown as two operations. Four input vectors first become four reduced shards through reduce-scatter, then all-gather replicates the complete summed vector on every rank. Four-rank all-reduce shown as two operations. Four input vectors first become four reduced shards through reduce-scatter, then all-gather replicates the complete summed vector on every rank.
Follow one column from four inputs to one reduced owner, then across every output rank. Reduce-scatter performs the arithmetic and partitions ownership; all-gather replicates the finished vector.

Start with the collective contract

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]Reference 2NCCL Collective Operationshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html

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

yi=∑r=0k−1xi(r),i∈{0,…,N−1}.\mathbf{y}_i = \sum_{r=0}^{k-1} \mathbf{x}^{(r)}_i, \qquad i \in \{0, \ldots, N-1\}.yi​=r=0∑k−1​xi(r)​,i∈{0,…,N−1}.

Here, kkk is rank count, NNN is elements per input vector, and iii selects one element. The operation combines values and puts the same answer on every rank. A training framework may divide the sum by kkk before applying a mean gradient, but that scaling convention sits above NCCL.

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.

OperationInput on each rankOutput ownershipCommon AI use
All-reduceSame-shaped tensorFull reduced tensor on every rankReplicated data-parallel gradients
Reduce-scatterSame-shaped tensorOne reduced chunk per rankSharded gradients in FSDP or ZeRO
All-gatherOne local chunkConcatenated chunks on every rankParameter or activation reconstruction
BroadcastFull tensor on rootRoot tensor copied to all ranksSeeds, metadata, or initial weights
All-to-allDestination chunksOne chunk from every sourceMoE token routing
Send / receivePoint-to-point bufferNamed peer onlyPipeline 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]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:

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

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 kkk chunks. Reduce-scatter circulates and combines chunks for k−1k-1k−1 steps; all-gather circulates the completed chunks for another k−1k-1k−1 steps.

For message size MMM, the payload sent by each rank in a ring all-reduce is approximately:

Vsend=2k−1kM.V_{send} = 2\frac{k-1}{k}M.Vsend​=2kk−1​M.

Each rank receives the same payload volume. As kkk grows, the sent volume approaches 2M2M2M, not kMkMkM. This bandwidth property is why ring all-reduce remains strong for large buffers when links can be kept busy.[3]Reference 3Bandwidth Optimal All-Reduce Algorithms for Clusters of Workstationshttps://doi.org/10.1016/j.jpdc.2008.09.002

Rings pay in latency. The algorithm needs 2(k−1)2(k-1)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⁡2k\log_2 klog2​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.[4]Reference 4Massively 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. 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]Reference 5NCCL Source Repositoryhttps://github.com/NVIDIA/nccl/tree/5067397c2676d5aed50042fc39e5c8ee96eb0027

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 shapeLikely pressureWhat to inspect
Large gradient bucketsLink bandwidthRing or fabric-offload path, channel count, bus bandwidth
Tiny tensor-parallel reductionsStartup latencyTree-like path, protocol, launch gaps
MoE all-to-allImbalance and network injectionPer-peer bytes, hot experts, NIC rails
Pipeline send/receiveBubble and orderingStage timeline, peer pairing, stream waits
FSDP all-gatherBurst overlapBucket timing, parameter prefetch, compute coverage

A communicator is an ordered program

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.

all_reduce_step.cu
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]Reference 6NCCL CUDA Stream Semanticshttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/streams.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: 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.

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.

NCCL topology map with a short intra-node GPU-to-NVSwitch path and a longer cross-node route through PCIe, network interface cards, and an InfiniBand or RoCE fabric. Labels mark where topology selection and transport bottlenecks appear. NCCL topology map with a short intra-node GPU-to-NVSwitch path and a longer cross-node route through PCIe, network interface cards, and an InfiniBand or RoCE fabric. Labels mark where topology selection and transport bottlenecks appear.
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.

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.

Read one call through the source tree

The pinned source snapshot makes the runtime boundary concrete.[5]Reference 5NCCL Source Repositoryhttps://github.com/NVIDIA/nccl/tree/5067397c2676d5aed50042fc39e5c8ee96eb0027 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.

Diagram showing Host planning, Device and transport, 1 · collectives.cc API + arguments, and 2 · enqueue.cc work planning. Diagram showing Host planning, Device and transport, 1 · collectives.cc API + arguments, and 2 · enqueue.cc work planning.
Host planning, Device and transport, 1 · collectives.cc API + arguments, and 2 · enqueue.cc work planning.

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

Source areaResponsibilityQuestion to carry
src/init.cc, bootstrap.ccCommunicator setup and peer discoveryWhich ranks and devices joined?
src/graph/topo.cc, search.ccPhysical topology and graph searchWhich links can form rings or trees?
src/graph/tuning.ccAlgorithm, protocol, and channel modelWhy did this message get this plan?
src/enqueue.cc, collectives.ccHost API validation and work creationWhat exactly was enqueued?
src/device/*.h, primitives.hDevice-side copy and reduction stepsWhich rank sends, receives, or reduces next?
src/transport/*P2P, shared memory, sockets, InfiniBand, pluginsWhich physical route carries bytes?
src/proxy.ccHost proxy progress for transports that need itIs progress waiting on host or network work?
src/debug.cc, src/ras/Logs and reliability signalsWhat 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.

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 patternNCCL trafficScaling failure when communication wins
Data-parallel trainingGradient all-reduceMore replicas add little tokens/sec
FSDP or ZeROReduce-scatter plus parameter all-gatherMemory fits, but layer gathers stall compute
Tensor parallelismFrequent activation reductions or gathersWide TP across slow links hurts every layer
Pipeline parallelismSend/receive activations and gradientsStage bubbles or peer waits grow
Expert parallelismAll-to-all token dispatch and combineHot experts create uneven peer traffic
Multi-GPU inferenceTensor-parallel reductions, KV or state exchangeInter-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.

Measure bytes, rounds, and exposed time

A simple communication estimate separates startup latency from payload time:

Tcollective≈Rα+VB,T_{collective} \approx R\alpha + \frac{V}{B},Tcollective​≈Rα+BV​,

where RRR is communication rounds, α\alphaα is per-round startup cost, VVV is bytes crossing the measured path, and BBB 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]Reference 7NCCL Testshttps://github.com/NVIDIA/nccl-tests

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.

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.

SymptomLikely classFirst evidence
One rank exited or OOMedMissing participantEarliest per-rank error, not peer timeouts
Same step hangs every runCollective order or shape mismatchPer-rank operation sequence and tensor metadata
Multi-node onlyInterface, firewall, NIC, or RDMA pathBootstrap and NET logs, host connectivity
Slow after scale-outTopology or stragglerTimeline by rank, link counters, nccl-tests
Container init failureShared memory or memlock/dev/shm, limits, NCCL warning
Wrong data without hangBuffer lifetime or stream dependencyCUDA 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]Reference 8NCCL 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 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:

  • 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 are powerful enough to 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.

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 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.[5]Reference 5NCCL Source Repositoryhttps://github.com/NVIDIA/nccl/tree/5067397c2676d5aed50042fc39e5c8ee96eb0027

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]Reference 9NVIDIA NCCL Contributorshttps://github.com/NVIDIA/nccl/graphs/contributors

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]Reference 10NCCL 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.[3]Reference 3Bandwidth 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.[4]Reference 4Massively 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. 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 if you deliberately skip the call on one process?

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.

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 minimize bandwidth cost for large messages; trees cut round count for smaller messages or large rank counts.
  • CUDA streams make collectives asynchronous, but explicit dependencies still protect buffer correctness.
  • Topology determines whether bytes use NVLink, PCIe, shared memory, sockets, InfiniBand, RoCE, or plugins.
  • Higher-level frameworks own sharding and scheduling. NCCL owns communication execution.
  • Debug the earliest rank divergence, then validate the fabric separately.
  • Current source is NVIDIA-led, publicly contributed, and carries mixed file-level license terms.

Mastery quiz

Complete the lesson

Mastery Check

Answer every question, then check your score. Score 75% or higher to mark this lesson complete.

1.Why can all-reduce be decomposed into reduce-scatter followed by all-gather?
2.Which operation best matches FSDP gradient ownership when each rank should retain only its reduced shard?
3.What is the main ring-versus-tree tradeoff for all-reduce?
4.An ncclAllReduce host call returns successfully. What does that prove?
5.All ranks except rank 5 are stuck inside the same collective. Rank 5 logged an OOM first. What should you investigate first?
6.Why can a wider tensor-parallel group reduce performance even when every GPU has spare compute?
7.What does a good nccl-tests result establish?
8.Which license statement matches the pinned NCCL source tree?

8 questions remaining.

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
XFacebookLinkedInBlueskyRedditHacker NewsEmail
References

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

Discussion

Questions and insights from fellow learners.

Discussion loads when you reach this section.