From 961922ea02a800447888d7dff2751002ab1e7912 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Thu, 3 Jun 2021 19:45:18 -0700 Subject: [PATCH] Add option to enable multiple SAT in SHARP (#380) * Add option to enable multiple SAT in SHARP * Extend number of NICs to 16 --- src/graph/rome_models.cc | 2 +- src/graph/search.cc | 3 ++ tools/TopoVisual/extract_topo.awk | 72 ++++++++++++------------------- tools/scripts/topo_val.sh | 5 ++- tools/topo_expl/topo_expl.cpp | 1 + 5 files changed, 36 insertions(+), 47 deletions(-) diff --git a/src/graph/rome_models.cc b/src/graph/rome_models.cc index f5b5229394..29a0df3207 100755 --- a/src/graph/rome_models.cc +++ b/src/graph/rome_models.cc @@ -29,7 +29,7 @@ THE SOFTWARE. #define MAX_ROME_CPUS 8 #define MAX_ROME_GPUS 16 -#define MAX_ROME_NICS 8 +#define MAX_ROME_NICS 16 struct rcclRomeModel { int nGpus; diff --git a/src/graph/search.cc b/src/graph/search.cc index 8465bf7124..98edf7f105 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -765,6 +765,7 @@ float speedArray[] = { 42.0, 30.0, 24.0, 21.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, #define NSPEEDS (sizeof(speedArray)/sizeof(float)) RCCL_PARAM(ModelMatchingDisable, "MODEL_MATCHING_DISABLE", 0); +RCCL_PARAM(EnableMultipleSAT, "ENABLE_MULTIPLE_SAT", 0); ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph) { int ngpus = system->nodes[GPU].count; @@ -809,6 +810,8 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph // limit single node max channels when searching ring graph on Rome graph->maxChannels = 2; } + if (graph->collNet && !rcclParamEnableMultipleSAT()) + graph->maxChannels = 1; if (ngpus == 1) if (graph->pattern != NCCL_TOPO_PATTERN_RING) graph->pattern = NCCL_TOPO_PATTERN_TREE; #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) diff --git a/tools/TopoVisual/extract_topo.awk b/tools/TopoVisual/extract_topo.awk index b9da532341..c6028666c0 100755 --- a/tools/TopoVisual/extract_topo.awk +++ b/tools/TopoVisual/extract_topo.awk @@ -29,6 +29,7 @@ BEGIN { has_collnet=0 max_collnet=0 max_collnet_rank=0 + max_collnet_channel=0 collnet[""]=0 collnet_conn[""]=0 collnet_conn_type[""]=0 @@ -111,17 +112,22 @@ BEGIN { } while ($col_1!="") } - if($col_start=="CollNet" && $col_p1=="Channel") { - chan=strtonum($col_p2) - rank=strtonum($col_p4) - up_rank=strtonum($col_p6) - collnet[up_rank "," rank "," chan]="1" - if(has_collnet==0) - has_collnet=1 - if(chan>max_collnet) - max_collnet=chan + if($col_start=="CollNet" && $col_p1=="channel" && $col_p5=="down") { + channel=strtonum($col_p2) + up_rank=strtonum($col_p4) if(up_rank>max_collnet_rank) max_collnet_rank=up_rank + for(s=col_p6;s<=NF;s++) { + if($s=="nDown") break; + rank=$s + collnet[up_rank "," rank]="1" + if(rank>max_collnet_rank) + max_collnet_rank=rank + } + if(has_collnet==0) + has_collnet=1 + if(channel>max_collnet_channel) + max_collnet_channel=channel } if($col_start=="Coll" && $col_p2==":") { @@ -134,6 +140,8 @@ BEGIN { else printf "Error!\n" collnet_conn_type[rank "," chan]=$col_p6 + if(chan>max_collnet) + max_collnet=chan } if($col_p6=="via") { @@ -229,8 +237,9 @@ END { for(r=0; has_collnet && r<=max_collnet; r++) { printf " subgraph collnet_%d {\n", r num_top_ranks=0 - for(s=0;s c%d_%d [dir=back label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, max_collnet_rank, r, rank, val, color, style - else - printf " c%d_%d -> c%d_%d [label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, max_collnet_rank, r, rank, val, color, style - while(1) { - for(s=0;s c%d_%d [dir=back label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, rank, r, s, val, color, style - else - printf " c%d_%d -> c%d_%d [label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, rank, r, s, val, color, style - rank=s - break; - } - } - if(s>=max_collnet_rank) { - break; + printf " c%d_%d -> c%d_%d [label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, rank_switch, r, rank, val, color, style + for(s=0;s<=max_collnet_rank;s++) { + if((rank "," s) in collnet) { + style="solid" + color="green" + printf " c%d_%d -> c%d_%d [label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, rank, r, s, "", color, style } } } @@ -285,6 +266,7 @@ END { for(s=0;s<=max_collnet_rank;s++) { printf " c%d_%d [label=\"%d\",fontsize=\"28\"];\n", r, s, s } + printf " c%d_%d [label=\"SHARP:%d\",fontsize=\"28\"];\n", r, rank_switch, r printf " }\n\n" } printf "}\n" diff --git a/tools/scripts/topo_val.sh b/tools/scripts/topo_val.sh index 3a152a4bd4..0f05a9fe56 100755 --- a/tools/scripts/topo_val.sh +++ b/tools/scripts/topo_val.sh @@ -21,11 +21,14 @@ DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -for i in {0..53} +for i in {0..54} do if [[ $i -eq 50 ]] || [[ $i -eq 51 ]] then NCCL_COLLNET_ENABLE=1 $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" + elif [[ $i -eq 54 ]] + then + RCCL_ENABLE_MULTIPLE_SAT=1 NCCL_COLLNET_ENABLE=1 $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" else $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" fi diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index 85ec4ed870..52796ee21f 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -123,6 +123,7 @@ NodeModelDesc model_descs[] = { {4, "topo_collnet_n4.xml", "4 nodes collnet 4 NICs"}, {1, "topo_8p_90a.xml", "single node gfx90a"}, {4, "topo_8p_rome_4n_1.xml", "4 nodes node 8 gfx908 Rome 4 NICs NPS=4"}, + {4, "topo_collnet_n4.xml", "4 nodes collnet 4 NICs for multiple SAT"}, }; int main(int argc,char* argv[])