yezhengmao/nccl
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #283 from lukeyeager/topo-trim-net-links

Browse Source 
Topo trim net links
This commit is contained in:
Sylvain Jeaugey 2020-01-16 15:40:36 -08:00 committed by GitHub
commit 44c34e5d10
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
2 changed files with 27 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
makefiles/common.mk
View File

@ -42,9 +42,10 @@ else
endif endif
#$(info NVCC_GENCODE is ${NVCC_GENCODE}) #$(info NVCC_GENCODE is ${NVCC_GENCODE})
CXXFLAGS := -DCUDA_MAJOR=$(CUDA_MAJOR) -DCUDA_MINOR=$(CUDA_MINOR) -fPIC -fvisibility=hidden CXXFLAGS := -DCUDA_MAJOR=$(CUDA_MAJOR) -DCUDA_MINOR=$(CUDA_MINOR) -fPIC -fvisibility=hidden \
CXXFLAGS += -Wall -Wno-unused-function -Wno-sign-compare -std=c++11 -Wvla -Wall -Wno-unused-function -Wno-sign-compare -std=c++11 -Wvla \
CXXFLAGS += -I $(CUDA_INC) -I $(CUDA_INC) \
$(CXXFLAGS)
# Maxrregcount needs to be set accordingly to NCCL_MAX_NTHREADS (otherwise it will cause kernel launch errors) # Maxrregcount needs to be set accordingly to NCCL_MAX_NTHREADS (otherwise it will cause kernel launch errors)
# 512 : 120, 640 : 96, 768 : 80, 1024 : 60 # 512 : 120, 640 : 96, 768 : 80, 1024 : 60
# We would not have to set this if we used __launch_bounds__, but this only works on kernels, not on functions. # We would not have to set this if we used __launch_bounds__, but this only works on kernels, not on functions.

23
src/graph/paths.cc
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 // Remove/free paths for a given type
static void ncclTopoRemovePathType(struct ncclTopoSystem* system, int nodeType) { static void ncclTopoRemovePathType(struct ncclTopoSystem* system, int nodeType) {
for (int t=0; t<NCCL_TOPO_NODE_TYPES; t++) { 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++) { for (int n=0; n<system->nodes[t].count; n++) {
struct ncclTopoNode* node = system->nodes[t].nodes+n; struct ncclTopoNode* node = system->nodes[t].nodes+n;
free(node->paths[nodeType]); free(node->paths[nodeType]);
node->paths[nodeType] = NULL; 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 // Trim network
ncclTopoRemovePathType(system, NET); ncclTopoRemovePathType(system, NET);
system->nodes[NET].count = 0; 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(domains);
free(ids); free(ids);