diff --git a/src/graph/paths.cc b/src/graph/paths.cc index a21f11327c..d4d1009ddc 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -576,6 +576,8 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclPeer return ncclSuccess; } +RCCL_PARAM(EnableIntranet, "ENABLE_INTRANET", 0); + ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* comm) { int *domains; int64_t *ids; @@ -669,6 +671,10 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* } } } + if (rcclParamEnableIntranet()) { + remove = 0; + system->type |= RCCL_TOPO_FORCE_INTRA; + } comm->localRanks = system->nodes[GPU].count; if (system->nodes[GPU].count == comm->nRanks && remove) { for (int n=system->nodes[NET].count-1; n>=0; n--) @@ -741,7 +747,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { else { // Round to next pow2 nChannelsPerPeer and nChannels comm->p2pnChannelsPerPeer = nextPow2(minChannels); - comm->p2pnChannels = nextPow2(comm->p2pnChannels); + comm->p2pnChannels = std::min(comm->topo->nodes[GPU].count == comm->topo->nRanks ? 2*comm->nRanks : comm->nRanks, nextPow2(comm->p2pnChannels)); } // Init channels that weren't used so far diff --git a/src/graph/rome_models.cc b/src/graph/rome_models.cc index 85a63e8df1..64fd7c94fa 100644 --- a/src/graph/rome_models.cc +++ b/src/graph/rome_models.cc @@ -573,8 +573,8 @@ static struct rcclRomeModel rome_model_76 = { .connMatrix = { 0, 4, 1, 0, 0, 0, 2, 0, 4, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 4, 2, 0, 0, 0, 0, 1, 4, 0, 0, 0, 0, 1, 0, 0, 2, 0, 0, 4, 1, 0, 0, 1, 0, 0, 4, 0, 0, 1, 2, 0, 0, 0, 1, 0, 0, 4, 0, 0, 0, 1, 0, 1, 4, 0, }, .gdrLevel = { PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, }, .pattern = "4444", - .ringBase = "N4 4 5 7 6 0 1 3 2 N2|N2 2 3 1 0 6 7 5 4 N4|N7 7 6 0 1 5 4 2 3 N3|N1 1 0 6 7 3 2 4 5 N5|N3 3 2 4 5 7 6 0 1 N1|N5 5 4 2 3 1 0 6 7 N7|N0 0 1 5 4 2 3 7 6 N6|N6 6 7 3 2 4 5 1 0 N0", - .options = "tuning=3", + .ringBase = "N0 0 1 3 2 4 5 7 6 N6|N2 2 3 1 0 6 7 5 4 N4|N5 5 4 2 3 7 6 0 1 N1|N1 1 0 6 7 3 2 4 5 N5|N4 4 5 7 6 0 1 3 2 N2|N2 2 3 1 0 6 7 5 4 N4|N0 0 1 5 4 2 3 7 6 N6|N3 3 2 4 5 1 0 6 7 N7|N4 4 5 7 6 0 1 3 2 N2|N6 6 7 5 4 2 3 1 0 N0|N7 7 6 0 1 5 4 2 3 N3|N6 6 7 3 2 4 5 1 0 N0|N3 3 2 0 1 5 4 6 7 N7|N1 1 0 2 3 7 6 4 5 N5|N5 5 4 6 7 3 2 0 1 N1|N7 7 6 4 5 1 0 2 3 N3", + .options = "disableNumaMatching=1,tuning=3", .treeBase = "", }; diff --git a/src/graph/topo.h b/src/graph/topo.h index 38e8c7cfe1..0aa08c3376 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -88,6 +88,7 @@ struct ncclTopoLinkList { #define RCCL_TOPO_4P2H_ROME 2 #define RCCL_TOPO_GDR_ALL 4 #define RCCL_TOPO_16P1H 8 +#define RCCL_TOPO_FORCE_INTRA 16 struct ncclTopoNode { int type; diff --git a/src/init.cc b/src/init.cc index fbbcf2ae34..ca542fcda4 100644 --- a/src/init.cc +++ b/src/init.cc @@ -817,7 +817,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm if ((comm->topo->type & RCCL_TOPO_4P2H_ROME) && (comm->topo->type & RCCL_TOPO_GDR_ALL)) { if (rcclParamP2pNetDisable() == 0) { - comm->p2pNet = 1; + if (!(comm->topo->type & RCCL_TOPO_FORCE_INTRA)) comm->p2pNet = 1; INFO(NCCL_INIT, "RCCL enabled same node P2P over network"); } else diff --git a/tools/scripts/topo_val.sh b/tools/scripts/topo_val.sh index f71879eefe..1dc06dd074 100755 --- a/tools/scripts/topo_val.sh +++ b/tools/scripts/topo_val.sh @@ -21,7 +21,7 @@ DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -for i in {0..76} +for i in {0..78} do if [[ $i -eq 50 ]] || [[ $i -eq 51 ]] then diff --git a/tools/topo_expl/models/topo_8p1h_5.xml b/tools/topo_expl/models/topo_8p1h_5.xml new file mode 100644 index 0000000000..5213f7ac8a --- /dev/null +++ b/tools/topo_expl/models/topo_8p1h_5.xml @@ -0,0 +1,146 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + \ No newline at end of file diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index d903ea3f8f..afbacc5c33 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -147,6 +147,8 @@ NodeModelDesc model_descs[] = { {4, "topo_8p1h_3.xml", "4 nodes 8P1H Alt."}, {1, "topo_8p1h_4.xml", "Single node 8P1H Alt."}, {2, "topo_8p1h_4.xml", "2 nodes 8P1H Alt."}, + {1, "topo_8p1h_5.xml", "Single node 8P1H Alt."}, + {2, "topo_8p1h_5.xml", "2 nodes 8P1H Alt."}, }; int main(int argc,char* argv[])