Implement NIC identification and remapping (#420)
* Add 1H16P GPU model * Implement NIC identification and remapping * Revert "Sort IB devices based on device name (#413)" This reverts commit2d0ed8dff6. * Fix permute and check order * Correction on IB speed reporting * Revert "Allow user to link layer with RCCL_IB_HCA_SKIP_LINK_LAYER (#361)" This reverts commitcaf5c9992a.
このコミットが含まれているのは:
Wenkai Du
GitHub
+290
-34
@@ -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;
|
||||
}
|
||||
|
||||
@@ -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
|
||||
+4
-1
@@ -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;
|
||||
|
||||
|
||||
@@ -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;
|
||||
|
||||
+1
-1
@@ -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;
|
||||
|
||||
@@ -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<ncclNIbDevs; d++) {
|
||||
|
||||
@@ -21,7 +21,7 @@
|
||||
|
||||
DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
|
||||
|
||||
for i in {0..55}
|
||||
for i in {0..57}
|
||||
do
|
||||
if [[ $i -eq 50 ]] || [[ $i -eq 51 ]]
|
||||
then
|
||||
|
||||
@@ -0,0 +1,191 @@
|
||||
<system version="2">
|
||||
<cpu numaid="0" affinity="00000000,00000000,00000000,ffffffff,00000000,00000000,00000000,ffffffff" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:41:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:4c:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:4e:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="0" sm="90" gcn="910" arch="38911" rank="0" gdr="0">
|
||||
<xgmi target="0000:51:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:56:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0e:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:4f:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:51:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="1" sm="90" gcn="910" arch="38911" rank="1" gdr="0">
|
||||
<xgmi target="0000:4e:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:56:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:59:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:11:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:52:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:54:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:56:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="2" sm="90" gcn="910" arch="38911" rank="2" gdr="0">
|
||||
<xgmi target="0000:4e:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:51:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:59:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:16:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:57:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:59:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="3" sm="90" gcn="910" arch="38911" rank="3" gdr="0">
|
||||
<xgmi target="0000:51:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:56:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:19:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="1" affinity="00000000,00000000,ffffffff,00000000,00000000,00000000,ffffffff,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:01:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:09:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:0c:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:0e:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="4" sm="90" gcn="910" arch="38911" rank="4" gdr="0">
|
||||
<xgmi target="0000:4e:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:11:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:8f:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:0f:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:11:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="5" sm="90" gcn="910" arch="38911" rank="5" gdr="0">
|
||||
<xgmi target="0000:51:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0e:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:92:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:97:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:0b:00.0" class="0x020000" vendor="0x8086" device="0x10d3" subsystem_vendor="0x8086" subsystem_device="0xa01f" link_speed="2.5 GT/s" link_width="1">
|
||||
<nic>
|
||||
<net name="enp11s0" dev="0" speed="1000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:12:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:14:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:16:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="6" sm="90" gcn="910" arch="38911" rank="6" gdr="0">
|
||||
<xgmi target="0000:56:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:19:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:92:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:97:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:17:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:19:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="7" sm="90" gcn="910" arch="38911" rank="7" gdr="0">
|
||||
<xgmi target="0000:59:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:16:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:9a:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="2" affinity="00000000,ffffffff,00000000,00000000,00000000,ffffffff,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:c2:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:ca:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:cd:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:cf:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="8" sm="90" gcn="910" arch="38911" rank="8" gdr="0">
|
||||
<xgmi target="0000:d2:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d7:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:8f:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:d0:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:d2:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="9" sm="90" gcn="910" arch="38911" rank="9" gdr="0">
|
||||
<xgmi target="0000:cf:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d7:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:da:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:92:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:d3:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:d5:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:d7:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="10" sm="90" gcn="910" arch="38911" rank="10" gdr="0">
|
||||
<xgmi target="0000:cf:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d2:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:da:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:97:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:d8:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:da:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="11" sm="90" gcn="910" arch="38911" rank="11" gdr="0">
|
||||
<xgmi target="0000:d2:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d7:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:9a:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="3" affinity="ffffffff,00000000,00000000,00000000,ffffffff,00000000,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:82:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:8a:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:8d:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:8f:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="12" sm="90" gcn="910" arch="38911" rank="12" gdr="0">
|
||||
<xgmi target="0000:0e:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:cf:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:92:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:90:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:92:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="13" sm="90" gcn="910" arch="38911" rank="13" gdr="0">
|
||||
<xgmi target="0000:11:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:16:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d2:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:8f:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:93:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:95:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:97:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="14" sm="90" gcn="910" arch="38911" rank="14" gdr="0">
|
||||
<xgmi target="0000:11:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:16:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d7:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:9a:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:98:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="Unknown speed" link_width="8">
|
||||
<pci busid="0000:9a:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="Unknown speed" link_width="16">
|
||||
<gpu dev="15" sm="90" gcn="910" arch="38911" rank="15" gdr="0">
|
||||
<xgmi target="0000:19:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:da:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:97:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,126 @@
|
||||
<system version="2">
|
||||
<cpu numaid="1" affinity="00000000,00000000,00000000,ffff0000,00000000,00000000,00000000,ffff0000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:43:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:48:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:4a:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="90" gcn="908" arch="38911" rank="0" gdr="1">
|
||||
<xgmi target="0000:50:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0f:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:4c:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:4e:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:50:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="90" gcn="908" arch="38911" rank="1" gdr="1">
|
||||
<xgmi target="0000:4a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0f:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:45:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x1590" subsystem_device="0x0345" link_speed="16 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_1" dev="3" speed="100000" port="1" guid="0x181515ffffa4e988" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="3" affinity="00000000,00000000,ffff0000,00000000,00000000,00000000,ffff0000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:03:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:05:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:08:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:0a:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="90" gcn="908" arch="38911" rank="2" gdr="1">
|
||||
<xgmi target="0000:4a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:50:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0f:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:0b:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:0d:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:0f:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="90" gcn="908" arch="38911" rank="3" gdr="1">
|
||||
<xgmi target="0000:4a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:50:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:0a:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:13:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x1590" subsystem_device="0x0345" link_speed="16 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_3" dev="2" speed="100000" port="1" guid="0x301515ffffa4e988" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="5" affinity="00000000,ffff0000,00000000,00000000,00000000,ffff0000,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:c3:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:c6:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:c8:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:ca:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="90" gcn="908" arch="38911" rank="4" gdr="1">
|
||||
<xgmi target="0000:d0:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:8a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:90:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:cc:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:ce:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:d0:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="90" gcn="908" arch="38911" rank="5" gdr="1">
|
||||
<xgmi target="0000:ca:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:8a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:90:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:c5:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x1590" subsystem_device="0x0345" link_speed="16 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_5" dev="1" speed="100000" port="1" guid="0x601515ffffa4e988" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="7" affinity="ffff0000,00000000,00000000,00000000,ffff0000,00000000,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:83:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:86:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:88:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:8a:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="90" gcn="908" arch="38911" rank="6" gdr="1">
|
||||
<xgmi target="0000:ca:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d0:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:90:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:8c:00.0" class="0x060400" vendor="0x11f8" device="0x4000" subsystem_vendor="0x11f8" subsystem_device="0xbeef" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:8e:00.0" class="0x060400" vendor="0x1002" device="0x14a0" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:90:00.0" class="0x038000" vendor="0x1002" device="0x738c" subsystem_vendor="0x1002" subsystem_device="0x0c34" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="90" gcn="908" arch="38911" rank="7" gdr="1">
|
||||
<xgmi target="0000:ca:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:d0:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:8a:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:85:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x1590" subsystem_device="0x0345" link_speed="16 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_7" dev="0" speed="100000" port="1" guid="0x64780affffa4e988" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -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[])
|
||||
|
||||
@@ -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;
|
||||
|
||||
新しいイシューから参照
ユーザーをブロックする