创建 Communicator

当创建 communicator 时，必须为属于该 communicator 的 n 个 CUDA 设备中的每一个分配一个 0 到 n-1 之间的唯一 rank。不支持将同一个 CUDA 设备多次用作同一个 NCCL communicator 的不同 rank，并且可能导致挂起。

给定 rank 到 CUDA 设备的静态映射，ncclCommInitRank()、ncclCommInitRankConfig() 和 ncclCommInitAll() 函数将创建 communicator 对象，每个 communicator 对象都与固定的 rank 和 CUDA 设备相关联。这些对象随后将用于启动通信操作。

在调用 ncclCommInitRank() 之前，您需要首先创建一个唯一对象，该对象将由所有进程和线程使用，以同步并了解它们是同一 communicator 的一部分。这通过调用 ncclGetUniqueId() 函数完成。

ncclGetUniqueId() 函数返回一个 ID，该 ID 必须使用任何 CPU 通信系统广播到所有参与的线程和进程，例如，将 ID 指针传递给多个线程，或使用 MPI 或另一个并行环境（例如，使用套接字）将其广播到其他进程。

您还可以调用 ncclCommInitAll 操作，在单个进程中一次创建 n 个 communicator 对象。由于它仅限于单个进程，因此此函数不允许节点间通信。 ncclCommInitAll 等效于调用 ncclGetUniqueId 和 ncclCommInitRank 的组合。

以下示例代码是 ncclCommInitAll 的简化实现。

ncclResult_t ncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist) {
  ncclUniqueId Id;
  ncclGetUniqueId(&Id);
  ncclGroupStart();
  for (int i=0; i<ndev; i++) {
    cudaSetDevice(devlist[i]);
    ncclCommInitRank(comm+i, ndev, Id, i);
  }
  ncclGroupEnd();
}

相关链接

• ncclCommInitAll()
• ncclGetUniqueId()
• ncclCommInitRank()

使用选项创建 communicator

ncclCommInitRankConfig() 函数允许使用特定选项创建 NCCL communicator。

NCCL 支持的 config 参数在此处列出：ncclConfig_t。

例如，“blocking”可以设置为 0，以要求 NCCL 永远不要在任何 NCCL 调用中阻塞，同时也可以设置其他 config 参数，以更精确地定义 communicator 行为。下面显示了一个简单的示例代码

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

相关链接：ncclCommGetAsyncError()

使用多个 ncclUniqueIds 创建 communicator

ncclCommInitRankScalable() 函数允许使用多个 ncclUniqueIds 创建 NCCL communicator。所有 NCCL ranks 都必须提供相同的 ncclUniqueIds 数组（相同的 ncclUniqueIds，并且顺序相同）。为了获得最佳性能，我们建议在 NCCL ranks 之间尽可能均匀地分配 ncclUniqueIds。

在内部，NCCL ranks 主要与单个 ncclUniqueId 通信。因此，为了获得最佳结果，我们建议在 ranks 之间均匀分配 ncclUniqueIds。

以下函数可用于确定 NCCL rank 是否应创建 ncclUniqueIds

bool rankHasRoot(const int rank, const int nRanks, const int nIds) {
  const int rmr = nRanks % nIds;
  const int rpr = nRanks / nIds;
  const int rlim = rmr * (rpr+1);
  if (rank < rlim) {
    return !(rank % (rpr + 1));
  } else {
    return !((rank - rlim) % rpr);
  }
}

例如，如果要将 3 个 ncclUniqueIds 分配给 7 个 NCCL ranks，则第一个 ncclUniqueId 将与 ranks 0-2 关联，而其他 ncclUniqueIds 将与 ranks 3-4 和 5-6 关联。因此，此函数在 rank 0、3 和 5 上将返回 true，否则返回 false。

注意：只有第一个 ncclUniqueId 将用于创建 communicator 哈希 ID，该哈希 ID 用于在日志文件和重放工具中标识 communicator。

创建更多 communicators

ncclCommSplit 函数可用于基于现有 communicator 创建 communicators。这允许将现有 communicator 拆分为多个子分区、复制现有 communicator，甚至创建 rank 较少的单个 communicator。

原始 communicator 中的所有 ranks 都需要调用 ncclCommSplit 函数。如果某些 ranks 不属于任何子组，它们仍然需要使用颜色 NCCL_SPLIT_NOCOLOR 调用 ncclCommSplit。

新创建的 communicators 将继承父 communicator 配置（例如，非阻塞）。如果父 communicator 在非阻塞模式下运行，则可以通过在父 communicator 上调用 ncclCommAbort，然后在任何新返回的 communicator 上调用 ncclCommAbort 来停止 ncclCommSplit 操作。这是因为在两个 communicators 中的任何一个上进行操作期间都可能发生挂起。

以下代码复制了现有 communicator

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

这将 communicator 拆分为两半

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

这将创建一个仅包含前 2 个 ranks 的 communicator

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

相关链接

• ncclCommSplit()

并发使用多个 NCCL communicators

使用多个 NCCL communicators 需要仔细同步，否则可能导致死锁。

NCCL kernels 是阻塞的（等待数据到达），任何 CUDA 操作都可能导致设备同步，这意味着它将等待所有 NCCL kernels 完成。由于 NCCL 操作本身执行 CUDA 调用，因此这可能会很快导致死锁。

因此，应在不同的 epoch 中使用不同 communicators 上的操作，并使用锁定机制，并且应用程序应确保跨 ranks 以相同的顺序提交操作。

启动多个通信操作（在不同的 streams 上）可能有效，前提是它们可以容纳在 GPU 中，但如果 NCCL 每个操作使用更多 CUDA blocks，或者如果 NCCL collectives 内部使用的某些调用要执行设备同步（例如，动态分配一些 CUDA 内存），则随时可能中断。

完成 communicator

ncclCommFinalize 将将 communicator 从 ncclSuccess 状态转换为 ncclInProgress 状态，开始在后台完成所有操作，并与其他 ranks 同步，这些 ranks 可能正在使用资源与其他 ranks 进行通信。与 communicator 关联的所有未完成操作和网络相关资源将通过 ncclCommFinalize 刷新和释放。一旦所有 NCCL 操作完成，communicator 将转换为 ncclSuccess 状态。用户可以使用 ncclCommGetAsyncError 查询该状态。如果 communicator 标记为非阻塞，则此操作是非阻塞的；否则，它是阻塞的。

相关链接：ncclCommFinalize()

销毁 communicator

一旦 communicator 完成，下一步是释放所有资源，包括 communicator 本身。可以使用 ncclCommDestroy 销毁与 communicator 关联的本地资源。如果在调用 ncclCommDestroy 时 communicator 的状态为 ncclSuccess，则保证调用是非阻塞的；否则，ncclCommDestroy 可能会阻塞。在所有情况下，ncclCommDestroy 调用都将释放 communicator 的资源并返回，并且在 ncclCommDestroy 返回后，不应再访问该 communicator。

相关链接：ncclCommDestroy()

错误处理和 communicator 中止

所有 NCCL 调用都返回一个 NCCL 错误代码，该代码在下表中进行了总结。如果 NCCL 调用返回的错误代码与 ncclSuccess 和 ncclInternalError 不同，并且如果 NCCL_DEBUG 设置为 WARN，则 NCCL 将打印一条人类可读的消息，解释发生了什么情况。如果 NCCL_DEBUG 设置为 INFO，则 NCCL 还将打印导致错误的调用堆栈。此消息旨在帮助用户解决问题。

下表总结了应如何理解和处理不同的错误。以下各节详细解释了每种情况。

NCCL 错误

错误	描述	解决方案	错误处理	Group 行为
ncclSuccess	无错误	无	无	无
ncclUnhandledCudaError	CUDA 调用期间出错 (1)	CUDA 配置/使用 (1)	Communicator 中止 (5)	全局 (6)
ncclSystemError	系统调用期间出错 (1)	系统配置/使用 (1)	Communicator 中止 (5)	全局 (6)
ncclInternalError	NCCL 内部错误 (2)	在 NCCL 中修复 (2)	Communicator 中止 (5)	全局 (6)
ncclInvalidArgument	NCCL 调用的参数无效 (3)	在应用程序中修复 (3)	无 (3)	单个 (3)
ncclInvalidUsage	NCCL 调用的使用无效 (4)	在应用程序中修复 (4)	Communicator 中止 (5)	全局 (6)
ncclInProgress	NCCL 调用仍在进行中	使用 ncclCommGetAsyncError 轮询完成情况	无	无

(1) ncclUnhandledCudaError 和 ncclSystemError 表示 NCCL 对外部组件的调用失败，这导致 NCCL 操作失败。错误消息应解释用户应查看并尝试修复哪个组件，可能需要系统管理员的帮助。

(2) ncclInternalError 表示 NCCL 错误。它可能不会报告 NCCL_DEBUG=WARN 的消息，因为它需要在 NCCL 源代码中进行修复。 NCCL_DEBUG=INFO 将打印导致错误的回溯。

(3) ncclInvalidArgument 表示参数值不正确，例如 NULL 指针或超出范围的值。当返回此错误时，NCCL 调用无效。组状态保持不变，communicator 仍然正常运行。应用程序可以调用 ncclCommAbort 或继续，就好像调用未发生一样。对于在组内发生的调用，将立即返回此错误，并应用于该特定的 NCCL 调用。它不会由 ncclGroupEnd 返回，因为 ncclGroupEnd 不接受任何参数。

(4) 当动态条件导致失败时，将返回 ncclInvalidUsage，这表示 NCCL API 的使用不正确。

(5) 这些错误对于 communicator 是致命的。要恢复，应用程序需要在 communicator 上调用 ncclCommAbort 并重新创建它。

(6) 组内操作的动态错误始终由 ncclGroupEnd 报告，并应用于组内的所有操作，这些操作可能已完成或未完成。应用程序必须在组内的所有 communicators 上调用 ncclCommAbort。

异步错误和错误处理

一些通信错误，特别是网络错误，通过 ncclCommGetAsyncError 函数报告。遇到异步错误的操作通常不会进行，并且永远不会完成。当发生异步错误时，应中止操作并使用 ncclCommAbort 销毁 communicator。在等待 NCCL 操作完成时，应用程序应调用 ncclCommGetAsyncError，并在发生错误时销毁 communicator。

以下代码显示了如何等待 NCCL 操作并轮询异步错误，而不是使用 cudaStreamSynchronize。

int ncclStreamSynchronize(cudaStream_t stream, ncclComm_t comm) {
  cudaError_t cudaErr;
  ncclResult_t ncclErr, ncclAsyncErr;
  while (1) {
   cudaErr = cudaStreamQuery(stream);
   if (cudaErr == cudaSuccess)
     return 0;

   if (cudaErr != cudaErrorNotReady) {
     printf("CUDA Error : cudaStreamQuery returned %d\n", cudaErr);
     return 1;
   }

   ncclErr = ncclCommGetAsyncError(comm, &ncclAsyncErr);
   if (ncclErr != ncclSuccess) {
     printf("NCCL Error : ncclCommGetAsyncError returned %d\n", ncclErr);
     return 1;
   }

   if (ncclAsyncErr != ncclSuccess) {
     // An asynchronous error happened. Stop the operation and destroy
     // the communicator
     ncclErr = ncclCommAbort(comm);
     if (ncclErr != ncclSuccess)
       printf("NCCL Error : ncclCommDestroy returned %d\n", ncclErr);
     // Caller may abort or try to create a new communicator.
     return 2;
   }

   // We might want to let other threads (including NCCL threads) use the CPU.
   sched_yield();
  }
}

相关链接

• ncclCommGetAsyncError()
• ncclCommAbort()

容错

NCCL 提供了一组功能，允许应用程序从致命错误（例如网络故障、节点故障或进程故障）中恢复。当发生此类错误时，应用程序应能够对 communicator 调用 ncclCommAbort 以释放所有资源，然后创建一个新的 communicator 以继续。所有 NCCL 调用都可以是非阻塞的，以确保可以在任何时候（在初始化、通信或完成 communicator 时）调用 ncclCommAbort。

要正确中止，当 communicator 中的任何 rank 失败时（例如，由于分段错误），所有其他 ranks 都需要调用 ncclCommAbort 以中止它们自己的 NCCL communicator。用户可以实现方法来决定何时以及是否中止 communicators 并重新启动 NCCL 操作。这是一个示例，展示了如何以非阻塞方式初始化和拆分 communicator，从而允许在任何时候中止

bool globalFlag;
bool abortFlag = false;
ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
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 */

用户需要提供 checkTimeout 函数，以确定应用程序应等待 NCCL 初始化的最长时间；同样，用户可以应用其他方法来检测错误，除了超时函数之外。类似的方法也可以应用于 NCCL 完成。