SWDEV-285388 - Bring back the inline asm till we have a proper compiler support

Change-Id: Ibafe33dd8298f6e19f8b6a6d2aa220b633b76862


[ROCm/clr commit: f4680d98a9]

projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h

@@ -281,16 +281,34 @@ __device__
 inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
 {
-    return __builtin_amdgcn_atomic_inc32(
-        address, val, __ATOMIC_RELAXED, "agent");
+    __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);
 }
 
 __device__
 inline
 unsigned int atomicDec(unsigned int* address, unsigned int val)
 {
-    return __builtin_amdgcn_atomic_dec32(
-        address, val, __ATOMIC_RELAXED, "agent");
+    __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);
 }
 
 __device__
@@ -580,16 +598,34 @@ __device__
 inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
 {
-   return __builtin_amdgcn_atomic_inc32(
-        address, val, __ATOMIC_RELAXED, "agent");
+    __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);
 }
 
 __device__
 inline
 unsigned int atomicDec(unsigned int* address, unsigned int val)
 {
-    return __builtin_amdgcn_atomic_dec32(
-        address, val, __ATOMIC_RELAXED, "agent");
+    __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);
 }
 
 __device__

projects/clr/hipamd/include/hip/amd_detail/llvm_intrinsics.h

@@ -22,7 +22,8 @@ THE SOFTWARE.
 
 /**
  *  @file  amd_detail/llvm_intrinsics.h
- *  @brief Contains declarations for wrapper functions for llvm intrinsics.
+ *  @brief Contains declarations for wrapper functions for llvm intrinsics
+ *         like llvm.amdgcn.s.barrier.
  */
 
 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_LLVM_INTRINSICS_H
@@ -30,23 +31,11 @@ THE SOFTWARE.
 
 #include "hip/amd_detail/host_defines.h"
 
-#if __has_builtin(__builtin_amdgcn_groupstaticsize)
-__device__
-inline
-unsigned __llvm_amdgcn_groupstaticsize() {
-  return __builtin_amdgcn_groupstaticsize();
-}
-#else
-// FIXME: This should be removed once the above builtin is available.
+// FIXME: These should all be removed and proper builtins used.
 __device__
 unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize");
-#endif
 
-template<typename __T>
 __device__
-inline
-__T __llvm_amdgcn_ds_swizzle(__T index, const int pattern) {
-  return __builtin_amdgcn_ds_swizzle(index, pattern);
-}
+int __llvm_amdgcn_ds_swizzle(int index, int pattern) __asm("llvm.amdgcn.ds.swizzle");
 
 #endif

projects/clr/hipamd/include/hip/amd_detail/math_fwd.h

@@ -388,29 +388,17 @@ __attribute__((const))
 float __ocml_fma_rtz_f32(float, float, float);
 
 __device__
-inline
 __attribute__((const))
-float __llvm_amdgcn_cos_f32(float __x) {
-  return __builtin_amdgcn_cosf(__x);
-}
+float __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32");
 __device__
-inline
 __attribute__((const))
-float __llvm_amdgcn_rcp_f32(float __x) {
-  return __builtin_amdgcn_rcpf(__x);
-}
+float __llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32");
 __device__
-inline
 __attribute__((const))
-float __llvm_amdgcn_rsq_f32(float __x) {
-  return __builtin_amdgcn_rsqf(__x);
-}
+float __llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32");
 __device__
-inline
 __attribute__((const))
-float __llvm_amdgcn_sin_f32(float __x) {
-  return __builtin_amdgcn_sinf(__x);
-}
+float __llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32");
 // END INTRINSICS
 // END FLOAT
 
@@ -711,17 +699,11 @@ __attribute__((const))
 double __ocml_fma_rtz_f64(double, double, double);
 
 __device__
-inline
 __attribute__((const))
-double __llvm_amdgcn_rcp_f64(double __x) {
-  return __builtin_amdgcn_rcp(__x);
-}
+double __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64");
 __device__
-inline
 __attribute__((const))
-double __llvm_amdgcn_rsq_f64(double __x) {
-  return __builtin_amdgcn_rsq(__x);
-}
+double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
 // END INTRINSICS
 // END DOUBLE