NCCL使用/调用步骤源码解读(单设备单进程为例):

步骤总结:

通过MPI获取本机rank(可理解为进程)数量localrank,用于rank绑定GPU;

rank0获取NCCL通信组ID,并通过MPI_Bcast广播给其它rank;

借助MPI获取的这些信息NCCL完成初始化,并进行集合通信。

核心步骤:

1、初试化和启动MPI通信。

2、计算主机名的哈希值,并MPI_allgather通信使得每个rank(进程)都获取其它rank的哈希值。

3、根据获取的哈希值,比较得到该rank所在的主机参与通信的rank总数localrank(哈希值相同的rank在同一主机上)。(哈希值就是主机名,其实可以用主机名来获取主机上参与通信的总rank数,只是主机命名五花八门,哈希值更容易比较)

4、rank0上获取NCCL的唯一ID,并MPI_Bcast广播给其它rank。(这个唯一的ID是用来标识通信组,因此所有通信组中的rank有相同的ID)

5、基于localrank绑定GPU,并分配发送接收缓冲区,创建CUDA流。

6、初始化NCCL通信器。

7、nccl allreduce通信。同步CUDA流,确保通信完成。

8、释放缓冲区。

9、销毁通信器。

10、终止MPI环境

视频教程

哈哈哈,感觉这期没必要做视频,后续有必要视频教程的在B站更新

 1.1 NCCL官网案例源码详解One Device per Process or Thread_哔哩哔哩_bilibili

对应源码

源码来源NCCL官方文档 Example 2: One Device per Process or Thread: Examples — NCCL 2.21.5 documentation (nvidia.com)

下期预告:

NCCL源码解读2:ncclGetUniqueId(&id)函数源码解读