From d2904ffe47f08116bf373eb29a450cccb494adb7 Mon Sep 17 00:00:00 2001 From: Ioannis Assiouras Date: Thu, 2 Feb 2023 12:08:59 +0000 Subject: [PATCH] SWDEV-380687 - Fixed typo in safe atomicAdd for gfx90a Change-Id: I87bc6d9e0ef1f564d679c6280c71c4633d3aa619 [ROCm/clr commit: 6704f3d0f0cc3eaadc80b24552f19a0110694a04] --- .../clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 7c2a329e4f..0100e99e71 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 @@ -310,7 +310,7 @@ __device__ inline double unsafeAtomicMin(double* addr, double val) { */ __device__ inline float safeAtomicAdd(float* addr, float value) { #if defined(__gfx908__) || \ - (defined(__gfx90a) && !__has_builtin(__hip_atomic_fetch_add)) + (defined(__gfx90a__) && !__has_builtin(__hip_atomic_fetch_add)) // On gfx908, we can generate unsafe FP32 atomic add that does not follow all // IEEE rules when -munsafe-fp-atomics is passed. Do a CAS loop emulation instead. // On gfx90a, if we do not have the __hip_atomic_fetch_add builtin, we need to