diff --git a/src/graph/rome_models.cc b/src/graph/rome_models.cc index c90677f332..b8abffa17e 100644 --- a/src/graph/rome_models.cc +++ b/src/graph/rome_models.cc @@ -37,6 +37,7 @@ struct rcclRomeModel { int64_t gpuNuma[NCCL_TOPO_MAX_NODES]; int64_t nicNuma[NCCL_TOPO_MAX_NODES]; uint8_t connMatrix[NCCL_TOPO_MAX_NODES*NCCL_TOPO_MAX_NODES]; + uint8_t gdrLevel[NCCL_TOPO_MAX_NODES*NCCL_TOPO_MAX_NODES]; const char *pattern; const char *ringBase; }; @@ -48,6 +49,7 @@ static struct rcclRomeModel rome_model_22 = { .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, }, + .gdrLevel = { 6, 6, 6, 5, 6, 6, 5, 6, }, .pattern = "10302120", .ringBase = "7 4 5 3 1 0 6 2|4 7 3 5 0 1 2 6", }; @@ -59,6 +61,7 @@ static struct rcclRomeModel rome_model_25 = { .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, }, + .gdrLevel = { 5, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 5, }, .pattern = "11303011", .ringBase = "2 1 0 3 6 7 5 4|7 6 4 5 1 2 3 0", }; @@ -70,6 +73,7 @@ static struct rcclRomeModel rome_model_27 = { .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, }, + .gdrLevel = { 5, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 5, }, .pattern = "11303011", .ringBase = "0 6 2 3 1 7 5 4|7 1 4 5 6 0 3 2", }; @@ -81,6 +85,7 @@ static struct rcclRomeModel rome_model_29 = { .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, }, + .gdrLevel = { 6, 6, 6, 6, 5, 5, 6, 6, }, .pattern = "10302120", .ringBase = "6 5 7 4 0 1 3 2|6 4 7 5 2 3 1 0", }; @@ -92,6 +97,7 @@ static struct rcclRomeModel rome_model_31 = { .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, }, + .gdrLevel = { 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, }, .pattern = "0110201010200110", .ringBase = "1 2 3 0 6 4 5 7|4 6 7 5 2 1 0 3", }; @@ -103,6 +109,7 @@ static struct rcclRomeModel rome_model_33 = { .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, }, + .gdrLevel = { 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, 6, }, .pattern = "0110201010200110", .ringBase = "1 4 5 7 0 3 2 6|4 1 7 5 6 2 3 0", }; @@ -114,6 +121,7 @@ static struct rcclRomeModel rome_model_30 = { .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, }, + .gdrLevel = { }, .pattern = "0010201010200010", .ringBase = "3 0 1 2 6 7 5 4|2 1 0 3 7 6 4 5", }; @@ -125,6 +133,7 @@ static struct rcclRomeModel rome_model_32 = { .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, }, + .gdrLevel = { }, .pattern = "0010201010200010", .ringBase = "0 6 2 3 4 5 7 1|3 2 6 0 1 7 5 4", }; @@ -136,6 +145,7 @@ static struct rcclRomeModel rome_model_24 = { .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, }, + .gdrLevel = { }, .pattern = "10303010", .ringBase = "0 1 2 3 5 7 6 4|1 0 3 2 7 5 4 6", }; @@ -147,6 +157,7 @@ static struct rcclRomeModel rome_model_26 = { .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, }, + .gdrLevel = { }, .pattern = "10303010", .ringBase = "4 5 7 1 0 3 2 6|3 0 6 2 1 7 5 4", }; @@ -158,6 +169,7 @@ static struct rcclRomeModel rome_model_23 = { .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, }, + .gdrLevel = { }, .pattern = "10302020", .ringBase = "1 7 6 4 5 2 0 3|2 5 3 0 4 6 7 1", }; @@ -169,6 +181,7 @@ static struct rcclRomeModel rome_model_38 = { .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, }, + .gdrLevel = { }, .pattern = "10201000201010", .ringBase = "6 7 1 4 3 5 2 0|0 2 5 3 4 1 7 6", }; @@ -180,6 +193,7 @@ static struct rcclRomeModel rome_model_28 = { .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, }, + .gdrLevel = { }, .pattern = "10302020", .ringBase = "0 3 2 1 4 5 6 7|7 6 5 4 1 2 3 0|0 2 5 7 4 6 3 1|1 3 6 4 7 5 2 0", }; @@ -191,6 +205,7 @@ static struct rcclRomeModel rome_model_40 = { .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, }, + .gdrLevel = { 6, 6, 6, 6, 5, 5, 6, 6, }, .pattern = "10302120", .ringBase = "6 7 1 4 0 5 3 2|7 6 4 1 0 2 3 5", }; @@ -202,6 +217,7 @@ static struct rcclRomeModel rome_model_42 = { .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, }, + .gdrLevel = { 6, 6, 6, 6, 6, 6, 6, 6, }, .pattern = "10201001201010", .ringBase = "7 4 6 1 3 0 2 5|6 4 7 1 3 2 5 0", }; @@ -213,6 +229,7 @@ static struct rcclRomeModel rome_model_44 = { .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, }, + .gdrLevel = { 6, 6, 6, 6, 5, 5, 6, 6, }, .pattern = "20202120", .ringBase = "5 4 7 6 2 1 3 0|5 6 7 4 1 0 2 3", }; @@ -224,6 +241,7 @@ static struct rcclRomeModel rome_model_45 = { .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, }, .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, }, + .gdrLevel = { }, .pattern = "10201000201010", .ringBase = "0 1 2 3 4 5 6 7|0 2 5 7 4 6 1 3|0 3 1 6 4 7 5 2|0 7 6 5 4 3 2 1", }; @@ -235,6 +253,7 @@ static struct rcclRomeModel rome_model_46 = { .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, }, .nicNuma = { 4, }, .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, }, + .gdrLevel = { 6, 6, 6, 6, 6, 6, 6, 6, }, .pattern = "10201001201010", .ringBase = "6 5 7 4 1 2 3 0|7 4 6 5 1 0 3 2", }; @@ -246,6 +265,7 @@ static struct rcclRomeModel rome_model_48 = { .gpuNuma = { 0, 0, 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, }, + .gdrLevel = { }, .pattern = "20202020", .ringBase = "0 1 2 3 4 5 6 7|7 6 5 4 3 2 1 0|0 1 2 3 4 5 6 7|7 6 5 4 3 2 1 0", }; @@ -257,6 +277,7 @@ static struct rcclRomeModel rome_model_49 = { .gpuNuma = { 0, 0, 1, 1, 2, 2, 3, 3, }, .nicNuma = { 0, 1, 2, 3, }, .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, }, + .gdrLevel = { 4, 4, 6, 6, 6, 6, 6, 6, 6, 6, 4, 4, 6, 6, 6, 6, 6, 6, 6, 6, 4, 4, 6, 6, 6, 6, 6, 6, 6, 6, 4, 4, }, .pattern = "21212121", .ringBase = "N0 0 1 2 3 4 5 6 7 N3|N3 7 6 5 4 3 2 1 0 N0|N1 2 3 0 1 6 7 4 5 N2|N2 5 4 7 6 1 0 3 2 N1", }; @@ -268,6 +289,7 @@ static struct rcclRomeModel rome_model_52 = { .gpuNuma = { 0, 0, 0, 0, 0, 0, 0, 0, }, .nicNuma = { }, .connMatrix = { 0, 1, 1, 0, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 1, 1, 0, }, + .gdrLevel = { }, .pattern = "80", .ringBase = "0 1 3 2 4 5 7 6|6 7 5 4 2 3 1 0|0 1 5 4 6 7 3 2|2 3 7 6 4 5 1 0", }; @@ -279,6 +301,7 @@ static struct rcclRomeModel rome_model_53 = { .gpuNuma = { 1, 1, 3, 3, 5, 5, 7, 7, }, .nicNuma = { 1, 3, 5, 7, }, .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, }, + .gdrLevel = { 4, 4, 6, 6, 6, 6, 6, 6, 6, 6, 4, 4, 6, 6, 6, 6, 6, 6, 6, 6, 4, 4, 6, 6, 6, 6, 6, 6, 6, 6, 4, 4, }, .pattern = "21212121", .ringBase = "N0 0 1 2 3 4 5 6 7 N3|N3 7 6 5 4 3 2 1 0 N0|N1 2 3 0 1 6 7 4 5 N2|N2 5 4 7 6 1 0 3 2 N1", }; @@ -290,6 +313,7 @@ static struct rcclRomeModel rome_model_43 = { .gpuNuma = { 0, 0, 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, }, + .gdrLevel = { }, .pattern = "20202020", .ringBase = "0 1 2 3 4 5 6 7|0 2 5 7 4 6 1 3|0 3 1 6 4 7 5 2|0 7 6 5 4 3 2 1", }; @@ -301,10 +325,23 @@ static struct rcclRomeModel rome_model_55 = { .gpuNuma = { 0, 0, 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, }, + .gdrLevel = { }, .pattern = "20202020", .ringBase = "0 1 2 3 4 5 6 7|7 6 5 4 3 2 1 0|2 3 0 1 6 7 4 5|5 4 7 6 1 0 3 2", }; +static struct rcclRomeModel rome_model_56 = { + .nGpus = 16, .nCpus = 4, .nNics = 0, .nLinks = 4, + .gpuIds = { 0x4e000, 0x51000, 0x56000, 0x59000, 0xe000, 0x11000, 0x16000, 0x19000, 0xcf000, 0xd2000, 0xd7000, 0xda000, 0x8f000, 0x92000, 0x97000, 0x9a000, }, + .nicIds = { }, + .gpuNuma = { 0, 0, 0, 0, 1, 1, 1, 1, 2, 2, 2, 2, 3, 3, 3, 3, }, + .nicNuma = { }, + .connMatrix = { 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, }, + .gdrLevel = { }, + .pattern = "40404040", + .ringBase = "0 1 3 2 6 7 15 14 10 11 9 8 12 13 5 4|0 1 2 3 7 6 13 12 8 9 10 11 15 14 5 4|0 2 3 7 6 14 15 11 10 8 9 13 12 4 5 1|4 5 13 12 8 9 11 10 14 15 7 6 2 3 1 0|4 5 14 15 11 10 9 8 12 13 6 7 3 2 1 0|1 5 4 12 13 9 8 10 11 15 14 6 7 3 2 0", +}; + static struct rcclRomeModel romeTopoModels[] = { rome_model_22, rome_model_25, @@ -330,6 +367,7 @@ static struct rcclRomeModel romeTopoModels[] = { rome_model_53, rome_model_43, rome_model_55, + rome_model_56, }; /* Parse user defined rings. Format is like : @@ -338,7 +376,7 @@ static struct rcclRomeModel romeTopoModels[] = { * Rings with a non-matching number of gpus are ignored so we can provide * rings for multiple cases. */ -ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int* gpu_map) { +ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int* gpu_map, int* net_map) { int gpus[NCCL_TOPO_MAX_NODES]; int nChannels = 0; int gpu = 0; @@ -410,6 +448,10 @@ ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct n } if (net_count) { + for (int i = 0; net_map && i < ngpus*2; i++) { + if (nets[i]-'N' < 0 || nets[i]-'N' >= nnets) continue; + nets[i] = net_map[nets[i]-'N']+'N'; + } memcpy(&graph->intraNets[ngpus*nChannels*2], nets, ngpus*2*sizeof(int)); graph->nIntraChannels++; if (nets[0]-'N' >= nnets || nets[ngpus*2-1]-'N' >= nnets) goto newchannel; @@ -485,7 +527,7 @@ ncclResult_t parseChordalRing(struct ncclTopoSystem* system, struct ncclTopoGrap } // create chordal ring based on reference and remapped ids system->type |= RCCL_TOPO_CR8G; - NCCLCHECK(parseGraph(ringBase, system, graph, id)); + NCCLCHECK(parseGraph(ringBase, system, graph, id, NULL)); if (system->nodes[NET].count && system->nodes[GPU].count != system->nRanks) { int *intra, *used; graph->nChannels = system->nodes[NET].count; @@ -517,27 +559,6 @@ ncclResult_t parseChordalRing(struct ncclTopoSystem* system, struct ncclTopoGrap return ncclSuccess; } -struct ncclGpuIdHIP { - int g; - int dev; -}; - -static int cmpIds(const void * g1, const void * g2) { - struct ncclGpuIdHIP *s1 = (struct ncclGpuIdHIP*)g1; - struct ncclGpuIdHIP *s2 = (struct ncclGpuIdHIP*)g2; - return s1->dev - s2->dev; -} - -struct ncclCpuNuma { - int c; - uint64_t numa; -}; - -static int cmpNuma(const void * g1, const void * g2) { - struct ncclCpuNuma *s1 = (struct ncclCpuNuma*)g1; - struct ncclCpuNuma *s2 = (struct ncclCpuNuma*)g2; - return s1->numa - s2->numa; -} static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRomeModel* romeTopo, char *pattern) { pattern[0] = 0; // pattern will be NULL for invalid topology @@ -545,13 +566,47 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo romeTopo->nCpus = system->nodes[CPU].count; romeTopo->nNics = system->nodes[NET].count; romeTopo->nLinks = 0; + + struct ncclGpuIdHIP { + int g; + int dev; + }; + + auto cmpIds = [](const void * g1, const void * g2) { + struct ncclGpuIdHIP *s1 = (struct ncclGpuIdHIP*)g1; + struct ncclGpuIdHIP *s2 = (struct ncclGpuIdHIP*)g2; + return s1->dev - s2->dev; + }; + + struct ncclCpuNuma { + int c; + uint64_t numa; + }; + + auto cmpNuma = [](const void * g1, const void * g2) { + struct ncclCpuNuma *s1 = (struct ncclCpuNuma*)g1; + struct ncclCpuNuma *s2 = (struct ncclCpuNuma*)g2; + return (int)(s1->numa - s2->numa); + }; + + struct ncclNetId { + int n; + uint64_t id; + }; + + auto cmpNets = [](const void * g1, const void * g2) { + struct ncclNetId *s1 = (struct ncclNetId*)g1; + struct ncclNetId *s2 = (struct ncclNetId*)g2; + return (int)(s1->id - s2->id); + }; + // sort GPU devices by HIP device ID - struct ncclGpuIdHIP scores[NCCL_TOPO_MAX_NODES]; + struct ncclGpuIdHIP gpu_scores[NCCL_TOPO_MAX_NODES]; for (int i = 0; i < romeTopo->nGpus; i ++) { - scores[i].g = i; - scores[i].dev = system->nodes[GPU].nodes[i].gpu.dev; + gpu_scores[i].g = i; + gpu_scores[i].dev = system->nodes[GPU].nodes[i].gpu.dev; } - qsort(scores, romeTopo->nGpus, sizeof(struct ncclGpuIdHIP), cmpIds); + qsort(gpu_scores, romeTopo->nGpus, sizeof(struct ncclGpuIdHIP), cmpIds); // sort CPU devices by NUMA id struct ncclCpuNuma cpu_scores[NCCL_TOPO_MAX_NODES]; for (int i = 0; i < romeTopo->nCpus; i ++) { @@ -559,10 +614,17 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo cpu_scores[i].numa = system->nodes[CPU].nodes[i].id; } qsort(cpu_scores, romeTopo->nCpus, sizeof(struct ncclCpuNuma), cmpNuma); + // sort NET devices by id + struct ncclNetId net_scores[NCCL_TOPO_MAX_NODES]; + for (int i = 0; i < romeTopo->nNics; i ++) { + net_scores[i].n = i; + net_scores[i].id = system->nodes[NET].nodes[i].id; + } + qsort(net_scores, romeTopo->nNics, sizeof(struct ncclNetId), cmpNets); for (int i = 0; i < romeTopo->nGpus; i ++) { int gpu, n, m, distance; - gpu = scores[i].g; + gpu = gpu_scores[i].g; romeTopo->gpuIds[i] = system->nodes[GPU].nodes[gpu].id; m = 0; distance = system->nodes[GPU].nodes[gpu].paths[CPU][m].count; @@ -591,17 +653,18 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo if (romeTopo->nLinks < count) romeTopo->nLinks = count; } - for (int net = 0; net < romeTopo->nNics; net++) { + for (int i = 0; i < romeTopo->nNics; i++) { int n, m, distance; m = 0; - romeTopo->nicIds[net] = system->nodes[NET].nodes[net].net.busId; + int net = net_scores[i].n; + romeTopo->nicIds[i] = system->nodes[NET].nodes[net].net.busId; distance = system->nodes[NET].nodes[net].paths[CPU][m].count; for (n = 0; n < romeTopo->nCpus; n++) if (system->nodes[NET].nodes[net].paths[CPU][n].count < distance) { distance = system->nodes[NET].nodes[net].paths[CPU][n].count; m = n; } - if (m < romeTopo->nCpus) romeTopo->nicNuma[net] = system->nodes[CPU].nodes[m].id; + if (m < romeTopo->nCpus) romeTopo->nicNuma[i] = system->nodes[CPU].nodes[m].id; else return ncclSuccess; } @@ -618,6 +681,15 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo } pattern[romeTopo->nCpus*2] = 0; + // compute gdr level matrix + for (int i = 0; i < romeTopo->nNics; i++) { + int n = net_scores[i].n; + for (int j = 0; j < romeTopo->nGpus; j++) { + int g = gpu_scores[j].g; + romeTopo->gdrLevel[i*romeTopo->nGpus+j] = system->nodes[GPU].nodes[g].paths[NET][n].type; + } + } + const char* romeModelFile = getenv("RCCL_DUMP_ROME_MODEL_FILE"); if (romeModelFile) { INFO(NCCL_ENV, "RCCL_DUMP_ROME_MODEL_FILE set by environment to %s", romeModelFile); @@ -644,6 +716,10 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo for (int i = 0; i < romeTopo->nGpus; i ++) for (int n = 0; n < romeTopo->nGpus; n++) fprintf(file, "%d, ", romeTopo->connMatrix[i*romeTopo->nGpus+n]); fprintf(file, "},\n"); + fprintf(file, " .gdrLevel = { "); + for (int i = 0; i < romeTopo->nNics; i ++) + for (int n = 0; n < romeTopo->nGpus; n++) fprintf(file, "%d, ", romeTopo->gdrLevel[i*romeTopo->nGpus+n]); + fprintf(file, "},\n"); fprintf(file, " .pattern = \"%s\",\n", pattern); fprintf(file, " .ringBase = \"\",\n"); fprintf(file, "};\n"); @@ -694,13 +770,44 @@ static bool permuteGpuIds(int *g, int n, int last, struct rcclRomeModel* ref, st return false; } +static bool permuteNetIds(int *n, int *g, int s, int last, struct rcclRomeModel* ref, struct rcclRomeModel* topo, int* time) { + (*time) ++; + if (s == last) { + int i, j; + // match NET numa + for (i = 0; i < ref->nNics; i++) { + if (ref->nicNuma[i] != topo->nicNuma[n[i]]) break; + } + if (i < ref->nNics) return false; + // match gdr level + for (i = 0; i < ref->nNics; i++) { + for (j = 0; j < ref->nGpus; j++) { + if (ref->gdrLevel[i*ref->nGpus+j] != topo->gdrLevel[n[i]*ref->nGpus+g[j]]) break; + } + if (j < ref->nGpus) break; + } + if (i < ref->nNics) return false; + return true; + } else { + for (int i = s; i <= last; i++) { + std::swap(n[s], n[i]); + if (permuteNetIds(n, g, s+1, last, ref, topo, time)) return true; + std::swap(n[s], n[i]); + } + } + return false; +} + + ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph* graph) { static char ringRemap[64]; int i; int ngpus = system->nodes[GPU].count; int ncpus = system->nodes[CPU].count; + int nnets = system->nodes[NET].count; + if (ngpus > 8) return ncclSuccess; // only valid on Rome int arch, vendor, model; NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model)); @@ -715,7 +822,7 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph* // recognize system as Rome 4P2H even if no matching model if (ngpus > 4 && romeTopo.nLinks) system->type |= RCCL_TOPO_4P2H_ROME; - int g[NCCL_TOPO_MAX_NODES]; + int g[NCCL_TOPO_MAX_NODES], n[NCCL_TOPO_MAX_NODES]; int time = 0; struct timeval tvs, tve; gettimeofday(&tvs, NULL); @@ -734,8 +841,14 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph* if (romeTopo.nCpus != romeTopoModels[i].nCpus || romeTopo.nGpus != romeTopoModels[i].nGpus || romeTopo.nNics != romeTopoModels[i].nNics || romeTopo.nLinks != romeTopoModels[i].nLinks) continue; if (strcmp(romeTopoModels[i].pattern, pattern)) continue; + // permute GPU IDs for (int j = 0; j < ngpus; j++) g[j] = (j+2)%ngpus; - if (permuteGpuIds(g, 0, ngpus-1, romeTopoModels+i, &romeTopo, &time, match_nbio)) break; + if (!permuteGpuIds(g, 0, ngpus-1, romeTopoModels+i, &romeTopo, &time, match_nbio)) continue; + if (nnets > 1) { + // permute NET IDs + for (int j = 0; j < nnets; j++) n[j] = (j+2)%nnets; + if (permuteNetIds(n, g, 0, nnets-1, romeTopoModels+i, &romeTopo, &time)) break; + } else break; } gettimeofday(&tve, NULL); float t = (tve.tv_sec - tvs.tv_sec)*1E3 + (tve.tv_usec - tvs.tv_usec)/1E3; @@ -752,9 +865,152 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph* sprintf(line+offset, "%d ", g[k]); offset = strlen(line); } + if (nnets > 1) { + sprintf(line+offset, "NET mapping: "); + offset = strlen(line); + for (int k = 0; k < nnets; k++) { + sprintf(line+offset, "%d ", n[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)); + NCCLCHECK(parseGraph(romeTopoModels[i].ringBase, system, graph, g, nnets > 1 ? n : NULL)); + return ncclSuccess; +} + +static bool permuteGpuIdsForNuma(int *r, int *g, int n, int last, int ngpusPerNuma, struct rcclRomeModel* ref, struct rcclRomeModel* topo, int* time) { + (*time) ++; + if (n == last) { + int i, j; + // match GPU numa + for (i = 0; i < ngpusPerNuma; i++) + if (ref->gpuNuma[r[i]] != topo->gpuNuma[g[i]]) break; + if (i < ngpusPerNuma) return false; + // match XGMI connection + for (i = 0; i < ngpusPerNuma; i++) { + for (j = 0; j < ngpusPerNuma; j++) { + if (ref->connMatrix[r[i]*ref->nGpus+r[j]] != topo->connMatrix[g[i]*ref->nGpus+g[j]]) break; + if ((ref->gpuIds[r[i]]-ref->gpuIds[r[j]])*(topo->gpuIds[g[i]]-topo->gpuIds[g[j]]) < 0) break; + } + if (j < ngpusPerNuma) break; + } + if (i < ngpusPerNuma) return false; + return true; + } else { + for (int i = n; i <= last; i++) { + std::swap(g[n], g[i]); + if (permuteGpuIdsForNuma(r, g, n+1, last, ngpusPerNuma, ref, topo, time)) return true; + std::swap(g[n], g[i]); + } + } + return false; +} + +ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* graph) { + static char ringRemap[256]; + int i; + + int ngpus = system->nodes[GPU].count; + int ncpus = system->nodes[CPU].count; + int nnets = system->nodes[NET].count; + + // only valid on Rome + 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) + return ncclSuccess; + + // number of GPUs and NICs on each numa node is used as first screening pattern + struct rcclRomeModel romeTopo; + char pattern[256]; + NCCLCHECK(parseRomeSystem(system, &romeTopo, pattern)); + + // only match for system with 16 GPUs + if (ngpus != 16) return ncclSuccess; + + int gcnt = 0, mcnt = 0; + int g16[NCCL_TOPO_MAX_NODES], n[NCCL_TOPO_MAX_NODES]; + struct timeval tvs, tve; + gettimeofday(&tvs, NULL); + for (i = 0; i < sizeof(romeTopoModels)/sizeof(romeTopoModels[0]); i++) { + if (romeTopo.nCpus != romeTopoModels[i].nCpus || romeTopo.nGpus != romeTopoModels[i].nGpus || + romeTopo.nNics != romeTopoModels[i].nNics || romeTopo.nLinks != romeTopoModels[i].nLinks) continue; + if (strcmp(romeTopoModels[i].pattern, pattern)) continue; + int j, r[ngpus], g[ngpus]; + // match GPUs for each CPU NUMA nodes + for (j = 0; j < ncpus; j++) { + int ngpusPerNuma = 0, cnt = 0; + for (int k = 0; k < ngpus; k++) { + if (romeTopoModels[i].gpuNuma[k] != j) continue; + r[ngpusPerNuma++] = k; + } + if (ngpusPerNuma == 0) continue; + gcnt++; + // init GPU mapping + for (int k = 0; k < ngpus; k++) { + if (romeTopo.gpuNuma[k] != j) continue; + g[(2+cnt++)%ngpusPerNuma] = k; + } + int time = 0; + if (permuteGpuIdsForNuma(r, g, 0, ngpusPerNuma-1, ngpusPerNuma, romeTopoModels+i, &romeTopo, &time)) { + //printf("g[%d] = ", j); for (int n = 0; n < ngpusPerNuma; n++) printf("%d ", g[n]); printf(" total %d\n", cnt16); + cnt = 0; + for (int k = 0; k < ngpus; k++) { + if (romeTopo.gpuNuma[k] != j) continue; + g16[k] = g[cnt++]; + } + mcnt++; + } + } + if (gcnt && gcnt == mcnt) { + // final check to match all GPUs' XGMI connection + int k; + for (k = 0; k < romeTopoModels[i].nGpus; k++) { + int m; + for (m = 0; m < romeTopoModels[i].nGpus; m++) { + if (romeTopoModels[i].connMatrix[k*romeTopoModels[i].nGpus+m] != romeTopo.connMatrix[g16[k]*romeTopoModels[i].nGpus+g16[m]]) break; + if ((romeTopoModels[i].gpuIds[k]-romeTopoModels[i].gpuIds[m])*(romeTopo.gpuIds[g16[k]]-romeTopo.gpuIds[g16[m]]) < 0) break; + } + if (m < romeTopoModels[i].nGpus) break; + } + if (k < romeTopoModels[i].nGpus) continue; + if (nnets > 1) { + // permute NET IDs + int time = 0; + for (int m = 0; m < nnets; m++) n[m] = (m+2)%nnets; + if (permuteNetIds(n, g16, 0, nnets-1, romeTopoModels+i, &romeTopo, &time)) break; + } else break; + } + } + gettimeofday(&tve, NULL); + float t = (tve.tv_sec - tvs.tv_sec)*1E3 + (tve.tv_usec - tvs.tv_usec)/1E3; + if (i >= sizeof(romeTopoModels)/sizeof(romeTopoModels[0])) { + //printf("No solution in %.2fms (%d iter)\n", t, time); + return ncclSuccess; + } + + char line[1024]; + //sprintf(line, "Found matching Rome model index %d in %.2fms (%d iter) with GPU mapping: ", i, t, time); + sprintf(line, "Found matching Rome model index %d with GPU mapping: ", i); + int offset = strlen(line); + for (int k = 0; k < ngpus; k++) { + sprintf(line+offset, "%d ", g16[k]); + offset = strlen(line); + } + if (nnets > 1) { + sprintf(line+offset, "NET mapping: "); + offset = strlen(line); + for (int k = 0; k < nnets; k++) { + sprintf(line+offset, "%d ", n[k]); + offset = strlen(line); + } + } + INFO(NCCL_GRAPH, "%s", line); + system->type |= RCCL_TOPO_16P1H; + + // create 16P1H based on reference and remapped ids + NCCLCHECK(parseGraph(romeTopoModels[i].ringBase, system, graph, g16, nnets > 1 ? n : NULL)); return ncclSuccess; } diff --git a/src/graph/rome_models.h b/src/graph/rome_models.h index 08526b953d..cb6548718f 100644 --- a/src/graph/rome_models.h +++ b/src/graph/rome_models.h @@ -22,8 +22,9 @@ THE SOFTWARE. #ifndef RCCL_ROME_MODELS_H_ #define RCCL_ROME_MODELS_H_ -ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int* gpu_map); +ncclResult_t parseGraph(const char* str, struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int* gpu_map, int* net_map); ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph* graph); ncclResult_t parseChordalRing(struct ncclTopoSystem* system, struct ncclTopoGraph* graph); +ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* graph); #endif \ No newline at end of file diff --git a/src/graph/search.cc b/src/graph/search.cc index 15726f4177..8e2265d34a 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -795,7 +795,7 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph str = getenv("NCCL_RINGS"); if (str) { // user supplied topo - NCCLCHECK(parseGraph(str, system, graph, NULL)); + NCCLCHECK(parseGraph(str, system, graph, NULL, NULL)); if (graph->nChannels) { system->type |= RCCL_TOPO_4P2H_ROME; } @@ -805,6 +805,9 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph if (graph->nChannels) return ncclSuccess; // try to match Rome 4P2H NCCLCHECK(parseRome4P2H(system, graph)); + if (graph->nChannels) return ncclSuccess; + // try to match 1H16P + NCCLCHECK(parse1H16P(system, graph)); } if (graph->nChannels) return ncclSuccess; diff --git a/src/graph/topo.h b/src/graph/topo.h index fdef9a0069..3a71609e15 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -84,6 +84,7 @@ struct ncclTopoLinkList { #define RCCL_TOPO_CR8G 1 #define RCCL_TOPO_4P2H_ROME 2 #define RCCL_TOPO_GDR_ALL 4 +#define RCCL_TOPO_16P1H 8 struct ncclTopoNode { int type; diff --git a/src/init.cc b/src/init.cc index 4021abbfae..73c34219b5 100644 --- a/src/init.cc +++ b/src/init.cc @@ -870,7 +870,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G)) allGather3Data[rank].nc = 4; if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910) - allGather3Data[rank].nc = 6; + allGather3Data[rank].nc = 4; allGather3Data[rank].tree.pattern = treeGraph.pattern; allGather3Data[rank].tree.nChannels = treeGraph.nChannels; allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels; diff --git a/src/transport/net_ib.cc b/src/transport/net_ib.cc index 2ca0b51b78..74f013e2be 100644 --- a/src/transport/net_ib.cc +++ b/src/transport/net_ib.cc @@ -99,7 +99,7 @@ static ncclResult_t ncclIbGetPciPath(char* devName, char** path, int* realPort) return ncclSuccess; } -static int ibvWidths[] = { 1, 4, 8, 12 }; +static int ibvWidths[] = { 1, 4, 8, 12, 2 }; static int ibvSpeeds[] = { 2500, 5000, 10000, 10000, 14000, 25000, 50000 }; static int firstBitSet(int val, int max) { int i = 0; @@ -113,8 +113,6 @@ static int ncclIbSpeed(int speed) { return ibvSpeeds[firstBitSet(speed, sizeof(ibvSpeeds)/sizeof(int)-1)]; } -RCCL_PARAM(IbHcaSkipLinkLayer, "IB_HCA_SKIP_LINK_LAYER", 0); - ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction) { static int shownIbHcaEnv = 0; if(wrap_ibv_symbols() != ncclSuccess) { return ncclInternalError; } @@ -167,8 +165,8 @@ ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction) { continue; } if (portAttr.state != IBV_PORT_ACTIVE) continue; - if ((rcclParamIbHcaSkipLinkLayer() == IBV_LINK_LAYER_INFINIBAND || portAttr.link_layer != IBV_LINK_LAYER_INFINIBAND) - && (rcclParamIbHcaSkipLinkLayer() == IBV_LINK_LAYER_ETHERNET || portAttr.link_layer != IBV_LINK_LAYER_ETHERNET)) continue; + if (portAttr.link_layer != IBV_LINK_LAYER_INFINIBAND + && portAttr.link_layer != IBV_LINK_LAYER_ETHERNET) continue; // check against user specified HCAs/ports if (! (matchIfList(devices[d]->name, port, userIfs, nUserIfs, searchExact) ^ searchNot)) { @@ -199,8 +197,6 @@ ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction) { if (ncclNIbDevs == 0) { INFO(NCCL_INIT|NCCL_NET, "NET/IB : No device found."); } else { - auto cmpIbDevs = [](const void* n1, const void* n2) { return strcmp(((struct ncclIbDev*)n1)->devName, ((struct ncclIbDev*)n2)->devName); }; - qsort(ncclIbDevs, ncclNIbDevs, sizeof(struct ncclIbDev), cmpIbDevs); char line[1024]; line[0] = '\0'; for (int d=0; d + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/models/topo_8p_rome_4n_2.xml b/tools/topo_expl/models/topo_8p_rome_4n_2.xml new file mode 100644 index 0000000000..031d62a585 --- /dev/null +++ b/tools/topo_expl/models/topo_8p_rome_4n_2.xml @@ -0,0 +1,126 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index 46675d6219..1c234e0159 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -124,7 +124,9 @@ NodeModelDesc model_descs[] = { {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"}, - {1, "topo_8p_rome_vm1.xml", "single node 8 gfx908 Rome VM"}, + {1, "topo_8p_rome_vm1.xml", "single node 8 gfx908 Rome VM"}, + {1, "topo_16p1h.xml", "single node 16P1H"}, + {4, "topo_8p_rome_4n_2.xml", "4 nodes 8 gfx908 Rome 4 NICs NPS=4 Alt. Model"}, }; int main(int argc,char* argv[]) diff --git a/tools/topo_expl/utils.cpp b/tools/topo_expl/utils.cpp index 9360df4eb4..e1f39b53f2 100644 --- a/tools/topo_expl/utils.cpp +++ b/tools/topo_expl/utils.cpp @@ -489,7 +489,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G)) allGather3Data[rank].nc = 4; if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910) - allGather3Data[rank].nc = 6; + allGather3Data[rank].nc = 4; allGather3Data[rank].tree.pattern = treeGraph.pattern; allGather3Data[rank].tree.nChannels = treeGraph.nChannels; allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels;