NCCL（NVIDIA聚合通信库）概述及示例

NCCL全称NVIDIA Collective Communication Library，是NVIDIA推出的用于GPU之间通信的库，包括聚合通信和点对点通信两部分。

NCCL的聚合通信接口采用异步调用的方式（通过stream来实现），接口调用后立即返回，用户需要调用cudaStreamSynchronize等待聚合通信完成。NCCL聚合通信包括以下接口：

• AllReduce

image

对所有GPU上的目标数据进行reduce操作（例如sum，max），并将结果写到所有的GPU上 。

使用示例：

    // NCCL使用组操作保证多个聚合通信操作同时下发到stream（在ncclGroupEnd时）
    ncclGroupStart()；
    // 单进程多GPU方式，nDev为GPU个数，对于每个GPU都调用1次AllReduce，
    //每个GPU上的sendbuf数据都参与reduce操作，然后将结果写到每个GPU上的recvbuf中
    for (int i = 0; i < nDev; ++i) {
        ncclAllReduce(sendbuf[i],recvbuf[i],size,ncclFloat, ncclSum, 0,comms[i], s[i]);
    }
    ncclGroupEnd()；
    // 调用stream同步方法等待每个GPU的broadcast操作完成
    for (int i = 0; i < nDev; ++i) {
        cudaSetDevice(i);
        cudaStreamSynchronize(s[i]);
    }

• Broadcast

image

将root GPU上的数据发送到所有GPU上，其中root可以由用户进行指定。

使用示例：

    // NCCL使用组操作保证多个聚合通信操作同时下发到stream（在ncclGroupEnd时）
    ncclGroupStart()；
    // 单进程多GPU方式，从root节点发送sendbuff中数据到所有GPU上的recvbuff中
    for (int i = 0; i < nDev; ++i) {
        ncclBroadcast(sendbuff[i], recvbuff[i], count, datatype, root, comm[i], stream[i]);
    }
    ncclGroupEnd()；
    // 调用stream同步方法等待每个GPU的broadcast操作完成
    for (int i = 0; i < nDev; ++i) {
        cudaSetDevice(i);
        cudaStreamSynchronize(s[i]);
    }

• Reduce

image

对所有GPU的数据进行reduce操作（例如sum，max），并将结果写到用户指定的GPU上。

• AllGather

image

集合所有GPU上的数据，并将集合后的数据写到所有的GPU上。

• ReduceScatter

image

对所有GPU上的数据进行reduce操作（如sum max），然后将结果切分到所有的GPU上。

从2.7版本起，NCCL实现了点对点通信，点对点通信也是异步方式，调用之后立即返回，用户需要调用cudaStreamSynchronize等待聚合通信完成。点对点通信包括send和recive 2个接口，send和recive需要配对使用，否则可能会导致阻塞。点对点通信给了用户更大的自由度，用户可以通过send/recive的组合可以实现诸如Scatter、Gather、All-to-All等多种多样的聚合通信操作。

• Scatter实现：

image

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

• Gather实现：

image

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

• All-to-All实现：

image

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

NCCL官方入口：https://docs.nvidia.com/deeplearning/nccl/