gfx11xx: disable LL protocol to workaround mtype issue (#840)

[ROCm/rccl commit: 8e58b65873]

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>

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)