From 67f6409eb1aaeeb743542e4cf8da75dc70abf3a8 Mon Sep 17 00:00:00 2001
From: "Sang, Tao" <Tao.Sang@amd.com>
Date: Wed, 11 Jun 2025 23:46:51 -0400
Subject: [PATCH] SWDEV-519446 - Part 1: update normal atomic function file
 header (#345)

Remove the scope of __has_builtin(__hip_atomic_compare_exchange_strong) == false
---
 .../include/hip/amd_detail/amd_hip_atomic.h   | 303 +-----------------
 1 file changed, 1 insertion(+), 302 deletions(-)

diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
index 0edff439de..d45c5a8b4c 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_atomic.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
@@ -26,8 +26,6 @@ THE SOFTWARE.
 #include "amd_device_functions.h"
 #endif
 
-#if __has_builtin(__hip_atomic_compare_exchange_strong)
-
 template<bool B, typename T, typename F> struct Cond_t;
 
 template<typename T, typename F> struct Cond_t<true, T, F> { using type = T; };
@@ -300,7 +298,7 @@ __device__
 inline
 void atomicAddNoRet(float* address, float val)
 {
-    __ockl_atomic_add_noret_f32(address, val);
+  unsafeAtomicAdd(address, val);
 }
 
 __device__
@@ -697,7 +695,6 @@ inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
 {
   return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
-
 }
 
 __device__
@@ -850,301 +847,3 @@ inline
 unsigned long long atomicXor_system(unsigned long long* address, unsigned long long val) {
   return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
-
-#else // __hip_atomic_compare_exchange_strong
-__device__
-inline
-unsigned short int atomicCAS(unsigned short int* address, unsigned short int compare,
-                             unsigned short int val)
-{
-    __atomic_compare_exchange_n(
-        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
-
-    return compare;
-}
-
-__device__
-inline
-int atomicCAS(int* address, int compare, int val)
-{
-    __atomic_compare_exchange_n(
-        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
-
-    return compare;
-}
-__device__
-inline
-unsigned int atomicCAS(
-    unsigned int* address, unsigned int compare, unsigned int val)
-{
-    __atomic_compare_exchange_n(
-        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
-
-    return compare;
-}
-__device__
-inline
-unsigned long long atomicCAS(
-    unsigned long long* address,
-    unsigned long long compare,
-    unsigned long long val)
-{
-    __atomic_compare_exchange_n(
-        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
-
-    return compare;
-}
-
-__device__
-inline
-int atomicAdd(int* address, int val)
-{
-    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicAdd(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicAdd(
-    unsigned long long* address, unsigned long long val)
-{
-    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-float atomicAdd(float* address, float val)
-{
-#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
-    return unsafeAtomicAdd(address, val);
-#else
-    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
-#endif
-}
-
-#if !defined(__HIPCC_RTC__)
-HIP_DEPRECATED("use atomicAdd instead")
-#endif // !defined(__HIPCC_RTC__)
-__device__
-inline
-void atomicAddNoRet(float* address, float val)
-{
-    __ockl_atomic_add_noret_f32(address, val);
-}
-
-__device__
-inline
-double atomicAdd(double* address, double val)
-{
-#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
-    return unsafeAtomicAdd(address, val);
-#else
-    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
-#endif
-}
-
-__device__
-inline
-int atomicSub(int* address, int val)
-{
-    return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicSub(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED);
-}
-
-__device__
-inline
-int atomicExch(int* address, int val)
-{
-    return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicExch(unsigned int* address, unsigned int val)
-{
-    return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicExch(unsigned long long* address, unsigned long long val)
-{
-    return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-float atomicExch(float* address, float val)
-{
-    return __uint_as_float(__atomic_exchange_n(
-        reinterpret_cast<unsigned int*>(address),
-        __float_as_uint(val),
-        __ATOMIC_RELAXED));
-}
-
-__device__
-inline
-int atomicMin(int* address, int val)
-{
-    return __atomic_fetch_min(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicMin(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_min(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicMin(
-    unsigned long long* address, unsigned long long val)
-{
-    unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
-    while (val < tmp) {
-        const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
-
-        if (tmp1 != tmp) { tmp = tmp1; continue; }
-
-        tmp = atomicCAS(address, tmp, val);
-    }
-
-    return tmp;
-}
-__device__ inline long long atomicMin(long long* address, long long val) {
-    long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
-    while (val < tmp) {
-        const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
-
-        if (tmp1 != tmp) {
-          tmp = tmp1;
-          continue;
-        }
-
-        tmp = atomicCAS(address, tmp, val);
-    }
-    return tmp;
-}
-
-__device__
-inline
-int atomicMax(int* address, int val)
-{
-    return __atomic_fetch_max(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicMax(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_max(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicMax(
-    unsigned long long* address, unsigned long long val)
-{
-    unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
-    while (tmp < val) {
-        const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
-
-        if (tmp1 != tmp) { tmp = tmp1; continue; }
-
-        tmp = atomicCAS(address, tmp, val);
-    }
-
-    return tmp;
-}
-__device__ inline long long atomicMax(long long* address, long long val) {
-    long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
-    while (tmp < val) {
-        const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
-
-        if (tmp1 != tmp) {
-          tmp = tmp1;
-          continue;
-        }
-
-        tmp = atomicCAS(address, tmp, val);
-    }
-    return tmp;
-}
-
-__device__
-inline
-unsigned int atomicInc(unsigned int* address, unsigned int val)
-{
-  return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
-}
-
-__device__
-inline
-unsigned int atomicDec(unsigned int* address, unsigned int val)
-{
-  return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
-}
-
-__device__
-inline
-int atomicAnd(int* address, int val)
-{
-    return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicAnd(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicAnd(
-    unsigned long long* address, unsigned long long val)
-{
-    return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
-}
-
-__device__
-inline
-int atomicOr(int* address, int val)
-{
-    return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicOr(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicOr(
-    unsigned long long* address, unsigned long long val)
-{
-    return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
-}
-
-__device__
-inline
-int atomicXor(int* address, int val)
-{
-    return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned int atomicXor(unsigned int* address, unsigned int val)
-{
-    return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
-}
-__device__
-inline
-unsigned long long atomicXor(
-    unsigned long long* address, unsigned long long val)
-{
-    return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
-}
-
-#endif // __hip_atomic_compare_exchange_strong