diff --git a/src/graph/paths.cc b/src/graph/paths.cc index 2aeb9fbb20..056a288489 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -414,14 +414,15 @@ ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* system, int64_t busId, int // Check if we are close enough that it makes sense to enable GDR int netGdrLevel = PATH_PXB; -#ifdef TOPO_EXPL - int arch, vendor, model; - NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model)); - if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_AMD && model == NCCL_TOPO_CPU_TYPE_ROME) - netGdrLevel = PATH_PHB; -#endif NCCLCHECK(ncclGetLevel(&ncclTopoUserGdrLevel, NULL, "NCCL_NET_GDR_LEVEL")); if (ncclTopoUserGdrLevel != -2) netGdrLevel = ncclTopoUserGdrLevel; + else { + int arch, vendor, model; + NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model)); + if((system->nodes[GPU].nodes[g].id & 0xf0000) == (system->nodes[NET].nodes[n].net.busId & 0xf0000)) + netGdrLevel = PATH_PHB; + } + int distance = gpu->paths[NET][n].type; if (distance > netGdrLevel) { INFO(NCCL_NET,"GPU Direct RDMA Disabled for GPU %lx / HCA %d (distance %d > %d)", busId, netDev, distance, netGdrLevel); diff --git a/src/graph/rome_models.h b/src/graph/rome_models.h index 65a572c917..286d860842 100755 --- a/src/graph/rome_models.h +++ b/src/graph/rome_models.h @@ -29,6 +29,7 @@ struct rcclRomeModel { int nNics; int nLinks; int64_t gpuIds[MAX_ROME_GPUS]; + int64_t nicIds[MAX_ROME_NICS]; int64_t gpuNuma[MAX_ROME_GPUS]; int64_t nicNuma[MAX_ROME_NICS]; int connMatrix[MAX_ROME_GPUS*MAX_ROME_GPUS]; @@ -39,6 +40,7 @@ struct rcclRomeModel { static struct rcclRomeModel rome_model_22 = { .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 2, .gpuIds = { 0x3000, 0x43000, 0x26000, 0xc3000, 0x83000, 0x23000, 0xc6000, 0xa3000, }, + .nicIds = { 0xe1000, }, .gpuNuma = { 1, 0, 1, 2, 3, 1, 2, 3, }, .nicNuma = { 2, }, .connMatrix = { 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, }, @@ -49,6 +51,7 @@ static struct rcclRomeModel rome_model_22 = { static struct rcclRomeModel rome_model_25 = { .nGpus = 8, .nCpus = 4, .nNics = 2, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { 0x61000, 0xa1000, }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, }, .nicNuma = { 0, 3, }, .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, }, @@ -59,6 +62,7 @@ static struct rcclRomeModel rome_model_25 = { static struct rcclRomeModel rome_model_27 = { .nGpus = 8, .nCpus = 4, .nNics = 2, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { 0x61000, 0xa1000, }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, }, .nicNuma = { 0, 3, }, .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, }, @@ -69,6 +73,7 @@ static struct rcclRomeModel rome_model_27 = { static struct rcclRomeModel rome_model_29 = { .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 3, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, }, + .nicIds = { 0xe1000, }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, }, .nicNuma = { 2, }, .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, }, @@ -79,6 +84,7 @@ static struct rcclRomeModel rome_model_29 = { static struct rcclRomeModel rome_model_31 = { .nGpus = 8, .nCpus = 8, .nNics = 2, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { 0x61000, 0xa1000, }, .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, }, .nicNuma = { 0, 6, }, .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, }, @@ -89,6 +95,7 @@ static struct rcclRomeModel rome_model_31 = { static struct rcclRomeModel rome_model_33 = { .nGpus = 8, .nCpus = 8, .nNics = 2, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { 0x61000, 0xa1000, }, .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, }, .nicNuma = { 0, 6, }, .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, }, @@ -99,6 +106,7 @@ static struct rcclRomeModel rome_model_33 = { static struct rcclRomeModel rome_model_30 = { .nGpus = 8, .nCpus = 8, .nNics = 0, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, }, .nicNuma = { }, .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, }, @@ -109,6 +117,7 @@ static struct rcclRomeModel rome_model_30 = { static struct rcclRomeModel rome_model_32 = { .nGpus = 8, .nCpus = 8, .nNics = 0, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, }, .nicNuma = { }, .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, }, @@ -119,6 +128,7 @@ static struct rcclRomeModel rome_model_32 = { static struct rcclRomeModel rome_model_24 = { .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, }, .nicNuma = { }, .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, }, @@ -129,6 +139,7 @@ static struct rcclRomeModel rome_model_24 = { static struct rcclRomeModel rome_model_26 = { .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, }, .nicNuma = { }, .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, }, @@ -139,6 +150,7 @@ static struct rcclRomeModel rome_model_26 = { static struct rcclRomeModel rome_model_23 = { .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, }, .nicNuma = { }, .connMatrix = { 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, }, @@ -149,6 +161,7 @@ static struct rcclRomeModel rome_model_23 = { static struct rcclRomeModel rome_model_38 = { .nGpus = 8, .nCpus = 7, .nNics = 0, .nLinks = 2, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, }, .nicNuma = { }, .connMatrix = { 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, }, @@ -159,6 +172,7 @@ static struct rcclRomeModel rome_model_38 = { static struct rcclRomeModel rome_model_28 = { .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 3, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, }, + .nicIds = { }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, }, .nicNuma = { }, .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, }, @@ -169,6 +183,7 @@ static struct rcclRomeModel rome_model_28 = { static struct rcclRomeModel rome_model_40 = { .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 3, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, }, + .nicIds = { 0xe1000, }, .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, }, .nicNuma = { 2, }, .connMatrix = { 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 1, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, }, @@ -179,6 +194,7 @@ static struct rcclRomeModel rome_model_40 = { static struct rcclRomeModel rome_model_42 = { .nGpus = 8, .nCpus = 7, .nNics = 1, .nLinks = 3, .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, }, + .nicIds = { 0xe1000, }, .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, }, .nicNuma = { 4, }, .connMatrix = { 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 1, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, }, @@ -189,11 +205,12 @@ static struct rcclRomeModel rome_model_42 = { static struct rcclRomeModel rome_model_44 = { .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 3, .gpuIds = { 0x63000, 0x43000, 0x27000, 0x3000, 0xe3000, 0xc3000, 0xa3000, 0x83000, }, + .nicIds = { 0xc4000, }, .gpuNuma = { 0, 0, 1, 1, 2, 2, 3, 3, }, .nicNuma = { 2, }, .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, }, .pattern = "20202120", - .ringBase = "4 5 7 6 2 1 3 0|7 4 5 6 2 0 1 3", + .ringBase = "5 4 7 6 2 1 3 0|5 6 7 4 1 0 2 3", }; static struct rcclRomeModel romeTopoModels[] = { diff --git a/src/graph/search.cc b/src/graph/search.cc index e6afe49a66..0a887b6a1e 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -905,6 +905,7 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo } if (j >= romeTopo->nNics) { net_map[j] = i; + romeTopo->nicIds[romeTopo->nNics] = system->nodes[NET].nodes[i].net.busId; (romeTopo->nNics)++; if (romeTopo->nNics >= MAX_ROME_NICS) break; } @@ -941,6 +942,9 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo fprintf(file, " .gpuIds = { "); for (int i = 0; i < romeTopo->nGpus; i ++) fprintf(file, "0x%lx, ", romeTopo->gpuIds[i]); fprintf(file, "},\n"); + fprintf(file, " .nicIds = { "); + for (int i = 0; i < romeTopo->nNics; i ++) fprintf(file, "0x%lx, ", romeTopo->nicIds[i]); + fprintf(file, "},\n"); fprintf(file, " .gpuNuma = { "); for (int i = 0; i < romeTopo->nGpus; i ++) fprintf(file, "%ld, ", romeTopo->gpuNuma[i]); fprintf(file, "},\n"); @@ -1038,13 +1042,18 @@ static ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopo } char line[1024]; +#ifdef ENABLE_TRACE sprintf(line, "Found matching Rome model index %d in %.2fms (%d iter) with GPU mapping: ", i, t, time); +#else + sprintf(line, "Found matching Rome model index %d with GPU mapping: ", i); +#endif int offset = strlen(line); for (int k = 0; k < ngpus; k++) { sprintf(line+offset, "%d ", g[k]); offset = strlen(line); } INFO(NCCL_GRAPH, "%s", line); + // create 4P2H based on reference and remapped ids NCCLCHECK(parseGraph(romeTopoModels[i].ringBase, system, graph, g, romeTopo.nNics, net_map)); return ncclSuccess; diff --git a/src/graph/topo.cc b/src/graph/topo.cc index 18fff02275..73463a0dfb 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -266,7 +266,7 @@ ncclResult_t ncclTopoSortSystem(struct ncclTopoSystem* system) { return ncclSuccess; } -ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* system, struct ncclTopoNode* nic) { +ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* system, struct ncclTopoNode* nic, int64_t busId) { int dev; NCCLCHECK(xmlGetAttrInt(xmlNet, "dev", &dev)); @@ -286,6 +286,7 @@ ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* s if (xmlGetAttrInt(xmlNet, "gdr", &net->net.gdrSupport) != ncclSuccess) net->net.gdrSupport = 0; if (xmlGetAttrInt(xmlNet, "maxconn", &net->net.maxChannels) != ncclSuccess) net->net.maxChannels = MAXCHANNELS; if (ncclCollNet && xmlGetAttrInt(xmlNet, "coll", &net->net.collSupport) != ncclSuccess) net->net.collSupport = 0; + net->net.busId = busId; ncclDebugNoWarn = 0; NCCLCHECK(ncclTopoConnectNodes(nic, net, LINK_NET, net->net.width)); @@ -293,14 +294,14 @@ ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* s return ncclSuccess; } -ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* system, struct ncclTopoNode* nic) { +ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* system, struct ncclTopoNode* nic, int64_t busId) { for (int s=0; snSubs; s++) { struct ncclXmlNode* xmlNet = xmlNic->subs[s]; if (strcmp(xmlNet->name, "net") != 0) continue; int index; NCCLCHECK(xmlGetAttrIndex(xmlNet, "dev", &index)); if (index == -1) continue; - NCCLCHECK(ncclTopoAddNet(xmlNet, system, nic)); + NCCLCHECK(ncclTopoAddNet(xmlNet, system, nic, busId)); } return ncclSuccess; } @@ -354,7 +355,7 @@ ncclResult_t ncclTopoAddPci(struct ncclXmlNode* xmlPci, struct ncclTopoSystem* s NCCLCHECK(ncclTopoCreateNode(system, &nicNode, type, busId)); node = nicNode; // Connect it to parent later on } - NCCLCHECK(ncclTopoAddNic(xmlNic, system, nicNode)); + NCCLCHECK(ncclTopoAddNic(xmlNic, system, nicNode, busId)); } else if (type == PCI) { NCCLCHECK(ncclTopoCreateNode(system, &node, type, busId)); for (int s=0; snSubs; s++) { @@ -421,7 +422,7 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s NCCLCHECK(ncclTopoConnectNodes(cpu, nic, LINK_PCI, LOC_WIDTH)); NCCLCHECK(ncclTopoConnectNodes(nic, cpu, LINK_PCI, LOC_WIDTH)); } - NCCLCHECK(ncclTopoAddNic(node, system, nic)); + NCCLCHECK(ncclTopoAddNic(node, system, nic, 0)); } } return ncclSuccess; diff --git a/src/graph/topo.h b/src/graph/topo.h index 4b97b00f9c..2ae0096e9e 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -100,6 +100,7 @@ struct ncclTopoNode { int gdrSupport; int collSupport; int maxChannels; + int64_t busId; }net; struct { int arch;