NCCL User Guide

https://docs.nvidia.com/deeplearning/nccl/user-guide/docs
Admin
NCCL is a library providing optimized inter-GPU collective communication primitives that are...
Tokens:357,892
Snippets:2,059
Trust Score:10
Update:1 day ago
Show doc for...
Tokens:
Raw
### Host-side setup for device communication kernel

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

Shows the host-side steps to prepare for a device communication kernel: initialize communicator, allocate and register a symmetric window, create a device communicator with requirements, and launch the kernel. The example specifies 16 CTAs and 512 threads per CTA, with one LSA barrier per CTA.

```cpp
int main() {
  [...]
  NCCLCHECK(ncclCommInitRank(&comm, nranks, id, rank));

  /* Buffer initialization and window creation */
  char* buffer;
  size_t size = 256*1048576;
  NCCLCHECK(ncclMemAlloc((void**)&buffer, size));
  ncclWindow_t win;
  NCCLCHECK(ncclCommWindowRegister(comm, buffer, size, &win, NCCL_WIN_COLL_SYMMETRIC));

  /* Get device communicator */
  ncclDevComm devComm;
  ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
  int nCTAs = 16;
  reqs.lsaBarrierCount = nCTAs;
  NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));

  /* Launch user kernel */
  customKernel<<<nCTAs, 512>>>(devComm, win);
  [...]
}
```

--------------------------------

### Example output for NCCL_DEBUG=VERSION

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

Example log output showing the NCCL version printed at startup when NCCL_DEBUG=VERSION is set.

```text
NCCL version 2.30.3+cuda13.0
```

--------------------------------

### cftPutKernel

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

Example kernel demonstrating CFT put operation: creates an ncclCft object, gets CFT local endpoint info, issues a put, then submits and flushes the work.

```APIDOC
## cftPutKernel

### Description
Example kernel that stages data through shared memory, creates an `ncclCft` object, issues a put operation, and then submits and flushes the work.

### Method
Kernel launch (CUDA __global__ function)

### Endpoint
`cftPutKernel(ncclDevComm devComm, ncclWindow_t win)`

### Parameters
- **devComm** (ncclDevComm) - Required - Device communicator
- **win** (ncclWindow_t) - Required - Window object for CFT operations

### Request Example
```cpp
__global__ void cftPutKernel(ncclDevComm devComm, ncclWindow_t win) {
  __shared__ ncclCftSmem cftSmem;
  __shared__ alignas(16) char smem[128];

  ncclCoopCta coop;
  ncclCft<ncclCoopCta> cft{coop, cftSmem};

  ncclTeam cftTeam = ncclTeamCft(devComm);
  int peer = (cftTeam.rank + 1) % cftTeam.nRanks;

  ncclCftLeId leId;
  size_t leOffset;
  ncclGetCftLeInfo(win, 0, peer, cftTeam, devComm, &leId, &leOffset);

  cft.put(coop, leId, leOffset, smem, sizeof(smem));
  cft.submit(coop);
  cft.flush(coop);
}
```

### Notes
- CFT source and destination shared memory pointers must be 16-byte aligned.
- Byte count must be a multiple of 16.
- See Device API - CFT for full list of put, get, multicast, reduction, and pull-reduce helpers.
```

--------------------------------

### NCCL configuration file example

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

Example content for /etc/nccl.conf or ${NCCL_CONF_FILE} that sets NCCL_DEBUG to WARN and restricts socket interface to ens1f0.

```plaintext
NCCL_DEBUG=WARN
NCCL_SOCKET_IFNAME==ens1f0
```

--------------------------------

### Example output for NCCL_DEBUG=INFO

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

Example log output showing detailed information about NCCL operations when NCCL_DEBUG=INFO is set.

```text
node-01:3873285:3873285 [0] NCCL INFO Initialized NET plugin IB
node-01:3873285:3873285 [0] NCCL INFO Assigned NET plugin IB to comm
node-01:3873285:3873285 [0] NCCL INFO Assigned GIN plugin GIN_XXXX to comm
node-01:3873285:3873285 [0] NCCL INFO Assigned RMA plugin RMA_XXXX to comm
node-01:3873285:3873285 [0] NCCL INFO Using network IB
node-01:3873285:3873285 [0] NCCL INFO DMA-BUF is available on GPU device 0
```

--------------------------------

### Pure GIN AlltoAll: Host setup and kernel

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

Host setup configures ncclDevCommRequirements with worldGinBarrierCount and ginSignalCount set to the grid width (NCCL_DEVICE_CTA_COUNT) and ginConnectionType set to NCCL_GIN_CONNECTION_FULL. The kernel uses ncclGin, ncclGinBarrierSession, gin.put, gin.waitSignal, and gin.flush to perform the AlltoAll. Requires NCCL 2.30 or later for the host setup; GPU-initiated networking is available since NCCL 2.28.7.

```CUDA C++
// Grid width (CTAs). Must match reqs.worldGinBarrierCount and reqs.ginSignalCount.
#define NCCL_DEVICE_CTA_COUNT 16
#define NCCL_DEVICE_THREADS_PER_CTA 512

int main() {
  [...]
  ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
  reqs.worldGinBarrierCount = NCCL_DEVICE_CTA_COUNT;
  reqs.ginSignalCount = NCCL_DEVICE_CTA_COUNT;
  reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL;
  NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));
  [...]
}

template <typename T>
__global__ void PureGinAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset,
                                      ncclWindow_t recvwin, size_t recvoffset,
                                      size_t count, struct ncclDevComm devComm) {
  int ginContext = 0; // single context for simplicity
  unsigned int signalIndex = blockIdx.x;
  ncclGin gin { devComm, ginContext };
  uint64_t signalValue = gin.readSignal(signalIndex);

  ncclGinBarrierSession<ncclCoopCta> bar { ncclCoopCta(), gin, ncclTeamTagWorld(), blockIdx.x };
  bar.sync(ncclCoopCta(), cuda::memory_order_acquire, ncclGinFenceLevel::None);

  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  int nthreads = blockDim.x * gridDim.x;

  const size_t size = count * sizeof(T);
  for (int r = tid; r < devComm.nRanks; r += nthreads) {
    gin.put(ncclTeamWorld(devComm), r,
        recvwin, recvoffset + devComm.rank * size,
        sendwin, sendoffset + r * size,
        size, ncclGin_WeakSignalInc{signalIndex});
  }

  // Wait only on the CTA whose blockIdx.x (signalIndex) accumulates all puts to this rank.
  int receivingCta = (devComm.rank % nthreads) / blockDim.x;
  if (blockIdx.x == receivingCta)
    gin.waitSignal(ncclCoopCta(), signalIndex, signalValue + devComm.nRanks);

  gin.flush(ncclCoopCta());
  bar.sync(ncclCoopCta(), cuda::memory_order_release, ncclGinFenceLevel::None);
}
```

--------------------------------

### Device-side endpoint query example

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

Example showing how to obtain a CFT team and logical endpoint information on the device side using ncclTeamCft and ncclGetCftLeInfo.

```APIDOC
## Device-side endpoint query example

### Description
Demonstrates obtaining a CFT team and logical endpoint information for a given window, byte offset, and CFT-team peer rank on the device side.

### Functions Used
- `ncclTeamCft(devComm)`
- `ncclGetCftLeInfo(win, byteOffset, peerCft, cftTeam, devComm, &leId, &leOffset)`

### Parameters
- `devComm` (ncclDevComm) - Device communicator.
- `win` (ncclWindow_t) - Registered window.
- `byteOffset` (size_t) - Byte offset within the window.
- `peerCft` (int) - CFT-team peer rank.
- `cftTeam` (ncclTeam) - CFT team obtained from `ncclTeamCft`.
- `leId` (ncclCftLeId*) - Output logical endpoint ID.
- `leOffset` (size_t*) - Output logical endpoint offset.

### Code Example
```c
ncclTeam cftTeam = ncclTeamCft(devComm);
ncclCftLeId leId;
size_t leOffset;

ncclGetCftLeInfo(win, byteOffset, peerCft, cftTeam, devComm,
                 &leId, &leOffset);
```
```

--------------------------------

### Example output for NCCL_DEBUG=WARN

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

Example log output showing a warning message when NCCL_DEBUG=WARN is set.

```text
[2026-05-05 06:16:47] node-01:189884:189884 [3] plugin/net.cc:334 NCCL WARN Failed to initialize any NET plugin
```

--------------------------------

### OpenSM QoS configuration example

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

Example OpenSM configuration defining one low-priority and one high-priority Virtual Lane (VL) mapped to Service Levels (SL) 0 and 1. The high-priority SL gets a larger share of network bandwidth. This configuration is used to map NCCL traffic classes to SLs on InfiniBand networks.

```text
...
qos_max_vls 2
qos_high_limit 255
qos_vlarb_high 1:4
qos_vlarb_low 0:1,1:4
qos_sl2vl 0,1

max_op_vls 2
....
```

--------------------------------

### Complete single-process multi-device NCCL example

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

A complete working example of a single process managing multiple devices. It includes error checking macros, device buffer allocation, NCCL initialization with ncclCommInitAll, an AllReduce operation using the group API, synchronization, cleanup, and finalization. The example uses 4 devices and performs an AllReduce on 32M floats.

```c
#include <stdlib.h>
#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"

#define CUDACHECK(cmd) do {                         \
  cudaError_t err = cmd;                            \
  if (err != cudaSuccess) {                         \
    printf("Failed: Cuda error %s:%d '%s'\n",       \
        __FILE__,__LINE__,cudaGetErrorString(err)); \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


#define NCCLCHECK(cmd) do {                         \
  ncclResult_t res = cmd;                           \
  if (res != ncclSuccess) {                         \
    printf("Failed, NCCL error %s:%d '%s'\n",       \
        __FILE__,__LINE__,ncclGetErrorString(res)); \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


int main(int argc, char* argv[])
{
  ncclComm_t comms[4];


  //managing 4 devices
  int nDev = 4;
  int size = 32*1024*1024;
  int devs[4] = { 0, 1, 2, 3 };


  //allocating and initializing device buffers
  float** sendbuff = (float**)malloc(nDev * sizeof(float*));
  float** recvbuff = (float**)malloc(nDev * sizeof(float*));
  cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);


  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));
    CUDACHECK(cudaMalloc((void**)sendbuff + i, size * sizeof(float)));
    CUDACHECK(cudaMalloc((void**)recvbuff + i, size * sizeof(float)));
    CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float)));
    CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float)));
    CUDACHECK(cudaStreamCreate(s+i));
  }


  //initializing NCCL
  NCCLCHECK(ncclCommInitAll(comms, nDev, devs));


   //calling NCCL communication API. Group API is required when using
   //multiple devices per thread
  NCCLCHECK(ncclGroupStart());
  for (int i = 0; i < nDev; ++i)
    NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum,
        comms[i], s[i]));
  NCCLCHECK(ncclGroupEnd());


  //synchronizing on CUDA streams to wait for completion of NCCL operation
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));
    CUDACHECK(cudaStreamSynchronize(s[i]));
  }


  //free device buffers
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(i));
    CUDACHECK(cudaFree(sendbuff[i]));
    CUDACHECK(cudaFree(recvbuff[i]));
  }


  //finalizing NCCL
  for(int i = 0; i < nDev; ++i)
      ncclCommDestroy(comms[i]);


  printf("Success \n");
  return 0;
}
```

--------------------------------

### Example output for NCCL_DEBUG=TRACE with NCCL_DEBUG_SUBSYS=CALL

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

Example log output showing function call tracing when NCCL_DEBUG=TRACE and NCCL_DEBUG_SUBSYS=CALL are set.

```text
node-01:986207:986207 NCCL CALL ncclGroupStart()
node-01:986207:986207 NCCL CALL ncclSend(0,0x...,8388608,7,0,2,0x...,0x...)
node-01:986207:986207 NCCL CALL ncclRecv(0,0x...,8388608,7,0,0,0x...,0x...)
node-01:986207:986207 NCCL CALL ncclGroupEnd()
```

--------------------------------

### Device API – Host-Side Setup

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

Provides host-side setup functions for the device API, including host-accessible device pointer functions.

```APIDOC
## Device API – Host-Side Setup

### Description
Provides host-side setup functions for the device API, including host-accessible device pointer functions.

### Host-Side Setup
- Functions to initialize and configure the device API from the host.

### Host-Accessible Device Pointer Functions
- Functions that return device pointers accessible from the host.
```

--------------------------------

### Example output for NCCL_DEBUG=ATTN

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

Example log output showing an ATTN message when NCCL_DEBUG=ATTN is set.

```text
[2026-05-05 06:16:47] node-01:189884:189884 [3] plugin/net.cc:208 (ncclNetPluginInit) NCCL ATTN Failed to initialize NET plugin IB
```

--------------------------------

### ncclParamDumpAll example output

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/param.html

Example output of ncclParamDumpAll when NCCL_PARAM_DUMP_ALL is set to true. The output shows the key, type, attributes, current value, set_by, default, and accepted value for each parameter. The information is not sorted.

```text
NCCL_NO_CACHE (const char *) [Private, Cached] Comma-separated list of param keys to disable caching (or ALL)
    Current value=<unset> set_by=Default default=<unset>
    Accepted value: String

NCCL_PARAM_DUMP_ALL (bool) [Private] Print all parameters including private ones
    Current value=TRUE set_by=EnvPlugin default=FALSE
    Accepted value: Boolean: 1/T/TRUE or 0/F/FALSE
```

--------------------------------

### ncclParamGetAllParameterKeys example

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/param.html

Example of using ncclParamGetAllParameterKeys to retrieve all registered parameter keys. By default, only parameters published in Environment Variables are returned. Setting NCCL_PARAM_DUMP_ALL to true forces the function to return all parameters. The returned table is valid until the next call on the same thread.

```c
ncclResult_t ret = ncclSuccess;
const char** keyTable = NULL;
int tableLen = 0;

ret = ncclParamGetAllParameterKeys(&keyTable, &tableLen);
if (ret == ncclSuccess) {
  for (int i = 0; i < tableLen; ++i) {
    printf("Key: %s\n", keyTable[i]); /* prints all published parameter keys */
  }
}

putenv("NCCL_PARAM_DUMP_ALL=true");
ret = ncclParamGetAllParameterKeys(&keyTable, &tableLen);
if (ret == ncclSuccess) {
  for (int i = 0; i < tableLen; ++i) {
    printf("Key: %s\n", keyTable[i]); /* prints all parameter keys */
  }
}
```

--------------------------------

### MPI-based NCCL AllReduce Example

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

Complete example of using NCCL with MPI to perform an all-reduce operation. It initializes MPI, calculates a local rank based on host hashes, initializes NCCL, performs the all-reduce, and finalizes. Requires MPI and CUDA.

```C
#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"
#include "mpi.h"
#include <unistd.h>
#include <stdint.h>
#include <stdlib.h>


#define MPICHECK(cmd) do {                          \
  int e = cmd;                                      \
  if( e != MPI_SUCCESS ) {                          \
    printf("Failed: MPI error %s:%d '%d'\n",        \
        __FILE__,__LINE__, e);   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


#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)


static uint64_t getHash(const char* string, size_t n) {
  // Based on DJB2a, result = result * 33 ^ char
  uint64_t result = 5381;
  for (size_t c = 0; c < n; c++){
    result = ((result << 5) + result) ^ string[c];
  }
  return result;
}

/* Generate a hash of the unique identifying string for this host
 * that will be unique for both bare-metal and container instances
 * Equivalent of a hash of;
 *
 * $(hostname)$(cat /proc/sys/kernel/random/boot_id)
 *
 */
#define HOSTID_FILE "/proc/sys/kernel/random/boot_id"
static uint64_t getHostHash(const char* hostname) {
  char hostHash[1024];

  // Fall back is the hostname if something fails
  (void) strncpy(hostHash, hostname, sizeof(hostHash));
  int offset = strlen(hostHash);

  FILE *file = fopen(HOSTID_FILE, "r");
  if (file != NULL) {
    char *p;
    if (fscanf(file, "%ms", &p) == 1) {
        strncpy(hostHash+offset, p, sizeof(hostHash)-offset-1);
        free(p);
    }
  }
  fclose(file);

  // Make sure the string is terminated
  hostHash[sizeof(hostHash)-1]='\0';

  return getHash(hostHash, strlen(hostHash));
}

static void getHostName(char* hostname, int maxlen) {
  gethostname(hostname, maxlen);
  for (int i=0; i< maxlen; i++) {
    if (hostname[i] == '.') {
        hostname[i] = '\0';
        return;
    }
  }
}


int main(int argc, char* argv[])
{
  int size = 32*1024*1024;


  int myRank, nRanks, localRank = 0;


  //initializing MPI
  MPICHECK(MPI_Init(&argc, &argv));
  MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
  MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));


  //calculating localRank based on hostname which is used in selecting a GPU
  uint64_t hostHashs[nRanks];
  char hostname[1024];
  getHostName(hostname, 1024);
  hostHashs[myRank] = getHostHash(hostname);
  MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));
  for (int p=0; p<nRanks; p++) {
     if (p == myRank) break;
     if (hostHashs[p] == hostHashs[myRank]) localRank++;
  }


  ncclUniqueId id;
  ncclComm_t comm;
  float *sendbuff, *recvbuff;
  cudaStream_t s;


  //get NCCL unique ID at rank 0 and broadcast it to all others
  if (myRank == 0) ncclGetUniqueId(&id);
  MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));


  //picking a GPU based on localRank, allocate device buffers
  CUDACHECK(cudaSetDevice(localRank));
  CUDACHECK(cudaMalloc(&sendbuff, size * sizeof(float)));
  CUDACHECK(cudaMalloc(&recvbuff, size * sizeof(float)));
  CUDACHECK(cudaStreamCreate(&s));


  //initializing NCCL
  NCCLCHECK(ncclCommInitRank(&comm, nRanks, id, myRank));


  //communicating using NCCL
  NCCLCHECK(ncclAllReduce((const void*)sendbuff, (void*)recvbuff, size, ncclFloat, ncclSum,
        comm, s));


  //completing NCCL operation by synchronizing on the CUDA stream
  CUDACHECK(cudaStreamSynchronize(s));


  //free device buffers
  CUDACHECK(cudaFree(sendbuff));
  CUDACHECK(cudaFree(recvbuff));


  //finalizing NCCL
  ncclCommDestroy(comm);


  //finalizing MPI
  MPICHECK(MPI_Finalize());


  printf("[MPI Rank %d] Success \n", myRank);
  return 0;
}
```

--------------------------------

### ncclGetLsaDevicePointer usage example

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/device_setup.html

Example of calling ncclGetLsaDevicePointer to obtain a pointer to LSA peer 1's memory. The window must be registered with a communicator that supports LSA, and the target peer must be load/store accessible (P2P connectivity required). On success, ncclSuccess is returned and the pointer is stored in lsaPtr.

```c
void* lsaPtr;
ncclResult_t result = ncclGetLsaDevicePointer(window, 0, 1, &lsaPtr);
if (result == ncclSuccess) {
    // Use lsaPtr to access LSA peer 1's memory
}
```

--------------------------------

### Complete working example with multiple MPI processes and multiple devices per process

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

Full example showing how to use multiple GPUs per MPI rank. Includes error checking macros, host hash generation for local rank calculation, and initialization of two GPUs per process.

```cpp
#include <stdio.h>
#include "cuda_runtime.h"
#include "nccl.h"
#include "mpi.h"
#include <unistd.h>
#include <stdint.h>


#define MPICHECK(cmd) do {                          \
  int e = cmd;                                      \
  if( e != MPI_SUCCESS ) {                          \
    printf("Failed: MPI error %s:%d '%d'\n",        \
        __FILE__,__LINE__, e);   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


#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)


static uint64_t getHash(const char* string) {
  // Based on DJB2a, result = result * 33 ^ char
  uint64_t result = 5381;
  for (int c = 0; string[c] != '\0'; c++){
    result = ((result << 5) + result) ^ string[c];
  }
  return result;
}

/* Generate a hash of the unique identifying string for this host
 * that will be unique for both bare-metal and container instances
 * Equivalent of a hash of;
 *
 * $(hostname)$(cat /proc/sys/kernel/random/boot_id)
 *
 */
#define HOSTID_FILE "/proc/sys/kernel/random/boot_id"
static uint64_t getHostHash(const char* hostname) {
  char hostHash[1024];

  // Fall back is the hostname if something fails
  (void) strncpy(hostHash, hostname, sizeof(hostHash));
  int offset = strlen(hostHash);

  FILE *file = fopen(HOSTID_FILE, "r");
  if (file != NULL) {
    char *p;
    if (fscanf(file, "%ms", &p) == 1) {
        strncpy(hostHash+offset, p, sizeof(hostHash)-offset-1);
        free(p);
    }
  }
  fclose(file);

  // Make sure the string is terminated
  hostHash[sizeof(hostHash)-1]='\0';

  return getHash(hostHash, strlen(hostHash));
}

static void getHostName(char* hostname, int maxlen) {
  gethostname(hostname, maxlen);
  for (int i=0; i< maxlen; i++) {
    if (hostname[i] == '.') {
        hostname[i] = '\0';
        return;
    }
  }
}


int main(int argc, char* argv[])
{
  int size = 32*1024*1024;


  int myRank, nRanks, localRank = 0;


  //initializing MPI
  MPICHECK(MPI_Init(&argc, &argv));
  MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
  MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));


  //calculating localRank which is used in selecting a GPU
  uint64_t hostHashs[nRanks];
  char hostname[1024];
  getHostName(hostname, 1024);
  hostHashs[myRank] = getHostHash(hostname);
  MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));
  for (int p=0; p<nRanks; p++) {
     if (p == myRank) break;
     if (hostHashs[p] == hostHashs[myRank]) localRank++;
  }


  //each process is using two GPUs
  int nDev = 2;


  float** sendbuff = (float**)malloc(nDev * sizeof(float*));
  float** recvbuff = (float**)malloc(nDev * sizeof(float*));
  cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
```

--------------------------------

### Example RAS diagnostics report with no issues

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

Shows the output format of an on-demand report with no reported issues. Each line starts with host:pid NCCL DIAG and includes result lines and a completion line.

```text
node042:12345 NCCL DIAG === RAS Diagnostics ===
node042:12345 NCCL DIAG [OK]   GPU inventory: 8x NVIDIA H100 per node consistent across 8 ranks in comm 0x1234
node042:12345 NCCL DIAG [OK]   CUDA driver version: 13000 consistent across 8 ranks in comm 0x1234
node042:12345 NCCL DIAG [OK]   ECC: no uncorrected volatile errors across 8 ranks in comm 0x1234
node042:12345 NCCL DIAG [OK]   Xid/SXid: no events requiring immediate attention across 2 hosts
node042:12345 NCCL DIAG [OK]   NVLink: 18 links per GPU, all active at consistent speed across 8 ranks in comm 0x1234
node042:12345 NCCL DIAG [OK]   NCCL environment: NCCL_* env vars consistent across 8 ranks in comm 0x1234
node042:12345 NCCL DIAG [OK]   rdma_topo check: passed on 8/8 ranks in comm 0x1234
node042:12345 NCCL DIAG [OK]   NVIDIA graphics driver version: 610.47.03 consistent across 8 ranks in comm 0x1234
node042:12345 NCCL DIAG RAS diagnostics completed in 12.3 ms across 8 RAS peers
```

--------------------------------

### CFT Put Kernel Example

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

Shows a CUDA kernel that performs a CFT put operation. It creates an ncclCft object, gets CFT team info, and issues put, submit, and flush operations. Shared memory pointers must be 16-byte aligned and byte count must be a multiple of 16.

```CUDA C++
__global__ void cftPutKernel(ncclDevComm devComm, ncclWindow_t win) {
  __shared__ ncclCftSmem cftSmem;
  __shared__ alignas(16) char smem[128];

  ncclCoopCta coop;
  ncclCft<ncclCoopCta> cft{coop, cftSmem};

  ncclTeam cftTeam = ncclTeamCft(devComm);
  int peer = (cftTeam.rank + 1) % cftTeam.nRanks;

  ncclCftLeId leId;
  size_t leOffset;
  ncclGetCftLeInfo(win, 0, peer, cftTeam, devComm, &leId, &leOffset);

  cft.put(coop, leId, leOffset, smem, sizeof(smem));
  cft.submit(coop);
  cft.flush(coop);
}
```

--------------------------------

### Initialize MPI and Run Simple Reduction

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

Download, compile, and run the MPI example cpi.c to verify MPI initialization and reduction operations. Replace <number of processes> and <processes per node> as needed.

```bash
wget https://raw.githubusercontent.com/pmodels/mpich/main/examples/cpi.c
mpicc -o cpi cpi.c
mpirun -np <number of processes> -N <processes per node> ./cpi
```

--------------------------------

### ncclLsaCopy broadcast example with barrier synchronization

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/device_reducecopy.html

Example of using ncclLsaCopy to copy from a local source to symmetric buffers on all LSA peers, including barrier synchronization before and after the copy. The example shows the broadcast phase of AllGather.

```cpp
ncclCoopCta ctaCoop;
ncclLsaBarrierSession<ncclCoopCta> bar { ctaCoop, devComm, ncclTeamLsa(devComm), devComm.lsaBarrier, blockIdx.x };
bar.sync(ctaCoop, cuda::memory_order_acquire);

size_t srcOffset = [...];  // byte offset into symmetric send buffer on each peer
size_t dstOffset = [...];  // byte offset into symmetric recv buffer on each peer
T* srcPtr = (T*)ncclGetLocalPointer(sendwin, srcOffset);
ncclLsaCopy<T, ncclCoopCta, size_t>(ctaCoop, srcPtr, recvwin, dstOffset, count, devComm);

bar.sync(ctaCoop, cuda::memory_order_release);
```

--------------------------------

### Usage example for host-accessible device pointer functions

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

Shows how to allocate symmetric memory, register a window, obtain multimem, LSA, and peer pointers, and clean up. Note that multimem pointer may be nullptr if not supported, and device pointers are invalidated after window deregistration.

```cpp
int main() {
  [...]
  // Allocate symmetric memory buffer
  char* buffer;
  size_t size = 256 * 1024 * 1024;  // 256 MB buffer
  NCCLCHECK(ncclMemAlloc((void**)&buffer, size));

  // Create window with the allocated buffer
  ncclWindow_t win;
  NCCLCHECK(ncclCommWindowRegister(comm, buffer, size, &win, NCCL_WIN_COLL_SYMMETRIC));

  // Get host-accessible pointers
  void* multimemPtr;
  void* lsaPtr;
  void* peerPtr;

  // Get multimem pointer (returns nullptr if multimem not supported)
  NCCLCHECK(ncclGetLsaMultimemDevicePointer(win, 0, &multimemPtr));
  if (multimemPtr == nullptr) {
      // Multimem not available, use fallback
  }

  // Get LSA pointer for peer 1
  NCCLCHECK(ncclGetLsaDevicePointer(win, 0, 1, &lsaPtr));

  // Get peer pointer for world rank 2
  NCCLCHECK(ncclGetPeerDevicePointer(win, 0, 2, &peerPtr));

  // Use pointers in custom kernels or legacy code
  customKernel<<<nCTAs, 256>>>(multimemPtr, lsaPtr, peerPtr);

  // Cleanup
  NCCLCHECK(ncclCommWindowDeregister(comm, &win));
  // Device pointers are invalidated after window deregistration
  NCCLCHECK(ncclMemFree(buffer));
  [...]
}
```

--------------------------------

### Example output for INIT subsystem

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

Typical NCCL INFO output when filtering by INIT subsystem, showing initialization and plugin assignment messages.

```text
node-01:1895902:1895913 [0] NCCL INFO Initialized NET plugin IB
node-01:1895902:1895913 [0] NCCL INFO Assigned NET plugin IB to comm
node-01:1895902:1895913 [0] NCCL INFO Assigned GIN plugin GIN_XXXX to comm
node-01:1895902:1895913 [0] NCCL INFO Assigned RMA plugin RMA_XXXX to comm
node-01:1895902:1895913 [0] NCCL INFO Using network IB
node-01:1895902:1895913 [0] NCCL INFO DMA-BUF is available on GPU device 0
node-01:1895902:1895913 [0] NCCL INFO [Rank 0] ncclCommInitRankConfig comm 0x... rank 0 nranks 2 cudaDev 0 nvmlDev 0 busId XXXX commId 0x... - Init START
node-01:1895902:1895913 [0] NCCL INFO ncclTopoGetCpuAffinity: Affinity for GPU 0 is X,X,X,X,X,X,X,X. (GPU affinity = X,X,X,X,X,X,X,X ; CPU affinity = X-X).
```

--------------------------------

### Pure GIN AlltoAll Kernel

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

Host setup and device kernel for a pure GIN AlltoAll collective. The host creates an ncclDevComm with GIN-specific resources and registers symmetric memory windows, then launches the kernel that performs the collective using GIN.

```APIDOC
## PureGinAlltoAllKernel

### Description
Performs an AlltoAll collective using GPU-initiated networking (GIN). All peer data moves over the network. The kernel uses GIN put operations with weak signals and barriers for synchronization.

### Host Setup

- Define grid width (CTAs) and threads per CTA.
- Initialize `ncclDevCommRequirements` with `NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER`.
- Set `worldGinBarrierCount` and `ginSignalCount` to the number of CTAs.
- Set `ginConnectionType` to `NCCL_GIN_CONNECTION_FULL`.
- Call `ncclDevCommCreate(comm, &reqs, &devComm)`.

### Kernel Parameters

- `sendwin` (ncclWindow_t) - Send window.
- `sendoffset` (size_t) - Offset into send window.
- `recvwin` (ncclWindow_t) - Receive window.
- `recvoffset` (size_t) - Offset into receive window.
- `count` (size_t) - Number of elements per rank.
- `devComm` (struct ncclDevComm) - Device communicator.

### Kernel Implementation

- Create a `ncclGin` object with a single context.
- Read a signal value using `gin.readSignal(signalIndex)`.
- Create a `ncclGinBarrierSession` for cooperative CTA synchronization.
- Perform `gin.put` operations to send data to all peers.
- Wait on the signal for the receiving CTA.
- Flush and synchronize with release semantics.

### Example Code

```cpp
// Grid width (CTAs). Must match reqs.worldGinBarrierCount and reqs.ginSignalCount.
#define NCCL_DEVICE_CTA_COUNT 16
#define NCCL_DEVICE_THREADS_PER_CTA 512

int main() {
  [...]
  ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
  reqs.worldGinBarrierCount = NCCL_DEVICE_CTA_COUNT;
  reqs.ginSignalCount = NCCL_DEVICE_CTA_COUNT;
  reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL;
  NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));
  [...]
}

template <typename T>
__global__ void PureGinAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset,
                                      ncclWindow_t recvwin, size_t recvoffset,
                                      size_t count, struct ncclDevComm devComm) {
  int ginContext = 0; // single context for simplicity
  unsigned int signalIndex = blockIdx.x;
  ncclGin gin { devComm, ginContext };
  uint64_t signalValue = gin.readSignal(signalIndex);

  ncclGinBarrierSession<ncclCoopCta> bar { ncclCoopCta(), gin, ncclTeamTagWorld(), blockIdx.x };
  bar.sync(ncclCoopCta(), cuda::memory_order_acquire, ncclGinFenceLevel::None);

  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  int nthreads = blockDim.x * gridDim.x;

  const size_t size = count * sizeof(T);
  for (int r = tid; r < devComm.nRanks; r += nthreads) {
    gin.put(ncclTeamWorld(devComm), r,
        recvwin, recvoffset + devComm.rank * size,
        sendwin, sendoffset + r * size,
        size, ncclGin_WeakSignalInc{signalIndex});
  }

  // Wait only on the CTA whose blockIdx.x (signalIndex) accumulates all puts to this rank.
  int receivingCta = (devComm.rank % nthreads) / blockDim.x;
  if (blockIdx.x == receivingCta)
    gin.waitSignal(ncclCoopCta(), signalIndex, signalValue + devComm.nRanks);

  gin.flush(ncclCoopCta());
  bar.sync(ncclCoopCta(), cuda::memory_order_release, ncclGinFenceLevel::None);
}
```

### Notes

- GPU-initiated networking is available since NCCL 2.28.7.
- Version-specific host and kernel changes for older NCCL builds are summarized under Compatibility adjustments.
- `ncclDevCommCreate()` fails if GIN cannot be provided.
```

--------------------------------

### ncclLsaCopyTma lambda-based custom layout example

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/device_reducecopy.html

Example of using ncclLsaCopyTma with a destination lambda for custom layouts. The lambda returns a T* for each destination index, and each destination pointer must be 16-byte aligned. The example uses ncclTeamLsa and ncclGetLsaPointer to obtain the destination pointers.

```cpp
ncclTeam team = ncclTeamLsa(devComm);
auto dstLambda = [=] __device__ (int i) -> T* {
  return (T*)ncclGetLsaPointer(recvwin, dstOffset, i);
};
ncclLsaCopyTma<T, ncclCoopCta, decltype(dstLambda), size_t, smemBytes>(
  ctaCoop, srcPtr, dstLambda, team.nRanks, count, smemScratch);
```

--------------------------------

### ncclSimInfo_t.NCCL_SIM_INFO_INITIALIZER (C macro)

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

C macro for initializing ncclSimInfo_t.

```APIDOC
## ncclSimInfo_t.NCCL_SIM_INFO_INITIALIZER (C macro)

### Description
C macro for initializing ncclSimInfo_t.

### Type
C macro

### Defined in
api/types.html
```

--------------------------------

### NcclDataType.numpy_dtype

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/nccl4py/types.html

Property that returns the equivalent NumPy dtype. For BFLOAT16 and float8 variants, ml-dtypes must be installed. Raises NcclInvalid if ml-dtypes is required but not installed.

```APIDOC
## NcclDataType.numpy_dtype

### Description
Returns the equivalent NumPy dtype. For `BFLOAT16` and the float8 variants, `ml-dtypes` must be installed.

### Property
`numpy_dtype: numpy.dtype`

### Returns
- **numpy.dtype** - NumPy dtype corresponding to this NCCL data type.

### Raises
- **NcclInvalid** - If `ml-dtypes` is required but not installed.

```

--------------------------------

### NCCL_OOB_NET_IFNAME example: prefix match

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

Uses the first interface found in the list of 'mlx5_1', 'mlx5_10', 'mlx5_11', etc. Requires NCCL_NET='IB' and NCCL_OOB_NET_ENABLE=1.

```bash
NCCL_NET="IB" NCCL_OOB_NET_ENABLE=1 NCCL_OOB_NET_IFNAME="mlx5_1"
```

--------------------------------

### Run nvloom pairwise and multicast benchmarks

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

Commands to benchmark local GPU communication with nvloom. Use for intra-node performance testing.

```bash
srun/mpirun -n <number of processes> ./nvloom -s pairwise --sizeMin 1M --sizeMax 4G
srun/mpirun -n <number of processes> ./nvloom -s multicast --sizeMin 1M --sizeMax 4G
```

--------------------------------

### Nonblocking Group Operation Example

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

Example code demonstrating how to issue multiple NCCL operations in a group with nonblocking communicators and handle the asynchronous completion of kernel issuance.

```APIDOC
## Nonblocking Group Operation

### Description
This example shows how to use `ncclGroupStart()` and `ncclGroupEnd()` with nonblocking communicators. When `ncclGroupEnd()` returns `ncclInProgress`, the user must wait for the communicator state to become `ncclSuccess` before synchronizing CUDA streams.

### Code Example
```c
ncclGroupStart();
  for (int g=0; g<ngpus; g++) {
    ncclAllReduce(sendbuffs[g]+offsets[i], recvbuffs[g]+offsets[i], counts[i], datatype[i], comms[g], streams[g]);
  }
ret = ncclGroupEnd();
if (ret == ncclInProgress) {
   for (int g=0; g<ngpus; g++) {
     do {
       ncclCommGetAsyncError(comms[g], &state);
     } while (state == ncclInProgress);
   }
} else if (ret == ncclSuccess) {
   /* Successfully issued */
   printf("NCCL kernel issue succeeded\n");
} else {
   /* Errors happen */
   reportErrorAndRestart();
}

for (int g=0; g<ngpus; g++) {
  cudaStreamSynchronize(streams[g]);
}
```

### Related Functions
- `ncclCommInitRankConfig()`
- `ncclCommGetAsyncError()`
```
Tokens:
Raw