From 5a402b9c99490e472c01e9c9cbb9e217e2b67eea Mon Sep 17 00:00:00 2001 From: Brian Sumner Date: Mon, 11 Jul 2022 13:57:28 -0700 Subject: [PATCH] SWDEV-333033 - add safe and unsafe atomic min and max including gfx940 and add missing nvidia support Change-Id: I829a0a5fd49c510e77eabbcb92d1a415ef6b5a4c [ROCm/clr commit: 4375b9f5b9b2322c45f7334c707f38011fda0644] --- .../include/hip/amd_detail/amd_hip_atomic.h | 136 ++++--- .../hip/amd_detail/amd_hip_unsafe_atomics.h | 356 +++++++++++++++++- .../hip/nvidia_detail/nvidia_hip_atomics.h | 75 ++++ .../hip/nvidia_detail/nvidia_hip_runtime.h | 1 + .../nvidia_detail/nvidia_hip_unsafe_atomics.h | 32 ++ 5 files changed, 548 insertions(+), 52 deletions(-) create mode 100644 projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_atomics.h diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h index 80e0b3f2a8..aff272e2e3 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -427,20 +427,30 @@ unsigned long long atomicMin_system(unsigned long long* address, unsigned long l __device__ inline -float atomicMin(float* address, float val) { - unsigned int* uaddr { reinterpret_cast(address) }; - #if __has_builtin(__hip_atomic_load) - unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)}; - #else - unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; - #endif - float value = __uint_as_float(tmp); - - while (val < value) { - value = atomicCAS(address, value, val); +float atomicMin(float* addr, float val) { +#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) + return unsafeAtomicMin(address, val); +#else + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value > val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); } - return value; + #else + unsigned int *uaddr = (unsigned int *)addr; + unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __uint_as_float(value) > val) { + done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __uint_as_float(value); + #endif +#endif } __device__ @@ -463,20 +473,30 @@ float atomicMin_system(float* address, float val) { __device__ inline -double atomicMin(double* address, double val) { - unsigned long long* uaddr { reinterpret_cast(address) }; - #if __has_builtin(__hip_atomic_load) - unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)}; - #else - unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; - #endif - double value = __longlong_as_double(tmp); - - while (val < value) { - value = atomicCAS(address, value, val); +double atomicMin(double* addr, double val) { +#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) + return unsafeAtomicMin(address, val); +#else + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value > val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); } - return value; + #else + unsigned long long *uaddr = (unsigned long long *)addr; + unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __longlong_as_double(value) > val) { + done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __longlong_as_double(value); + #endif +#endif } __device__ @@ -547,20 +567,30 @@ unsigned long long atomicMax_system(unsigned long long* address, unsigned long l __device__ inline -float atomicMax(float* address, float val) { - unsigned int* uaddr { reinterpret_cast(address) }; - #if __has_builtin(__hip_atomic_load) - unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)}; - #else - unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; - #endif - float value = __uint_as_float(tmp); - - while (value < val) { - value = atomicCAS(address, value, val); +float atomicMax(float* addr, float val) { +#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) + return unsafeAtomicMax(addr, val); +#else + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value < val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); } - return value; + #else + unsigned int *uaddr = (unsigned int *)addr; + unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __uint_as_float(value) < val) { + done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __uint_as_float(value); + #endif +#endif } __device__ @@ -583,20 +613,30 @@ float atomicMax_system(float* address, float val) { __device__ inline -double atomicMax(double* address, double val) { - unsigned long long* uaddr { reinterpret_cast(address) }; - #if __has_builtin(__hip_atomic_load) - unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT)}; - #else - unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; - #endif - double value = __longlong_as_double(tmp); - - while (value < val) { - value = atomicCAS(address, value, val); +double atomicMax(double* addr, double val) { +#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__) + return unsafeAtomicMax(addr, val); +#else + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value < val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); } - return value; + #else + unsigned long long *uaddr = (unsigned long long *)addr; + unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __longlong_as_double(value) < val) { + done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __longlong_as_double(value); + #endif +#endif } __device__ diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h index ba5795897b..243b5a64d9 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h @@ -50,7 +50,10 @@ THE SOFTWARE. * @return Original value contained in \p addr. */ __device__ inline float unsafeAtomicAdd(float* addr, float value) { -#if defined(__gfx90a__) && \ +#if defined(__gfx940__) && \ + __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f32) + return __builtin_amdgcn_flat_atomic_fadd_f32(addr, value); +#elif defined(__gfx90a__) && \ __has_builtin(__builtin_amdgcn_is_shared) && \ __has_builtin(__builtin_amdgcn_is_private) && \ __has_builtin(__builtin_amdgcn_ds_atomic_fadd_f32) && \ @@ -73,6 +76,78 @@ __device__ inline float unsafeAtomicAdd(float* addr, float value) { #endif } +/** + * @brief Unsafe floating point rmw atomic max. + * + * Performs a relaxed read-modify-write floating point atomic max with + * device memory scope. The original value at \p addr is returned and + * the value at \p addr is replaced by \p val if greater. + * + * @note This operation is currently identical to that performed by + * atomicMax and is included for completeness. + * + * @param [in,out] addr Pointer to value to be updated + * @param [in] val Value used to update the value at \p addr. + * @return Original value contained in \p addr. + */ +__device__ inline float unsafeAtomicMax(float* addr, float val) { + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value < val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned int *uaddr = (unsigned int *)addr; + unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __uint_as_float(value) < val) { + done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __uint_as_float(value); + #endif +} + +/** + * @brief Unsafe floating point rmw atomic min. + * + * Performs a relaxed read-modify-write floating point atomic min with + * device memory scope. The original value at \p addr is returned and + * the value at \p addr is replaced by \p val if lesser. + * + * @note This operation is currently identical to that performed by + * atomicMin and is included for completeness. + * + * @param [in,out] addr Pointer to value to be updated + * @param [in] val Value used to update the value at \p addr. + * @return Original value contained in \p addr. + */ +__device__ inline float unsafeAtomicMin(float* addr, float val) { + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value > val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned int *uaddr = (unsigned int *)addr; + unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __uint_as_float(value) > val) { + done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __uint_as_float(value); + #endif +} + /** * @brief Unsafe double precision rmw atomic add. * @@ -95,12 +170,12 @@ __device__ inline float unsafeAtomicAdd(float* addr, float value) { * Passing in global segment addresses in fine grain allocations will result in * undefined behavior and are not supported. * - * @param [in,out] addr Pointer to value to be increment by \p value. + * @param [in,out] addr Pointer to value to be updated. * @param [in] value Value by \p addr is to be incremented. * @return Original value contained in \p addr. */ __device__ inline double unsafeAtomicAdd(double* addr, double value) { -#if defined(__gfx90a__) && \ +#if (defined(__gfx90a__) || defined(__gfx940_)) && \ __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f64) return __builtin_amdgcn_flat_atomic_fadd_f64(addr, value); #elif defined (__hip_atomic_fetch_add) @@ -110,6 +185,112 @@ __device__ inline double unsafeAtomicAdd(double* addr, double value) { #endif } +/** + * @brief Unsafe double precision rmw atomic max. + * + * Performs a relaxed read-modify-write double precision atomic max with + * device memory scope. Original value at \p addr is returned and + * the value of \p addr is updated with \p val if greater. + * + * @note This operation currently only performs different operations for + * the gfx90a target. Other devices continue to use safe atomics. + * + * It can be used to generate code that uses fast hardware floating point atomic + * operations which may handle rounding and subnormal values differently than + * non-atomic floating point operations. + * + * The operation is not always safe and can have undefined behavior unless + * following condition are met: + * + * - \p addr is at least 8 byte aligned + * - If \p addr is a global segment address, it is in a coarse grain allocation. + * Passing in global segment addresses in fine grain allocations will result in + * undefined behavior and are not supported. + * + * @param [in,out] addr Pointer to value to be updated. + * @param [in] val Value used to updated the contents at \p addr + * @return Original value contained at \p addr. + */ +__device__ inline double unsafeAtomicMax(double* addr, double val) { +#if (defined(__gfx90a__) || defined(__gfx940__)) && \ + __has_builtin(__builtin_amdgcn_flat_atomic_fmax_f64) + return __builtin_amdgcn_flat_atomic_fmax_f64(addr, val); +#else + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value < val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned long long *uaddr = (unsigned long long *)addr; + unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __longlong_as_double(value) < val) { + done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __longlong_as_double(value); + #endif +#endif +} + +/** + * @brief Unsafe double precision rmw atomic min. + * + * Performs a relaxed read-modify-write double precision atomic min with + * device memory scope. Original value at \p addr is returned and + * the value of \p addr is updated with \p val if lesser. + * + * @note This operation currently only performs different operations for + * the gfx90a target. Other devices continue to use safe atomics. + * + * It can be used to generate code that uses fast hardware floating point atomic + * operations which may handle rounding and subnormal values differently than + * non-atomic floating point operations. + * + * The operation is not always safe and can have undefined behavior unless + * following condition are met: + * + * - \p addr is at least 8 byte aligned + * - If \p addr is a global segment address, it is in a coarse grain allocation. + * Passing in global segment addresses in fine grain allocations will result in + * undefined behavior and are not supported. + * + * @param [in,out] addr Pointer to value to be updated. + * @param [in] val Value used to updated the contents at \p addr + * @return Original value contained at \p addr. + */ +__device__ inline double unsafeAtomicMin(double* addr, double val) { +#if (defined(__gfx90a__) || defined(__gfx940__)) && \ + __has_builtin(__builtin_amdgcn_flat_atomic_fmin_f64) + return __builtin_amdgcn_flat_atomic_fmin_f64(addr, val); +#else + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value > val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned long long *uaddr = (unsigned long long *)addr; + unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __longlong_as_double(value) > val) { + done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __longlong_as_double(value); + #endif +#endif +} + /** * @brief Safe floating point rmw atomic add. * @@ -163,6 +344,78 @@ __device__ inline float safeAtomicAdd(float* addr, float value) { #endif } +/** + * @brief Safe floating point rmw atomic max. + * + * Performs a relaxed read-modify-write floating point atomic max with + * device memory scope. The original value at \p addr is returned and + * the value at \p addr is replaced by \p val if greater. + * + * @note This operation ensures that, on all targets, we produce safe atomics. + * This will be the case even when -munsafe-fp-atomics is passed into the compiler. + * + * @param [in,out] addr Pointer to value to be updated + * @param [in] val Value used to update the value at \p addr. + * @return Original value contained in \p addr. + */ +__device__ inline float safeAtomicMax(float* addr, float val) { + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value < val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned int *uaddr = (unsigned int *)addr; + unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __uint_as_float(value) < val) { + done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __uint_as_float(value); + #endif +} + +/** + * @brief Safe floating point rmw atomic min. + * + * Performs a relaxed read-modify-write floating point atomic min with + * device memory scope. The original value at \p addr is returned and + * the value at \p addr is replaced by \p val if lesser. + * + * @note This operation ensures that, on all targets, we produce safe atomics. + * This will be the case even when -munsafe-fp-atomics is passed into the compiler. + * + * @param [in,out] addr Pointer to value to be updated + * @param [in] val Value used to update the value at \p addr. + * @return Original value contained in \p addr. + */ +__device__ inline float safeAtomicMin(float* addr, float val) { + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value > val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned int *uaddr = (unsigned int *)addr; + unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __uint_as_float(value) > val) { + done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __uint_as_float(value); + #endif +} + /** * @brief Safe double precision rmw atomic add. * @@ -178,7 +431,7 @@ __device__ inline float safeAtomicAdd(float* addr, float value) { * @return Original value contained in \p addr. */ __device__ inline double safeAtomicAdd(double* addr, double value) { -#if defined(__gfx90a__) && \ +#if (defined(__gfx90a__) || defined(__gfx940__)) && \ __has_builtin(__hip_atomic_fetch_add) // On gfx90a, with the __hip_atomic_fetch_add builtin, relaxed system-scope // atomics will produce safe CAS loops, but are otherwise not different than @@ -215,4 +468,99 @@ __device__ inline double safeAtomicAdd(double* addr, double value) { #endif // __has_builtin(__hip_atomic_fetch_add) #endif } + +/** + * @brief Safe double precision rmw atomic max. + * + * Performs a relaxed read-modify-write double precision atomic max with + * device memory scope. Original value at \p addr is returned and + * the value of \p addr is updated with \p val if greater. + * + * @note This operation ensures that, on all targets, we produce safe atomics. + * This will be the case even when -munsafe-fp-atomics is passed into the compiler. + * + * @param [in,out] addr Pointer to value to be updated. + * @param [in] val Value used to updated the contents at \p addr + * @return Original value contained at \p addr. + */ +__device__ inline double safeAtomicMax(double* addr, double val) { + #if __has_builtin(__builtin_amdgcn_is_private) + if (__builtin_amdgcn_is_private( + (const __attribute__((address_space(0))) void*)addr)) { + double old = *addr; + *addr = __builtin_fmax(old, val); + return old; + } else { + #endif + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value < val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned long long *uaddr = (unsigned long long *)addr; + unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __longlong_as_double(value) < val) { + done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __longlong_as_double(value); + #endif + #if __has_builtin(__builtin_amdgcn_is_private) + } + #endif +} + +/** + * @brief Safe double precision rmw atomic min. + * + * Performs a relaxed read-modify-write double precision atomic min with + * device memory scope. Original value at \p addr is returned and + * the value of \p addr is updated with \p val if lesser. + * + * @note This operation ensures that, on all targets, we produce safe atomics. + * This will be the case even when -munsafe-fp-atomics is passed into the compiler. + * + * @param [in,out] addr Pointer to value to be updated. + * @param [in] val Value used to updated the contents at \p addr + * @return Original value contained at \p addr. + */ +__device__ inline double safeAtomicMin(double* addr, double val) { + #if __has_builtin(__builtin_amdgcn_is_private) + if (__builtin_amdgcn_is_private( + (const __attribute__((address_space(0))) void*)addr)) { + double old = *addr; + *addr = __builtin_fmin(old, val); + return old; + } else { + #endif + #if __has_builtin(__hip_atomic_load) && \ + __has_builtin(__hip_atomic_compare_exchange_strong) + double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + bool done = false; + while (!done && value > val) { + done = __hip_atomic_compare_exchange_strong(addr, &value, val, + __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + return value; + #else + unsigned long long *uaddr = (unsigned long long *)addr; + unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED); + bool done = false; + while (!done && __longlong_as_double(value) > val) { + done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false, + __ATOMIC_RELAXED, __ATOMIC_RELAXED); + } + return __longlong_as_double(value); + #endif + #if __has_builtin(__builtin_amdgcn_is_private) + } + #endif +} + #endif diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_atomics.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_atomics.h new file mode 100644 index 0000000000..f9a92d582a --- /dev/null +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_atomics.h @@ -0,0 +1,75 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_ATOMICS_H +#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_ATOMICS_H + + +__device__ inline float atomicMax(float* addr, float val) { + unsigned int *uaddr = (unsigned int *)addr; + float value = __uint_as_float(*uaddr); + + while (value < val) { + value = __uint_as_float(atomicCAS(uaddr, __float_as_uint(value), + __float_as_uint(val))); + } + return value; +} + +__device__ inline double atomicMax(double* addr, double val) { + unsigned long long* uaddr = (unsigned long long *)addr; + double value = __longlong_as_double(*uaddr); + + while (value < val) { + value = __longlong_as_double(atomicCAS(uaddr, + __double_as_longlong(value), + __double_as_longlong(val))); + } + + return value; +} + +__device__ inline float atomicMin(float* addr, float val) { + unsigned int *uaddr = (unsigned int *)addr; + float value = __uint_as_float(*uaddr); + + while (value > val) { + value = __uint_as_float(atomicCAS(uaddr, __float_as_uint(value), + __float_as_uint(val))); + } + return value; +} + +__device__ inline double atomicMin(double* addr, double val) { + unsigned long long* uaddr = (unsigned long long *)addr; + double value = __longlong_as_double(*uaddr); + + while (value > val) { + value = __longlong_as_double(atomicCAS(uaddr, + __double_as_longlong(value), + __double_as_longlong(val))); + } + + return value; +} + +#endif diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime.h index 19be62c6f6..c63e35700b 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime.h +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime.h @@ -76,6 +76,7 @@ typedef int hipLaunchParm; #ifdef __CUDACC__ +#include "nvidia_hip_atomics.h" #include "nvidia_hip_unsafe_atomics.h" #define hipThreadIdx_x threadIdx.x diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_unsafe_atomics.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_unsafe_atomics.h index 919353129a..993f17507b 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_unsafe_atomics.h +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_unsafe_atomics.h @@ -44,6 +44,22 @@ __device__ inline double unsafeAtomicAdd(double* addr, double value) { #endif } +__device__ inline float unsafeAtomicMax(float* addr, float value) { + return atomicMax(addr, value); +} + +__device__ inline double unsafeAtomicMax(double* addr, double val) { + return atomicMax(addr, val); +} + +__device__ inline float unsafeAtomicMin(float* addr, float value) { + return atomicMin(addr, value); +} + +__device__ inline double unsafeAtomicMin(double* addr, double val) { + return atomicMin(addr, val); +} + __device__ inline float safeAtomicAdd(float* addr, float value) { return atomicAdd(addr, value); } @@ -65,4 +81,20 @@ __device__ inline double safeAtomicAdd(double* addr, double value) { #endif } +__device__ inline float safeAtomicMax(float* addr, float value) { + return atomicMax(addr, value); +} + +__device__ inline double safeAtomicMax(double* addr, double val) { + return atomicMax(addr, val); +} + +__device__ inline float safeAtomicMin(float* addr, float value) { + return atomicMin(addr, value); +} + +__device__ inline double safeAtomicMin(double* addr, double val) { + return atomicMin(addr, val); +} + #endif