From b0919dc46c4f15e70fe547e923c41f10cc339797 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Tue, 25 Aug 2020 13:16:34 -0700
Subject: [PATCH] Only use software barrier for synchronization (#258)

---
 src/collectives/device/primitives.h | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h
index 4269615a6d..ef8c753a1f 100644
--- a/src/collectives/device/primitives.h
+++ b/src/collectives/device/primitives.h
@@ -84,9 +84,7 @@ class ncclPrimitives {
 
   inline __device__ void barrier() {
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
-    if (nthreads == gridDim.x) {
-      __syncthreads();
-    } else if (wid == 0) {
+    if (wid == 0) {
       if (NRECV < NSEND) barrier_by_id(0);
       else barrier_by_id(1);
     }