SWDEV-346579 - atomics test cases for non MI2xx/MI3xx (#3060)

Replace the internal __builtin__ functions with external APIs

Change-Id: Ia3cf37af004b5b9a4833b18a8771a8864772beb4


[ROCm/hip commit: 0acac9c7db]

projects/hip/tests/catch/unit/deviceLib/BuiltIns_fadd.cc

@@ -34,14 +34,14 @@ This testfile verifies __builtin_amdgcn_global_atomic_fadd_f64 API scenarios
 #define INITIAL_VAL 5
 __global__ void AtomicAdd_GlobalMem(double* addr, double* result) {
   double inc_val = 10;
-  *result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val);
+  *result = unsafeAtomicAdd(addr, inc_val);
 }
 static constexpr auto AtomicAddGlobalMem{
 R"(
 extern "C"
 __global__ void AtomicAdd_GlobalMem(double* addr, double* result) {
   double inc_val = 10;
-  *result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val);
+  *result = unsafeAtomicAdd(addr, inc_val);
 }
 )"};
 /*

projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmax.cc

@@ -39,20 +39,12 @@ __global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) {
   __shared__ double int_val;
   int_val = 5;
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-        (const __attribute__((address_space(0))) void*)(&int_val)))
-    *result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp);
-  else
-    *result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp);
+  *result = unsafeAtomicMax(&int_val, comp);
   *addr = int_val;
 }
 __global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) {
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-        (const __attribute__((address_space(0))) void*)(addr)))
-    *result =  __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp);
-  else
-    *result =  __builtin_amdgcn_global_atomic_fmax_f64(addr, comp);
+  *result = unsafeAtomicMax(addr, comp);
 }
 static constexpr auto fmaxFlatMem {
 R"(
@@ -61,11 +53,7 @@ __global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) {
   __shared__ double int_val;
   int_val = 5;
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-        (const __attribute__((address_space(0))) void*)(&int_val)))
-    *result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp);
-  else
-    *result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp);
+  *result = unsafeAtomicMax(&int_val, comp);
   *addr = int_val;
 }
 )"};
@@ -75,11 +63,7 @@ R"(
 extern "C"
 __global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) {
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-        (const __attribute__((address_space(0))) void*)(addr)))
-    *result =  __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp);
-  else
-    *result =  __builtin_amdgcn_global_atomic_fmax_f64(addr, comp);
+  *result = unsafeAtomicMax(addr, comp);
 }
 )"};
 

projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmin.cc

@@ -41,12 +41,8 @@ __global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) {
   __shared__ double int_val;
   int_val = 5;
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-          (const __attribute__((address_space(0))) void*)(&int_val)))
-     *result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp);
-  else
-     *result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp);
-     *addr = int_val;
+  *result = unsafeAtomicMin(&int_val, comp);
+  *addr = int_val;
 }
 )"};
 
@@ -55,11 +51,7 @@ R"(
 extern "C"
 __global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) {
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-          (const __attribute__((address_space(0))) void*)(addr)))
-     *result =  __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp);
-  else
-     *result =  __builtin_amdgcn_global_atomic_fmin_f64(addr, comp);
+  *result = unsafeAtomicMin(addr, comp);   
 }
 )"};
 
@@ -67,20 +59,12 @@ __global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) {
   __shared__ double int_val;
   int_val = 5;
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-          (const __attribute__((address_space(0))) void*)(&int_val)))
-     *result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp);
-  else
-     *result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp);
+  *result = unsafeAtomicMin(&int_val, comp);
   *addr = int_val;
 }
 __global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) {
   double comp = 10;
-  if (__builtin_amdgcn_is_shared(
-          (const __attribute__((address_space(0))) void*)(addr)))
-     *result =  __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp);
-  else
-     *result =  __builtin_amdgcn_global_atomic_fmin_f64(addr, comp);
+  *result = unsafeAtomicMin(addr, comp);   
 }
 
 /*