diff --git a/CMakeLists.txt b/CMakeLists.txt index 42f87beb35..d0a316ca69 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) @@ -281,7 +282,7 @@ set(SRC_FILES src/msccl.cc src/proxy.cc src/transport.cc -# src/init_nvtx.cc + src/init_nvtx.cc # src/clique/AllReduceCliqueKernel.h # src/clique/CliqueCommon.h # src/clique/CliqueManager.cc @@ -413,6 +414,26 @@ set(SRC_FILES src/include/nvtx3/nvtxExtDetail/nvtxExtInit.h src/include/nvtx3/nvtxExtDetail/nvtxExtPayloadTypeInfo.h src/include/nvtx3/nvtxExtDetail/nvtxExtTypes.h + src/include/nvtx.h + src/include/nvtx_stub.h + src/include/p2p.h + src/include/param.h + src/include/profiler.h + src/include/proxy.h + src/include/rccl_bfloat16.h + src/include/rccl_vars.h + src/include/rocm_smi_wrap.h + src/include/rocmwrap.h + src/include/shm.h + src/include/signals.h + src/include/socket.h + src/include/strongstream.h + src/include/timer.h + src/include/transport.h + src/include/trees.h + src/include/utils.h + src/init.cc + src/init_nvtx.cc src/misc/archinfo.cc src/misc/argcheck.cc # src/misc/cudawrap.cc @@ -483,6 +504,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}) @@ -524,7 +547,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 1c62143047..fecf166b62 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/docs/sphinx/requirements.in b/docs/sphinx/requirements.in index 46fe1144f1..170ba424e7 100644 --- a/docs/sphinx/requirements.in +++ b/docs/sphinx/requirements.in @@ -1 +1 @@ -rocm-docs-core==0.33.0 +rocm-docs-core==0.34.0 diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index ea9683ca9d..e5fec3f896 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 @@ -100,7 +100,7 @@ requests==2.31.0 # via # pygithub # sphinx -rocm-docs-core==0.33.0 +rocm-docs-core==0.34.0 # via -r requirements.in smmap==5.0.0 # via gitdb 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/device/msccl_kernel_impl.h b/src/device/msccl_kernel_impl.h index ac9a57c8d3..f3e172c9cc 100644 --- a/src/device/msccl_kernel_impl.h +++ b/src/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++; } } diff --git a/src/graph/paths.cc b/src/graph/paths.cc index f7d8704a8a..e185bec511 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -863,7 +863,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 && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94")) comm->p2pnChannelsPerPeer *= 2; comm->p2pnChannels = nextPow2(comm->p2pnChannels); } @@ -874,9 +875,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 e0cbde1a64..1f8bec0faa 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -467,14 +467,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; inodes[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 d236c45566..fc6336cb5d 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -242,4 +242,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 e81d85eeda..2e3b8d51dd 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -602,7 +602,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/include/core.h b/src/include/core.h index 84c65f313f..a8d858bb10 100644 --- a/src/include/core.h +++ b/src/include/core.h @@ -37,6 +37,10 @@ #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 diff --git a/src/transport/net_ib.cc b/src/transport/net_ib.cc index 0a80859651..25183bda9d 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;