AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge remote-tracking branch 'remotes/nccl/master'

Browse Source 
[ROCm/rccl commit: 8e73a2ad60]
This commit is contained in:
Wenkai Du
2020-02-27 11:45:50 -08:00
parent a8f784e269 6034c27655
commit b750defc28
7 changed files with 48 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/graph/paths.cc
+23
View File
@@ -179,11 +179,18 @@ static ncclResult_t addCpuStep(struct ncclTopoSystem* system, int c, int t1, int
// Remove/free paths for a given type
static void ncclTopoRemovePathType(struct ncclTopoSystem* system, int nodeType) {
for (int t=0; t<NCCL_TOPO_NODE_TYPES; t++) {
// Remove links _to_ the given type
for (int n=0; n<system->nodes[t].count; n++) {
struct ncclTopoNode* node = system->nodes[t].nodes+n;
free(node->paths[nodeType]);
node->paths[nodeType] = NULL;
}
// Remove links _from_ the given type
for (int n=0; n<system->nodes[nodeType].count; n++) {
struct ncclTopoNode* node = system->nodes[nodeType].nodes+n;
free(node->paths[t]);
node->paths[t] = NULL;
}
}
}
@@ -309,6 +316,22 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm*
// Trim network
ncclTopoRemovePathType(system, NET);
system->nodes[NET].count = 0;
for (int t=0; t<NCCL_TOPO_NODE_TYPES; t++) {
for (int n=0; n<system->nodes[t].count; n++) {
struct ncclTopoNode* node = system->nodes[t].nodes+n;
for (int l=0; l<node->nlinks; l++) {
struct ncclTopoLink* link = &(node->links[l]);
if (link->remNode->type == NET) {
// Remove the link
for (int i=l; i<(node->nlinks-1); i++) {
memcpy(&(node->links[i]), &(node->links[i+1]), sizeof(ncclTopoLink));
}
node->nlinks--;
l--; // revisit the same value of "l" for the next iteration, since we edited the list in the middle of the loop
}
}
}
}
}
free(domains);
free(ids);
projects/rccl/src/graph/tuning.cc
+1 -1
View File
@@ -193,7 +193,7 @@ ncclResult_t ncclSetThresholds(struct ncclComm* comm, int minCompCap, int maxCom
// Override defaults with user env
char* str = getenv("NCCL_THREAD_THRESHOLDS");
if (str) {
ssize_t t[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { -2 };
ssize_t t[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = {{ -2, -2, -2 }, { -2, -2, -2}};
sscanf(str, "%ld %ld %ld %ld %ld %ld", t[0], t[0]+1, t[0]+2, t[1], t[1]+1, t[1]+2);
for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) {
for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {