Skip to content
Snippets Groups Projects
Commit 77a22268 authored by Richard Angersbach's avatar Richard Angersbach
Browse files

Avoid duplicate definition of atomicMin/Max for HIP

parent 0c40ed63
No related branches found
No related tags found
1 merge request!438Reduction Support
Pipeline #74237 passed
......@@ -13,10 +13,11 @@ typedef __hip_uint16_t uint16_t;
typedef __hip_int16_t int16_t;
#endif
#ifdef __CUDA_ARCH__
// No direct implementation of atomic multiplication, minimum and maximum available
// No direct implementation for all atomic operations available
// -> add support by custom implementations using a CAS mechanism
#if defined(__CUDA_ARCH__) || defined(__HIPCC_RTC__)
// - atomicMul (double/float)
// see https://stackoverflow.com/questions/43354798/atomic-multiplication-and-division
__device__ double atomicMul(double* address, double val) {
......@@ -43,6 +44,10 @@ __device__ float atomicMul(float* address, float val) {
return __int_as_float(old);
}
#endif
#ifdef __CUDA_ARCH__
// - atomicMin (double/float)
// see https://stackoverflow.com/questions/17399119/how-do-i-use-atomicmax-on-floating-point-values-in-cuda
__device__ __forceinline__ double atomicMin(double *address, double val)
......@@ -94,4 +99,5 @@ __device__ __forceinline__ float atomicMax(float *address, float val)
}
return __int_as_float(ret);
}
#endif
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment