AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

Add option to enable multiple SAT in SHARP (#380)

Quellcode durchsuchen 
* Add option to enable multiple SAT in SHARP

* Extend number of NICs to 16
Dieser Commit ist enthalten in:
Wenkai Du
2021-06-03 19:45:18 -07:00
committet von GitHub
Ursprung 903c84050d
Commit 961922ea02
5 geänderte Dateien mit 36 neuen und 47 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
src/graph/rome_models.cc
+1 -1
Datei anzeigen
@@ -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;
src/graph/search.cc
+3
Datei anzeigen
@@ -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__)
tools/TopoVisual/extract_topo.awk
+27 -45
Datei anzeigen
@@ -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<max_collnet_rank;s++) {
if((max_collnet_rank "," s "," r) in collnet)
rank_switch=max_collnet_rank+1
for(s=0;s<=max_collnet_rank;s++) {
if((s "," r) in collnet_conn_type)
top_ranks[num_top_ranks++]=s
}
for(d=0; d<num_top_ranks; d++) {
@@ -244,40 +253,12 @@ END {
if (match(val,"GDRDMA"))
color="green"
}
if (match(val,"P2P")) {
color="green"
}
if(send)
printf " c%d_%d -> 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<max_collnet_rank;s++) {
if((rank "," s "," r) in collnet) {
if(send)
val=conn[s "," rank "," r]
else
val=conn[rank "," s "," r]
style="solid"
color="red"
if (match(val,"NET")) {
style="dashed"
if (match(val,"GDRDMA"))
color="green"
}
if (match(val,"P2P")) {
color="green"
}
if(send)
printf " c%d_%d -> 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"
tools/scripts/topo_val.sh
+4 -1
Datei anzeigen
@@ -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
tools/topo_expl/topo_expl.cpp
+1
Datei anzeigen
@@ -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[])