From c419755576745b9f4d67fcf761d6dcc745d90793 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Thu, 3 Aug 2023 13:46:40 -0700
Subject: [PATCH] Fix merge error and replace inline asm (#838)

[ROCm/rccl commit: 60efe265494bc6a7893cc110c3e53095176e65ee]
---
 projects/rccl/src/collectives/device/prims_ll128.h | 6 +-----
 1 file changed, 1 insertion(+), 5 deletions(-)

diff --git a/projects/rccl/src/collectives/device/prims_ll128.h b/projects/rccl/src/collectives/device/prims_ll128.h
index 469b843890..4d8ef8f0ed 100644
--- a/projects/rccl/src/collectives/device/prims_ll128.h
+++ b/projects/rccl/src/collectives/device/prims_ll128.h
@@ -292,10 +292,6 @@ private:
           needReload &= (0 == checkAbort(spins, i, 0));
         } while (__any(needReload));
 
-        #pragma unroll
-        for (int u=0; u<ELEMS_PER_THREAD; u+=2)
-          load128(ptr+u*WARP_SIZE, vr[u], vr[u+1]);
-
         #pragma unroll
         for (int u=0; u<ELEMS_PER_THREAD; u+=2) {
           v[u]   = applyReduce(redOp, vr[u], v[u]);
@@ -314,7 +310,7 @@ private:
     }
 
 #if RCCL_USE_WBINVL1_VOL
-    if (tid == 0) __asm__ __volatile__("buffer_wbinvl1_vol");
+    if (tid == 0) __builtin_amdgcn_buffer_wbinvl1();
 #endif
     /************************ Send **************************/
     if (SEND) {