01What is NCCL?

NCCL (NVIDIA Collective Communications Library) is a topology-aware library that provides high-performance collective communication primitives for multi-GPU and multi-node systems. It is the de facto standard for distributed deep-learning frameworks — PyTorch, TensorFlow, JAX, and Horovod all build on top of NCCL for AllReduce, Broadcast, and friends.

NCCL uses a ring or tree algorithm (or a hybrid) chosen at runtime based on the topology. It leverages high-speed interconnects like NVLink, PCIe, and InfiniBand, achieving near-peak bandwidth with minimal latency. The API is written in C, with first-class support for CUDA streams and ncclComm_t communicator objects.

02Getting Started

To compile an NCCL program, link against -lnccl and include the header. The library ships with the CUDA Toolkit or can be built from source for custom interconnect support.

Build Command

# Compile with NCCL and CUDA runtime
nvcc -O3 -arch=native -I/usr/local/cuda/include \
     -L/usr/local/cuda/lib64 -lnccl -lcudart \
     my_program.cu -o my_program

Minimal Header

#include <nccl.h>
#include <cuda_runtime.h>
#include <stdio.h>

03Core Concepts

NCCL is organised around a communicator (ncclComm_t) that represents a group of GPUs. Each GPU is identified by a rank (0 to size-1). All collective operations must be called collectively by every rank in the communicator.

The basic lifecycle is: getUniqueId → CommInitRank → collectives → CommDestroy. All ranks must call ncclCommInitRank with the same unique ID.

04Usage Examples

The following examples demonstrate the essential NCCL primitives. Each assumes rank and size are known (e.g. from MPI or environment variables) and that a CUDA device has already been selected with cudaSetDevice(rank).

4.1 — AllReduce (Sum)

The most common collective in deep learning. Every rank contributes a buffer of n floats; after the call, every rank has the element-wise sum across all ranks.

#include <nccl.h>

void allreduce_example(ncclComm_t comm, float* data,
                         int n, cudaStream_t stream, int rank)
{
    // data is a device pointer of n floats
    ncclResult_t res = ncclAllReduce(
        data,        // sendbuff (in-place)
        data,        // recvbuff (in-place)
        n,           // count
        ncclFloat,   // datatype
        ncclSum,     // reduction op
        comm,
        stream
    );

    if (res != ncclSuccess) {
        fprintf(stderr, "Rank %d: ncclAllReduce failed: %s\n",
                rank, ncclGetErrorString(res));
    }
}

4.2 — Broadcast

Rank 0 sends its buffer to all other ranks. After the call, every rank has a copy of the data from rank 0.

void broadcast_example(ncclComm_t comm, float* buf,
                            int n, cudaStream_t stream, int root)
{
    // All ranks pass the same root rank
    ncclResult_t res = ncclBroadcast(
        buf,        // buffer (same on all ranks)
        buf,        // in-place
        n,
        ncclFloat,
        root,       // source rank
        comm,
        stream
    );
}

4.3 — AllGather

Each rank contributes n elements; after the call, every rank has the concatenation of all contributions, ordered by rank.

void allgather_example(ncclComm_t comm, float* sendbuf,
                            float* recvbuf, int n, cudaStream_t stream)
{
    // recvbuf must be size * n floats
    ncclResult_t res = ncclAllGather(
        sendbuf,    // local contribution (n elements)
        recvbuf,    // full buffer (size*n elements)
        n,
        ncclFloat,
        comm,
        stream
    );
}

4.4 — ReduceScatter

The inverse of AllGather: each rank ends up with a slice of the reduced result. Commonly used in ZeRO-style optimisers.

void reduce_scatter_example(ncclComm_t comm, float* sendbuf,
                                  float* recvbuf, int n, cudaStream_t stream)
{
    // sendbuf: size*n elements; recvbuf: n elements (local slice)
    ncclResult_t res = ncclReduceScatter(
        sendbuf,
        recvbuf,
        n,
        ncclFloat,
        ncclSum,
        comm,
        stream
    );
}

4.5 — Point-to-Point (Send / Recv)

NCCL also supports point-to-point communication. Useful for pipeline parallelism where rank i sends to rank i+1.

void p2p_example(ncclComm_t comm, float* buf, int n,
                      cudaStream_t stream, int rank, int size)
{
    if (rank > 0) {
        // Receive from previous rank
        ncclRecv(buf, n, ncclFloat, rank - 1, comm, stream);
    }
    if (rank < size - 1) {
        // Send to next rank
        ncclSend(buf, n, ncclFloat, rank + 1, comm, stream);
    }
}

4.6 — Full Communicator Setup

A complete skeleton showing communicator initialisation. In practice, you'd exchange the ncclUniqueId between processes using MPI, shared memory, or a file.

int main(int argc, char** argv)
{
    int rank = get_rank_from_env();  // e.g. via MPI or SLURM
    int size = get_size_from_env();

    cudaSetDevice(rank);

    ncclComm_t comm;
    ncclUniqueId id;

    if (rank == 0) {
        ncclGetUniqueId(&id);
        broadcast_id_to_all_ranks(&id);  // user-defined
    } else {
        receive_id_from_rank0(&id);       // user-defined
    }

    ncclResult_t res = ncclCommInitRank(&comm, size, id, rank);
    if (res != ncclSuccess) {
        fprintf(stderr, "CommInitRank failed: %s\n",
                ncclGetErrorString(res));
        return 1;
    }

    // ... use collectives ...

    ncclCommDestroy(comm);
    return 0;
}

05Environment Variables

NCCL behaviour can be tuned via environment variables. These are especially useful for debugging or squeezing out extra performance on specific hardware.

06Best Practices

These guidelines help you avoid common pitfalls and get the most out of NCCL.

1. Use one communicator per process group. Creating communicators is expensive. Reuse them across iterations and destroy them only at program exit.
2. Keep buffers in device memory. NCCL operates on GPU pointers. Avoid host-to-device copies in the critical path — overlap them with computation using CUDA streams.
3. Use in-place operations when possible. sendbuff == recvbuff halves memory traffic and simplifies buffer management.
4. Match collectives across all ranks. Every rank in a communicator must call the same collective with the same arguments. Mismatched calls cause deadlocks or undefined behaviour.
5. Be mindful of data types. NCCL supports ncclFloat, ncclHalf, ncclInt, ncclInt64, etc. Use the smallest type that fits your data to maximise bandwidth.
6. Set NCCL_DEBUG=WARN in production. It catches silent issues without the overhead of full tracing.
7. Use ncclGroupStart/End for multiple collectives. Batching independent collectives into a group lets NCCL fuse operations and reduce latency.

// Fuse multiple collectives into a single group
ncclGroupStart();
ncclAllReduce(grad_a, grad_a, n, ncclFloat, ncclSum, comm, stream);
ncclAllReduce(grad_b, grad_b, n, ncclFloat, ncclSum, comm, stream);
ncclAllReduce(grad_c, grad_c, n, ncclFloat, ncclSum, comm, stream);
ncclGroupEnd();

07Troubleshooting

Common issues and how to resolve them.

For deeper dives, the official documentation at docs.nvidia.com covers the full API and hardware-specific tuning guides.