From d257970ad105ecbdd4145211a6a00d8c3e351e3f Mon Sep 17 00:00:00 2001
From: Gilbert Lee <gilbert.lee@amd.com>
Date: Thu, 10 Oct 2019 22:39:45 -0700
Subject: [PATCH 1/3] Fix for GenericOp device primitive bug

[ROCm/rccl commit: 8ae1bce3bbcfaaa8d4aa743b75469d61f36c79e7]
---
 projects/rccl/src/collectives/device/primitives.h | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h
index 81a4d4cb7f..8c63a6a4f1 100644
--- a/projects/rccl/src/collectives/device/primitives.h
+++ b/projects/rccl/src/collectives/device/primitives.h
@@ -193,10 +193,11 @@ class ncclPrimitives {
       if (SEND) __threadfence_system();
       if (tid == 0) FOR_SEND(postSend);
       if (tid == 0) FOR_RECV(postRecv);
+
+      for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
+      for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;
+      offset += sliceSize;
     }
-    for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
-    for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;
-    offset += sliceSize;
   }
 
   __device__ void loadRecvConn(struct ncclConnInfo* conn, int i, T* directBuff) {

From cf597ff257ad7b54dbc6324bc0f63e60b6408c06 Mon Sep 17 00:00:00 2001
From: Gilbert Lee <gilbert.lee@amd.com>
Date: Fri, 11 Oct 2019 09:16:19 -0700
Subject: [PATCH 2/3] Performing __threadfence_system() with only first thread

[ROCm/rccl commit: 1392dd29977fe2ae6eddd64d6d80609ced56de9f]
---
 projects/rccl/src/collectives/device/primitives.h | 11 +++++++----
 1 file changed, 7 insertions(+), 4 deletions(-)

diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h
index 8c63a6a4f1..b0d4a2938d 100644
--- a/projects/rccl/src/collectives/device/primitives.h
+++ b/projects/rccl/src/collectives/device/primitives.h
@@ -189,10 +189,13 @@ class ncclPrimitives {
         }
       }
       exitIfAbortBarrier(abort, abortCount);
-      if (tid == 0) FOR_SEND(postSendSize, realSize*sizeof(T));
-      if (SEND) __threadfence_system();
-      if (tid == 0) FOR_SEND(postSend);
-      if (tid == 0) FOR_RECV(postRecv);
+      if (tid == 0)
+      {
+        FOR_SEND(postSendSize, realSize*sizeof(T));
+        __threadfence_system();
+        FOR_SEND(postSend);
+        FOR_RECV(postRecv);
+      }
 
       for (int i=0; i<RECV*NRECV+SRC; i++) srcs[i] += sliceSize;
       for (int i=0; i<SEND*NSEND+DST; i++) dsts[i] += sliceSize;

From 7560929bd7895cbd39a34b03b04eb616bc1c5250 Mon Sep 17 00:00:00 2001
From: Gilbert Lee <gilbert.lee@amd.com>
Date: Fri, 11 Oct 2019 09:20:10 -0700
Subject: [PATCH 3/3] Reverting GenericOp bug workaround modifications to
 slice/chunk steps

[ROCm/rccl commit: 37603ae6cb734e52ac94c95748f1a532d54b962b]
---
 projects/rccl/src/collectives/collectives.h | 18 ++++++------------
 1 file changed, 6 insertions(+), 12 deletions(-)

diff --git a/projects/rccl/src/collectives/collectives.h b/projects/rccl/src/collectives/collectives.h
index 63fcfd2017..27de1f7540 100644
--- a/projects/rccl/src/collectives/collectives.h
+++ b/projects/rccl/src/collectives/collectives.h
@@ -57,18 +57,12 @@
 DECL_ALL_COLLS
 
 // CHUNKSIZE must be a multiple of SLICESIZE
-//#define ALLREDUCE_SLICESTEPS (NCCL_STEPS/4)
-//#define ALLREDUCE_CHUNKSTEPS (NCCL_STEPS/2)
-//#define ALLGATHER_SLICESTEPS (NCCL_STEPS/4)
-//#define ALLGATHER_CHUNKSTEPS (NCCL_STEPS/2)
-//#define REDUCESCATTER_SLICESTEPS (NCCL_STEPS/4)
-//#define REDUCESCATTER_CHUNKSTEPS (NCCL_STEPS/2)
-#define ALLREDUCE_SLICESTEPS 4
-#define ALLREDUCE_CHUNKSTEPS 4
-#define ALLGATHER_SLICESTEPS 4
-#define ALLGATHER_CHUNKSTEPS 4
-#define REDUCESCATTER_SLICESTEPS 4
-#define REDUCESCATTER_CHUNKSTEPS 4
+#define ALLREDUCE_SLICESTEPS (NCCL_STEPS/4)
+#define ALLREDUCE_CHUNKSTEPS (NCCL_STEPS/2)
+#define ALLGATHER_SLICESTEPS (NCCL_STEPS/4)
+#define ALLGATHER_CHUNKSTEPS (NCCL_STEPS/2)
+#define REDUCESCATTER_SLICESTEPS (NCCL_STEPS/4)
+#define REDUCESCATTER_CHUNKSTEPS (NCCL_STEPS/2)
 #define BROADCAST_SLICESTEPS 1
 #define BROADCAST_CHUNKSTEPS 1
 #define REDUCE_SLICESTEPS 1