组调用

组函数 (ncclGroupStart/ncclGroupEnd) 可用于将多个调用合并为一个。这对于三个目的来说是必要的：从一个线程管理多个 GPU（以避免死锁）、聚合通信操作以提高性能，或合并多个发送/接收点对点操作（请参阅 点对点通信 部分）。所有三种用法可以组合在一起，但有一个例外：对 ncclCommInitRank() 的调用不能与其他调用合并。

从单个线程管理多个 GPU

当单个线程管理多个设备时，必须使用组语义。这是因为每个 NCCL 调用都可能需要阻塞，等待其他线程/rank 到达，然后才能有效地在给定流上发布 NCCL 操作。因此，像下面显示的在多个设备上的简单循环可能会在第一个调用上阻塞，等待其他调用。

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

为了定义这些调用是同一集体操作的一部分，应使用 ncclGroupStart 和 ncclGroupEnd。

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

这将告诉 NCCL 将 ncclGroupStart 和 ncclGroupEnd 之间的所有调用视为对多个设备的单个调用。

注意：当在组内调用时，流操作（如 ncclAllReduce）可能在没有将操作排队到流上的情况下返回。因此，只有在 ncclGroupEnd 返回后才能调用像 cudaStreamSynchronize 这样的流操作。

当一个线程管理多个设备时，也必须使用组调用来创建通信器。

ncclGroupStart();
for (int i=0; i<nLocalDevs; i++) {
  cudaSetDevice(device[i]);
  ncclCommInitRank(comms+i, nranks, commId, rank[i]);
}
ncclGroupEnd();

注意：与 NCCL 1.x 相反，在组内每次 NCCL 通信调用之前无需设置 CUDA 设备，但当在组内调用 ncclCommInitRank 时仍然需要。

相关链接

• ncclGroupStart()
• ncclGroupEnd()

聚合操作 (2.2 及更高版本)

组语义还可以用于在单个 NCCL 启动中执行多个集体操作。这对于减少启动开销（换句话说，延迟）非常有用，因为它仅在多个操作中发生一次。初始化函数不能与其他的初始化函数聚合，也不能与通信函数聚合。

集体操作的聚合可以通过在 ncclGroupStart / ncclGroupEnd 部分中进行多次 NCCL 调用来简单地完成。

在以下示例中，我们一起启动一个 broadcast 和两个 allReduce 操作作为一个单独的 NCCL 启动。

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

允许将聚合与多 GPU 启动相结合，并在组启动中使用不同的通信器，如“从单个线程管理多个 GPU”主题中所示。当组合多 GPU 启动和聚合时，ncclGroupStart 和 ncclGroupEnd 可以使用一次或在每个级别使用。以下示例对来自不同层和多个 CUDA 设备上的 allReduce 操作进行分组。

ncclGroupStart();
for (int i=0; i<nlayers; i++) {
  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]);
  }
  ncclGroupEnd();
}
ncclGroupEnd();

注意：NCCL 操作仅在最后一次调用 ncclGroupEnd 期间作为一个整体启动。for 循环内的 ncclGroupStart 和 ncclGroupEnd 调用不是必需的，也不执行任何操作。

相关链接

• ncclGroupStart()
• ncclGroupEnd()

非阻塞组操作

如果通过 ncclCommInitRankConfig 将通信器标记为非阻塞，则组函数将相应地变为异步。在这种情况下，如果用户在一个组中发出多个 NCCL 操作，从 ncclGroupEnd() 返回可能并不意味着 NCCL 通信内核已发布到 CUDA 流。如果 ncclGroupEnd() 返回 ncclSuccess，则表示 NCCL 内核已发布到流；如果返回 ncclInProgress，则表示 NCCL 内核正在后台发布到流。用户有责任确保通信器的状态变为 ncclSuccess，然后再调用相关的 CUDA 调用（例如 cudaStreamSynchronize）。

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

相关链接

• ncclCommInitRankConfig()
• ncclCommGetAsyncError()