diff --git a/projects/rccl/.github/CODEOWNERS b/projects/rccl/.github/CODEOWNERS old mode 100644 new mode 100755 index 35fef7d910..2b5b09701b --- a/projects/rccl/.github/CODEOWNERS +++ b/projects/rccl/.github/CODEOWNERS @@ -1 +1,4 @@ -* @wenkaidu @gilbertlee-amd @akolliasAMD @edgargabriel @PedramAlizadeh @BertanDogancay @nusislam @nileshnegi @KawtharShafie @AtlantaPepsi @mberenjk @haripriya-amd \ No newline at end of file +* @wenkaidu @gilbertlee-amd @akolliasAMD @edgargabriel @PedramAlizadeh @BertanDogancay @nusislam @nileshnegi @KawtharShafie @AtlantaPepsi @mberenjk @haripriya-amd# Documentation files +docs/* @ROCm/rocm-documentation +*.md @ROCm/rocm-documentation +*.rst @ROCm/rocm-documentation diff --git a/projects/rccl/cmake/Generator.cmake b/projects/rccl/cmake/Generator.cmake index fe4ca52b39..1dc0a5391f 100644 --- a/projects/rccl/cmake/Generator.cmake +++ b/projects/rccl/cmake/Generator.cmake @@ -154,7 +154,7 @@ function(gen_device_table) foreach(func IN LISTS FUNC_LIST) string(FIND "${func}" "LL128" IS_LL128) if(NOT IS_LL128 EQUAL -1) - file(APPEND ${DEVICE_TABLE_H_FILE} "#if defined(__gfx90a__)\n") + file(APPEND ${DEVICE_TABLE_H_FILE} "#if defined(__gfx90a__) && defined(ENABLE_LL128)\n") endif() if(ENABLE_IFC) file(APPEND ${DEVICE_TABLE_H_FILE} "__device__ void ${func}();\n") @@ -173,7 +173,7 @@ function(gen_device_table) foreach(func ${FUNC_LIST}) string(FIND "${func}" "LL128" IS_LL128) if(NOT IS_LL128 EQUAL -1) - file(APPEND ${DEVICE_TABLE_H_FILE} "#if defined(__gfx90a__)\n") + file(APPEND ${DEVICE_TABLE_H_FILE} "#if defined(__gfx90a__) && defined(ENABLE_LL128)\n") file(APPEND ${DEVICE_TABLE_H_FILE} " ${func},\n") file(APPEND ${DEVICE_TABLE_H_FILE} "#else\n") string(REPLACE "LL128" "LL" func "${func}") @@ -284,11 +284,9 @@ endfunction() ## Function to generate MSCCL Kernels ########################################################################################################### function(gen_msccl_kernels) - foreach(REDOP_CURRENT IN LISTS ALL_REDOPS) + set(MSCCL_REDOP Sum Prod MinMax) + foreach(REDOP_CURRENT IN LISTS MSCCL_REDOP) foreach(DATA_TYPE ${ALL_TYPES}) - if (REDOP_CURRENT STREQUAL "SumPostDiv" AND DATA_TYPE IN_LIST FLOATS_LIST) - continue() # Skip the iteration for FLOATS_LIST when REDOP_CURRENT is SumPostDiv - endif() set(FILE_NAME "${HIPIFY_DIR}/src/device/msccl_kernel_${REDOP_CURRENT}_${DATA_TYPE}.cpp") message(STATUS "Generating ${FILE_NAME}") file(WRITE ${FILE_NAME} @@ -296,8 +294,7 @@ function(gen_msccl_kernels) #include \"primitives.h\" #include \"nccl_common.h\" #include \"device.h\" - MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP_TYPE(${REDOP_CURRENT}, ${DATA_TYPE}, false); - MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP_TYPE(${REDOP_CURRENT}, ${DATA_TYPE}, true);") + MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP_TYPE(${REDOP_CURRENT}, ${DATA_TYPE}, false);") list(APPEND HIP_SOURCES ${FILE_NAME}) endforeach() endforeach() @@ -365,7 +362,7 @@ function(gen_collectives) file(WRITE ${FILE_PATH} "#include \"${COLL_LOWER}.h\"\n#include \"common.h\"\n#include \"collectives.h\"\n") string(FIND "${list_name}" "LL128" IS_LL128) if(NOT IS_LL128 EQUAL -1) - file(APPEND ${FILE_PATH} "#if defined(__gfx90a__)\n") + file(APPEND ${FILE_PATH} "#if defined(__gfx90a__) && defined(ENABLE_LL128)\n") endif() foreach(IMPL IN LISTS IMPL_LIST) file(APPEND ${FILE_PATH} "${IMPL}") diff --git a/projects/rccl/docs/sphinx/requirements.in b/projects/rccl/docs/sphinx/requirements.in index 2dc317aca6..46fe1144f1 100644 --- a/projects/rccl/docs/sphinx/requirements.in +++ b/projects/rccl/docs/sphinx/requirements.in @@ -1 +1 @@ -rocm-docs-core==0.31.0 +rocm-docs-core==0.33.0 diff --git a/projects/rccl/docs/sphinx/requirements.txt b/projects/rccl/docs/sphinx/requirements.txt index 279bb23e11..ea9683ca9d 100644 --- a/projects/rccl/docs/sphinx/requirements.txt +++ b/projects/rccl/docs/sphinx/requirements.txt @@ -100,7 +100,7 @@ requests==2.31.0 # via # pygithub # sphinx -rocm-docs-core==0.31.0 +rocm-docs-core==0.33.0 # via -r requirements.in smmap==5.0.0 # via gitdb diff --git a/projects/rccl/src/device/msccl_kernel_impl.h b/projects/rccl/src/device/msccl_kernel_impl.h index 5716946b82..ac9a57c8d3 100644 --- a/projects/rccl/src/device/msccl_kernel_impl.h +++ b/projects/rccl/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 ? __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) ? __ATOMIC_RELEASE : __ATOMIC_RELAXED); step++; } } @@ -427,8 +427,6 @@ __global__ void MSCCL_KERNEL_ENTRY_NAME(devredop, type, Simple, fullOps)(struct #define MSCCL_IMPL_KERNEL_ENTRY_FUNC() \ MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP(Sum, false) \ MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP(Prod, false) \ - MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP(MinMax, false) \ - MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP(PreMulSum, false) \ - MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP_NOFLOAT(SumPostDiv, false) + MSCCL_IMPL_KERNEL_ENTRY_FUNC_DEVREDOP(MinMax, false) #endif diff --git a/projects/rccl/src/enqueue.cc b/projects/rccl/src/enqueue.cc index 3657771d98..2094fd1f7f 100644 --- a/projects/rccl/src/enqueue.cc +++ b/projects/rccl/src/enqueue.cc @@ -18,6 +18,7 @@ #include "rocmwrap.h" #include "rccl_vars.h" #include "transport.h" +#include "common.h" #include // std::memcpy #include // PRIx64 diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index 5ff42163b4..f7d8704a8a 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -863,6 +863,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { } else { // Round to next pow2 nChannelsPerPeer and nChannels comm->p2pnChannelsPerPeer = (ncclParamNChannelsPerPeer() == -2 ? nextPow2(minChannels) : ncclParamNChannelsPerPeer()); + comm->p2pnChannelsPerPeer *= 2; comm->p2pnChannels = nextPow2(comm->p2pnChannels); } diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc index db92476c4d..48631546da 100644 --- a/projects/rccl/src/graph/tuning.cc +++ b/projects/rccl/src/graph/tuning.cc @@ -362,6 +362,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom if (coll == ncclFuncAllGather && a != NCCL_ALGO_RING && a != NCCL_ALGO_NVLS) continue; for (int p=0; ptopo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") && comm->topo->nodes[GPU].count == comm->topo->nRanks) continue; if ((a == NCCL_ALGO_NVLS || a == NCCL_ALGO_NVLS_TREE) && p != NCCL_PROTO_SIMPLE) continue; int collnet = (a == NCCL_ALGO_COLLNET_DIRECT || a == NCCL_ALGO_COLLNET_CHAIN) ? 1 : 0; float bw = nNodes <= 2 || collnet ? graphs[a]->bwIntra : graphs[a]->bwInter; diff --git a/projects/rccl/src/include/enqueue.h b/projects/rccl/src/include/enqueue.h index 94e331e194..634f037cb3 100644 --- a/projects/rccl/src/include/enqueue.h +++ b/projects/rccl/src/include/enqueue.h @@ -10,7 +10,6 @@ #include "comm.h" #include "group.h" #include "collectives.h" -#include "common.h" #include "utils.h" #define NCCL_MIN_CHANNEL_SIZE (NCCL_LL_THREAD_THRESHOLD*64) diff --git a/projects/rccl/src/include/msccl/msccl_kernel.h b/projects/rccl/src/include/msccl/msccl_kernel.h index 97975f6e8d..7fe496539d 100644 --- a/projects/rccl/src/include/msccl/msccl_kernel.h +++ b/projects/rccl/src/include/msccl/msccl_kernel.h @@ -39,9 +39,7 @@ __global__ void MSCCL_KERNEL_ENTRY_NAME(devredop, type, proto, fullOps)(struct n #define MSCCL_DECL_KERNEL_ENTRY_FUNC() \ MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP(Sum, false) \ MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP(Prod, false) \ - MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP(MinMax, false) \ - MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP(PreMulSum, false) \ - MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP_NOFLOAT(SumPostDiv, false) + MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP(MinMax, false) MSCCL_DECL_KERNEL_ENTRY_FUNC() diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 6e29f35ee8..eaf1e70a40 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -191,6 +191,7 @@ void *ncclCommThreadMain(void *arg) { memset(head, 0, sizeof(int)*MAXCHANNELS); vega_gpu_rtc_freq = GetDeviceWallClockRateInKhz(comm->cudaDev) * 1.0E3; do { + int numActiveChans = MAXCHANNELS; for (int channel = 0; channel < MAXCHANNELS; channel++) { int tail = comm->collTraceTail[channel].tail%COLLTRACE_NUM_ITEMS; int count; @@ -198,8 +199,8 @@ void *ncclCommThreadMain(void *arg) { count = tail - head[channel]; else count = COLLTRACE_NUM_ITEMS + head[channel] - tail; - if (!count) { - usleep(1000); //sleep 1ms + if (count == 0) { + numActiveChans--; continue; } for (int i = 0; i < count; i++) { @@ -260,7 +261,10 @@ void *ncclCommThreadMain(void *arg) { head[channel] %= COLLTRACE_NUM_ITEMS; } } - } while(!comm->collTraceExit); + if (comm->collTraceExit && numActiveChans == 0) + break; + usleep(1000); //sleep 1ms + } while(true); pthread_exit(NULL); } #endif diff --git a/projects/rccl/src/misc/msccl/msccl_lifecycle.cc b/projects/rccl/src/misc/msccl/msccl_lifecycle.cc index 105d093716..034a597769 100644 --- a/projects/rccl/src/misc/msccl/msccl_lifecycle.cc +++ b/projects/rccl/src/misc/msccl/msccl_lifecycle.cc @@ -270,11 +270,10 @@ static ncclResult_t mscclInternalSchedulerSelectAlgo(struct mscclSchedulerParam* mscclStatus& status = mscclGetStatus(); param->scheduled = false; - - /*// Current MSCCL doesn't support pre/post op + // Current MSCCL doesn't support pre/post op if (param->op >= ncclAvg) { return ncclSuccess; - }*/ + } // Whether the algorithm is in-place bool isInPlace = false; diff --git a/projects/rccl/src/misc/msccl/msccl_setup.cc b/projects/rccl/src/misc/msccl/msccl_setup.cc index d49ec91228..21be019361 100644 --- a/projects/rccl/src/misc/msccl/msccl_setup.cc +++ b/projects/rccl/src/misc/msccl/msccl_setup.cc @@ -332,12 +332,10 @@ static ncclResult_t hostToDevRedOp( #define MSCCL_KERNEL_ENTRY() \ MSCCL_KERNEL_ENTRY_DEVREDOP(Sum, false), \ MSCCL_KERNEL_ENTRY_DEVREDOP(Prod, false), \ - MSCCL_KERNEL_ENTRY_DEVREDOP(MinMax, false), \ - MSCCL_KERNEL_ENTRY_DEVREDOP(PreMulSum, false), \ - MSCCL_KERNEL_ENTRY_DEVREDOP_NOFLOAT(SumPostDiv, false) + MSCCL_KERNEL_ENTRY_DEVREDOP(MinMax, false) // Except for ncclDevPreMulSum and ncclDevSumPostDiv required by ncclAvg -void* mscclKernelEntries[ncclNumDevRedOps * ncclNumTypes * NCCL_NUM_PROTOCOLS] = { +void* mscclKernelEntries[(ncclNumDevRedOps-2) * ncclNumTypes * NCCL_NUM_PROTOCOLS] = { #ifdef COMPILE_MSCCL_KERNEL MSCCL_KERNEL_ENTRY() #endif diff --git a/projects/rccl/test/common/PtrUnion.hpp b/projects/rccl/test/common/PtrUnion.hpp index f29d41a0bc..d905648942 100644 --- a/projects/rccl/test/common/PtrUnion.hpp +++ b/projects/rccl/test/common/PtrUnion.hpp @@ -47,6 +47,8 @@ namespace RcclUnitTesting double* F8; // ncclFloat64 rccl_bfloat16* B2; // ncclBfloat16 + constexpr PtrUnion() : ptr(nullptr) {} + ErrCode Attach(void *ptr); ErrCode Attach(PtrUnion ptrUnion);