### NCCL and MPI Example Setup

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/examples.rst.txt

Includes necessary headers and defines macros for checking MPI, CUDA, and NCCL calls. Also includes utility functions for hashing hostnames and generating unique IDs.

```C
#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*));

```

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

### Complete MPI and NCCL Integration Example

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

A full implementation demonstrating MPI initialization, local rank calculation based on hostnames, and NCCL communicator setup.

```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));
```

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

### Get Host-Accessible Device Pointers with NCCL

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

Example demonstrating how to allocate memory, register a window, and obtain host-accessible pointers to device memory using NCCL functions. Ensure pointers are used within the lifetime of the window and communicator. Multimem support should be checked.

```c++
int main() {
  [...] // Assume comm is initialized
  // 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));
  [...] // Assume comm is destroyed
}
```

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

### NCCL PutSignal and WaitSignal Setup

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/p2p.rst.txt

Demonstrates the setup for one-sided communication using ncclPutSignal and ncclWaitSignal. This includes memory allocation and window registration for remote memory access (RMA).

```C
// Allocate symmetric memory for RMA operations
 void *sendbuff, *recvbuff;
 NCCLCHECK(ncclMemAlloc((void**)&sendbuff, size));
 NCCLCHECK(ncclMemAlloc((void**)&recvbuff, size));

 // Register buffers as symmetric windows
 ncclWindow_t sendWindow, recvWindow;
 NCCLCHECK(ncclCommWindowRegister(comm, sendbuff, size, &sendWindow, NCCL_WIN_COLL_SYMMETRIC));
 NCCLCHECK(ncclCommWindowRegister(comm, recvbuff, size, &recvWindow, NCCL_WIN_COLL_SYMMETRIC));

 int peer = (rank == 0) ? 1 : 0;
 ncclWaitSignalDesc_t waitDesc = {.opCnt = 1, .peer = peer, .sigIdx = 0, .ctx = ctx};
```

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

### Complete NCCL Example: Single Process, Multiple Devices

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/examples.rst.txt

A full C example demonstrating NCCL initialization, communication (AllReduce), synchronization, and cleanup for 4 devices within a single process. Includes CUDA and NCCL error checking macros.

```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;
 }
```

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

### Configure NCCL via Configuration File

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/env.rst.txt

Example content for /etc/nccl.conf or ${NCCL_CONF_FILE} to set system-wide environment variables.

```C
NCCL_DEBUG=WARN
NCCL_SOCKET_IFNAME==ens1f0
```

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

### Complete NCCL Example with MPI and CUDA

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/examples.rst.txt

A comprehensive example demonstrating NCCL operations within a multi-process MPI environment, utilizing CUDA for device management and synchronization.

```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;
 }

```

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

### Multimem Device Kernel Setup

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

Host-side setup for enabling multimem support by setting `lsaMultimem` in `ncclDevCommRequirements`. The `lsaBarrierCount` should match the number of CTAs.

```C++
int main() {
  [...] 
  reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
  int nCTAs = 16;
  reqs.lsaBarrierCount = nCTAs;
  reqs.lsaMultimem = true;
  NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));
  [...] 
}

```

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

### Initialize Communicator with Config

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

Example of initializing a communicator with specific configuration parameters such as blocking mode and CTA limits.

```c
ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
config.blocking = 0;
config.minCTAs = 4;
config.maxCTAs = 16;
config.cgaClusterSize = 2;
config.netName = "Socket";
CHECK(ncclCommInitRankConfig(&comm, nranks, id, rank, &config));
do {
  CHECK(ncclCommGetAsyncError(comm, &state));
  // Handle outside events, timeouts, progress, ...
} while(state == ncclInProgress);
```

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

### ncclLsaCopy Usage Example

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

Example demonstrating the broadcast phase of an AllGather operation using ncclLsaCopy with barrier synchronization.

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

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);
```

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

### Infiniband Subnet Manager QoS Configuration

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

Example configuration for defining Virtual Lane mapping and priority levels in the subnet manager.

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

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

### Register Buffers into NCCL Window for Communication

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/bufferreg.rst.txt

Example demonstrating how to register local buffers into an NCCL window for low-latency and high-bandwidth communication. This requires buffers allocated via VMM-based allocators or `ncclMemAlloc`.

```c
void* src;
void* dst;
ncclWindow_t src_win;
ncclWindow_t dst_win;

CHECK(ncclMemAlloc(&src, src_size));
CHECK(ncclMemAlloc(&dst, dst_size));
```

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

### NCCL AllReduce Loop With Grouping

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

This example shows how to use ncclGroupStart and ncclGroupEnd to group multiple NCCL calls for managing multiple devices from one thread. This prevents deadlocks by ensuring calls are treated as a single collective operation.

```c
ncclGroupStart();
for (int i=0; i<nLocalDevs; i++) {
  ncclAllReduce(..., comm[i], stream[i]);
}
ncclGroupEnd();

```

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

### ncclGroupStart

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

Starts a group call to prevent blocking during subsequent NCCL operations.

```APIDOC
## ncclGroupStart

### Description
Start a group call. All subsequent calls to NCCL until ncclGroupEnd will not block due to inter-CPU synchronization.

### Method
C Function

### Endpoint
ncclGroupStart()

### Response
- **ncclResult_t** - Returns the result of the operation.
```

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

### Enable Zero-CTA Optimization with NCCL

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

Configure NCCL to use Zero-CTA optimization by setting the CTAPolicy to NCCL_CTA_POLICY_ZERO. This example demonstrates buffer allocation, symmetric window registration, performing an AllGather collective, and subsequent cleanup.

```c
ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
// NCCL_CTA_POLICY_ZERO to enable zero-CTA optimization whenever possible
config.CTAPolicy = NCCL_CTA_POLICY_ZERO;
CHECK(ncclCommInitRankConfig(&comm, nranks, id, rank, &config));

void* src;
void* dst;
ncclWindow_t src_win;
ncclWindow_t dst_win;

CHECK(ncclMemAlloc(&src, src_size));
CHECK(ncclMemAlloc(&dst, dst_size));

// Register the buffers into NCCL symmetric window
CHECK(ncclCommWindowRegister(comm, src, src_size, &src_win, NCCL_WIN_COLL_SYMMETRIC));
CHECK(ncclCommWindowRegister(comm, dst, dst_size, &dst_win, NCCL_WIN_COLL_SYMMETRIC));

CHECK(ncclAllGather(src, dst, 1, ncclInt8, comm, stream));
CHECK(cudaStreamSynchronize(stream));

CHECK(ncclCommWindowDeregister(comm, src_win));
CHECK(ncclCommWindowDeregister(comm, dst_win));

CHECK(ncclMemFree(src));
CHECK(ncclMemFree(dst));

```

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

### Compile and Run MPI Pi Calculation Example

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting.rst.txt

Download, compile, and run a simple MPI application (pi calculation) to ensure MPI initialization and basic communication are working correctly. This requires MPI compiler (mpicc) and launcher (mpirun).

```shell
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
```

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

### RAS Job Summary Output

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting/ras.rst.txt

Example output showing the job configuration summary for a healthy job.

```text
  Job summary
  ===========

    Nodes  Processes         GPUs  Processes     GPUs
  (total)   per node  per process    (total)  (total)
        4          8            1         32       32
```

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

### Aggregating Broadcast and AllReduce Operations

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

This example shows how to aggregate multiple collective operations, including a broadcast and two allReduce calls, within a single ncclGroupStart/ncclGroupEnd section to reduce launch overhead.

```c
ncclGroupStart();
ncclBroadcast(sendbuff1, recvbuff1, count1, datatype, root, comm, stream);
ncclAllReduce(sendbuff2, recvbuff2, count2, datatype, comm, stream);
ncclAllReduce(sendbuff3, recvbuff3, count3, datatype, comm, stream);
ncclGroupEnd();

```

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

### Initialize NCCL Communicator with MPI

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/examples.rst.txt

When each process or thread manages at most one GPU, ncclCommInitRank can be used. This example shows communicator creation within an MPI context, where each MPI rank uses one device.

```C
int myRank, nRanks;
MPI_Comm_rank(MPI_COMM_WORLD, &myRank);
MPI_Comm_size(MPI_COMM_WORLD, &nRanks);
```

```C
ncclUniqueId id;
 if (myRank == 0) ncclGetUniqueId(&id);
 MPI_Bcast(&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);
```

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

### Initialize NCCL Communicator and Device Communication

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/deviceapi.rst.txt

This C code demonstrates the host-side setup for device-initiated communication. It initializes an NCCL communicator, allocates a buffer, registers a symmetric memory window, and creates a device communicator with specified requirements before launching a custom CUDA kernel.

```C
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);
    [...]
  }
```

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

### Ensure Operation Order in Groups

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/groups.rst.txt

Users must guarantee the same issuing order of operations among different GPUs, even when using different communicators within a group. This example shows the correct ordering for multiple operations on two independent communicators.

```C
ncclGroupStart();
ncclBroadcast(sendbuff1, recvbuff1, count1, datatype, root, comm0, stream);
ncclAllReduce(sendbuff2, recvbuff2, count2, datatype, comm0, stream);
ncclAllReduce(sendbuff3, recvbuff3, count3, datatype, comm0, stream);
ncclAllReduce(sendbuff4, recvbuff4, count4, datatype, comm1, stream);
ncclGroupEnd();
```

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

### Configure Host and Implement GIN AlltoAll Kernel

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

Demonstrates host-side communicator requirements for GIN and a device-side AlltoAll kernel using ncclGin for one-sided puts and signal-based synchronization.

```cpp
int main() {
  [...]
  reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
  int nCTAs = 1;
  reqs.railGinBarrierCount = nCTAs;
  reqs.ginSignalCount = 1;
  NCCLCHECK(ncclDevCommCreate(comm, &reqs, &devComm));
  [...]
}

template <typename T>
__global__ void ginAlltoAllKernel(ncclDevComm devComm, ncclWindow_t win,
                                  size_t inputOffset, size_t outputOffset, size_t count) {
  int ginContext = 0;
  ncclGinSignal_t signalIndex = 0;
  ncclGin gin { devComm, ginContext };
  uint64_t signalValue = gin.readSignal(signalIndex);

  ncclGinBarrierSession<ncclCoopCta> bar { ncclCoopCta(), gin, ncclTeamWorld(devComm),
                                           devComm.railGinBarrier, blockIdx.x };
  bar.sync(ncclCoopCta(), cuda::memory_order_relaxed, ncclGinFenceLevel::Relaxed);

  const int rank = devComm.rank, nRanks = devComm.nRanks;
  const int tid = threadIdx.x + blockIdx.x * blockDim.x;
  const int nThreads = blockDim.x * gridDim.x;

  const size_t size = count * sizeof(T);
  for (int peer = tid; peer < nRanks; peer += nThreads) {
    gin.put(ncclTeamWorld(devComm), peer, win, outputOffset + rank * size,
            win, inputOffset + peer * size, size, ncclGin_SignalInc{signalIndex});
  }

  gin.waitSignal(ncclCoopCta(), signalIndex, signalValue + nRanks);
  gin.flush(ncclCoopCta());
}
```

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

### Initialize Device Communicator and Launch Kernel

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

Demonstrates the host-side workflow for registering symmetric memory windows, creating a device communicator, and launching a custom CUDA kernel.

```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);
  [...]
}
```

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

### Get local LSA pointer

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

Returns a load-store accessible pointer to the current device's memory buffer within a window. Offset is in bytes. This is a shortcut for getting the local peer's pointer.

```cpp
void *ncclGetLocalPointer(ncclWindow_t w, size_t offset)
```

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

### ReduceSumCopy API Example

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

This example demonstrates the usage of the ncclLsaReduceSumCopy function, showcasing common template parameters like element type (T), cooperation level (Coop), integer count type (IntCount), and unroll factor (UNROLL).

```APIDOC
## ncclLsaReduceSumCopy

### Description
Fused reduce-then-copy building block.

### Method
Not Applicable (Function Signature)

### Endpoint
Not Applicable (Function Signature)

### Parameters
#### Template Parameters
- **T** (type) - Required - Element type. Supported types: `float`, `double`, `half`, `int8`, `int16`, `int32`, `int64`; and `__nv_bfloat16`, `__nv_fp8_e4m3`, `__nv_fp8_e5m2` where available.
- **Coop** (type) - Required - Cooperation level (e.g., `ncclCoopCta` or `ncclCoopThread`).
- **IntCount** (type) - Required - Type for the element count (e.g., `unsigned int` or `size_t`).
- **UNROLL** (integer) - Optional; default `4*16/sizeof(T)` - Represents the tradeoff between register usage and achievable peak bandwidth.

#### Function Parameters
- **ctaCoop** (type) - Description not provided.
- **sendwin** (type) - Description not provided.
- **srcOffset** (type) - Required - Byte offset into symmetric send buffer on each peer.
- **recvwin** (type) - Description not provided.
- **dstOffset** (type) - Required - Byte offset into symmetric receive buffer on each peer.
- **count** (type) - Description not provided.
- **team** (type) - Description not provided.

### Request Example
```c
size_t srcOffset = [...];  // byte offset into symmetric send buffer on each peer
size_t dstOffset = [...];  // byte offset into symmetric recv buffer on each peer
ncclLsaReduceSumCopy<float, ncclCoopCta, size_t, 16>(ctaCoop, sendwin, srcOffset, recvwin, dstOffset, count, team);
```

### Response
#### Success Response (200)
Not Applicable (Function Signature)

#### Response Example
Not Applicable (Function Signature)
```

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

### Initialize Communicator with Options

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/communicators.rst.txt

Demonstrates configuring a communicator with specific parameters like blocking behavior and CTA limits using ncclCommInitRankConfig.

```C
  ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
  config.blocking = 0;
  config.minCTAs = 4;
  config.maxCTAs = 16;
  config.cgaClusterSize = 2;
  config.netName = "Socket";
  CHECK(ncclCommInitRankConfig(&comm, nranks, id, rank, &config));
  do {
    CHECK(ncclCommGetAsyncError(comm, &state));
    // Handle outside events, timeouts, progress, ...
  } while(state == ncclInProgress);
```

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

### Split NCCL Communicators

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/communicators.rst.txt

Examples of using ncclCommSplit to duplicate, partition, or filter communicators.

```C
 int rank;
 ncclCommUserRank(comm, &rank);
 ncclCommSplit(comm, 0, rank, &newcomm, NULL);
```

```C
 int rank, nranks;
 ncclCommUserRank(comm, &rank);
 ncclCommCount(comm, &nranks);
 ncclCommSplit(comm, rank/(nranks/2), rank%(nranks/2), &newcomm, NULL);
```

```C
 int rank;
 ncclCommUserRank(comm, &rank);
 ncclCommSplit(comm, rank<2 ? 0 : NCCL_SPLIT_NOCOLOR, rank, &newcomm, NULL);
```

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

### Register and Use NCCL Windows for Communication

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

Demonstrates allocating memory with ncclMemAlloc, registering buffers with ncclCommWindowRegister using symmetric flags, and performing an AllGather operation.

```cpp
void* src;
void* dst;
ncclWindow_t src_win;
ncclWindow_t dst_win;

CHECK(ncclMemAlloc(&src, src_size));
CHECK(ncclMemAlloc(&dst, dst_size));
// Passing NCCL_WIN_COLL_SYMMETRIC requires users to provide the symmetric buffers among all ranks in collectives.
// Every rank needs to call ncclCommWindowRegister to register its buffers.
CHECK(ncclCommWindowRegister(comm, src, src_size, &src_win, NCCL_WIN_COLL_SYMMETRIC));
CHECK(ncclCommWindowRegister(comm, dst, dst_size, &dst_win, NCCL_WIN_COLL_SYMMETRIC));
// Use the registered buffers for communication to enable symmetric communication benefits.
// In this example, every rank has 0x1000 offset and 0x2000 offset from the head address of
// src and dst respectively, which satisfies the symmetric buffer requirement.
CHECK(ncclAllGather((uint8_t*)src + 0x1000, (uint8_t*)dst + 0x2000, 1, ncclInt8, comm, stream));
CHECK(cudaStreamSynchronize(stream));

CHECK(ncclCommWindowDeregister(comm, src_win));
CHECK(ncclCommWindowDeregister(comm, dst_win));

CHECK(ncclMemFree(src));
CHECK(ncclMemFree(dst));
```

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

### RAS Warning Output

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting/ras.rst.txt

Example output showing warning details when a mismatch is detected.

```text
  Warnings
  ========

  #0-0 (27a079b828ff1a75) MISMATCH
    Communicator ranks have different collective operation counts
```

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

### RAS Communicator Status Output

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting/ras.rst.txt

Example output showing the status of communicators in a healthy job.

```text
  Communicators... (0.00s)
  =============

  Group     Comms     Nodes     Ranks     Ranks     Ranks    Status  Errors
      #  in group  per comm  per node  per comm  in group
      0         8         4         1         4        32   RUNNING      OK
```

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

### Initialize and split a non-blocking communicator

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/communicators.rst.txt

Demonstrates setting up a non-blocking communicator to allow for safe aborts and restarts during initialization or splitting.

```C
bool globalFlag;
bool abortFlag = false;
ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
/* set communicator as nonblocking */
config.blocking = 0;
CHECK(ncclCommInitRankConfig(&comm, nRanks, id, myRank, &config));
do {
  CHECK(ncclCommGetAsyncError(comm, &state));
} while(state == ncclInProgress && checkTimeout() != true);

if (checkTimeout() == true || state != ncclSuccess) abortFlag = true;

/* sync abortFlag among all healthy ranks. */
reportErrorGlobally(abortFlag, &globalFlag);

if (globalFlag) {
  /* time is out or initialization failed: every rank needs to abort and restart. */
ncclCommAbort(comm);
  /* restart NCCL; this is a user implemented function, it might include
   * resource cleanup and ncclCommInitRankConfig() to create new communicators. */
  restartNCCL(&comm);
}

/* nonblocking communicator split. */
CHECK(ncclCommSplit(comm, color, key, &childComm, &config));
do {
  CHECK(ncclCommGetAsyncError(comm, &state));
} while(state == ncclInProgress && checkTimeout() != true);

if (checkTimeout() == true || state != ncclSuccess) abortFlag = true;

/* sync abortFlag among all healthy ranks. */
reportErrorGlobally(abortFlag, &globalFlag);

if (globalFlag) {
  ncclCommAbort(comm);
  /* if chilComm is not NCCL_COMM_NULL, user should abort child communicator
   * here as well for resource reclamation. */
  if (childComm != NCCL_COMM_NULL) ncclCommAbort(childComm);
  restartNCCL(&comm);
}
/* application workload */
```

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

### Retrieve MPI Information

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/examples.rst.txt

Get the rank and size of the MPI communicator. This is a prerequisite for distributed operations.

```C
 int myRank, nRanks;
 MPI_Comm_rank(MPI_COMM_WORLD, &myRank);
 MPI_Comm_size(MPI_COMM_WORLD, &nRanks);
```

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

### RAS Mismatch Status Output

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting/ras.rst.txt

Example output indicating a communicator rank mismatch during an active query.

```text
  Group     Comms     Nodes     Ranks     Ranks     Ranks    Status  Errors
      #  in group  per comm  per node  per comm  in group
      0         1         4         8        32        32   RUNNING  MISMATCH
```

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

### Implement In-Place AllReduce Kernel with LSA

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/usage/deviceapi.rst.txt

A basic device kernel demonstrating an in-place AllReduce using LSA memory load/store instructions. It requires manual synchronization via ncclLsaBarrierSession and iteration over communicator ranks.

```C
    ncclLsaBarrierSession<ncclCoopCta> bar { ncclCoopCta(), devComm, ncclTeamTagLsa(), blockIdx.x };
    bar.sync(ncclCoopCta(), cuda::memory_order_relaxed);

    const int rank = devComm.lsaRank, nRanks = devComm.lsaSize;
    const int globalTid = threadIdx.x + blockDim.x * (rank + blockIdx.x * nRanks);
    const int globalNthreads = blockDim.x * gridDim.x * nRanks;

    for (size_t o = globalTid; o < count; o += globalNthreads) {
      T v = 0;
      for (int peer = 0; peer < nRanks; peer++) {
        T* inputPtr = (T*)ncclGetLsaPointer(win, offset, peer);
        v += inputPtr[o];
      }
      for (int peer = 0; peer < nRanks; peer++) {
        T* outputPtr = (T*)ncclGetLsaPointer(win, offset, peer);
        outputPtr[o] = v;
      }
    }

    bar.sync(ncclCoopCta(), cuda::memory_order_release);
  }
```

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

### Run p2pBandwidthLatencyTest

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting.rst.txt

Use this CUDA sample to test GPU-to-GPU communication. Ensure it completes successfully and reports good performance.

```shell
cd cuda-samples/Samples/5_Domain_Specific/p2pBandwidthLatencyTest
make
./p2pBandwidthLatencyTest
```

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

### Debug Shared Memory Errors

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting.rst.txt

Example of a warning message generated by NCCL when shared memory allocation fails.

```shell
NCCL WARN Error: failed to extend /dev/shm/nccl-03v824 to 4194660 bytes
```

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

### Identify Communicator Errors and Timeouts

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/troubleshooting/ras.rst.txt

Example output demonstrating how RAS reports missing communicator data and communication timeouts.

```text
  Communicators... (2.05s)
  =============

  Group     Comms     Nodes     Ranks     Ranks     Ranks    Status  Errors
      #  in group  per comm  per node  per comm  in group
      0         1         4       7-8        32        32   RUNNING  INCOMPLETE

  Errors
  ======

  INCOMPLETE
    Missing communicator data from 1 job process
    Process 3487984 on node 172.16.64.213 managing GPU 5

  #0-0 (cf264af53edbe986) INCOMPLETE
    Missing communicator data from 1 rank
    The missing rank: 21

  Warnings
  ========

  TIMEOUT
    Encountered 2 communication timeouts while gathering communicator data
```

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

### Get Multimem LSA Pointer

Source: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/_sources/api/device_memory.rst.txt

Returns a multicast memory pointer associated with a window and device communicator. Availability is hardware-dependent.

```cpp
void* ncclGetLsaMultimemPointer(ncclWindow_t w, size_t offset, ncclDevComm const& devComm)
```