Use relaxed atomics for LL on GFX11 (#859)

[ROCm/rccl commit: 6a0a6a37d9]
2023-08-21 16:28:39 -07:00
parent 75e3927f50
commit 5983f0e371
3 changed files with 30 additions and 18 deletions
@@ -449,7 +449,7 @@ 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++) {
    // Disable LL protocol on gfx11xx
-    int pEnable = (p == NCCL_PROTO_LL && comm->topo->nodes[GPU].nodes[0].gpu.gcn/100 == 11) ? 0 : protoEnable[p];
+    int pEnable = protoEnable[p];
    if (pEnable == 2 && p == NCCL_PROTO_LL128) {
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
 #if defined(ENABLE_LL128)