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;
}
4556
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
