Tune AUTOUNROLL for better performance

Also remove all unused UNROLL defines


[ROCm/rccl commit: 1faededc03]

projects/rccl/src/collectives/device/all_gather.cu

@@ -9,6 +9,4 @@
 #include "all_gather.h"
 #include "collectives.h"
 
-#define UNROLL 4
-
 IMPL_COLL3(ncclAllGather, copy, FuncSum, i8, int8_t, ncclCollAllGather, ncclSum, ncclInt8);

projects/rccl/src/collectives/device/all_reduce.cu

@@ -9,8 +9,6 @@
 #include "all_reduce.h"
 #include "collectives.h"
 
-#define UNROLL 4
-
 IMPL_COLL2(ncclAllReduce, sum,  FuncSum,  ncclCollAllReduce, ncclSum);
 IMPL_COLL2(ncclAllReduce, prod, FuncProd, ncclCollAllReduce, ncclProd);
 IMPL_COLL2(ncclAllReduce, min,  FuncMin,  ncclCollAllReduce, ncclMin);

projects/rccl/src/collectives/device/broadcast.cu

@@ -9,6 +9,4 @@
 #include "broadcast.h"
 #include "collectives.h"
 
-#define UNROLL 4
-
 IMPL_COLL3(ncclBroadcast, copy, FuncSum, i8, int8_t, ncclCollBroadcast, ncclSum, ncclInt8);

projects/rccl/src/collectives/device/common_kernel.h

@@ -336,9 +336,14 @@ __device__ void ReduceCopy128bMulti( const int w, const int nw, const int t,
 template <typename T>
 __device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(Pack128); }
 
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
+// Use UNROLL 4 for 2 SRCs, 2 for the rest
+#define AUTOUNROLL (UNROLL*(2/MINSRCS))
+#else
 // Try to limit consecutive load/stores to 8.
 // Use UNROLL 8 when we have a single source and a single destination, 4 otherwise
 #define AUTOUNROLL (UNROLL*(4/(MINDSTS+MINSRCS)))
+#endif
 
 template<int UNROLL, class FUNC, typename T, int MINSRCS, int MAXSRCS, int MINDSTS, int MAXDSTS>
 __device__ void ReduceOrCopyMulti(const int tid, const int nthreads,

projects/rccl/src/collectives/device/reduce.cu

@@ -9,8 +9,6 @@
 #include "reduce.h"
 #include "collectives.h"
 
-#define UNROLL 4
-
 IMPL_COLL2(ncclReduce, sum,  FuncSum,  ncclCollReduce, ncclSum);
 IMPL_COLL2(ncclReduce, prod, FuncProd, ncclCollReduce, ncclProd);
 IMPL_COLL2(ncclReduce, min,  FuncMin,  ncclCollReduce, ncclMin);

projects/rccl/src/collectives/device/reduce_scatter.cu

@@ -9,8 +9,6 @@
 #include "reduce_scatter.h"
 #include "collectives.h"
 
-#define UNROLL 4
-
 IMPL_COLL2(ncclReduceScatter, sum,  FuncSum,  ncclCollReduceScatter, ncclSum);
 IMPL_COLL2(ncclReduceScatter, prod, FuncProd, ncclCollReduceScatter, ncclProd);
 IMPL_COLL2(ncclReduceScatter, min,  FuncMin,  ncclCollReduceScatter, ncclMin);