From f70e3e569bb2d8b9e43e75aabfe528d6feeb2485 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Thu, 17 Aug 2023 08:50:51 -0700
Subject: [PATCH] gfx11: don't use LL for sendrecv (#853)

* gfx11: don't use LL for sendrecv

* Use builtin instead of inline asm
---
 src/collectives/device/primitives.h   |  1 +
 src/collectives/device/prims_simple.h | 13 +++++++------
 src/enqueue.cc                        |  3 ++-
 3 files changed, 10 insertions(+), 7 deletions(-)

diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h
index 170a9cad43..a1a5d173cb 100644
--- a/src/collectives/device/primitives.h
+++ b/src/collectives/device/primitives.h
@@ -21,6 +21,7 @@
   } else { \
     const int w = threadIdx.x/WARP_SIZE; \
     const int wid = threadIdx.x%WARP_SIZE; \
+    __threadfence(); \
     if (wid == 0) { \
       barrier_next[w] += nthreads/WARP_SIZE; \
       atomicAdd((unsigned long long *)barriers, 1); \
diff --git a/src/collectives/device/prims_simple.h b/src/collectives/device/prims_simple.h
index 3aa21d882c..9bcac9e88e 100644
--- a/src/collectives/device/prims_simple.h
+++ b/src/collectives/device/prims_simple.h
@@ -159,17 +159,18 @@ private:
 
   template<int Recv, int Send>
   inline __device__ void postPeer(bool dataStored) {
+    if (Send && (flags & RolePostSend) && dataStored)
+#ifdef __GFX9__
+      __builtin_amdgcn_buffer_wbinvl1();
+#else
+      __threadfence_system();
+#endif
+
     if ((flags & Send*RolePostSend) && next_hdp_reg)
       STORE((unsigned int *)next_hdp_reg, 0x1);
 
     if (flags & (Recv*RolePostRecv | Send*RolePostSend)) {
       step += StepPerSlice;
-      if (Send && (flags & RolePostSend) && dataStored)
-#ifdef __GFX9__
-        __asm__ __volatile__("buffer_wbinvl1_vol");
-#else
-        __threadfence_system();
-#endif
       STORE(connStepPtr, step);
     }
   }
diff --git a/src/enqueue.cc b/src/enqueue.cc
index 856bc57cff..801f4c8036 100644
--- a/src/enqueue.cc
+++ b/src/enqueue.cc
@@ -318,7 +318,8 @@ static ncclResult_t addP2pToPlan(
   // 1 is connIndex
   struct ncclConnInfo* conn = isSendNotRecv ?
     &comm->channels[channelId].peers[peer]->send[1].conn : &comm->channels[channelId].peers[peer]->recv[1].conn;
-  info.protocol = ((conn->buffs[NCCL_PROTO_LL] != nullptr) && bytes <= ncclParamP2pLLThreshold()) ? NCCL_PROTO_LL : NCCL_PROTO_SIMPLE;
+  // do not use LL on gfx11
+  info.protocol = ((conn->buffs[NCCL_PROTO_LL] != nullptr) && bytes <= ncclParamP2pLLThreshold() && comm->topo->nodes[GPU].nodes[0].gpu.gcn/100 != 11) ? NCCL_PROTO_LL : NCCL_PROTO_SIMPLE;
 
   struct ncclProxyOp proxyOp = {};
   NCCLCHECK(ncclProxyComputeP2p(&info, &proxyOp));