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

Nvidia NCCL如何构建物理拓扑

视频教程在这：

2.2 NCCL源码分析：物理拓扑识别感知xml通信topo构建 ncclTopoGetSystem()_哔哩哔哩_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;  
}