Using the NCCL Library: A Practical Guide

Using the NCCL Library: A Practical Guide

NCCL (NVIDIA Collective Communications Library) is a library of multi-GPU collective communication primitives that are topology-aware and can be easily integrated into applications. Here’s a practical guide to using NCCL with example code.

Prerequisites

• NVIDIA GPUs with CUDA support
• NCCL library installed (comes with CUDA or can be installed separately)
• Basic understanding of MPI or multi-GPU programming

Basic NCCL Concepts

NCCL provides optimized implementations of:

• AllReduce
• Broadcast
• Reduce
• AllGather
• ReduceScatter
• Point-to-point send/receive

Practical Demo Code

Here’s a complete example demonstrating NCCL AllReduce across multiple GPUs:

#include <nccl.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <mpi.h>#define CUDACHECK(cmd) do {                         \cudaError_t e = cmd;                              \if( e != cudaSuccess ) {                          \printf("Failed: Cuda error %s:%d '%s'\n",       \__FILE__,__LINE__,cudaGetErrorString(e));   \exit(EXIT_FAILURE);                             \}                                                 \
} while(0)#define NCCLCHECK(cmd) do {                         \ncclResult_t r = cmd;                             \if (r!= ncclSuccess) {                            \printf("Failed, NCCL error %s:%d '%s'\n",       \__FILE__,__LINE__,ncclGetErrorString(r));   \exit(EXIT_FAILURE);                             \}                                                 \
} while(0)int main(int argc, char* argv[]) {// Initialize MPIMPI_Init(&argc, &argv);int mpi_rank, mpi_size;MPI_Comm_rank(MPI_COMM_WORLD, &mpi_rank);MPI_Comm_size(MPI_COMM_WORLD, &mpi_size);// Make sure exactly 2 GPUs are available per node for this demoint local_rank = -1;char* local_rank_str = getenv("LOCAL_RANK");if (local_rank_str != NULL) {local_rank = atoi(local_rank_str);} else {// Fallback: use MPI rank if LOCAL_RANK not setlocal_rank = mpi_rank;}// Assign GPU to this processCUDACHECK(cudaSetDevice(local_rank));// NCCL variablesncclUniqueId id;ncclComm_t comm;float *sendbuff, *recvbuff;const size_t count = 32 * 1024 * 1024; // 32M elements// Allocate buffersCUDACHECK(cudaMalloc(&sendbuff, count * sizeof(float)));CUDACHECK(cudaMalloc(&recvbuff, count * sizeof(float)));// Initialize send buffer with some valuesfloat* hostBuff = (float*)malloc(count * sizeof(float));for (size_t i = 0; i < count; i++) {hostBuff[i] = 1.0f * (local_rank + 1); // Different value per rank}CUDACHECK(cudaMemcpy(sendbuff, hostBuff, count * sizeof(float), cudaMemcpyHostToDevice));free(hostBuff);// Generate NCCL unique ID at rank 0 and broadcast it to all othersif (mpi_rank == 0) {ncclGetUniqueId(&id);}MPI_Bcast(&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);// Initialize NCCL communicatorNCCLCHECK(ncclCommInitRank(&comm, mpi_size, id, mpi_rank));// Perform AllReduce operationNCCLCHECK(ncclAllReduce((const void*)sendbuff, (void*)recvbuff, count, ncclFloat, ncclSum, comm, cudaStreamDefault));// Synchronize to make sure the operation is completeCUDACHECK(cudaStreamSynchronize(cudaStreamDefault));// Verify results (only rank 0 for demo purposes)if (mpi_rank == 0) {float* verifyBuff = (float*)malloc(count * sizeof(float));CUDACHECK(cudaMemcpy(verifyBuff, recvbuff, count * sizeof(float), cudaMemcpyDeviceToHost));// Expected sum is (1 + 2 + ... + mpi_size) for each elementfloat expected = mpi_size * (mpi_size + 1) / 2.0f;for (size_t i = 0; i < 10; i++) { // Check first 10 elementsif (verifyBuff[i] != expected) {printf("ERROR: Expected %f, got %f at index %zu\n", expected, verifyBuff[i], i);break;}}printf("Rank %d: NCCL AllReduce test completed successfully\n", mpi_rank);free(verifyBuff);}// CleanupCUDACHECK(cudaFree(sendbuff));CUDACHECK(cudaFree(recvbuff));NCCLCHECK(ncclCommDestroy(comm));MPI_Finalize();return 0;
}

Compilation and Execution

To compile and run this code:

1. Compile with:

nvcc -o nccl_demo nccl_demo.cu -lnccl -lmpi

2. Run with MPI (example for 4 processes):

mpirun -np 4 ./nccl_demo

Key Steps Explained

1. Initialization:

   • Initialize MPI to get rank and size
   • Determine local rank for GPU assignment
   • Set CUDA device based on local rank

2. Memory Allocation:

   • Allocate device buffers for sending and receiving data
   • Initialize send buffer with rank-specific values

3. NCCL Setup:

   • Generate a unique NCCL ID at rank 0 and broadcast it
   • Initialize NCCL communicator with this ID

4. Collective Operation:

   • Perform AllReduce operation (sum in this case)
   • Synchronize to ensure completion

5. Verification:

   • Check results (only on rank 0 for simplicity)
   • Expected result is the sum of all ranks’ values

6. Cleanup:

   • Free device memory
   • Destroy NCCL communicator
   • Finalize MPI

Common Patterns

1. Point-to-Point Communication

// Send from rank 0 to rank 1
if (rank == 0) {NCCLCHECK(ncclSend(sendbuff, count, ncclFloat, 1, comm, cudaStreamDefault));
} 
else if (rank == 1) {NCCLCHECK(ncclRecv(recvbuff, count, ncclFloat, 0, comm, cudaStreamDefault));
}

2. Broadcast

// Broadcast from rank 0 to all others
NCCLCHECK(ncclBroadcast(sendbuff, recvbuff, count, ncclFloat, 0, comm, cudaStreamDefault));

3. Using Streams

cudaStream_t stream;
cudaStreamCreate(&stream);// Perform operation with custom stream
NCCLCHECK(ncclAllReduce(sendbuff, recvbuff, count, ncclFloat, ncclSum, comm, stream));// Synchronize the specific stream
cudaStreamSynchronize(stream);
cudaStreamDestroy(stream);

Best Practices

1. Topology Awareness: NCCL automatically optimizes for the system topology. Ensure proper GPU affinity.

2. Stream Management: Use separate streams for computation and communication to overlap them.

3. Buffer Reuse: Reuse communication buffers when possible to avoid allocation overhead.

4. Error Checking: Always check NCCL and CUDA return codes as shown in the example.

5. Multi-Node: For multi-node setups, ensure proper network configuration (InfiniBand, NVLink, etc.).

This example provides a foundation for using NCCL in your applications. The library offers significant performance benefits for multi-GPU communication compared to naive implementations.