SWDEV-364628 - 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

Change-Id: I45cc8a5ed62b73643dd3b74024c8ec835c632659

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, "agent");
 }
 
 __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, "agent");
 }
 
 __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, "agent");
 }
 
 __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, "agent");
 }
 
 __device__