Fully enable cache bypass access for Simple/LL/LL128

projects/rccl/src/device/op128.h

@@ -12,13 +12,27 @@
 #include "device/rccl_ptr.h"
 
 inline __device__ void load128(const uint64_t* ptr, uint64_t &v0, uint64_t &v1) {
+#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  union { v4u v; uint64_t u64[2]; } u;
+  u.v = __builtin_amdgcn_global_load_b128((v4u_gptr) ptr, RCCL_SYSTEM_SYNCSCOPE);
+  v0 = u.u64[0];
+  v1 = u.u64[1];
+#else
   v0 = __builtin_nontemporal_load((u64_gptr) ptr);
   v1 = __builtin_nontemporal_load((u64_gptr) ptr+1);
+#endif
 }
 
 inline __device__ void store128(uint64_t* ptr, uint64_t v0, uint64_t v1) {
+#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  union { v4u v; uint64_t u64[2]; } u;
+  u.u64[0] = v0;
+  u.u64[1] = v1;
+  __builtin_amdgcn_global_store_b128((v4u_gptr) ptr, u.v, RCCL_SYSTEM_SYNCSCOPE);
+#else
   *((u64_gptr) ptr) = v0;
   *((u64_gptr) ptr + 1) = v1;
+#endif
 }
 
 inline __device__ uint64_t* shmemCvtPtr(volatile uint64_t* shmemGenericPtr) {
@@ -240,7 +254,29 @@ template<> __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack
 // template<> __device__ __forceinline__ void st_relaxed_gpu_global<0>(uintptr_t addr, BytePack<0> value) {}
 
 // Used to define implementations for above prototypes.
-#define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
+#define DEFINE_ld_st__size_space_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
+  template<> \
+  __device__ __forceinline__ BytePack<bytes> ld_##space<bytes>(addr_cxx_ty addr) { \
+    data_cxx_ty tmp; \
+    tmp = *((data_cxx_ty *)addr); \
+    BytePack<bytes> ans; \
+    ans.native = tmp; \
+    return ans; \
+  } \
+  template<> \
+  __device__ __forceinline__ BytePack<bytes> ld_volatile_##space<bytes>(addr_cxx_ty addr) { \
+    data_cxx_ty tmp; \
+    tmp =  __hip_atomic_load((data_cxx_ty *)addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); \
+    BytePack<bytes> ans; \
+    ans.native = tmp; \
+    return ans; \
+  } \
+  template<> \
+  __device__ __forceinline__ void st_##space<bytes>(addr_cxx_ty addr, BytePack<bytes> value) { \
+    __hip_atomic_store((data_cxx_ty *)addr, value.native, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); \
+  }
+
+#define DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
   template<> \
   __device__ __forceinline__ BytePack<bytes> ld_##space<bytes>(addr_cxx_ty addr) { \
     data_cxx_ty tmp; \
@@ -252,7 +288,7 @@ template<> __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack
   template<> \
   __device__ __forceinline__ BytePack<bytes> ld_volatile_##space<bytes>(addr_cxx_ty addr) { \
     data_cxx_ty tmp; \
-    tmp =  __builtin_nontemporal_load((data_cxx_ty *)addr); \
+    tmp = __builtin_nontemporal_load((data_cxx_ty *)addr); \
     BytePack<bytes> ans; \
     ans.native = tmp; \
     return ans; \
@@ -262,6 +298,14 @@ template<> __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack
     __builtin_nontemporal_store(value.native, (data_cxx_ty *)addr); \
   }
 
+#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+#define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
+  DEFINE_ld_st__size_space_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty)
+#else
+#define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \
+  DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty)
+#endif
+
 // #if __CUDA_ARCH__ >= 700
 //   #define PTX_relaxed_gpu "relaxed.gpu"
 // #else
@@ -296,6 +340,8 @@ DEFINE_ld_st__size(4, uint32_t, b32, r)
 DEFINE_ld_st__size(8, uint64_t, b64, l)
 
 #undef DEFINE_ld_st__size_space
+#undef DEFINE_ld_st__size_space_hip_atomic
+#undef DEFINE_ld_st__size_space_fallback
 #undef DEFINE_ld_st__size
 
 
@@ -362,12 +408,20 @@ DEFINE_ld_st_16__space(global, uintptr_t, l)
 
 __device__ __forceinline__ uint64_t ld_volatile_global(uint64_t *ptr) {
   uint64_t ans;
+  #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  ans = __hip_atomic_load(ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  #else
   ans = __builtin_nontemporal_load(ptr);
+  #endif
   return ans;
 }
 __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t *ptr) {
   uint64_t ans;
+  #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  ans = __hip_atomic_load(ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  #else
   ans = __builtin_nontemporal_load(ptr);
+  #endif
   return ans;
 }
 
@@ -383,18 +437,34 @@ __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t *ptr) {
 
 __device__ __forceinline__ uint64_t ld_acquire_sys_global(uint64_t *ptr) {
   uint64_t ans;
+  #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  ans = __hip_atomic_load(ptr, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SYSTEM);
+  #else
   ans = __atomic_load_n(ptr ,__ATOMIC_SEQ_CST);
+  #endif
   return ans;
 }
 
 __device__ __forceinline__ void st_volatile_global(uint64_t *ptr, uint64_t val) {
+  #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  __hip_atomic_store(ptr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  #else
   __builtin_nontemporal_store(val, ptr);
+  #endif
 }
 __device__ __forceinline__ void st_relaxed_sys_global(uint64_t *ptr, uint64_t val) {
+  #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  __hip_atomic_store(ptr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  #else
   __builtin_nontemporal_store(val, ptr);
+  #endif
 }
 __device__ __forceinline__ void st_release_sys_global(uint64_t *ptr, uint64_t val) {
+  #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  __hip_atomic_store(ptr, val, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
+  #else
   __atomic_store_n(ptr, val, __ATOMIC_SEQ_CST);
+  #endif
 }
 
 __device__ __forceinline__ void fence_acq_rel_sys() {

projects/rccl/src/device/prims_ll.h

@@ -167,6 +167,8 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #ifdef __GFX11__
       asm volatile ("global_load_b128 %0, %1, off glc slc dlc\n"
         "s_waitcnt vmcnt(0)\n" : "=v"(i4.i4) : "v"(&src->i4));
+#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+      i4.v4u = __builtin_amdgcn_global_load_b128((v4u_gptr)src->v, RCCL_SYSTEM_SYNCSCOPE);
 #else
       *((u64_gptr)i4.v) = __builtin_nontemporal_load((u64_gptr)src->v);
       *((u64_gptr)i4.v + 1) = __builtin_nontemporal_load((u64_gptr)src->v+1);
@@ -210,6 +212,8 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #ifdef __GFX11__
         asm volatile ("global_load_b128 %0, %1, off glc slc dlc\n"
           "s_waitcnt vmcnt(0)\n" : "=v"(line[i].i4) : "v"(&src->i4));
+#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+        line[i].v4u = __builtin_amdgcn_global_load_b128((v4u_gptr)src->v, RCCL_SYSTEM_SYNCSCOPE);
 #else
         line[i].v[0] = __builtin_nontemporal_load(src->v);
         line[i].v[1] = __builtin_nontemporal_load(src->v+1);
@@ -237,6 +241,8 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #ifdef __GFX11__
       asm volatile ("global_load_b128 %0, %1, off glc slc dlc\n"
         "s_waitcnt vmcnt(0)\n" : "=v"(line[i].i4) : "v"(&src->i4));
+#elif RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+      line[i].v4u = __builtin_amdgcn_global_load_b128((v4u_gptr)src->v, RCCL_SYSTEM_SYNCSCOPE);
 #else
       line[i].v[0] = __builtin_nontemporal_load((u64_gptr)src->v);
       line[i].v[1] = __builtin_nontemporal_load((u64_gptr)src->v+1);
@@ -268,8 +274,13 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
     i4.flag1 = flag;
     i4.data2 = (val >> 32);
     i4.flag2 = flag;
+    #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+    // System scope store that bypasses the hardware caches, should generate global_store_dwordx4 instruction with sc0 and sc1 bits set to 1 on gfx942/gfx950.
+    __builtin_amdgcn_global_store_b128((v4u_gptr) dst->v, i4.v4u, RCCL_SYSTEM_SYNCSCOPE);
+    #else
     *((u64_gptr) dst->v) = *((u64_gptr) i4.v);
-    *((u64_gptr) dst->v+1) = *((u64_gptr) i4.v+1); 
+    *((u64_gptr) dst->v+1) = *((u64_gptr) i4.v+1);
+    #endif
 #if defined(__gfx950__) && ROCM_VERSION < 70002
     __builtin_amdgcn_fence(__ATOMIC_RELEASE, ""); // flush cache on gfx950 if ROCr fix for hipHostMallocUncached is not available (ROCm version < 7.0.2)
 #endif
@@ -290,6 +301,16 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
       uint64_t u8;
     };
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+    if(sizeof(U) == 1)
+      u1 = __hip_atomic_load((uint8_t*)src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    else if(sizeof(U) == 2)
+      u2 = __hip_atomic_load((uint16_t*)src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    else if(sizeof(U) == 4)
+      u4 = __hip_atomic_load((uint32_t*)src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    else
+      u8 = __hip_atomic_load((uint64_t*)src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#else
     if(sizeof(U) == 1)
 #ifdef __GFX11__
       u1 = __atomic_load_n((uint8_t*)src, __ATOMIC_RELAXED);
@@ -314,6 +335,7 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #else
       u8 = __builtin_nontemporal_load((u64_gptr)src);
 #endif
+#endif
 #else
     if(sizeof(U) == 1)
       asm volatile("ld.volatile.global.b8 %0,[%1];" : "=r"(u4) : "l"(src) : "memory");
@@ -338,6 +360,16 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
     };
     elt = val;
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
+#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+    if(sizeof(U) == 1)
+      __hip_atomic_store((uint8_t*)dst, u1, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    else if(sizeof(U) == 2)
+      __hip_atomic_store((uint16_t*)dst, u2, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    else if(sizeof(U) == 4)
+      __hip_atomic_store((uint32_t*)dst, u4, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    else
+      __hip_atomic_store((uint64_t*)dst, u8, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#else
     if(sizeof(U) == 1)
       __builtin_nontemporal_store(u1, (uint8_t*)dst);
     else if(sizeof(U) == 2)
@@ -349,6 +381,7 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #if defined(__gfx950__) && ROCM_VERSION < 70002
     __builtin_amdgcn_fence(__ATOMIC_RELEASE, ""); // flush cache on gfx950 if ROCr fix for hipHostMallocUncached is not available (ROCm version < 7.0.2)
 #endif
+#endif
 #else
     if(sizeof(U) == 1)
       asm volatile("st.volatile.global.b8 [%0],%1;" :: "l"(dst), "r"(u4) : "memory");

projects/rccl/src/device/rccl_ptr.h

@@ -37,6 +37,7 @@ using u8_gptr = __attribute__((address_space(1))) uint8_t*;
 
 #ifdef __HIP_DEVICE_COMPILE__
 #if (defined(__gfx942__) || defined(__gfx950__)) && __has_builtin(__builtin_amdgcn_global_load_b128) && __has_builtin(__builtin_amdgcn_global_store_b128) && !defined(DWORDX4_INTRINSICS_FORCE_OFF)
+#define RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS 1
 #pragma message "RCCL DWORDX4 Builtins Enabled on GFX942/GFX950"
 #else
 #define RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS 0

projects/rccl/src/include/device.h

@@ -96,6 +96,10 @@ struct ncclDevRedOpFull {
   uint64_t scalarArg;
 };
 
+#ifdef __HIP_DEVICE_COMPILE__
+#include "device/rccl_ptr.h"
+#endif
+
 union ncclLLFifoLine {
   /* Flags have to be *after* data, because otherwise, an incomplete receive
     from the network may receive the flag but not the data.
@@ -109,6 +113,9 @@ union ncclLLFifoLine {
   };
   uint64_t v[2];
   int4 i4;
+#if defined(__HIP_DEVICE_COMPILE__) && RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
+  v4u v4u;  /* same layout as data1,flag1,data2,flag2 for b128 load/store */
+#endif
 };
 
 #if __HIP_DEVICE_COMPILE__

projects/rccl/tools/topo_expl/Makefile

@@ -77,6 +77,7 @@ hipify:
 	hipify-perl -inplace -quiet-warnings hipify_rccl/src/device/include/*.h
 	# PATCH 1: Fix device includes - remove "device/" prefix since files are in same directory
 	sed -i 's|#include "device/|#include "|g' hipify_rccl/src/device/include/*.h
+	sed -i 's|#include "device/|#include "|g' hipify_rccl/include/*.h
 	# PATCH 2: Use NVTX stubs - topo_expl doesn't need full NVTX implementation
 	cp hipify_rccl/include/nvtx_stub.h hipify_rccl/include/nvtx.h
 	bash ../../cmake/scripts/add_unroll.sh "hipify_rccl/src/device/include/common.h"