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: dcb5d29f1e]

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__