点对点通信

(自 NCCL 2.7 起) 点对点通信可用于表达 ranks 之间的任何通信模式。任何点对点通信都需要两次 NCCL 调用:在一个 rank 上调用 ncclSend(),并在另一个 rank 上调用相应的 ncclRecv(),且计数和数据类型相同。

可以结合 ncclGroupStart()ncclGroupEnd() 将针对不同对等方的多个 ncclSend()ncclRecv() 调用融合在一起,以形成更复杂的通信模式,例如一对多(分散)、多对一(收集)、全对全或与 N 维空间中的邻居通信。

组内的点对点调用将阻塞,直到该组调用完成,但组内的调用可以被视为独立进行,因此永远不应相互阻塞。因此,合并需要并发进行的调用以避免死锁非常重要。

以下是并行应用程序使用的经典点对点通信模式的一些示例。NCCL 语义允许每 rank 使用不同的大小、数据类型和缓冲区的所有变体。

Sendrecv (发送接收)

在 MPI 术语中,sendrecv 操作是指两个 ranks 交换数据,同时发送和接收。这可以通过将 ncclSend 和 ncclRecv 调用合并为一个来完成

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

一对多 (scatter) (分散)

来自 root rank 的一对多操作可以通过将所有发送和接收操作合并到一个组中来表达

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

多对一 (gather) (收集)

类似地,到 root rank 的多对一操作将以这种方式实现

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