From a785d554ae4c0f26f0bcfbc05fac31491a60761b Mon Sep 17 00:00:00 2001
From: Aaron En Ye Shi <enye.shi@gmail.com>
Date: Tue, 11 May 2021 18:04:39 +0000
Subject: [PATCH] SWDEV-285388 - Clean up llvm intrinsics using __asm

Instead of using inline asm, use clang builtins
for llvm intrinsics.

Change-Id: I30287f5a8de035ccd7e48d10e559a8a2e1d389f7


[ROCm/clr commit: 3fa21c1d49c78c36947e5c7d86a195ca5a2b028c]
---
 .../include/hip/amd_detail/amd_hip_atomic.h   | 52 +++----------------
 .../include/hip/amd_detail/llvm_intrinsics.h  | 15 ++++--
 .../hipamd/include/hip/amd_detail/math_fwd.h  | 30 ++++++++---
 3 files changed, 42 insertions(+), 55 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..72c177d3dd 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,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__
@@ -598,34 +580,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__
diff --git a/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h b/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h
index c100b846c0..8c626f7616 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h
@@ -22,8 +22,7 @@ THE SOFTWARE.
 
 /**
  *  @file  amd_detail/llvm_intrinsics.h
- *  @brief Contains declarations for wrapper functions for llvm intrinsics
- *         like llvm.amdgcn.s.barrier.
+ *  @brief Contains declarations for wrapper functions for llvm intrinsics.
  */
 
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_LLVM_INTRINSICS_H
@@ -31,11 +30,17 @@ THE SOFTWARE.
 
 #include "hip/amd_detail/host_defines.h"
 
-// FIXME: These should all be removed and proper builtins used.
 __device__
-unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize");
+inline
+unsigned __llvm_amdgcn_groupstaticsize() {
+  return __builtin_amdgcn_groupstaticsize();
+}
 
+template<typename __T>
 __device__
-int __llvm_amdgcn_ds_swizzle(int index, int pattern) __asm("llvm.amdgcn.ds.swizzle");
+inline
+__T __llvm_amdgcn_ds_swizzle(__T index, const int pattern) {
+  return __builtin_amdgcn_ds_swizzle(index, pattern);
+}
 
 #endif
diff --git a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h
index ac46d537a8..b299698c2f 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h
@@ -388,17 +388,29 @@ __attribute__((const))
 float __ocml_fma_rtz_f32(float, float, float);
 
 __device__
+inline
 __attribute__((const))
-float __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32");
+float __llvm_amdgcn_cos_f32(float __x) {
+  return __builtin_amdgcn_cosf(__x);
+}
 __device__
+inline
 __attribute__((const))
-float __llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32");
+float __llvm_amdgcn_rcp_f32(float __x) {
+  return __builtin_amdgcn_rcpf(__x);
+}
 __device__
+inline
 __attribute__((const))
-float __llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32");
+float __llvm_amdgcn_rsq_f32(float __x) {
+  return __builtin_amdgcn_rsqf(__x);
+}
 __device__
+inline
 __attribute__((const))
-float __llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32");
+float __llvm_amdgcn_sin_f32(float __x) {
+  return __builtin_amdgcn_sinf(__x);
+}
 // END INTRINSICS
 // END FLOAT
 
@@ -699,11 +711,17 @@ __attribute__((const))
 double __ocml_fma_rtz_f64(double, double, double);
 
 __device__
+inline
 __attribute__((const))
-double __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64");
+double __llvm_amdgcn_rcp_f64(double __x) {
+  return __builtin_amdgcn_rcp(__x);
+}
 __device__
+inline
 __attribute__((const))
-double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
+double __llvm_amdgcn_rsq_f64(double __x) {
+  return __builtin_amdgcn_rsq(__x);
+}
 // END INTRINSICS
 // END DOUBLE