Merge remote-tracking branch 'rccl/develop' into 2.19.4

[ROCm/rccl commit: 1a134b283b]

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

projects/rccl/src/enqueue.cc

@@ -18,6 +18,7 @@
 #include "rocmwrap.h"
 #include "rccl_vars.h"
 #include "transport.h"
+#include "common.h"
 #include <cstring> // std::memcpy
 #include <cinttypes> // PRIx64
 

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);
   }
 

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; p<NCCL_NUM_PROTOCOLS; p++) {
+        if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_SIMPLE && IsArchMatch(comm->topo->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;

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)

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()
 

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

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;

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