From 5b28e386f6b5bcacc2ea4c9606e87abdba936891 Mon Sep 17 00:00:00 2001
From: kjayapra-amd <karthik.jayaprakash@amd.com>
Date: Fri, 22 Mar 2024 15:13:14 -0400
Subject: [PATCH] SWDEV-448288 - Remove references to deprecated llvm
 references.

Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72
---
 .../include/hip/amd_detail/amd_hip_atomic.h   | 22 ++-----------------
 1 file changed, 2 insertions(+), 20 deletions(-)

diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
index e35a79abdc..d6e4d81869 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_atomic.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
@@ -973,22 +973,13 @@ inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
 {
 #if defined(__gfx941__)
-  __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 hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
     address,
     val,
     [](unsigned int& x, unsigned int y) { x = (x >= y) ? 0 : (x + 1); },
     [=]() {
     return
-      __builtin_amdgcn_atomic_inc(address, val, __ATOMIC_RELAXED, 1, false);
+      __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
   });
 #else
     return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
@@ -1001,22 +992,13 @@ inline
 unsigned int atomicDec(unsigned int* address, unsigned int val)
 {
 #if defined(__gfx941__)
-  __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 hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
     address,
     val,
     [](unsigned int& x, unsigned int y) { x = (!x || x > y) ? y : (x - 1); },
     [=]() {
     return
-      __builtin_amdgcn_atomic_dec(address, val, __ATOMIC_RELAXED, 1, false);
+      __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
   });
 #else
   return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");