From 2b5806542f04e7aa94004a3dfffa9a007682c513 Mon Sep 17 00:00:00 2001
From: kjayapra-amd <karthik.jayaprakash@amd.com>
Date: Wed, 23 Jun 2021 16:26:49 -0500
Subject: [PATCH] SWDEV-278500 - Implement missing atomic APIs and types to
 match CUDA.

Change-Id: I5f5bd94af9b68853bf354c8f90da217434934ddd


[ROCm/clr commit: 4aab5c29ff5dee39678cb84b157deaa3f13961a9]
---
 .../include/hip/amd_detail/amd_hip_atomic.h   | 241 ++++++++++++++++++
 1 file changed, 241 insertions(+)

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..bba9988035 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
@@ -45,6 +45,22 @@ unsigned int atomicCAS_system(unsigned int* address, unsigned int compare, unsig
   return compare;
 }
 
+__device__
+inline
+unsigned long atomicCAS(unsigned long* address, unsigned long compare, unsigned long val) {
+  __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+                                       __HIP_MEMORY_SCOPE_AGENT);
+  return compare;
+}
+
+__device__
+inline
+unsigned long atomicCAS_system(unsigned long* address, unsigned long compare, unsigned long val) {
+  __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+                                       __HIP_MEMORY_SCOPE_SYSTEM);
+  return compare;
+}
+
 __device__
 inline
 unsigned long long atomicCAS(unsigned long long* address, unsigned long long compare,
@@ -63,6 +79,38 @@ unsigned long long atomicCAS_system(unsigned long long* address, unsigned long l
   return compare;
 }
 
+__device__
+inline
+float atomicCAS(float* address, float compare, float val) {
+  __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+                                       __HIP_MEMORY_SCOPE_AGENT);
+    return compare;
+}
+
+__device__
+inline
+float atomicCAS_system(float* address, float compare, float val) {
+  __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+                                       __HIP_MEMORY_SCOPE_SYSTEM);
+    return compare;
+}
+
+__device__
+inline
+double atomicCAS(double* address, double compare, double val) {
+  __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+                                       __HIP_MEMORY_SCOPE_AGENT);
+    return compare;
+}
+
+__device__
+inline
+double atomicCAS_system(double* address, double compare, double val) {
+  __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+                                       __HIP_MEMORY_SCOPE_SYSTEM);
+    return compare;
+}
+
 __device__
 inline
 int atomicAdd(int* address, int val) {
@@ -87,6 +135,18 @@ unsigned int atomicAdd_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long atomicAdd(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicAdd_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicAdd(unsigned long long* address, unsigned long long val) {
@@ -157,6 +217,54 @@ unsigned int atomicSub_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long atomicSub(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicSub_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long long atomicSub(unsigned long long* address, unsigned long long val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long long atomicSub_system(unsigned long long* address, unsigned long long val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+float atomicSub(float* address, float val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+float atomicSub_system(float* address, float val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+double atomicSub(double* address, double val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+double atomicSub_system(double* address, double val) {
+  return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 int atomicExch(int* address, int val) {
@@ -181,6 +289,18 @@ unsigned int atomicExch_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long atomicExch(unsigned long* address, unsigned long val) {
+  return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicExch_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicExch(unsigned long long* address, unsigned long long val) {
@@ -205,6 +325,18 @@ float atomicExch_system(float* address, float val) {
   return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+double atomicExch(double* address, double val) {
+  return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+double atomicExch_system(double* address, double val) {
+  return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 int atomicMin(int* address, int val) {
@@ -229,6 +361,18 @@ unsigned int atomicMin_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long long atomicMin(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicMin_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicMin(unsigned long long* address, unsigned long long val) {
@@ -241,6 +385,30 @@ unsigned long long atomicMin_system(unsigned long long* address, unsigned long l
   return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+float atomicMin(float* address, float val) {
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+float atomicMin_system(float* address, float val) {
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+double atomicMin(double* address, double val) {
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+double atomicMin_system(double* address, double val) {
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 int atomicMax(int* address, int val) {
@@ -265,6 +433,18 @@ unsigned int atomicMax_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long atomicMax(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicMax_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicMax(unsigned long long* address, unsigned long long val) {
@@ -277,6 +457,30 @@ unsigned long long atomicMax_system(unsigned long long* address, unsigned long l
   return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+float atomicMax(float* address, float val) {
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+float atomicMax_system(float* address, float val) {
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+double atomicMax(double* address, double val) {
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+double atomicMax_system(double* address, double val) {
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
@@ -334,6 +538,19 @@ inline
 unsigned int atomicAnd_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
+
+__device__
+inline
+unsigned long atomicAnd(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicAnd_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicAnd(unsigned long long* address, unsigned long long val) {
@@ -370,6 +587,18 @@ unsigned int atomicOr_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long atomicOr(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicOr_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicOr(unsigned long long* address, unsigned long long val) {
@@ -406,6 +635,18 @@ unsigned int atomicXor_system(unsigned int* address, unsigned int val) {
   return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
 }
 
+__device__
+inline
+unsigned long atomicXor(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicXor_system(unsigned long* address, unsigned long val) {
+  return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
 __device__
 inline
 unsigned long long atomicXor(unsigned long long* address, unsigned long long val) {