From dcb5d29f1efabd4eb16b7aca1240907aa29e2246 Mon Sep 17 00:00:00 2001 From: Pierre van Houtryve Date: Wed, 28 Sep 2022 14:15:25 +0000 Subject: [PATCH] Replace LLVM `atomic.inc` intrinsics with Clang builtins With opaque pointers, the suffix of those intrinsics changed. This caused build failures that should be solved by just using the corresponding Clang builtins instead of using intrinsics directly See SWDEV-356581 Change-Id: Icd1d9b9438cac4bef0f7c52d4cd341ac76500890 --- .../include/hip/amd_detail/amd_hip_atomic.h | 52 +++---------------- 1 file changed, 8 insertions(+), 44 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h index 8cd4265647..0f516364f0 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -661,34 +661,16 @@ __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); + return __builtin_amdgcn_atomic_inc32( + address, val, __ATOMIC_RELAXED, "workgroup"); } __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); + return __builtin_amdgcn_atomic_dec32( + address, val, __ATOMIC_RELAXED, "workgroup"); } __device__ @@ -1023,34 +1005,16 @@ __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); + return __builtin_amdgcn_atomic_inc32( + address, val, __ATOMIC_RELAXED, "workgroup"); } __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); + return __builtin_amdgcn_atomic_dec32( + address, val, __ATOMIC_RELAXED, "workgroup"); } __device__