From 6d160e5a695de99be2951542b2e641ab18f5e5ba 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 [ROCm/clr commit: dcb5d29f1efabd4eb16b7aca1240907aa29e2246] --- .../include/hip/amd_detail/amd_hip_atomic.h | 52 +++---------------- 1 file changed, 8 insertions(+), 44 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h index 8cd4265647..0f516364f0 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/projects/clr/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__