diff --git a/projects/rccl/src/collectives/device/all_gather.cu b/projects/rccl/src/collectives/device/all_gather.cu
index 3fd3e0c63e..f9859e94d8 100644
--- a/projects/rccl/src/collectives/device/all_gather.cu
+++ b/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);
diff --git a/projects/rccl/src/collectives/device/all_reduce.cu b/projects/rccl/src/collectives/device/all_reduce.cu
index 704197160e..b028161158 100644
--- a/projects/rccl/src/collectives/device/all_reduce.cu
+++ b/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);
diff --git a/projects/rccl/src/collectives/device/broadcast.cu b/projects/rccl/src/collectives/device/broadcast.cu
index c4b1cbc5e9..fa42d34d37 100644
--- a/projects/rccl/src/collectives/device/broadcast.cu
+++ b/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);
\ No newline at end of file
diff --git a/projects/rccl/src/collectives/device/common_kernel.h b/projects/rccl/src/collectives/device/common_kernel.h
index 7cf85671a3..b05e06e7d0 100644
--- a/projects/rccl/src/collectives/device/common_kernel.h
+++ b/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,
diff --git a/projects/rccl/src/collectives/device/reduce.cu b/projects/rccl/src/collectives/device/reduce.cu
index dbfa1b7fad..8f748616d4 100644
--- a/projects/rccl/src/collectives/device/reduce.cu
+++ b/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);
diff --git a/projects/rccl/src/collectives/device/reduce_scatter.cu b/projects/rccl/src/collectives/device/reduce_scatter.cu
index 82cb408a16..0bb26f88e8 100644
--- a/projects/rccl/src/collectives/device/reduce_scatter.cu
+++ b/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);