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

SWDEV-399522 - Fix the issue #2 raised on github clr repo. Adds signed long long atomics

Browse Source 
Change-Id: I09c386d45e219e472d2da1449f678b9d17fe3e2e
This commit is contained in:
Jatin Chaudhary
2023-05-01 16:32:40 +01:00
committed by Jatin Jaikishan Chaudhary
parent a6e72de83d
commit be02e1e524
1 changed files with 84 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_hip_atomic.h
+84
View File
@@ -576,6 +576,34 @@ unsigned long long atomicMin_system(unsigned long long* address, unsigned long l
#endif // __gfx941__
}
__device__
inline
long long atomicMin(long long* address, long long val) {
#if defined(__gfx941__)
return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
address, val, [](long long x, long long y) { return x < y; },
[=]() {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
});
#else
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
#endif // __gfx941__
}
__device__
inline
long long atomicMin_system(long long* address, long long val) {
#if defined(__gfx941__)
return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
address, val, [](long long x, long long y) { return x < y; },
[=]() {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
});
#else
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
#endif // __gfx941__
}
__device__
inline
float atomicMin(float* addr, float val) {
@@ -792,6 +820,34 @@ unsigned long long atomicMax_system(unsigned long long* address, unsigned long l
#endif // __gfx941__
}
__device__
inline
long long atomicMax(long long* address, long long val) {
#if defined(__gfx941__)
return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
address, val, [](long long x, long long y) { return y < x; },
[=]() {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
});
#else
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
#endif // __gfx941__
}
__device__
inline
long long atomicMax_system(long long* address, long long val) {
#if defined(__gfx941__)
return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
address, val, [](long long x, long long y) { return y < x; },
[=]() {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
});
#else
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
#endif // __gfx941__
}
__device__
inline
float atomicMax(float* addr, float val) {
@@ -1437,6 +1493,20 @@ unsigned long long atomicMin(
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
@@ -1466,6 +1536,20 @@ unsigned long long atomicMax(
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