From 68289d1dfe4735b98fca45fde7b438effbc57210 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 9 Jul 2021 13:47:08 -0400 Subject: [PATCH] Revert "SWDEV-293408 Add atomicInc_system and atomicDec_system" This reverts commit 45205c5e5399795fe569cc158ba90348a9a3409a. Change-Id: I808b52758cc08882daa7fc889e53ca4c2fe64a61 --- .../include/hip/amd_detail/amd_hip_atomic.h | 36 +++++++++++-------- 1 file changed, 21 insertions(+), 15 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h index bb2f4d4ba9..6a25533beb 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -281,28 +281,34 @@ __device__ inline unsigned int atomicInc(unsigned int* address, unsigned int val) { - return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent"); + __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); } __device__ inline unsigned int atomicDec(unsigned int* address, unsigned int val) { - return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent"); -} + __device__ + extern + unsigned int __builtin_amdgcn_atomic_dec( + unsigned int*, + unsigned int, + unsigned int, + unsigned int, + bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32"); -__device__ -inline -unsigned int atomicInc_system(unsigned int* address, unsigned int val) -{ - return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "one-as"); -} - -__device__ -inline -unsigned int atomicDec_system(unsigned int* address, unsigned int val) -{ - return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "one-as"); + return __builtin_amdgcn_atomic_dec( + address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false); } __device__