From 1d989f65249c3cb602a64cce1dd666464dc058ad Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Fri, 2 Feb 2024 12:41:00 -0800 Subject: [PATCH 1/9] Doubling P2P channels per peer on single node only (#1069) --- src/graph/paths.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/graph/paths.cc b/src/graph/paths.cc index 4180f3e660..bb1e33a1c0 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -861,7 +861,8 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { } else { // Round to next pow2 nChannelsPerPeer and nChannels comm->p2pnChannelsPerPeer = (ncclParamNChannelsPerPeer() == -2 ? nextPow2(minChannels) : ncclParamNChannelsPerPeer()); - comm->p2pnChannelsPerPeer *= 2; + // Doubling P2P channels per peer on single node + if (comm->topo->nodes[GPU].count == comm->topo->nRanks) comm->p2pnChannelsPerPeer *= 2; comm->p2pnChannels = nextPow2(comm->p2pnChannels); } From ca007ddad3148fa5b36395051132be575a6e9fe7 Mon Sep 17 00:00:00 2001 From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com> Date: Tue, 6 Feb 2024 15:59:52 -0700 Subject: [PATCH 2/9] Bump cryptography from 41.0.6 to 42.0.0 in /docs/sphinx (#1070) Bumps [cryptography](https://github.com/pyca/cryptography) from 41.0.6 to 42.0.0. - [Changelog](https://github.com/pyca/cryptography/blob/main/CHANGELOG.rst) - [Commits](https://github.com/pyca/cryptography/compare/41.0.6...42.0.0) --- updated-dependencies: - dependency-name: cryptography dependency-type: indirect ... Signed-off-by: dependabot[bot] Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> --- docs/sphinx/requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index ea9683ca9d..a9e52a8cc1 100644 --- a/docs/sphinx/requirements.txt +++ b/docs/sphinx/requirements.txt @@ -26,7 +26,7 @@ charset-normalizer==3.1.0 # via requests click==8.1.3 # via sphinx-external-toc -cryptography==41.0.6 +cryptography==42.0.0 # via pyjwt deprecated==1.2.13 # via pygithub From a9214032fc4c6603395219579ad51267f6d205c9 Mon Sep 17 00:00:00 2001 From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com> Date: Tue, 6 Feb 2024 16:00:30 -0700 Subject: [PATCH 3/9] Bump rocm-docs-core from 0.33.0 to 0.33.1 in /docs/sphinx (#1071) Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.33.0 to 0.33.1. - [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases) - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md) - [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.33.0...v0.33.1) --- updated-dependencies: - dependency-name: rocm-docs-core dependency-type: direct:production update-type: version-update:semver-patch ... Signed-off-by: dependabot[bot] Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> --- docs/sphinx/requirements.in | 2 +- docs/sphinx/requirements.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/sphinx/requirements.in b/docs/sphinx/requirements.in index 46fe1144f1..1a04e80e03 100644 --- a/docs/sphinx/requirements.in +++ b/docs/sphinx/requirements.in @@ -1 +1 @@ -rocm-docs-core==0.33.0 +rocm-docs-core==0.33.1 diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index a9e52a8cc1..2f9b385a13 100644 --- a/docs/sphinx/requirements.txt +++ b/docs/sphinx/requirements.txt @@ -100,7 +100,7 @@ requests==2.31.0 # via # pygithub # sphinx -rocm-docs-core==0.33.0 +rocm-docs-core==0.33.1 # via -r requirements.in smmap==5.0.0 # via gitdb From 704c9ef0d195827946f5c2d40542ae1d276f716d Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Wed, 7 Feb 2024 14:05:57 -0800 Subject: [PATCH 4/9] Doubling P2P channels per peer on single node gfx94x only (#1074) --- src/graph/paths.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/graph/paths.cc b/src/graph/paths.cc index bb1e33a1c0..c4b2c06e87 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -862,7 +862,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { // Round to next pow2 nChannelsPerPeer and nChannels comm->p2pnChannelsPerPeer = (ncclParamNChannelsPerPeer() == -2 ? nextPow2(minChannels) : ncclParamNChannelsPerPeer()); // Doubling P2P channels per peer on single node - if (comm->topo->nodes[GPU].count == comm->topo->nRanks) comm->p2pnChannelsPerPeer *= 2; + if (comm->topo->nodes[GPU].count == comm->topo->nRanks && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94")) comm->p2pnChannelsPerPeer *= 2; comm->p2pnChannels = nextPow2(comm->p2pnChannels); } From be45f0effdbf65625d78010cfe252270117ac6bc Mon Sep 17 00:00:00 2001 From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com> Date: Thu, 8 Feb 2024 09:26:47 -0700 Subject: [PATCH 5/9] Bump rocm-docs-core from 0.33.1 to 0.33.2 in /docs/sphinx (#1073) Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.33.1 to 0.33.2. - [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases) - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md) - [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.33.1...v0.33.2) --- updated-dependencies: - dependency-name: rocm-docs-core dependency-type: direct:production update-type: version-update:semver-patch ... Signed-off-by: dependabot[bot] Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> --- docs/sphinx/requirements.in | 2 +- docs/sphinx/requirements.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/sphinx/requirements.in b/docs/sphinx/requirements.in index 1a04e80e03..3e94264700 100644 --- a/docs/sphinx/requirements.in +++ b/docs/sphinx/requirements.in @@ -1 +1 @@ -rocm-docs-core==0.33.1 +rocm-docs-core==0.33.2 diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index 2f9b385a13..2f9c10444a 100644 --- a/docs/sphinx/requirements.txt +++ b/docs/sphinx/requirements.txt @@ -100,7 +100,7 @@ requests==2.31.0 # via # pygithub # sphinx -rocm-docs-core==0.33.1 +rocm-docs-core==0.33.2 # via -r requirements.in smmap==5.0.0 # via gitdb From 5257c753c59ee708c96aa8c0eb296e6f9e2f98a1 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Thu, 8 Feb 2024 12:09:56 -0800 Subject: [PATCH 6/9] msccl: use relaxed atomics on scratch buffer (#1075) --- src/collectives/device/msccl_kernel_impl.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/collectives/device/msccl_kernel_impl.h b/src/collectives/device/msccl_kernel_impl.h index 3cbad93175..5e9b992e0e 100644 --- a/src/collectives/device/msccl_kernel_impl.h +++ b/src/collectives/device/msccl_kernel_impl.h @@ -372,7 +372,7 @@ __device__ __forceinline__ void mscclRunInterpreter( return; } if (t->hasDependence && tid == nthreads-1) - __atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), (t->type == MSCCL_REDUCE || t->type == MSCCL_RECV) ? __ATOMIC_RELEASE : __ATOMIC_RELAXED); + __atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), ((t->type == MSCCL_REDUCE || t->type == MSCCL_RECV) && (t->dstBuffer != MSCCL_SCRATCH_BUFFER)) ? __ATOMIC_RELEASE : __ATOMIC_RELAXED); step++; } } From 8a442faa12b4dcbc89cd9af2fba886c08349c22e Mon Sep 17 00:00:00 2001 From: Bertan Dogancay <111835151+BertanDogancay@users.noreply.github.com> Date: Thu, 8 Feb 2024 14:08:24 -0700 Subject: [PATCH 7/9] Nvtx support (#1076) * NVTX support --- CMakeLists.txt | 9 +++++++-- cmake/Generator.cmake | 1 - install.sh | 10 +++++++++- src/include/core.h | 4 ++++ 4 files changed, 20 insertions(+), 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 12efe8fc8f..3202833350 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -19,6 +19,7 @@ option(COLLTRACE "Collective Trace Option" option(ENABLE_MSCCL_KERNEL "Enable MSCCL while compiling" ON) option(ENABLE_IFC "Enable indirect function call" OFF) option(INSTALL_DEPENDENCIES "Force install dependencies" OFF) +option(NVTX "Enable NVTX" OFF) option(PROFILE "Enable profiling" OFF) option(TIMETRACE "Enable time-trace during compilation" OFF) option(TRACE "Enable additional tracing" OFF) @@ -411,7 +412,7 @@ set(SRC_FILES src/include/trees.h src/include/utils.h src/init.cc -# src/init_nvtx.cc + src/init_nvtx.cc src/misc/archinfo.cc src/misc/argcheck.cc # src/misc/cudawrap.cc @@ -484,6 +485,8 @@ foreach(SRC_FILE ${SRC_FILES}) ) endforeach() +# Generate device/host tables and all the collective functions that are going to be in librccl.so +#================================================================================================== if(ONLY_FUNCS) ## Generate only the specified functions gen_functions(${ONLY_FUNCS}) @@ -525,7 +528,9 @@ if(DEMANGLE_DIR) endif() ## Set RCCL compile definitions -target_compile_definitions(rccl PRIVATE NVTX_NO_IMPL) # NVTX is not supported +if(NOT NVTX) + target_compile_definitions(rccl PRIVATE NVTX_NO_IMPL) +endif() if(COLLTRACE) target_compile_definitions(rccl PRIVATE ENABLE_COLLTRACE) endif() diff --git a/cmake/Generator.cmake b/cmake/Generator.cmake index a38cd80690..9b842a6fac 100644 --- a/cmake/Generator.cmake +++ b/cmake/Generator.cmake @@ -21,7 +21,6 @@ # SOFTWARE. set(ALL_PARAMS "ALL_COLLS" "ALL_ALGOS" "ALL_PROTOS" "ALL_REDOPS" "ALL_TYPES") - set(ALL_COLLS "AllGather" "AllReduce" "AllToAllPivot" "Broadcast" "Reduce" "ReduceScatter" "SendRecv") set(ALL_ALGOS "TREE" "RING" "COLLNET_DIRECT" "COLLNET_CHAIN") set(ALL_PROTOS "LL" "LL128" "SIMPLE") diff --git a/install.sh b/install.sh index 3cc42a3649..018c7abb33 100755 --- a/install.sh +++ b/install.sh @@ -25,6 +25,7 @@ install_library=false msccl_kernel_enabled=true num_parallel_jobs=$(nproc) npkit_enabled=false +nvtx_enabled=false run_tests=false run_tests_all=false time_trace=false @@ -50,6 +51,7 @@ function display_help() echo " --amdgpu_targets Only compile for specified GPU architecture(s). For multiple targets, seperate by ';' (builds for all supported GPU architectures by default)" echo " --no_clean Don't delete files if they already exist" echo " --npkit-enable Compile with npkit enabled" + echo " --nvtx-enable Compile with nvtx enabled" echo " -p|--package_build Build RCCL package" echo " --prefix Specify custom directory to install RCCL to (default: /opt/rocm)" echo " --rm-legacy-include-dir Remove legacy include dir Packaging added for file/folder reorg backward compatibility" @@ -68,7 +70,7 @@ function display_help() # check if we have a modern version of getopt that can handle whitespace and long parameters getopt -T if [[ $? -eq 4 ]]; then - GETOPT_PARSE=$(getopt --name "${0}" --options dfhij:lprt --longoptions address-sanitizer,dependencies,debug,enable_backtrace,disable-colltrace,disable-msccl-kernel,fast,help,install,jobs:,local_gpu_only,amdgpu_targets:,no_clean,npkit-enable,package_build,prefix:,rm-legacy-include-dir,run_tests_all,run_tests_quick,static,tests_build,time-trace,verbose -- "$@") + GETOPT_PARSE=$(getopt --name "${0}" --options dfhij:lprt --longoptions address-sanitizer,dependencies,debug,enable_backtrace,disable-colltrace,disable-msccl-kernel,fast,help,install,jobs:,local_gpu_only,amdgpu_targets:,no_clean,npkit-enable,nvtx-enable,package_build,prefix:,rm-legacy-include-dir,run_tests_all,run_tests_quick,static,tests_build,time-trace,verbose -- "$@") else echo "Need a new version of getopt" exit 1 @@ -97,6 +99,7 @@ while true; do --amdgpu_targets) build_amdgpu_targets=${2}; shift 2 ;; --no_clean) clean_build=false; shift ;; --npkit-enable) npkit_enabled=true; shift ;; + --nvtx-enable) nvtx_enabled=true; shift ;; -p | --package_build) build_package=true; shift ;; --prefix) install_prefix=${2}; shift 2 ;; --rm-legacy-include-dir) build_freorg_bkwdcomp=false; shift ;; @@ -220,6 +223,11 @@ if ($install_dependencies); then cmake_common_options="${cmake_common_options} -DINSTALL_DEPENDENCIES=ON" fi +# Enable NVTX +if [[ "${nvtx_enabled}" == true ]]; then + cmake_common_options="${cmake_common_options} -DNVTX=ON" +fi + cmake_executable=cmake case "${OS_ID}" in centos|rhel) diff --git a/src/include/core.h b/src/include/core.h index a1d644d295..8ab92765a5 100644 --- a/src/include/core.h +++ b/src/include/core.h @@ -60,6 +60,10 @@ static __inline__ int ncclTypeSize(ncclDataType_t type) { #include "alloc.h" #include "utils.h" #include "param.h" +#ifdef NVTX_NO_IMPL #include "nvtx_stub.h" +#else +#include "nvtx.h" +#endif #endif // end include guard From 3e505a991c8ded73932f652bbc1f4e83ff07b98a Mon Sep 17 00:00:00 2001 From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com> Date: Fri, 9 Feb 2024 10:12:07 -0700 Subject: [PATCH 8/9] Bump rocm-docs-core from 0.33.2 to 0.34.0 in /docs/sphinx (#1078) Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.33.2 to 0.34.0. - [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases) - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md) - [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.33.2...v0.34.0) --- updated-dependencies: - dependency-name: rocm-docs-core dependency-type: direct:production update-type: version-update:semver-minor ... Signed-off-by: dependabot[bot] Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> --- docs/sphinx/requirements.in | 2 +- docs/sphinx/requirements.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/sphinx/requirements.in b/docs/sphinx/requirements.in index 3e94264700..170ba424e7 100644 --- a/docs/sphinx/requirements.in +++ b/docs/sphinx/requirements.in @@ -1 +1 @@ -rocm-docs-core==0.33.2 +rocm-docs-core==0.34.0 diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index 2f9c10444a..e5fec3f896 100644 --- a/docs/sphinx/requirements.txt +++ b/docs/sphinx/requirements.txt @@ -100,7 +100,7 @@ requests==2.31.0 # via # pygithub # sphinx -rocm-docs-core==0.33.2 +rocm-docs-core==0.34.0 # via -r requirements.in smmap==5.0.0 # via gitdb From 5669b0d7b6c6eb437129d2966ec85d807e5546bd Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Fri, 9 Feb 2024 09:18:38 -0800 Subject: [PATCH 9/9] 2.18.5 fix (#1077) * Revert "Revert "2.18.5-1"" This reverts commit 767fde8210160ced688f494e20c37873d23bcaed. * Fix initial net device value --- src/graph/paths.cc | 4 +- src/graph/search.cc | 32 +++++++--- src/graph/topo.cc | 137 ++++++++++++---------------------------- src/graph/topo.h | 9 +++ src/graph/xml.cc | 13 +++- src/transport/net_ib.cc | 2 +- 6 files changed, 85 insertions(+), 112 deletions(-) diff --git a/src/graph/paths.cc b/src/graph/paths.cc index c4b2c06e87..a9a1b43d02 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -873,9 +873,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { // fill the whole space of nChannels. To do so we mirror the bits in the // nChannels space. for (int c=0; cp2pnChannels; c++) { - int mirror = 0; - for (int b=1, mb=(comm->p2pnChannels>>1); bp2pnChannels; b<<=1, mb>>=1) if (c & b) mirror |= mb; - comm->p2pChannels[c] = mirror; + comm->p2pChannels[c] = mirrorBits(c, comm->p2pnChannels); } return ncclSuccess; } diff --git a/src/graph/search.cc b/src/graph/search.cc index 23cb9dbc59..406383bef9 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -412,6 +412,28 @@ ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, in int* localNets; NCCLCHECK(ncclCalloc(&localNets, system->nodes[NET].count)); + // First add the preferred NICs + for (int g=0; gnodes[GPU].count; g++) { + if (gpu != -1 && gpu != g) continue; + localNetCount = 0; + struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g; + for (int c = 0;; c++) { + int netId; + NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &netId)); + NCCLCHECK(ncclTopoIdToIndex(system, NET, netId, localNets+localNetCount)); + if (localNetCount > 0 && localNets[localNetCount] == localNets[0]) break; + localNetCount++; + } + // Append NICs to list + for (int i=0; inodes[GPU].count; g++) { if (gpu != -1 && gpu != g) continue; @@ -421,14 +443,6 @@ ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, in for (int n=0; nnodes[NET].count; n++) { if (paths[n].type == t) localNets[localNetCount++] = n; } - if (localNetCount == 0) continue; - // Shuffle by gpu NVML device number so that GPUs on the same PCI switch - // with multiple NICs don't use the same one as first choice. - for (int r=0; rnodes[GPU].nodes[g].gpu.dev % localNetCount; r++) { - int net0 = localNets[0]; - for (int i=0; ipattern == NCCL_TOPO_PATTERN_NVLS) { if (graph->nChannels < netcount) { int gpu; - NCCLCHECK(ncclTopoGetLocalGpu(system, nets[graph->nChannels], &gpu)); + NCCLCHECK(ncclTopoGetLocalGpu(system, system->nodes[NET].nodes[nets[graph->nChannels]].id, &gpu)); if (gpu != -1) NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, -1, -1, gpu)); } } else { diff --git a/src/graph/topo.cc b/src/graph/topo.cc index dd6d01d3c0..e645de978e 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -758,126 +758,67 @@ ncclResult_t ncclTopoGetSystem(struct ncclComm* comm, struct ncclTopoSystem** sy return ncclSuccess; } -static ncclResult_t getLocalNetMask(struct ncclTopoSystem* system, int g, uint64_t* localNetMask, int* type) { +ncclResult_t ncclTopoGetLocal(struct ncclTopoSystem* system, int type, int index, int resultType, int** locals, int* localCount, int* pathType) { int minType = PATH_DIS; float maxBw = 0; int count = 0; - int* nets; - NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count)); - for (int n=0; nnodes[NET].count; n++) { - struct ncclTopoLinkList* path = system->nodes[NET].nodes[n].paths[GPU]+g; - if (path->bw > maxBw || (path->bw == maxBw && path->type < minType)) { - maxBw = path->bw; - minType = path->type; - if (type) *type = minType; + NCCLCHECK(ncclCalloc(locals, system->nodes[resultType].count)); + struct ncclTopoLinkList* paths = system->nodes[type].nodes[index].paths[resultType]; + for (int i=0; inodes[resultType].count; i++) { + if (paths[i].bw > maxBw || (paths[i].bw == maxBw && paths[i].type < minType)) { + maxBw = paths[i].bw; + minType = paths[i].type; + if (pathType) *pathType = minType; count = 0; } - if (path->bw == maxBw && path->type == minType) nets[count++] = system->nodes[NET].nodes[n].id; + if (paths[i].bw == maxBw && paths[i].type == minType) (*locals)[count++] = i; } - - *localNetMask = 0ULL; - for (int n=0; n= 64) return ncclInternalError; - *localNetMask |= 1ULL<nodes[GPU].count; - NCCLCHECK(ncclCalloc(&localNetMasks, ngpus)); - - // Fill localNetMasks for all GPUs. - for (int g=0; gnodes[NET].nodes[localNets[net%localNetCount]].id; } + free(localNets); + free(localGpus); + return ncclSuccess; } ncclResult_t ncclTopoGetLocalGpu(struct ncclTopoSystem* system, int net, int* gpuIndex) { - int ngpus = system->nodes[GPU].count; - int* gpus; - NCCLCHECK(ncclCalloc(&gpus, ngpus)); - - // Find localNetMask which includes net with the most local GPUs. - int netLocalGpus = 0, minType = PATH_DIS; - uint64_t localNetMask = 0ULL; - for (int g=0; gnodes[GPU].count; g++) { + struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g; + int id; + NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &id)); + if (net == id) { + *gpuIndex = g; return ncclSuccess; } - gIndex++; - if (gIndex == netLocalGpus) { - gIndex = 0; - cId++; - } } - n = (n+1) % 64; } + *gpuIndex = -1; + return ncclSuccess; } /****************************/ diff --git a/src/graph/topo.h b/src/graph/topo.h index 5cd973e9c8..445ffb4747 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -241,4 +241,13 @@ static float ncclTopoXGMISpeed(const char* gcn) { #define ncclGetKernelIndex(p_comm) (0) #endif +// Mirror bits +static bool isPow2(int val) { + return (val & (val-1)) == 0; +} +static int mirrorBits(int val, int pow2) { + int mirror = 0; + for (int b=1, mb=(pow2>>1); b>=1) if (val & b) mirror |= mb; + return mirror; +} #endif diff --git a/src/graph/xml.cc b/src/graph/xml.cc index 902477a275..c72d6c94f1 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -592,7 +592,18 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml* } } pciNode->parent = parent; - parent->subs[parent->nSubs++] = pciNode; + // Keep PCI sub devices ordered by PCI Bus ID (Issue #820) + int subIndex = parent->nSubs; + const char* newBusId; + NCCLCHECK(xmlGetAttrStr(pciNode, "busid", &newBusId)); + for (int s=0; snSubs; s++) { + const char* busId; + NCCLCHECK(xmlGetAttrStr(parent->subs[s], "busid", &busId)); + if (strcmp(newBusId, busId) < 0) { subIndex = s; break; } + } + for (int s = parent->nSubs; s > subIndex; s--) parent->subs[s] = parent->subs[s-1]; + parent->subs[subIndex] = pciNode; + parent->nSubs++; } if (strcmp(parent->name, "pci") == 0) { NCCLCHECK(ncclTopoGetXmlFromSys(parent, xml)); diff --git a/src/transport/net_ib.cc b/src/transport/net_ib.cc index d8e7c217c7..c54418f608 100644 --- a/src/transport/net_ib.cc +++ b/src/transport/net_ib.cc @@ -68,7 +68,7 @@ struct userIbDev { uint16_t port_en; }; -#define MAX_IB_DEVS 16 +#define MAX_IB_DEVS 32 struct ncclIbDev ncclIbDevs[MAX_IB_DEVS]; struct userIbDev userIbDevs[MAX_IB_DEVS]; pthread_mutex_t ncclIbLock = PTHREAD_MUTEX_INITIALIZER;