NCCL源码详解6：通信拓扑识别感知构建物理拓扑xml文件 ncclTopoGetSystem() 视频教程

爱串门的小马驹

已于 2025-02-19 10:49:53 修改

阅读量2.2k

点赞数 22

分类专栏：集合通信文章标签：网络分布式人工智能信息与通信拓扑学

于 2024-07-08 21:05:39 首次发布

本文链接：https://blog.csdn.net/lianghuaju/article/details/139090676

版权

集合通信专栏收录该内容

15 篇文章

订阅专栏

Nvidia NCCL如何构建物理拓扑

视频教程在这：

NCCL集合通信源码解读、案例、任务调度、拓扑_哔哩哔哩_bilibili

一、ncclTopoGetSystem()拓扑构建

1.1 ncclTopoGetSystem()拓扑构建核心逻辑：

1、尝试从文件加载已有拓扑信息。

2、如果没有已有拓扑信息，创建一个名为"system"的根节点；

3、遍历本服务器所有GPU，拓扑树中添加GPU节点和NVlink；

4、遍历所有网络设备，拓扑树中添加网络拓扑节点；

5、移除不可用的节点；

6、Multi-Node NVLink (MNNVL) 跨服务器NVLink支持；

7、保存拓扑文件；

最最核心的就下面这三步

一、创建根节点

二、遍历并插入GPU节点和NVlink

三、遍历并插入网卡节点

哈哈哈，原来如此简单！

ncclTopoGetSystem()中调用的最核心函数是啥，那必然是添加节点的函数，后面我们就来看看添加GPU节点ncclTopoFillGpu()，这里我们先过一下ncclTopoGetSystem()源码。

1.2 ncclTopoGetSystem()源码速递：

源码位置：nccl-master\src\graph\topo.cc

ncclResult_t ncclTopoGetSystem(struct ncclComm* comm, struct ncclTopoSystem** system) {  

  // 分配一个XML结构，用于存储拓扑信息  
  struct ncclXml* xml;  
  NCCLCHECK(xmlAlloc(&xml, NCCL_TOPO_XML_MAX_NODES));  

  ///
  1、 尝试从文件加载已有拓扑信息。
  // 尝试从环境变量中获取XML拓扑文件的路径  
  const char* xmlTopoFile = ncclGetEnv("NCCL_TOPO_FILE");  

  if (xmlTopoFile) {  
    // 如果环境变量设置了，则打印信息并加载该文件到xml结构中  
    INFO(NCCL_ENV, "NCCL_TOPO_FILE set by environment to %s", xmlTopoFile);  
    NCCLCHECK(ncclTopoGetXmlFromFile(xmlTopoFile, xml, 1));  
  } else {  
    // 如果没有设置环境变量，则尝试从默认位置加载XML拓扑文件  
    // Try default XML topology location  
    NCCLCHECK(ncclTopoGetXmlFromFile("/var/run/nvidia-topologyd/virtualTopology.xml", xml, 0));  
  }  

  
  /2、如果没有已有拓扑信息，创建一个名为"system"的根节点；//
  // 如果xml结构中没有任何节点（即没有加载到任何拓扑信息）  
  if (xml->maxIndex == 0) {  
    // 创建一个名为"system"的根节点，并设置其版本属性  
    // Create top tag  
    struct ncclXmlNode* top;  
    NCCLCHECK(xmlAddNode(xml, NULL, "system", &top));  
    NCCLCHECK(xmlSetAttrInt(top, "version", NCCL_TOPO_XML_VERSION));  
  }  

  
  /
  3、遍历本服务器所有GPU，拓扑树中添加GPU节点和NVlink；
  // 如果需要，自动检测GPU设备  
  // Auto-detect GPUs if needed  
  for (int r=0; r<comm->nRanks; r++) {  
    // 如果当前排名r对应的hostHash与当前rank的hostHash相同（可能是同一台机器上的不同GPU）  
    if (comm->peerInfo[r].hostHash == comm->peerInfo[comm->rank].hostHash) {  
      // 将busId转换为可读的PCI总线ID格式  
      char busId[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];  
      NCCLCHECK(int64ToBusId(comm->peerInfo[r].busId, busId));   

      // 填充一个表示GPU的XML节点  
      struct ncclXmlNode* node;  
      NCCLCHECK(ncclTopoFillGpu(xml, busId, &node));  

      // 如果没有成功创建节点，则继续下一次循环  
      if (node == NULL) continue;  
      // 设置该GPU节点的"keep"属性为1，表示需要保留这个节点  

      NCCLCHECK(xmlSetAttrInt(node, "keep", 1));  
      // 设置该GPU节点的"rank"属性为当前排名r  
      NCCLCHECK(xmlSetAttrInt(node, "rank", r));  

      // 设置该GPU节点的"gdr"属性，表示是否支持GPU Direct RDMA  
      NCCLCHECK(xmlInitAttrInt(node, "gdr", comm->peerInfo[r].gdrSupport));  

    }  

  }  


/
4、遍历所有网络设备，拓扑树中添加网络拓扑节点；/
// 如果需要的话，自动检测NICs（网络接口卡）。net和collnet共享相同的xml/graph节点，  
// 所以我们先从collnet开始，以便它有更高的优先级。  
// Auto-detect NICs if needed. net/collnet share the same xml/graph nodes,  
// so we start with collnet so that it has precedence.  
int netDevCount = 0; // 初始化网络设备计数为0  
// 如果comm支持collNet  
if (collNetSupport(comm)) {  
    // 获取comm支持的网络设备数量  
    NCCLCHECK(collNetDevices(comm, &netDevCount));  

    // 遍历每个网络设备  
    for (int n=0; n<netDevCount; n++) {  
        ncclNetProperties_t props; // 定义一个ncclNetProperties_t类型的变量props，用于存储设备属性  
        // 获取第n个网络设备的属性  
        NCCLCHECK(collNetGetProperties(comm, n, &props));  
        // 创建一个XML节点来表示这个网络设备  
        struct ncclXmlNode* netNode;  
        // 使用设备的pci路径和名称来填充XML节点  
        NCCLCHECK(ncclTopoFillNet(xml, props.pciPath, props.name, &netNode));  
        // 将"keep"属性设置为1，可能表示这个节点需要被保留  
        NCCLCHECK(xmlSetAttrInt(netNode, "keep", 1));  
        // 将"dev"属性设置为n，表示这是第n个设备  
        NCCLCHECK(xmlSetAttrInt(netNode, "dev", n));  
        // 将速度、端口、GUID等属性添加到XML节点中  
        NCCLCHECK(xmlInitAttrInt(netNode, "speed", props.speed));  
        NCCLCHECK(xmlInitAttrInt(netNode, "port", props.port));  
        NCCLCHECK(xmlInitAttrUint64(netNode, "guid", props.guid));  
        NCCLCHECK(xmlInitAttrInt(netNode, "maxconn", props.maxComms));  
        // 检查是否支持GPU Direct RDMA（GDR）  
        bool gdrSupport = (props.ptrSupport & NCCL_PTR_CUDA) || (comm->dmaBufSupport && (props.ptrSupport & NCCL_PTR_DMABUF));  
        // 打印GDR支持状态和设备信息  
        INFO(NCCL_NET,"NET/%s : GPU Direct RDMA %s for HCA %d '%s'", comm->ncclNet->name, gdrSupport ? "Enabled" : "Disabled", n, props.name);  
        // 将GDR支持状态添加到XML节点中  
        NCCLCHECK(xmlInitAttrInt(netNode, "gdr", gdrSupport));  
        // 将"coll"属性设置为1，可能表示这是一个集合通信网络接口  
        NCCLCHECK(xmlInitAttrInt(netNode, "coll", 1));  
    }  
} 

// 循环遍历所有的网络设备，其中 netDevCount 是网络设备的总数  
for (int n=0; n<netDevCount; n++) {   
    // 定义一个 ncclNetProperties_t 类型的变量 props，用于存储网络设备的属性  
    ncclNetProperties_t props;  
    // 调用 getProperties 函数获取网络设备的属性，并检查调用是否成功  
    // 参数 n 是当前网络设备的索引，&props 是用于存储属性的指针  
    NCCLCHECK(comm->ncclNet->getProperties(n, &props));    
    // 定义一个指向 ncclXmlNode 结构的指针 netNode，该结构将用于表示 XML 中的节点  
    struct ncclXmlNode* netNode;   
    // 调用 ncclTopoFillNet 函数在 XML 结构中创建一个新的节点，并检查调用是否成功  
    // 参数 xml 是 XML 结构的指针，props.pciPath 和 props.name 是网络设备的 PCI 路径和名称  
    // &netNode 是用于存储新节点指针的指针  
    NCCLCHECK(ncclTopoFillNet(xml, props.pciPath, props.name, &netNode));  
    // 设置新节点的 keep 属性为 1，表示该节点应该被保留  
    NCCLCHECK(xmlSetAttrInt(netNode, "keep", 1));  
    // 设置新节点的 dev 属性为当前网络设备的索引 n  
    NCCLCHECK(xmlSetAttrInt(netNode, "dev", n));    
    // 设置新节点的 speed 属性为网络设备的速度  
    NCCLCHECK(xmlInitAttrInt(netNode, "speed", props.speed));  
    // 设置新节点的 port 属性为网络设备的端口号  
    // 并检查设置属性是否成功  
    NCCLCHECK(xmlInitAttrInt(netNode, "port", props.port));  
    // 设置新节点的 latency 属性为网络设备的延迟  
    NCCLCHECK(xmlInitAttrFloat(netNode, "latency", props.latency));    
    // 设置新节点的 guid 属性为网络设备的全局唯一标识符  
    NCCLCHECK(xmlInitAttrUint64(netNode, "guid", props.guid));   
    // 设置新节点的 maxconn 属性为网络设备支持的最大并发通信数  
    NCCLCHECK(xmlInitAttrInt(netNode, "maxconn", props.maxComms));  
    // 检查网络设备是否支持 GPU Direct RDMA  
    // 如果 props.ptrSupport 包含 NCCL_PTR_CUDA 或者如果 comm->dmaBufSupport 为真且 props.ptrSupport 包含 NCCL_PTR_DMABUF，则 gdrSupport 为真  
    bool gdrSupport = (props.ptrSupport & NCCL_PTR_CUDA) || (comm->dmaBufSupport && (props.ptrSupport & NCCL_PTR_DMABUF));   
    // 打印日志信息，显示网络设备是否支持 GPU Direct RDMA  
    // 其中 comm->ncclNet->name 是网络设备的名称，n 是设备的索引，props.name 是设备的名字  
    INFO(NCCL_NET,"NET/%s : GPU Direct RDMA %s for HCA %d '%s'", comm->ncclNet->name, gdrSupport ? "Enabled" : "Disabled", n, props.name);  
    // 设置新节点的 gdr 属性，表示是否支持 GPU Direct RDMA  
    NCCLCHECK(xmlInitAttrInt(netNode, "gdr", gdrSupport));  
}  


5、移除不可用的节点；/
// 移除 XML 中不包含 keep="1" 节点的分支
NCCLCHECK(ncclTopoTrimXml(xml));  

/
6、Multi-Node NVLink (MNNVL) 跨服务器NVLink支持；  
// 如果 MNNVL被启用  
if (comm->MNNVL) {  
    // MNNVL 集群支持  
    // 分配内存来存储所有集群成员的网络拓扑数据  
    char* mem;  
    // 为每个集群成员分配足够的内存空间来存储 XML 数据  
    // 假设每个成员的 XML 数据不超过 NCCL_TOPO_XML_MAX_NODES 大小  
    NCCLCHECK(ncclCalloc(&mem, comm->clique.size * xmlMemSize(NCCL_TOPO_XML_MAX_NODES)));   
    // 获取当前集群成员的 XML 数据区域  
    struct ncclXml* rankXml = (struct ncclXml*)(mem + xmlMemSize(NCCL_TOPO_XML_MAX_NODES) * comm->cliqueRank);  
    // 复制当前集群成员的 XML 数据  
    memcpy(rankXml, xml, xmlMemSize(NCCL_TOPO_XML_MAX_NODES));   
    // 将当前集群成员的 XML 数据转换为内部表示形式（可能是为了更高效的通信）  
    NCCLCHECK(ncclTopoConvertXml(rankXml, (uintptr_t)xml->nodes, 1));    
    // 在集群内所有成员间收集各自的 XML 数据  
    // bootstrapIntraNodeAllGather 可能是某种集群内收集数据的函数  
    NCCLCHECK(bootstrapIntraNodeAllGather(comm->bootstrap, comm->clique.ranks, comm->cliqueRank, comm->clique.size, mem, xmlMemSize(NCCL_TOPO_XML_MAX_NODES)));   
    // 分配一个新的 XML 结构来存储融合后的集群拓扑数据  
    struct ncclXml* cliqueXml;  
    NCCLCHECK(xmlAlloc(&cliqueXml, comm->clique.size * NCCL_TOPO_XML_MAX_NODES));   
    // 融合集群内所有成员的 XML 数据  
    for (int i = 0; i < comm->clique.size; i++) {  
        // 获取集群中每个成员的 XML 数据  
        struct ncclXml* peerXml = (struct ncclXml*)(mem + xmlMemSize(NCCL_TOPO_XML_MAX_NODES) * i);  
        // 将 XML 数据转换为内部表示形式（这次可能为了融合做准备）  
        NCCLCHECK(ncclTopoConvertXml(peerXml, (uintptr_t)peerXml->nodes, 0));  
        // 将当前成员的 XML 数据融合到 cliqueXml 中  
        NCCLCHECK(ncclTopoFuseXml(cliqueXml, peerXml));  
    }   
    // 释放原来的 XML 数据  
    free(xml);  
    // 更新 xml 指针以指向融合后的集群 XML 数据  
    xml = cliqueXml;  
}  

//
7、保持拓扑文件；/  
// 获取环境变量 NCCL_TOPO_DUMP_FILE 的值，用于存储 XML 拓扑数据  
xmlTopoFile = ncclGetEnv("NCCL_TOPO_DUMP_FILE");  
// 如果环境变量被设置，并且当前进程是负责输出拓扑数据的进程（由 ncclParamTopoDumpFileRank() 确定）  
if (xmlTopoFile && comm->rank == ncclParamTopoDumpFileRank()) {  
    // 输出环境变量 NCCL_TOPO_DUMP_FILE 的值  
    INFO(NCCL_ENV, "NCCL_TOPO_DUMP_FILE set by environment to %s", xmlTopoFile);  
    // 将融合后的 XML 拓扑数据写入到指定的文件中  
    NCCLCHECK(ncclTopoDumpXmlToFile(xmlTopoFile, xml));  
}    
// 从 XML 数据中提取系统信息，并存储在 system 中  
// comm->peerInfo[comm->rank].hostHash 可能用于区分不同主机的哈希值  
NCCLCHECK(ncclTopoGetSystemFromXml(xml, system, comm->peerInfo[comm->rank].hostHash));  
  
// 释放 XML 数据的内存  
free(xml);   
// 返回成功状态  
return ncclSuccess;

}

ncclTopoGetSystem()中调用的最核心函数是添加节点的函数，下面我们就来看看添加GPU节点ncclTopoFillGpu()。

二、ncclTopoFillGpu核心逻辑：

2.1 ncclTopoFillGpu核心逻辑：

1、ncclTopoGetPciNode()确定当前GPU卡是否已创建xml node，没有就创建。

2、ncclTopoGetXmlFromSys()获取GPU到cpu的路径，路径信息获取，生成xml树。

3、GPU相关信息获取，设置NVlink信息。

2.2 源码速递：

源码位置：nccl-master\src\graph\xml.cc

// ncclTopoFillGpu，它接受一个ncclXml指针、一个char指针（busId）和一个ncclXmlNode指针的指针（gpuNode）作为参数  
ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct ncclXmlNode** gpuNode) {  
  
    // 定义一个ncclXmlNode指针node，用于临时存储找到的PCI节点的指针  
    struct ncclXmlNode* node;  
  
    
    1、ncclTopoGetPciNode()确定当前GPU卡是否已创建xml node，没有就创建。
    // 调用ncclTopoGetPciNode函数，根据busId在xml结构中查找匹配的PCI节点，并将找到的节点指针存储在node中   
    NCCLCHECK(ncclTopoGetPciNode(xml, busId, &node));  
  
    // 调用xmlSetAttrIfUnset函数，确保node节点有"class"属性，并且其值为"0x03"  
    // 如果该属性已存在且值不是"0x03"，则不会进行任何操作  
    NCCLCHECK(xmlSetAttrIfUnset(node, "class", "0x03"));  
  
    
    //
    2、ncclTopoGetXmlFromSys()获取GPU到cpu的路径，路径信息获取，生成xml树。
    // 调用ncclTopoGetXmlFromSys函数，从系统中获取与node节点相关的更多XML信息，并添加到xml结构中  
    NCCLCHECK(ncclTopoGetXmlFromSys(node, xml));  
  
    // 定义一个nvmlDevice_t类型的变量nvmlDev，用于存储从busId获取的NVML设备句柄  
    nvmlDevice_t nvmlDev;  
  
    // 调用ncclNvmlDeviceGetHandleByPciBusId函数，根据busId获取对应的NVML设备句柄，并存储在nvmlDev中  
    NCCLCHECK(ncclNvmlDeviceGetHandleByPciBusId(busId, &nvmlDev));  
  

    
    3、GPU相关信息获取，设置NVlink。
    // 调用ncclTopoGetXmlFromGpu函数，从nvmlDev设备句柄中获取与GPU相关的XML信息，并添加到xml结构中  
    // 同时，将找到的GPU节点的指针存储在gpuNode中  
    NCCLCHECK(ncclTopoGetXmlFromGpu(node, nvmlDev, xml, gpuNode));  
  
    // 如果以上所有操作都成功完成，则返回ncclSuccess表示成功  
    return ncclSuccess;  
}

2.3 ncclTopoGetPciNode()

我们过一下第一步中的函数ncclTopoGetPciNode()

1、ncclTopoGetPciNode()确定当前GPU卡是否已创建xml node，没有就创建。

源码位置：nccl-master\src\graph\xml.cc

// 定义一个函数ncclTopoGetPciNode，它接受一个ncclXml结构体指针xml，一个字符串指针busId用于指定PCI节点的busid，  
// 以及一个指向ncclXmlNode指针的指针pciNode，用于返回找到的或新创建的PCI节点的地址。  
ncclResult_t ncclTopoGetPciNode(struct ncclXml* xml, const char* busId, struct ncclXmlNode** pciNode) {  
  
    // 调用xmlFindTagKv函数在xml中查找标签为"pci"且属性"busid"等于busId的节点。  
    // 如果找到，将找到的节点的地址存储在*pciNode中。   
    NCCLCHECK(xmlFindTagKv(xml, "pci", pciNode, "busid", busId));  
  
    // 如果*pciNode是NULL，表示没有找到与busId相对应的PCI节点。  
    if (*pciNode == NULL) {  
        // 调用xmlAddNode函数在xml中添加一个新的"pci"节点，并将其地址存储在*pciNode中。  
        // 这里的NULL作为父节点参数，意味着新节点将被添加到XML树的根目录下。  
        NCCLCHECK(xmlAddNode(xml, NULL, "pci", pciNode));  
  
        // 调用xmlSetAttr函数设置新创建的PCI节点的"busid"属性为busId。  
        NCCLCHECK(xmlSetAttr(*pciNode, "busid", busId));  
    }  
  
    // 函数成功完成，返回ncclSuccess表示操作成功。  
    return ncclSuccess;  
}

2.4 ncclTopoGetXmlFromSys()

当然ncclTopoFillGpu中调用的最最核心的，还是这个函数ncclTopoGetXmlFromSys()。

2.4.1 ncclTopoGetXmlFromSys()核心逻辑

1、getPciPath()获取GPU到cpu的路径；

2、获取link_width，link_speed等属性；

3、根据路径查找父节点，查找不到就创建父节点，继续查找父节点的父节点（爷爷节点），就这样循环查找和创建，构建xml树，直到找到父节点；

4、插入节点GPU节点。

2.4.2 ncclTopoGetXmlFromSys()源码速递

ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml* xml) {  

  // 从pciNode节点中获取busid属性  
  const char* busId;  
  NCCLCHECK(xmlGetAttr(pciNode, "busid", &busId)); 
  // 分配一个字符指针path，用于存储PCI设备的路径（初始化为NULL）  
  char* path = NULL;  
  // 临时设置ncclDebugNoWarn为NCCL_GRAPH，可能是为了关闭某些与图相关的警告  
  ncclDebugNoWarn = NCCL_GRAPH;  

  
  ///
  1、getPciPath()获取GPU到cpu的路径；
  // 调用getPciPath函数，根据busId获取PCI设备的路径，并存储在path中  
  getPciPath(busId, &path);  
  // 恢复ncclDebugNoWarn为0，关闭之前的特殊调试设置  
  ncclDebugNoWarn = 0;   
  // 如果path不为空（即成功获取了PCI设备的路径）  
  if (path) {  
    // 调用ncclTopoSetAttrFromSys函数，从系统中获取与path对应的"class"属性，并设置到pciNode节点中  
    NCCLCHECK(ncclTopoSetAttrFromSys(pciNode, path, "class", "class"));  
  }  

  /
  2、获取link_width，link_speed等属性；/
  // 以下代码块类似地处理其他PCI设备属性：vendor, device, subsystem_vendor, subsystem_device  
  int index;  
  ncclDebugNoWarn = NCCL_GRAPH;  
  // 尝试从pciNode节点中获取vendor属性的索引  
  NCCLCHECK(xmlGetAttrIndex(pciNode, "vendor", &index));  
  // 如果索引为-1（即属性不存在）  
  if (index == -1) {  
    // 如果path不为空，则从系统中获取与path对应的"vendor"属性，并设置到pciNode节点中  
    if (path) ncclTopoSetAttrFromSys(pciNode, path, "vendor", "vendor");  
  }    
  // ...（类似地处理device, subsystem_vendor, subsystem_device属性）  

  // 尝试从pciNode节点中获取link_speed属性的索引  
  NCCLCHECK(xmlGetAttrIndex(pciNode, "link_speed", &index));  
  // 如果索引为-1（即属性不存在）  
  if (index == -1) {  
    // 如果path不为空  
    if (path) {  
      // 定义两个字符串数组用于存储从系统中读取的最大链路速度  
      char deviceSpeedStr[MAX_STR_LEN];  
      char portSpeedStr[MAX_STR_LEN];  
      // 初始化设备速度和端口速度为FLT_MAX（一个表示无穷大的浮点数）  
      float deviceSpeed = FLT_MAX;  
      float portSpeed = FLT_MAX;  
      // 从系统中获取与path对应的"max_link_speed"属性，并存储在deviceSpeedStr中  
      NCCLCHECK(ncclTopoGetStrFromSys(path, "max_link_speed", deviceSpeedStr));  
      // 将字符串转换为浮点数  
      sscanf(deviceSpeedStr, "%f GT/s", &deviceSpeed);  
      // 类似地，从系统的父路径（可能是端口）中获取"max_link_speed"属性  
      NCCLCHECK(ncclTopoGetStrFromSys(path, "../max_link_speed", portSpeedStr));  
      sscanf(portSpeedStr, "%f GT/s", &portSpeed);  
      // 设置link_speed属性为设备速度和端口速度中较小的一个  
      NCCLCHECK(xmlSetAttr(pciNode, "link_speed", portSpeed < deviceSpeed ? portSpeedStr : deviceSpeedStr));  
    } else {  
      // 如果path为空，则设置link_speed属性为空字符串  
      NCCLCHECK(xmlSetAttr(pciNode, "link_speed", ""));  
    }  
  } 
  NCCLCHECK(xmlGetAttrIndex(pciNode, "link_width", &index));
  if (index == -1) {
    if (path) {
      char strValue[MAX_STR_LEN];
      NCCLCHECK(ncclTopoGetStrFromSys(path, "max_link_width", strValue));
      int deviceWidth = strtol(strValue, NULL, 0);
      NCCLCHECK(ncclTopoGetStrFromSys(path, "../max_link_width", strValue));
      int portWidth = strtol(strValue, NULL, 0);
      NCCLCHECK(xmlSetAttrInt(pciNode, "link_width", std::min(deviceWidth,portWidth)));
    } else {
      NCCLCHECK(xmlSetAttr(pciNode, "link_width", ""));
    }
  }

/
3、根据路径查找父节点，查找不到就创建父节点，继续查找父节点的父节点（爷爷节点），就这样循环查找和创建，构建xml树，直到找到父节点；//
struct ncclXmlNode* parent = pciNode->parent; // 获取当前pciNode的父节点  
if (parent == NULL) { // 如果当前pciNode没有父节点  
    if (path) { // 如果path（PCI设备的路径）不为空  
        // Save that for later in case next step is a CPU  
        char numaIdStr[MAX_STR_LEN]; // 用于存储NUMA节点ID的字符串
        // 从系统中根据PCI路径获取NUMA节点ID    
        NCCLCHECK(ncclTopoGetStrFromSys(path, "numa_node", numaIdStr)); 
  
        // Go up one level in the PCI tree. Rewind two "/" and follow the upper PCI  
        // switch, or stop if we reach a CPU root complex.  
        int slashCount = 0; // 用于计数'/'的数量  
        int parentOffset; // 用于遍历path的偏移量  
        for (parentOffset = strlen(path)-1; parentOffset>0; parentOffset--) { // 从path的末尾开始遍历  
            if (path[parentOffset] == '/') { // 如果遇到'/'  
                slashCount++; // '/'计数加1  
                path[parentOffset] = '\0'; // 临时将'/'替换为'\0'，以便获取上级路径  
                int start = parentOffset - 1; // 从'/'前一个字符开始寻找下一个'/'  
                while (start>0 && path[start] != '/') start--; // 找到上一个'/'  

                // 检查上级路径是否符合PCI设备的BDF格式（总线:设备:功能）
                // Check whether the parent path looks like "BBBB:BB:DD.F" or not.  
                if (checkBDFFormat(path+start+1) == 0) {   
                    // This a CPU root complex. Create a CPU tag and stop there.  
                    struct ncclXmlNode* topNode; // 定义一个指向系统顶层的指针  
                    NCCLCHECK(xmlFindTag(xml, "system", &topNode)); // 查找系统标签  
                    NCCLCHECK(xmlGetSubKv(topNode, "cpu", &parent, "numaid", numaIdStr)); // 在系统标签下查找numaid为numaIdStr的cpu标签  
                    if (parent == NULL) { // 如果没有找到  
                        NCCLCHECK(xmlAddNode(xml, topNode, "cpu", &parent)); // 在系统标签下添加一个新的cpu标签  
                        NCCLCHECK(xmlSetAttrLong(parent, "host_hash", getHostHash())); // 设置host_hash属性  
                        NCCLCHECK(xmlSetAttr(parent, "numaid", numaIdStr)); // 设置numaid属性  
                    }  
                } else if (slashCount == 2) { //   
                    // Continue on the upper PCI switch  
                    for (int i = strlen(path)-1; i>0; i--) { // 从path的末尾开始遍历  
                        if (path[i] == '/') { // 如果遇到'/'  
                            NCCLCHECK(xmlFindTagKv(xml, "pci", &parent, "busid", path+i+1)); // 查找busid为path+i+1的pci标签  
                            if (parent == NULL) { // 如果没有找到  
                                NCCLCHECK(xmlAddNode(xml, NULL, "pci", &parent)); // 在XML树的根下添加一个新的pci标签  
                                NCCLCHECK(xmlSetAttr(parent, "busid", path+i+1)); // 设置busid属性  
                            }  
                            break; // 找到后跳出循环  
                        }  
                    }  
                }  
            }  
            if (parent) break; // 如果找到了父节点，则跳出循环  
        }  
    } else {  
      // 如果没有在 /sys 目录下找到相关信息，那么将 GPU 附加到一个未知的 CPU 上  
      // No information on /sys, attach GPU to unknown CPU    
      // 尝试在 XML 中查找名为 "cpu" 的标签，并查找具有 "numaid" 属性的标签，如果找不到，则 "numaid" 设置为 "-1"  
      NCCLCHECK(xmlFindTagKv(xml, "cpu", &parent, "numaid", "-1"));  
      // 如果 parent 指针为空，说明没有找到符合条件的 "cpu" 标签  
      if (parent == NULL) {  
        // 定义一个指向 XML 节点的指针 topNode  
        struct ncclXmlNode* topNode;   
        // 在 XML 中查找名为 "system" 的标签，并将找到的节点赋值给 topNode  
        NCCLCHECK(xmlFindTag(xml, "system", &topNode));   
        // 在 topNode 下添加一个名为 "cpu" 的子节点，并将新节点的指针赋值给 parent  
        NCCLCHECK(xmlAddNode(xml, topNode, "cpu", &parent));  
        // 设置新节点的 "host_hash" 属性为当前主机的哈希值  
        NCCLCHECK(xmlSetAttrLong(parent, "host_hash", getHostHash()));  
        // 设置新节点的 "numaid" 属性为 "-1"，表示这是一个未知的 CPU  
        NCCLCHECK(xmlSetAttr(parent, "numaid", "-1"));  
        // 调用函数从 CPU 获取 XML 数据并填充到 parent 节点中  
        NCCLCHECK(ncclTopoGetXmlFromCpu(parent, xml));  
      }  
      
      ///
      4、插入节点GPU节点。/
      // 将 pciNode 的 parent 指针设置为前面找到的或创建的 parent 节点  
      pciNode->parent = parent;  
  
      // 为了保持 PCI 子设备按 PCI 总线 ID 排序（解决 Issue #820）  
      // 首先获取新 PCI 设备的 busid  
      int subIndex = parent->nSubs; // 假设新设备被添加到子节点的末尾  
      const char* newBusId;  
      NCCLCHECK(xmlGetAttrStr(pciNode, "busid", &newBusId)); // 获取新 PCI 设备的 busid  
  
      // 遍历 parent 的所有子节点  
      for (int s=0; s<parent->nSubs; s++) {  
        // 获取当前子节点的 busid  
        const char* busId;  
        NCCLCHECK(xmlGetAttr(parent->subs[s], "busid", &busId));  
  
        // 如果当前子节点的 busid 不为空，并且新设备的 busid 比当前子节点的 busid 小  
        // 则将 subIndex 设置为当前索引 s，并跳出循环  
        if (busId != NULL && strcmp(newBusId, busId) < 0) { subIndex = s; break; }  
      }  
  
      // 如果 parent 的子节点数量已经达到了最大值 MAX_SUBS  
      if (parent->nSubs == MAX_SUBS) {  
        // 发出警告，并返回 ncclInternalError 错误  
        WARN("Error : XML parser is limited to %d subnodes", MAX_SUBS);  
        return ncclInternalError;  
      }  
  
      // 将 parent 的子节点从 subIndex 开始向后移动一个位置，为新节点腾出空间  
      for (int s = parent->nSubs; s > subIndex; s--) parent->subs[s] = parent->subs[s-1];  
  
      // 将 pciNode 插入到正确的位置 subIndex  
      parent->subs[subIndex] = pciNode;  
  
      // 增加 parent 的子节点数量  
      parent->nSubs++;  
    }  
  
    // 如果 parent 节点的名称是 "pci"，则从系统获取 XML 数据并填充到 parent 节点中  
    if (strcmp(parent->name, "pci") == 0) {  
      NCCLCHECK(ncclTopoGetXmlFromSys(parent, xml));  
    }   
    // 如果 parent 节点的名称是 "cpu"，则从 CPU 获取 XML 数据并填充到 parent 节点中  
    else if (strcmp(parent->name, "cpu") == 0) {  
      NCCLCHECK(ncclTopoGetXmlFromCpu(parent, xml));  
    }  
  
    // 释放之前分配给 path 的内存  
    free(path);  
  
    // 返回 ncclSuccess 表示操作成功  
    return ncclSuccess;  
}