fix(rccl): replace __hip_atomic builtins with Clang scoped atomics

projects/rccl/CMakeLists.txt

@@ -110,6 +110,7 @@ option(TRACE                                   "Enable additional tracing"
 option(QP_TRACKING                             "Enable per-device QP tracking"                 OFF)
 option(FAULT_INJECTION                         "Enable fault injection"                        ON)
 option(QUIET_WARNINGS                          "Supress compiler warnings"                     OFF)
+option(RCCL_POISON_HIP_ATOMICS                 "Reject __hip_atomic_* builtins in RCCL sources" ON)
 option(DWORDX4_INTRINSICS                      "Turn off compilation with dwordx4 intrinsics"  ON)
 if(ENABLE_TDM_SIMPLE)
   add_compile_definitions(ENABLE_TDM_SIMPLE=1)

projects/rccl/README.md

@@ -85,6 +85,7 @@ RCCL build & installation helper script
     -DENABLE_COMPRESS=OFF                 Disable GPU code compression (default: ON)
     -DENABLE_IFC=ON                       Enable indirect function call (default: OFF)
     -DFAULT_INJECTION=OFF                 Disable fault injection (default: ON)
+    -DRCCL_POISON_HIP_ATOMICS=OFF         Allow __hip_atomic_* builtins in RCCL sources (default: ON)
     -DRCCL_ROCPROFILER_REGISTER=OFF       Disable rocprofiler-register support (default: ON)
     -DTIMETRACE=ON                        Enable time-trace during compilation (default: OFF)
 

projects/rccl/install.sh

@@ -118,6 +118,7 @@ function display_help()
     echo "    -DENABLE_COMPRESS=OFF                 Disable GPU code compression (default: ON)"
     echo "    -DENABLE_IFC=ON                       Enable indirect function call (default: OFF)"
     echo "    -DFAULT_INJECTION=OFF                 Disable fault injection (default: ON)"
+    echo "    -DRCCL_POISON_HIP_ATOMICS=OFF         Allow __hip_atomic_* builtins in RCCL sources (default: ON)"
     echo "    -DRCCL_ROCPROFILER_REGISTER=OFF       Disable rocprofiler-register support (default: ON)"
     echo "    -DTIMETRACE=ON                        Enable time-trace during compilation (default: OFF)"
     echo ""

projects/rccl/src/CMakeLists.txt

@@ -1209,6 +1209,26 @@ if(RCCL_ROCPROFILER_REGISTER)
   target_link_libraries(
       rccl PRIVATE rocprofiler-register::rocprofiler-register)
 endif()
+
+# Poison the __hip_atomic_* builtins
+#==================================================================================================
+# RCCL spells device atomics with the Clang scoped-atomic builtins so that memory
+# scope is explicit at every call site. Force-include (rather than #include) the
+# poison header so it also covers hipified and generated sources, and set it here,
+# before DeviceLinker.cmake snapshots COMPILE_OPTIONS, so the per-kernel device
+# compiles inherit it too. See src/include/poison_hip_atomics.h for the coverage
+# this does and does not give.
+if(RCCL_POISON_HIP_ATOMICS)
+  # Joined --include=<file> rather than separate-argument -include <file>: as a
+  # single token it needs no cooperation from anything that forwards flags one
+  # at a time, so it survives rccl-device-compile's catch-all for unrecognized
+  # `-*` tokens as-is. That script also parses the separate-argument spelling
+  # (see the -include/-imacros branch in parse_compiler_flags), so either form
+  # works today; the joined one is used because it does not depend on that.
+  target_compile_options(rccl PRIVATE --include=${RCCL_SOURCE_DIR}/src/include/poison_hip_atomics.h)
+  message(STATUS "RCCL_POISON_HIP_ATOMICS enabled - __hip_atomic_* builtins rejected in RCCL sources")
+endif()
+
 # Device Linker: include pipeline and wire up dependencies
 # (placed after all target_compile_definitions so DeviceLinker.cmake
 #  can read the rccl target's definitions directly)

projects/rccl/src/algorithms/dda/device/CollCommon_ll128.h

@@ -65,10 +65,10 @@ __host__ __device__ __forceinline__ size_t ddaLL128NumLines(size_t nWords) {
 // comes from gfx1250 preserving per-thread program order of system-scope stores,
 // not from an atomic fence. Matches the validated microbenchmark primitives.
 __device__ __forceinline__ void ddaLL128StoreWord(uint64_t* p, uint64_t v) {
-  __hip_atomic_store((u64_gptr)p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  __scoped_atomic_store_n((u64_gptr)p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
 }
 __device__ __forceinline__ uint64_t ddaLL128LoadWord(const uint64_t* p) {
-  return __hip_atomic_load((u64_gptr) const_cast<uint64_t*>(p), __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+  return __scoped_atomic_load_n((u64_gptr) const_cast<uint64_t*>(p), __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
 }
 
 // Element-wise add of the T-elements packed into two 8B payload words. An 8B

projects/rccl/src/device/all_gather.h

@@ -195,7 +195,7 @@ struct RunWorkColl<ncclFuncAllGather, T, RedOp, NCCL_ALGO_PAT, NCCL_PROTO_SIMPLE
       while (1) {
         struct ncclPatStep* ps = shmem->patSteps + (step % NCCL_SHMEM_PAT_STEPS);
         int* poll = &ps->flags;
-        while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) != 0) {
+        while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) != 0) {
           pollCount++;// Wait for workers to be done with step 'step-NCCL_SHMEM_PAT_STEPS'
         }
         patAlgo.getNextOp(ps);
@@ -224,14 +224,14 @@ struct RunWorkColl<ncclFuncAllGather, T, RedOp, NCCL_ALGO_PAT, NCCL_PROTO_SIMPLE
       while (1) {
         struct ncclPatStep* ps = shmem->patSteps + (step % NCCL_SHMEM_PAT_STEPS);
         int* poll = &ps->flags;
-        while (__hip_atomic_load(poll, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) == 0) {
+        while (__scoped_atomic_load_n(poll, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) == 0) {
           pollCount++; // Wait for compute thread
         }
         int last = ps->last;
         prims.patCopy(ps, shmem);
         if (tidInGroup == 0)
-          __hip_atomic_store(poll, 0, __ATOMIC_RELEASE,
-                             __HIP_MEMORY_SCOPE_WORKGROUP); // Return element to compute thread
+          __scoped_atomic_store_n(poll, 0, __ATOMIC_RELEASE,
+                                  __MEMORY_SCOPE_WRKGRP); // Return element to compute thread
         if (last) break;
         step += nGroups;
       }

projects/rccl/src/device/ce_reduce/ce_reduce_impl.h.in

@@ -38,12 +38,12 @@ __device__ __forceinline__ void ncclCeGlobalBlockBarrier(uint32_t* d_barrierSync
   uint32_t* doneGen = d_barrierSync + 1;
   __syncthreads();
   if (threadIdx.x == 0) {
-    uint32_t prev = __hip_atomic_fetch_add(arrival, 1, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+    uint32_t prev = __scoped_atomic_fetch_add(arrival, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
     if (prev + 1 == (uint32_t)numBlocks) {
-      __hip_atomic_store(arrival, 0, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
-      __hip_atomic_store(doneGen, gen, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_AGENT);
+      __scoped_atomic_store_n(arrival, 0, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
+      __scoped_atomic_store_n(doneGen, gen, __ATOMIC_RELEASE, __MEMORY_SCOPE_DEVICE);
     } else {
-      while (__hip_atomic_load(doneGen, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_AGENT) != gen) {
+      while (__scoped_atomic_load_n(doneGen, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_DEVICE) != gen) {
         __builtin_amdgcn_s_sleep(1);
       }
     }
@@ -249,8 +249,8 @@ __global__ __launch_bounds__(256) void ncclCeLocalReduceKernelVec(
         while (true) {
           int ready = 0;
           for (int r = 0; r < nRanks; r++) {
-            uint32_t sv = __hip_atomic_load((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], __ATOMIC_ACQUIRE,
-                                            __HIP_MEMORY_SCOPE_SYSTEM);
+            uint32_t sv = __scoped_atomic_load_n((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], __ATOMIC_ACQUIRE,
+                                                 __MEMORY_SCOPE_SYSTEM);
             if (sv == 1) {
               ready++;
             }
@@ -354,8 +354,8 @@ __global__ __launch_bounds__(256) void ncclCeLocalReduceKernelVec(
       }
       if (blockIdx.x == 0) {
         for (int r = (int)threadIdx.x; r < nRanks; r += (int)blockDim.x) {
-          __hip_atomic_store((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], 0, __ATOMIC_RELEASE,
-                             __HIP_MEMORY_SCOPE_SYSTEM);
+          __scoped_atomic_store_n((uint32_t*)&signalBuffer[(size_t)slot * nRanks + r], 0, __ATOMIC_RELEASE,
+                                  __MEMORY_SCOPE_SYSTEM);
         }
       }
     }

projects/rccl/src/device/common.h

@@ -21,8 +21,8 @@
 #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
 #define STORE(DST, SRC) \
   { \
-    __hip_atomic_store((__attribute__((address_space(1))) __typeof__(*(DST))*)(DST), (SRC), __ATOMIC_RELAXED, \
-                       __HIP_MEMORY_SCOPE_SYSTEM); \
+    __scoped_atomic_store_n((__attribute__((address_space(1))) __typeof__(*(DST))*)(DST), (SRC), __ATOMIC_RELAXED, \
+                            __MEMORY_SCOPE_SYSTEM); \
   }
 #elif defined(__GFX9__)
 #define STORE(DST, SRC) \

projects/rccl/src/device/op128.h

@@ -272,8 +272,8 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
 // 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_hip_atomic(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \
-                                            addr_reg_ty) \
+#define DEFINE_ld_st__size_space_scoped_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; \
@@ -285,16 +285,16 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
   template <> \
   __device__ __forceinline__ BytePack<bytes> ld_volatile_##space<bytes>(addr_cxx_ty addr) { \
     data_cxx_ty tmp; \
-    tmp = __hip_atomic_load((__attribute__((address_space(1))) data_cxx_ty*)addr, __ATOMIC_RELAXED, \
-                            __HIP_MEMORY_SCOPE_SYSTEM); \
+    tmp = __scoped_atomic_load_n((__attribute__((address_space(1))) data_cxx_ty*)addr, __ATOMIC_RELAXED, \
+                                 __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((__attribute__((address_space(1))) data_cxx_ty*)addr, value.native, __ATOMIC_RELAXED, \
-                       __HIP_MEMORY_SCOPE_SYSTEM); \
+    __scoped_atomic_store_n((__attribute__((address_space(1))) data_cxx_ty*)addr, value.native, __ATOMIC_RELAXED, \
+                            __MEMORY_SCOPE_SYSTEM); \
   }
 
 #define DEFINE_ld_st__size_space_fallback(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, \
@@ -322,7 +322,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
 
 #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)
+  DEFINE_ld_st__size_space_scoped_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)
@@ -359,7 +359,7 @@ __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value)
 DEFINE_ld_st__size(1, uint8_t, b8, r) DEFINE_ld_st__size(2, uint16_t, b16, h) 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_scoped_atomic
 #undef DEFINE_ld_st__size_space_fallback
 #undef DEFINE_ld_st__size
 
@@ -427,13 +427,13 @@ DEFINE_ld_st_16__space(global, uintptr_t, l)
   __device__ __forceinline__ uint64_t ld_volatile_global(uint64_t* ptr) {
   uint64_t ans;
   ans =
-    __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
   return ans;
 }
 __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t* ptr) {
   uint64_t ans;
   ans =
-    __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+    __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
   return ans;
 }
 
@@ -450,21 +450,21 @@ __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;
   ans =
-    __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SYSTEM);
+    __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)ptr, __ATOMIC_ACQUIRE, __MEMORY_SCOPE_SYSTEM);
   return ans;
 }
 
 __device__ __forceinline__ void st_volatile_global(uint64_t* ptr, uint64_t val) {
-  __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
-                     __HIP_MEMORY_SCOPE_SYSTEM);
+  __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
+                          __MEMORY_SCOPE_SYSTEM);
 }
 __device__ __forceinline__ void st_relaxed_sys_global(uint64_t* ptr, uint64_t val) {
-  __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
-                     __HIP_MEMORY_SCOPE_SYSTEM);
+  __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELAXED,
+                          __MEMORY_SCOPE_SYSTEM);
 }
 __device__ __forceinline__ void st_release_sys_global(uint64_t* ptr, uint64_t val) {
-  __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELEASE,
-                     __HIP_MEMORY_SCOPE_SYSTEM);
+  __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)ptr, val, __ATOMIC_RELEASE,
+                          __MEMORY_SCOPE_SYSTEM);
 }
 
 __device__ __forceinline__ void fence_acq_rel_sys() {

projects/rccl/src/device/primitives.h

@@ -27,10 +27,10 @@
       if (wid == 0) { \
         (BARRIER_NEXT) += (NWORKERS) / WARP_SIZE; \
         __THREAD_FENCE; \
-        __hip_atomic_fetch_add((BARRIERS_PTR), 1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_WORKGROUP); \
+        __scoped_atomic_fetch_add((BARRIERS_PTR), 1, __ATOMIC_RELEASE, __MEMORY_SCOPE_WRKGRP); \
         int spins = 0; \
         int rate_limit = 50; \
-        while (__hip_atomic_load((BARRIERS_PTR), __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP) < (BARRIER_NEXT)) { \
+        while (__scoped_atomic_load_n((BARRIERS_PTR), __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP) < (BARRIER_NEXT)) { \
           spins++; \
           if (spins == NCCL_SPINS_BEFORE_CHECK_ABORT) { \
             if (__atomic_load_n(ncclShmem.comm.abortFlag, __ATOMIC_SEQ_CST)) { \
@@ -42,7 +42,7 @@
           if (spins == 0 && rate_limit > 0) { \
             rate_limit--; \
             traceData(__LINE__, threadIdx.x, \
-                      __hip_atomic_load((BARRIERS_PTR), __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_WORKGROUP), \
+                      __scoped_atomic_load_n((BARRIERS_PTR), __ATOMIC_ACQUIRE, __MEMORY_SCOPE_WRKGRP), \
                       (BARRIER_NEXT)); \
           } \
           __builtin_amdgcn_s_sleep(1); \

projects/rccl/src/device/prims_ll.h

@@ -288,8 +288,8 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
     *((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1);
 #else
 #if defined(__gfx1200__) || defined(__gfx1201__)
-    __hip_atomic_store((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
-    __hip_atomic_store((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
+    __scoped_atomic_store_n((u64_gptr)dst->v, i4.v[0], __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
+    __scoped_atomic_store_n((u64_gptr)dst->v + 1, i4.v[1], __ATOMIC_RELEASE, __MEMORY_SCOPE_SYSTEM);
 #else
     *((u64_gptr)dst->v) = *((u64_gptr)i4.v);
     *((u64_gptr)dst->v + 1) = *((u64_gptr)i4.v + 1);
@@ -321,17 +321,17 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
 #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
     if (sizeof(U) == 1)
-      u1 =
-        __hip_atomic_load((__attribute__((address_space(1))) uint8_t*)src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+      u1 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint8_t*)src, __ATOMIC_RELAXED,
+                                  __MEMORY_SCOPE_SYSTEM);
     else if (sizeof(U) == 2)
-      u2 = __hip_atomic_load((__attribute__((address_space(1))) uint16_t*)src, __ATOMIC_RELAXED,
-                             __HIP_MEMORY_SCOPE_SYSTEM);
+      u2 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint16_t*)src, __ATOMIC_RELAXED,
+                                  __MEMORY_SCOPE_SYSTEM);
     else if (sizeof(U) == 4)
-      u4 = __hip_atomic_load((__attribute__((address_space(1))) uint32_t*)src, __ATOMIC_RELAXED,
-                             __HIP_MEMORY_SCOPE_SYSTEM);
+      u4 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint32_t*)src, __ATOMIC_RELAXED,
+                                  __MEMORY_SCOPE_SYSTEM);
     else
-      u8 = __hip_atomic_load((__attribute__((address_space(1))) uint64_t*)src, __ATOMIC_RELAXED,
-                             __HIP_MEMORY_SCOPE_SYSTEM);
+      u8 = __scoped_atomic_load_n((__attribute__((address_space(1))) uint64_t*)src, __ATOMIC_RELAXED,
+                                  __MEMORY_SCOPE_SYSTEM);
 #else
     if (sizeof(U) == 1)
 #ifdef __GFX11__
@@ -380,17 +380,17 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
 #if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS
     if (sizeof(U) == 1)
-      __hip_atomic_store((__attribute__((address_space(1))) uint8_t*)dst, u1, __ATOMIC_RELAXED,
-                         __HIP_MEMORY_SCOPE_SYSTEM);
+      __scoped_atomic_store_n((__attribute__((address_space(1))) uint8_t*)dst, u1, __ATOMIC_RELAXED,
+                              __MEMORY_SCOPE_SYSTEM);
     else if (sizeof(U) == 2)
-      __hip_atomic_store((__attribute__((address_space(1))) uint16_t*)dst, u2, __ATOMIC_RELAXED,
-                         __HIP_MEMORY_SCOPE_SYSTEM);
+      __scoped_atomic_store_n((__attribute__((address_space(1))) uint16_t*)dst, u2, __ATOMIC_RELAXED,
+                              __MEMORY_SCOPE_SYSTEM);
     else if (sizeof(U) == 4)
-      __hip_atomic_store((__attribute__((address_space(1))) uint32_t*)dst, u4, __ATOMIC_RELAXED,
-                         __HIP_MEMORY_SCOPE_SYSTEM);
+      __scoped_atomic_store_n((__attribute__((address_space(1))) uint32_t*)dst, u4, __ATOMIC_RELAXED,
+                              __MEMORY_SCOPE_SYSTEM);
     else
-      __hip_atomic_store((__attribute__((address_space(1))) uint64_t*)dst, u8, __ATOMIC_RELAXED,
-                         __HIP_MEMORY_SCOPE_SYSTEM);
+      __scoped_atomic_store_n((__attribute__((address_space(1))) uint64_t*)dst, u8, __ATOMIC_RELAXED,
+                              __MEMORY_SCOPE_SYSTEM);
 #else
     if (sizeof(U) == 1) __builtin_nontemporal_store(u1, (uint8_t*)dst);
     else if (sizeof(U) == 2) __builtin_nontemporal_store(u2, (uint16_t*)dst);