diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index 7bdeb1ec8c..67f4df4462 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -510,8 +510,22 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* } if (ret) { system->type |= RCCL_TOPO_GDR_ALL; - remove = 0; - INFO(NCCL_GRAPH, "GDR is available on all GPUs"); + bool allXgmi = true; + // don't trim NICs unless all GPUs are connected by XGMI + for (int i = 0; i < system->nodes[GPU].count && allXgmi; i++) { + int cudaDev1 = system->nodes[GPU].nodes[i].gpu.dev; + for (int j = 0; j < system->nodes[GPU].count && allXgmi; j++) { + if (i == j) continue; + int cudaDev2 = system->nodes[GPU].nodes[j].gpu.dev; + bool isXGMI; + NCCLCHECK(ncclTopoGetLinkType(comm->topo, cudaDev1, cudaDev2, &isXGMI)); + allXgmi &= isXGMI; + } + } + if (!allXgmi) { + remove = 0; + INFO(NCCL_GRAPH, "GDR is available on all GPUs"); + } } } comm->localRanks = system->nodes[GPU].count; diff --git a/projects/rccl/src/graph/search.cc b/projects/rccl/src/graph/search.cc index dbd38c8d61..ed71511f1d 100644 --- a/projects/rccl/src/graph/search.cc +++ b/projects/rccl/src/graph/search.cc @@ -1046,7 +1046,8 @@ ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, str return ncclSuccess; } -ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, bool direct_only) { +ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, int maxInter, int nInter, int *inter) { + int interGpus[MAX_XGMI_INTER_GPUS+1]; int ngpus = system->nodes[GPU].count; *isXGMI = false; // check for direct XGMI connection @@ -1065,15 +1066,39 @@ ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, in } } } - if (direct_only) return ncclSuccess; - // check if there is intermediate GPU that is connected to both + if (maxInter == 0) return ncclSuccess; + // check if there are intermediate GPUs that are connected to both + bool res1, res2, res3; + int j; + for (j=0; j 0 && inter != nullptr) { + ncclTopoGetLinkType(system, inter[nInter], cudaDev2, &res2, 0); + if (res2) { + *isXGMI = true; + return ncclSuccess; + } + memcpy(interGpus+1, inter+1, sizeof(int)*nInter); + } + interGpus[0] = cudaDev1; + // add one more intermediate GPU recursively util reaching max depth + nInter++; + if (nInter+2 > ngpus || nInter > MAX_XGMI_INTER_GPUS || nInter > maxInter) return ncclSuccess; for (int i=0; inodes[GPU].nodes[i].gpu.dev == cudaDev1 || system->nodes[GPU].nodes[i].gpu.dev == cudaDev2) - continue; - bool res1, res2; - ncclTopoGetLinkType(system, system->nodes[GPU].nodes[i].gpu.dev, cudaDev1, &res1, true); - ncclTopoGetLinkType(system, system->nodes[GPU].nodes[i].gpu.dev, cudaDev2, &res2, true); - if (res1 && res2) { + int dev = system->nodes[GPU].nodes[i].gpu.dev; + // skip duplicated GPU + if (dev == cudaDev2) continue; + for (j=0; jlocalRanks == comm->nRanks) + if (comm->localRanks == comm->nRanks && comm->topo->nodes[GPU].nodes[0].gpu.gcn != 910) { // Check that all the GPUs have peer access to one another and are XGMI connected bool hasPeerAccess = true; @@ -855,7 +855,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm } bool isXGMI; - NCCLCHECK(ncclTopoGetLinkType(comm->topo, i, j, &isXGMI)); + // Limit to single intermediate GPU for enabling clique + NCCLCHECK(ncclTopoGetLinkType(comm->topo, i, j, &isXGMI, 1)); allXgmi &= isXGMI; } } diff --git a/projects/rccl/tools/topo_expl/models/topo_16p1h.xml b/projects/rccl/tools/topo_expl/models/topo_16p1h.xml index 4c6a229c0a..9225986e4a 100644 --- a/projects/rccl/tools/topo_expl/models/topo_16p1h.xml +++ b/projects/rccl/tools/topo_expl/models/topo_16p1h.xml @@ -1,19 +1,19 @@ - - - - - + + + + + - - - + + + @@ -21,11 +21,16 @@ + + + + + - - - - + + + + @@ -33,33 +38,38 @@ - - - + + + + + + + + - - - - - + + + + + - - - + + + @@ -67,16 +77,16 @@ - + - + - - - - + + + + @@ -84,33 +94,38 @@ - - - + + + + + + + + - - - - - + + + + + - - - + + + @@ -118,11 +133,16 @@ + + + + + - - - - + + + + @@ -130,33 +150,38 @@ - - - + + + + + + + + - - - - - + + + + + - - - + + + @@ -164,11 +189,16 @@ + + + + + - - - - + + + + @@ -176,16 +206,26 @@ - - - + + + + + + + + + + + + + diff --git a/projects/rccl/tools/topo_expl/utils.cpp b/projects/rccl/tools/topo_expl/utils.cpp index 0dd240a03c..95d3702cef 100644 --- a/projects/rccl/tools/topo_expl/utils.cpp +++ b/projects/rccl/tools/topo_expl/utils.cpp @@ -576,7 +576,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t bool allXgmi = true; { // [RCCL] Check if clique-based kernels can be enabled and initialize CliqueManager //CliqueManager::cliqueMode_t cliqueMode = CliqueManager::CLIQUE_DISABLED; - if (comm->localRanks == comm->nRanks) + if (comm->localRanks == comm->nRanks && comm->topo->nodes[GPU].nodes[0].gpu.gcn != 910) { // Check that all the GPUs have peer access to one another and are XGMI connected bool hasPeerAccess = true; @@ -595,7 +595,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t //} bool isXGMI; - NCCLCHECK(ncclTopoGetLinkType(comm->topo, i, j, &isXGMI)); + NCCLCHECK(ncclTopoGetLinkType(comm->topo, i, j, &isXGMI, 1)); allXgmi &= isXGMI; } }