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.

21 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

Animated four-rank all-reduce. Four colored payloads move into a reduction point; the completed sum then moves out to every rank.
Final frame5.2s loop
Follow one tensor position: four rank-local values move into the reduction, then the completed sum moves to every rank. This is logical ownership flow, not a claim that NCCL sends everything through one physical center.

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.

Rank 0 calls all-reduce, while rank 1 calls broadcast on the same communicator. Why can neither operation finish?

Answer

Collectives are ordered group agreements. Each rank is waiting for peers to join the same operation with compatible arguments, but the two ranks entered different protocols.

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

These contracts become easier to distinguish when payloads retain identity while moving. Each animation below follows one operation from initial to final ownership.

Animated NCCL broadcast. Root rank zero retains its payload while three copies move to ranks one, two, and three.
Final frame5.2s loop
Broadcast starts with one root-owned buffer. Copies fan out to peers; values from non-root ranks never enter the operation.
Animated NCCL reduce-scatter. Rank-local chunks move into four destination-specific reductions, then one completed shard moves smoothly to each owning rank.
Final frame5.2s loop
Reduce-scatter groups matching chunk positions across sources, reduces each group, and moves only the completed shard to its owner. Every rank finishes with one shard, not a full tensor.
Animated NCCL all-gather. One labeled shard from each rank moves into a concatenated set, then copies move to every rank.
Final frame5.2s loop
All-gather starts where reduce-scatter ends: each rank owns one finished shard. Those shards circulate until every rank reconstructs the ordered tensor.
Animated NCCL all-to-all. Sixteen labeled payloads move from source-rank rows to destination-rank rows, preserving source identity while ownership changes.
Final frame5.2s loop
Read color as destination and label as source-to-destination address. Every payload takes its own path; destination rank R2 finishes with one R2-bound payload from every source. This transpose-like ownership change is core MoE dispatch and combine traffic.

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.

After reduce-scatter in the four-rank example, rank 2 owns 3333. What extra operation is needed if rank 2 alone will update only that shard?

Answer

None for gradient ownership. All-gather is needed only when later computation requires every rank to reconstruct the full tensor or a full parameter bucket.

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.

Animated ring all-reduce following tensor chunk x2. A partial sum travels through four neighboring ranks during reduce-scatter, then the completed shard travels through the ring during all-gather.
Final frame8.5s loop
Follow only chunk `x2` to keep the choreography legible. Three neighbor steps turn `3` into `3333`; three more distribute that completed shard. Other chunks move through the same ring concurrently, offset in the pipeline.

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));
Animated CUDA stream timeline. Backward compute produces a gradient bucket, NCCL reduce-scatter overlaps later compute on a communication stream, and an optimizer waits for a completion event before reading its shard.
Final frame6.7s loop
Overlap shortens exposed communication time, but doesn't weaken dependencies. NCCL may start after a bucket-ready event while later backward kernels continue. The optimizer must still wait for the reduced-shard event before reading that buffer.

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

The same all-reduce is fast within one server but slows sharply across two servers. What changed if tensor size and rank count stayed fixed?

Answer

The physical route changed. Cross-node traffic adds PCIe, network interfaces, and fabric links, so placement, rail balance, or network bandwidth can become the limit even though the collective API is unchanged.

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

Correct answer: Reduce-scatter computes and partitions the reduced chunks, then all-gather replicates those finished chunks

The first operation produces one reduced chunk per rank. The second distributes those completed chunks so every rank owns the full reduced tensor.

2.Which operation best matches FSDP gradient ownership when each rank should retain only its reduced shard?

Correct answer: Reduce-scatter

Reduce-scatter combines values across ranks and leaves one output partition on each rank, matching sharded gradient ownership.

3.What is the main ring-versus-tree tradeoff for all-reduce?

Correct answer: Rings favor large-message bandwidth, while trees reduce communication rounds

Ring traffic approaches the bandwidth lower bound for large payloads but needs linearly many neighbor steps. Tree depth grows logarithmically.

4.An ncclAllReduce host call returns successfully. What does that prove?

Correct answer: The operation was enqueued on its CUDA stream, not necessarily completed

NCCL follows CUDA stream semantics. Consumers need correct stream ordering, events, or synchronization before reading the result.

5.All ranks except rank 5 are stuck inside the same collective. Rank 5 logged an OOM first. What should you investigate first?

Correct answer: The rank 5 OOM that removed a collective participant

Peer stacks show where survivors wait. The earliest rank failure usually explains why the distributed agreement could no longer finish.

6.Why can a wider tensor-parallel group reduce performance even when every GPU has spare compute?

Correct answer: Frequent layer collectives can cross slower links and expose communication latency

Tensor parallelism communicates within many layers. Crossing a topology boundary can make those small, frequent collectives dominate token latency.

7.What does a good nccl-tests result establish?

Correct answer: The tested fabric and message shape can achieve that standalone communication result

The benchmark isolates communication capability. Application overlap, rank balance, operation ordering, and compute scheduling still need separate evidence.

8.Which license statement matches the pinned NCCL source tree?

Correct answer: Mostly Apache-2.0 source, retained BSD sections, and separate terms for borrowed files

The root license explicitly preserves file-level boundaries. Check headers and third-party notices when redistributing or modifying source.

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

https://docs.nvidia.com/deeplearning/nccl/user-guide/index.html

NCCL Collective Operations

NVIDIA · 2026

https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html

Bandwidth Optimal All-Reduce Algorithms for Clusters of Workstations

Patarasuk, P., & Yuan, X. · 2009 · Journal of Parallel and Distributed Computing

https://doi.org/10.1016/j.jpdc.2008.09.002

Massively Scale Your Deep Learning Training with NCCL 2.4

Jeaugey, S. · 2019

https://developer.nvidia.com/blog/massively-scale-deep-learning-training-nccl-2-4/

NCCL Source Repository

NVIDIA and NCCL Contributors · 2026

https://github.com/NVIDIA/nccl/tree/5067397c2676d5aed50042fc39e5c8ee96eb0027

NCCL CUDA Stream Semantics

NVIDIA · 2026

https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/streams.html

NCCL Tests

NVIDIA and NCCL Contributors · 2026

https://github.com/NVIDIA/nccl-tests

NCCL Troubleshooting

NVIDIA · 2026

https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting.html

NVIDIA NCCL Contributors

NVIDIA and NCCL Contributors · 2026

https://github.com/NVIDIA/nccl/graphs/contributors

NCCL Source License

NVIDIA · 2026

https://github.com/NVIDIA/nccl/blob/5067397c2676d5aed50042fc39e5c8ee96eb0027/LICENSE.txt

Discussion

Questions and insights from fellow learners.

Discussion loads when you reach this section.