From 5ae3f3f9543d6e7abc6066a0c07e0a8d1161c57d Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Tue, 14 Sep 2021 18:30:16 -0700
Subject: [PATCH] Remove extra L1 cache invalidate and restore __ATOMIC_SEQ_CST
 atomics (#426)

---
 src/collectives/device/primitives.h | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h
index 395dae600a..596abdf618 100644
--- a/src/collectives/device/primitives.h
+++ b/src/collectives/device/primitives.h
@@ -136,7 +136,7 @@ class ncclPrimitives {
     spins = 0;
     while (connHeadCache + NCCL_STEPS < step + SLICESTEPS) {
       __builtin_amdgcn_s_sleep(8);
-      connHeadCache = atomicAdd_system((unsigned long long *)connHeadPtr, 0);
+      connHeadCache = LOAD(connHeadPtr);
       if (checkAbort()) break;
     }
     __asm__ __volatile__("s_wakeup");
@@ -158,7 +158,7 @@ class ncclPrimitives {
 #endif
     while (connTailCache < step + SLICESTEPS) {
       __builtin_amdgcn_s_sleep(8);
-      connTailCache = atomicAdd_system((unsigned long long *)connTailPtr, 0);
+      connTailCache =LOAD(connTailPtr);
       if (checkAbort()) break;
     }
     __asm__ __volatile__("s_wakeup");
@@ -207,7 +207,11 @@ class ncclPrimitives {
         }
       }
       barrier();
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
+      // L1 cache invalidation by __threadfence_system() is already done in above barrier()
+#else
       if (SEND && (role & ROLE_POST_SEND) && realSize > 0 && index == 0) __threadfence_system();
+#endif
       __syncwarp();
       if (SEND && (role & ROLE_POST_SEND)) postSend();
       if (RECV && (role & ROLE_POST_RECV)) postRecv();