AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-519446 - Part 1: update normal atomic function file header (#345)

Browse Source 
Remove the scope of __has_builtin(__hip_atomic_compare_exchange_strong) == false
This commit is contained in:
Sang, Tao
2025-06-11 23:46:51 -04:00
committed by GitHub
parent b7778ecd65
commit 67f6409eb1
1 changed files with 1 additions and 302 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_hip_atomic.h
+1 -302
View File
@@ -26,8 +26,6 @@ THE SOFTWARE.
#include "amd_device_functions.h" #include "amd_device_functions.h"
#endif #endif
#if __has_builtin(__hip_atomic_compare_exchange_strong)
template<bool B, typename T, typename F> struct Cond_t; template<bool B, typename T, typename F> struct Cond_t;
template<typename T, typename F> struct Cond_t<true, T, F> { using type = T; }; template<typename T, typename F> struct Cond_t<true, T, F> { using type = T; };
@@ -300,7 +298,7 @@ __device__
inline inline
void atomicAddNoRet(float* address, float val) void atomicAddNoRet(float* address, float val)
{ {
__ockl_atomic_add_noret_f32(address, val); unsafeAtomicAdd(address, val);
} }
__device__ __device__
@@ -697,7 +695,6 @@ inline
unsigned int atomicInc(unsigned int* address, unsigned int val) unsigned int atomicInc(unsigned int* address, unsigned int val)
{ {
return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent"); return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
} }
__device__ __device__
@@ -850,301 +847,3 @@ inline
unsigned long long atomicXor_system(unsigned long long* address, unsigned long long val) { 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); 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