NCCL源码详解2:通信初始化如何获取唯一ID UniqueId,ncclGetUniqueId()中ncclInit()、bootstrapGetUniqueId()包含视频教程

上节课我们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博客

1.3 NCCL通信器初始化ncclCommInitRank()源码解读_哔哩哔哩_bilibili

  • 10
    点赞
  • 12
    收藏
    觉得还不错? 一键收藏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值