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/