From c3282948205c9cf761c1e9fe2abe2dc5f1bcc8d1 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Fri, 4 Aug 2023 07:53:07 -0700
Subject: [PATCH] gfx11xx: disable LL protocol to workaround mtype issue (#840)

[ROCm/rccl commit: 8e58b6587317ea9e851760d739d2c078777ca065]
---
 projects/rccl/src/collectives/device/prims_simple.h | 4 ++++
 projects/rccl/src/graph/tuning.cc                   | 3 ++-
 2 files changed, 6 insertions(+), 1 deletion(-)

diff --git a/projects/rccl/src/collectives/device/prims_simple.h b/projects/rccl/src/collectives/device/prims_simple.h
index 27c02bf0bd..39f6e46ed6 100644
--- a/projects/rccl/src/collectives/device/prims_simple.h
+++ b/projects/rccl/src/collectives/device/prims_simple.h
@@ -100,7 +100,11 @@ private:
     #endif
     // volatile is faster than acquire but not as correct. Make sure reduceCopy
     // loads data using volatile so it doesn't see stale data in L1.
+#ifdef __GFX9__
     return atomicAdd((unsigned long long *)ptr, 0);
+#else
+    return __atomic_load_n(ptr, __ATOMIC_SEQ_CST);
+#endif
   }
 
   template <int DirectRecv, int DirectSend, int Recv, int Send, int Src, int Dst>
diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc
index 9b1f2fb853..ce32a05a18 100644
--- a/projects/rccl/src/graph/tuning.cc
+++ b/projects/rccl/src/graph/tuning.cc
@@ -448,7 +448,8 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
   }
 
   for (int c=0; c<NCCL_NUM_FUNCTIONS; c++) for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
-    int pEnable = protoEnable[p];
+    // Disable LL protocol on gfx11xx
+    int pEnable = (p == NCCL_PROTO_LL && comm->topo->nodes[GPU].nodes[0].gpu.gcn/100 == 11) ? 0 : protoEnable[p];
     if (pEnable == 2 && p == NCCL_PROTO_LL128) {
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
 #if defined(ENABLE_LL128)