上节课我们Nvidia NCCL的官方案例中
NCCL源码解读1:NCCL使用/调用案例 Example : One Device per Process or Thread-CSDN博客
讲到了,第四步是:4、ncclGetUniqueId():rank0上获取NCCL的唯一ID,并MPI_Bcast广播给其它rank。(这个唯一的ID是用来标识通信组,因此所有通信组中的rank有相同的ID)
这节课我们讲一下ncclGetUniqueId()源码和其调用的ncclInit()和bootstrapGetUniqueId()
视频教程请见:
1.2 NCCL通信初始化源码详解 ncclGetUniqueId()中ncclInit()、bootstrapGetUniqueId()_哔哩哔哩_bilibili
一、ncclGetUniqueId()源码详解
ncclGetUniqueId()核心逻辑
1、调用ncclInit()进行nccl库初始化(初步的初始化,基本函数可调用,但还难以通信)。
2、调用bootstrapGetUniqueId()函数来获取一个唯一的ID。
核心就这两步,其它与主逻辑无关。so easy!
源码速递
源码位置:nccl-master\src\init.cc
ncclResult_t ncclGetUniqueId(ncclUniqueId* out) {
1、确保NCCL库已经初始化。如果尚未初始化,则进行初始化。
NCCLCHECK(ncclInit());
// 检查传入的out指针是否为非空。
NCCLCHECK(PtrCheck(out, "GetUniqueId", "out"));
2、调用bootstrapGetUniqueId函数来获取一个唯一的ID,并将这个ID存储在传入的out指针所指向的内存位置。
ncclResult_t res = bootstrapGetUniqueId((struct ncclBootstrapHandle*)out);
// TRACE_CALL是一个用于日志记录或跟踪的宏。
TRACE_CALL("ncclGetUniqueId(0x%llx)", (unsigned long long)hashUniqueId(*out));
// 返回bootstrapGetUniqueId函数的结果。表示操作是否成功
return res;
}
二、ncclInit()源码详解
ncclInit()核心逻辑
1、initEnv();//初始化环境设置
A、setEnvFile(confFilePath);//根据配置文件初始化设置
2、initGdrCopy()//初始化 GPU Direct RDMA (GDR)
A、ncclGdrInit()//初始化与GPU Direct RDMA Copy相关的库和驱动程序
3、bootstrapNetInit()//
初始化引导网络,为NCCL网络通信的顺利进行奠定了基础
4、ncclNetPluginInit()//NCCL网络插件初始化,抽象和封装底层网络细节,方便NCCL灵活应用
备注:“A、setEnvFile(confFilePath);//根据配置文件初始化设置” 的首行缩进表示initEnv()调用了setEnvFile(confFilePath)函数。缩进表示调用,或者该代码片段中使用。
备注:bootstrap引导网络主要在初始化时完成一些小数据量的信息交换,例如ip地址。
源码速递
源码位置:nccl-master\src\init.cc
static ncclResult_t ncclInit() {
// 检查 initialized 是否已经被设置为 true,表示库是否已经被初始化。
if (__atomic_load_n(&initialized, __ATOMIC_ACQUIRE)) return ncclSuccess;
// 锁定互斥锁 initLock,确保在初始化过程中只有一个线程可以执行。
pthread_mutex_lock(&initLock);
// 再次检查 initialized,因为在获取锁的过程中可能已经有其他线程完成了初始化。
if (!initialized) {
//1、初始化库的环境设置。会调用setEnvFile(confFilePath);//根据配置文件初始化设置
initEnv();
//2、初始化GDR(GPU Direct RDMA)拷贝功能。会调用ncclGdrInit()//初始化与GDR相关的库和驱动程序
initGdrCopy();
//3、始终初始化引导网络,通过解析环境变量、查找网络接口和设置网络通信参数等步骤,
//bootstrapNetInit函数为NCCL网络通信的顺利进行奠定了基础。
NCCLCHECK(bootstrapNetInit());
//4、尝试加载并初始化外部的网络插件,抽象和封装底层网络细节,方便NCCL灵活应用。
NCCLCHECK(ncclNetPluginInit());
// 初始化与 NVTX 相关的枚举类型,以便在性能分析工具中使用。
initNvtxRegisteredEnums();
// 将 initialized 设置为 true,表示初始化已完成。
__atomic_store_n(&initialized, true, __ATOMIC_RELEASE);
}
// 解锁互斥锁 initLock,允许其他线程进入临界区。
pthread_mutex_unlock(&initLock);
// 返回初始化成功的结果。
return ncclSuccess;
}
三、bootstrapGetUniqueId()源码详解
bootstrapGetUniqueId()核心逻辑
1、生成一个随机数,填充ncclUniqueId的前半部分。
2、如果环境变量中有NCCL_COMM_ID的值,将环境变量解析为网络地址,赋值给ncclUniqueId的后半部分。
3、如果环境变量中没有NCCL_COMM_ID的值,将bootstrap网络地址,赋值给ncclUniqueId的后半部分。
注意,标识通信组的唯一ID ncclUniqueId本质上由两部分组成,前半部分是随机数,后半部分是网络地址。
源码速递
源码位置:nccl-master\src\bootstrap.cc
ncclResult_t bootstrapGetUniqueId(struct ncclBootstrapHandle* handle) {
// 初始化传入的ncclBootstrapHandle结构体,将其内存区域全部置为0
memset(handle, 0, sizeof(ncclBootstrapHandle));
1、生成一个随机数,填充ncclUniqueId的前半部分。
// 生成一个随机的magic值,并存储在handle的magic字段中
NCCLCHECK(getRandomData(&handle->magic, sizeof(handle->magic)));
// 尝试从环境变量中获取NCCL_COMM_ID的值
const char* env = ncclGetEnv("NCCL_COMM_ID");
// 如果环境变量存在(即env不为NULL)
if (env) {
2、如果环境变量中有NCCL_COMM_ID的值,将环境变量解析为网络地址,赋值给ncclUniqueId的后半部分。
// 打印一条信息,表明NCCL_COMM_ID是由环境变量设置的,并显示其值
INFO(NCCL_ENV, "NCCL_COMM_ID set by environment to %s", env);
// 尝试将环境变量中的字符串解析为一个网络地址,存储在handle的addr字段中
if (ncclSocketGetAddrFromString(&handle->addr, env) != ncclSuccess) {
// 如果解析失败,打印一条警告信息,指出NCCL_COMM_ID的格式不正确
WARN("Invalid NCCL_COMM_ID, please use format: <ipv4>:<port> or [<ipv6>]:<port> or <hostname>:<port>");
// 并返回ncclInvalidArgument错误码
return ncclInvalidArgument;
}
} else {
3、如果环境变量中没有NCCL_COMM_ID的值,将bootstrap网络地址,赋值给ncclUniqueId的后半部分。
// 如果环境变量不存在,则将默认的bootstrap网络地址复制到handle的addr字段中
memcpy(&handle->addr, &bootstrapNetIfAddr, sizeof(union ncclSocketAddress));
// 尝试创建一个根节点(可能是启动NCCL通信的服务器或领导者节点)
NCCLCHECK(bootstrapCreateRoot(handle, false));
}
// 如果所有操作都成功完成,则返回ncclSuccess表示成功
return ncclSuccess;
}
下期预告
通信器初始化ncclCommInitRank() 含视频教程
NCCL源码详解3:通信器初始化ncclCommInitRank() 含视频教程_succlcomminitrank-CSDN博客