From 5ea09794fc2c839faa9caf4b3594b2fe6b724334 Mon Sep 17 00:00:00 2001
From: Michael LIAO <michaelhongbin.liao@amd.com>
Date: Tue, 6 Jul 2021 16:08:12 -0400
Subject: [PATCH] SWDEV-293408 Add atomicInc_system and atomicDec_system

- Also revise the implementation of atomicInc/atomicDec based on clang
  builtins.

Change-Id: Ia5e06e88c7be1da5f8fc7ac70037a120c604e343


[ROCm/clr commit: 45205c5e5399795fe569cc158ba90348a9a3409a]
---
 .../include/hip/amd_detail/amd_hip_atomic.h   | 36 ++++++++-----------
 1 file changed, 15 insertions(+), 21 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 6a25533beb..bb2f4d4ba9 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
@@ -281,34 +281,28 @@ __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_dec32(address, val, __ATOMIC_RELAXED, "agent");
+}
 
-    return __builtin_amdgcn_atomic_dec(
-        address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
+__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");
 }
 
 __device__