AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
cf07633dc88fb0d024f9db02607d4da672a83557
rocm-systems/hipamd/include/hip/amd_detail/amd_hip_atomic.h
T
Joseph Greathouse f68149eafd SWDEV-332811 - Clean up and extend HIP unsafe atomic add 
Update HIP's unsafeAtomicAdd to:
 - Compile properly even when not compiling for gfx90a
 - Fall back to safe atomic add on non-gfx90a architectures
 - use flat atomic add for FP64 on gfx90a, instead of dynamically
   checking memory spaces.

In addition, when the compiler is passed -munsafe-fp-atomics, it
will define __AMDGCN_UNSAFE_FP_ATOMICS__. When this happens, the
compiler is requesting that the HIP headers force all HIP
atomicAdd() calls on floats or doubles to use their unsafe versions.

This patch thus causes unsafeAtomicAdd() calls when that define
is seen. This call to unsafeAtomicAdd() is also done for atomicSub(),
since that calls atomicAdd underneath. This is not done for
system-scope atomicAdd because, on gfx90a, system-scope atomic FP
add instructions would need to target fine-grained memory, which is
always unsafe.

This patch also creates safeAtomicAdd() functions for float and double.
These functions will create a standalone safe atomic, even when the
application is compiled with -munsafe-fp-atomics.

Finally, this patch adds wrappers in the Nvidia path of HIP so that
these HIP functions call through to atomicAdd there as well.

Change-Id: I8af0621d3d28ea30c9278bfeea7393d03bbdac6d
2022-04-29 07:19:58 -04:00

1077 lines
30 KiB
C
Raw Blame History

/*
Copyright (c) 2015 - Present 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.
*/
#pragma once
#include "amd_device_functions.h"
#if __has_builtin(__hip_atomic_compare_exchange_strong)
#if !__HIP_DEVICE_COMPILE__
//TODO: Remove this after compiler pre-defines the following Macros.
#define __HIP_MEMORY_SCOPE_SINGLETHREAD 1
#define __HIP_MEMORY_SCOPE_WAVEFRONT 2
#define __HIP_MEMORY_SCOPE_WORKGROUP 3
#define __HIP_MEMORY_SCOPE_AGENT 4
#define __HIP_MEMORY_SCOPE_SYSTEM 5
#endif
#include "amd_hip_unsafe_atomics.h"
__device__
inline
int atomicCAS(int* address, int compare, int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return compare;
}
__device__
inline
int atomicCAS_system(int* address, int compare, int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
return compare;
}
__device__
inline
unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return compare;
}
__device__
inline
unsigned int atomicCAS_system(unsigned int* address, unsigned int compare, unsigned int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
return compare;
}
__device__
inline
unsigned long atomicCAS(unsigned long* address, unsigned long compare, unsigned long val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return compare;
}
__device__
inline
unsigned long atomicCAS_system(unsigned long* address, unsigned long compare, unsigned long val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
return compare;
}
__device__
inline
unsigned long long atomicCAS(unsigned long long* address, unsigned long long compare,
unsigned long long val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return compare;
}
__device__
inline
unsigned long long atomicCAS_system(unsigned long long* address, unsigned long long compare,
unsigned long long val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
return compare;
}
__device__
inline
float atomicCAS(float* address, float compare, float val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return compare;
}
__device__
inline
float atomicCAS_system(float* address, float compare, float val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
return compare;
}
__device__
inline
double atomicCAS(double* address, double compare, double val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return compare;
}
__device__
inline
double atomicCAS_system(double* address, double compare, double val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
return compare;
}
__device__
inline
int atomicAdd(int* address, int val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicAdd_system(int* address, int val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicAdd(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicAdd_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicAdd(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicAdd_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicAdd(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicAdd_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__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);
#endif
}
__device__
inline
float atomicAdd_system(float* address, float val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
#if !defined(__HIPCC_RTC__)
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 __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
#endif
}
__device__
inline
double atomicAdd_system(double* address, double val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
int atomicSub(int* address, int val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicSub_system(int* address, int val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicSub(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicSub_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicSub(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicSub_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicSub(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicSub_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
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);
#endif
}
__device__
inline
float atomicSub_system(float* address, float val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
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);
#endif
}
__device__
inline
double atomicSub_system(double* address, double val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
int atomicExch(int* address, int val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicExch_system(int* address, int val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicExch(unsigned int* address, unsigned int val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicExch_system(unsigned int* address, unsigned int val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicExch(unsigned long* address, unsigned long val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicExch_system(unsigned long* address, unsigned long val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicExch(unsigned long long* address, unsigned long long val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicExch_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
float atomicExch(float* address, float val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
float atomicExch_system(float* address, float val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
double atomicExch(double* address, double val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
double atomicExch_system(double* address, double val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
int atomicMin(int* address, int val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicMin_system(int* address, int val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicMin(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicMin_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicMin(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicMin_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicMin(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicMin_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
float atomicMin(float* address, float val) {
unsigned int* uaddr { reinterpret_cast<unsigned int*>(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);
}
return value;
}
__device__
inline
float atomicMin_system(float* address, float val) {
unsigned int* uaddr { reinterpret_cast<unsigned int*>(address) };
#if __has_builtin(__hip_atomic_load)
unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
#else
unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
float value = __uint_as_float(tmp);
while (val < value) {
value = atomicCAS_system(address, value, val);
}
return value;
}
__device__
inline
double atomicMin(double* address, double val) {
unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(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);
}
return value;
}
__device__
inline
double atomicMin_system(double* address, double val) {
unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(address) };
#if __has_builtin(__hip_atomic_load)
unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
#else
unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
double value = __longlong_as_double(tmp);
while (val < value) {
value = atomicCAS_system(address, value, val);
}
return value;
}
__device__
inline
int atomicMax(int* address, int val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicMax_system(int* address, int val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicMax(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicMax_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicMax(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicMax_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicMax(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicMax_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
float atomicMax(float* address, float val) {
unsigned int* uaddr { reinterpret_cast<unsigned int*>(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);
}
return value;
}
__device__
inline
float atomicMax_system(float* address, float val) {
unsigned int* uaddr { reinterpret_cast<unsigned int*>(address) };
#if __has_builtin(__hip_atomic_load)
unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
#else
unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
float value = __uint_as_float(tmp);
while (value < val) {
value = atomicCAS_system(address, value, val);
}
return value;
}
__device__
inline
double atomicMax(double* address, double val) {
unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(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);
}
return value;
}
__device__
inline
double atomicMax_system(double* address, double val) {
unsigned long long* uaddr { reinterpret_cast<unsigned long long*>(address) };
#if __has_builtin(__hip_atomic_load)
unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
#else
unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
double value = __longlong_as_double(tmp);
while (value < val) {
value = atomicCAS_system(address, value, val);
}
return value;
}
__device__
inline
unsigned int atomicInc(unsigned int* address, unsigned int val)
{
__device__
extern
unsigned int __builtin_amdgcn_atomic_inc(
unsigned int*,
unsigned int,
unsigned int,
unsigned int,
bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32");
return __builtin_amdgcn_atomic_inc(
address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
}
__device__
inline
unsigned int atomicDec(unsigned int* address, unsigned int val)
{
__device__
extern
unsigned int __builtin_amdgcn_atomic_dec(
unsigned int*,
unsigned int,
unsigned int,
unsigned int,
bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32");
return __builtin_amdgcn_atomic_dec(
address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
}
__device__
inline
int atomicAnd(int* address, int val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicAnd_system(int* address, int val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicAnd(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicAnd_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicAnd(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicAnd_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicAnd(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicAnd_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
int atomicOr(int* address, int val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicOr_system(int* address, int val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicOr(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicOr_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicOr(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicOr_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicOr(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long long atomicOr_system(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
int atomicXor(int* address, int val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
int atomicXor_system(int* address, int val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned int atomicXor(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned int atomicXor_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long atomicXor(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
inline
unsigned long atomicXor_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__
inline
unsigned long long atomicXor(unsigned long long* address, unsigned long long val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__
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
__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__)
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
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
unsigned int atomicInc(unsigned int* address, unsigned int val)
{
__device__
extern
unsigned int __builtin_amdgcn_atomic_inc(
unsigned int*,
unsigned int,
unsigned int,
unsigned int,
bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32");
return __builtin_amdgcn_atomic_inc(
address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
}
__device__
inline
unsigned int atomicDec(unsigned int* address, unsigned int val)
{
__device__
extern
unsigned int __builtin_amdgcn_atomic_dec(
unsigned int*,
unsigned int,
unsigned int,
unsigned int,
bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32");
return __builtin_amdgcn_atomic_dec(
address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
}
__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
Reference in New Issue View Git Blame Copy Permalink