From feed2de479c4c4af445ff8a6c4a8ef608b43e44b Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 23 Aug 2021 11:12:30 -0700 Subject: [PATCH] SWDEV-293749 - Add cast before calling the builtin Change-Id: I6224a3693f91a5b4e040bd76918ab04a56a6454b --- hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h b/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h index 63f59b0d22..bf36a3f6ae 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h +++ b/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h @@ -22,6 +22,7 @@ THE SOFTWARE. #pragma once +#ifdef __cplusplus /** * @brief Unsafe floating point rmw atomic add for gfx90a. * @@ -52,7 +53,8 @@ THE SOFTWARE. __has_builtin(__builtin_amdgcn_ds_atomic_fadd_f32) && \ __has_builtin(__builtin_amdgcn_global_atomic_fadd_f32) __device__ inline float unsafeAtomicAdd(float* addr, float value) { - if (__builtin_amdgcn_is_shared(addr)) + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)addr)) return __builtin_amdgcn_ds_atomic_fadd_f32(addr, value); else return __builtin_amdgcn_global_atomic_fadd_f32(addr, value); @@ -89,9 +91,11 @@ __device__ inline float unsafeAtomicAdd(float* addr, float value) { __has_builtin(__builtin_amdgcn_ds_atomic_fadd_f64) && \ __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f64) __device__ inline double unsafeAtomicAdd(double* addr, double value) { - if (__builtin_amdgcn_is_shared(addr)) + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)addr)) return __builtin_amdgcn_ds_atomic_fadd_f64(addr, value); else return __builtin_amdgcn_flat_atomic_fadd_f64(addr, value); } -#endif \ No newline at end of file +#endif +#endif