点对点通信

(自NCCL 2.7起) 点对点通信可以用于在进程之间表达任意的通信模式。

任何点对点通信都需要两个NCCL调用：在一个进程上调用 ncclSend，在另一个进程上调用相应的 ncclRecv，传递相同的数量和数据类型。

可以使用 ncclGroupStart和ncclGroupEnd将多个针对不同进程的 ncclSend和ncclRecv调用合并成更复杂的通信模式，例如单对多（散射）、多对一（聚集）、多对多或N维空间中与邻居的通信。

在一个组内的点对点调用将会阻塞，直到该组内的所有调用完成，但是组内的调用可以认为是并行进行的，因此它们不应该阻塞彼此。因此，重要的是合并需要并行进行的调用，以避免死锁。

以下是并行应用程序中使用的几个经典点对点通信模式的示例。NCCL语义允许每个进程以不同的大小、数据类型和缓冲区来实现这些变体。

Sendrecv

在MPI术语中，sendrecv操作是指两个进程同时交换数据，既发送又接收。可以通过将ncclSend和ncclRecv调用合并为一个来实现：

ncclGroupStart();
ncclSend(sendbuff, sendcount, sendtype, peer, comm, stream);
ncclRecv(recvbuff, recvcount, recvtype, peer, comm, stream);
ncclGroupEnd();

单对多 (散射)

从一个root进程进行的单对多操作可以通过将所有发送和接收操作合并在一个组内来表达：

ncclGroupStart();
if (rank == root) {
  for (int r=0; r<nranks; r++)
    ncclSend(sendbuff[r], size, type, r, comm, stream);
}
ncclRecv(recvbuff, size, type, root, comm, stream);
ncclGroupEnd();

多对一 (聚集)

类似地，对于一个root进程的多对一操作可以通过以下方式实现：

ncclGroupStart();
if (rank == root) {
  for (int r=0; r<nranks; r++)
    ncclRecv(recvbuff[r], size, type, r, comm, stream);
}
ncclSend(sendbuff, size, type, root, comm, stream);
ncclGroupEnd();

多对多

多对多操作将会是一个向所有进程发送/接收的循环合并：

ncclGroupStart();
for (int r=0; r<nranks; r++) {
  ncclSend(sendbuff[r], sendcount, sendtype, r, comm, stream);
  ncclRecv(recvbuff[r], recvcount, recvtype, r, comm, stream);
}
ncclGroupEnd();

邻居交换

最后，可以通过以下方式与N维空间中的邻居交换数据：

ncclGroupStart();
for (int d=0; d<ndims; d++) {
  ncclSend(sendbuff[d], sendcount, sendtype, next[d], comm, stream);
  ncclRecv(recvbuff[d], recvcount, recvtype, prev[d], comm, stream);
}
ncclGroupEnd();