Skip to content
This repository has been archived by the owner on Jan 26, 2024. It is now read-only.

amd_detail: fix atomicMax and atomicMin for floating point types #40

Open
wants to merge 1 commit into
base: develop
Choose a base branch
from
Open

amd_detail: fix atomicMax and atomicMin for floating point types #40

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
84 changes: 50 additions & 34 deletions include/hip/amd_detail/amd_hip_atomic.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -434,13 +434,15 @@ float atomicMin(float* address, float val) {
#else
unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
float value = __uint_as_float(tmp);
float assumed = __uint_as_float(tmp);
float old = assumed;

while (val < value) {
value = atomicCAS(address, value, val);
while (val < assumed && __float_as_uint(assumed) != __float_as_uint(old)) {
assumed = old;
old = atomicCAS(address, assumed, val);
}

return value;
return old;
}

__device__
Expand All @@ -452,13 +454,15 @@ float atomicMin_system(float* address, float val) {
#else
unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
float value = __uint_as_float(tmp);
float assumed = __uint_as_float(tmp);
float old = assumed;

while (val < value) {
value = atomicCAS_system(address, value, val);
while (val < assumed && __float_as_uint(assumed) != __float_as_uint(old)) {
assumed = old;
old = atomicCAS_system(address, assumed, val);
}

return value;
return old;
}

__device__
Expand All @@ -470,13 +474,15 @@ double atomicMin(double* address, double val) {
#else
unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
double value = __longlong_as_double(tmp);
double assumed = __longlong_as_double(tmp);
double old = assumed;

while (val < value) {
value = atomicCAS(address, value, val);
while (val < assumed && __double_as_longlong(assumed) != __double_as_longlong(old)) {
assumed = old;
old = atomicCAS(address, assumed, val);
}

return value;
return old;
}

__device__
Expand All @@ -488,13 +494,15 @@ double atomicMin_system(double* address, double val) {
#else
unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
double value = __longlong_as_double(tmp);
double assumed = __longlong_as_double(tmp);
double old = assumed;

while (val < value) {
value = atomicCAS_system(address, value, val);
while (val < assumed && __double_as_longlong(assumed) != __double_as_longlong(old)) {
assumed = old;
old = atomicCAS_system(address, assumed, val);
}

return value;
return old;
}

__device__
Expand Down Expand Up @@ -554,13 +562,15 @@ float atomicMax(float* address, float val) {
#else
unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
float value = __uint_as_float(tmp);
float assumed = __uint_as_float(tmp);
float old = assumed;

while (value < val) {
value = atomicCAS(address, value, val);
}
while (val > assumed && __float_as_uint(assumed) != __float_as_uint(old)) {
assumed = old;
old = atomicCAS(address, assumed, val);
}

return value;
return old;
}

__device__
Expand All @@ -572,13 +582,15 @@ float atomicMax_system(float* address, float val) {
#else
unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
float value = __uint_as_float(tmp);
float assumed = __uint_as_float(tmp);
float old = assumed;

while (value < val) {
value = atomicCAS_system(address, value, val);
}
while (val > assumed && __float_as_uint(assumed) != __float_as_uint(old)) {
assumed = old;
old = atomicCAS_system(address, assumed, val);
}

return value;
return old;
}

__device__
Expand All @@ -590,13 +602,15 @@ double atomicMax(double* address, double val) {
#else
unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
double value = __longlong_as_double(tmp);
double assumed = __longlong_as_double(tmp);
double old = assumed;

while (value < val) {
value = atomicCAS(address, value, val);
while (val > assumed && __double_as_longlong(assumed) != __double_as_longlong(old)) {
assumed = old;
old = atomicCAS(address, assumed, val);
}

return value;
return old;
}

__device__
Expand All @@ -608,13 +622,15 @@ double atomicMax_system(double* address, double val) {
#else
unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
#endif
double value = __longlong_as_double(tmp);
double assumed = __longlong_as_double(tmp);
double old = assumed;

while (value < val) {
value = atomicCAS_system(address, value, val);
while (val > assumed && __double_as_longlong(assumed) != __double_as_longlong(old)) {
assumed = old;
old = atomicCAS_system(address, assumed, val);
}

return value;
return old;
}

__device__
Expand Down