diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h index d45c5a8b4c..c3e689109d 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -275,13 +275,21 @@ unsigned long long atomicAdd_system(unsigned long long* address, unsigned long l return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); } +#if defined(__has_extension) && __has_extension(clang_atomic_attributes) +#define __HIP_FINE_GRAINED_MEMORY [[clang::atomic(fine_grained_memory)]] +#else +#define __HIP_FINE_GRAINED_MEMORY +#endif + __device__ inline float atomicAdd(float* address, float val) { #if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) return unsafeAtomicAdd(address, val); #else - return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __HIP_FINE_GRAINED_MEMORY { + return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } #endif } @@ -307,7 +315,9 @@ double atomicAdd(double* address, double val) { #if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) return unsafeAtomicAdd(address, val); #else - return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __HIP_FINE_GRAINED_MEMORY { + return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } #endif } @@ -371,7 +381,9 @@ float atomicSub(float* address, float val) { #if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) return unsafeAtomicAdd(address, -val); #else - return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __HIP_FINE_GRAINED_MEMORY { + return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } #endif } @@ -387,7 +399,9 @@ double atomicSub(double* address, double val) { #if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) return unsafeAtomicAdd(address, -val); #else - return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __HIP_FINE_GRAINED_MEMORY { + return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } #endif } @@ -529,12 +543,6 @@ long long atomicMin_system(long long* address, long long val) { return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); } -#if defined(__has_extension) && __has_extension(clang_atomic_attributes) -#define __HIP_FINE_GRAINED_MEMORY [[clang::atomic(fine_grained_memory)]] -#else -#define __HIP_FINE_GRAINED_MEMORY -#endif - __device__ inline float atomicMin(float* addr, float val) {