From 02a1e1141d29ffacb69de5bbbe20316883d5cd87 Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Thu, 17 Nov 2022 03:31:05 +0530
Subject: [PATCH] SWDEV-346579 - atomics test cases for non MI2xx/MI3xx (#3060)

Replace the internal __builtin__ functions with external APIs

Change-Id: Ia3cf37af004b5b9a4833b18a8771a8864772beb4


[ROCm/hip commit: 0acac9c7db45c8e3c107198eebfa83cb031715c4]
---
 .../catch/unit/deviceLib/BuiltIns_fadd.cc     |  4 +--
 .../catch/unit/deviceLib/BuiltIns_fmax.cc     | 24 +++--------------
 .../catch/unit/deviceLib/BuiltIns_fmin.cc     | 26 ++++---------------
 3 files changed, 11 insertions(+), 43 deletions(-)

diff --git a/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fadd.cc b/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fadd.cc
index 2eaaa94baa..0819f90b2c 100644
--- a/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fadd.cc
+++ b/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);
 }
 )"};
 /*
diff --git a/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmax.cc b/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmax.cc
index 5049348c61..9dcddc0596 100644
--- a/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmax.cc
+++ b/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);
 }
 )"};
 
diff --git a/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmin.cc b/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmin.cc
index 25e52eff29..d669bee9aa 100644
--- a/projects/hip/tests/catch/unit/deviceLib/BuiltIns_fmin.cc
+++ b/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);   
 }
 
 /*