SWDEV-399522 - Fix the issue #2 raised on github clr repo. Adds signe…

…d long long atomics Change-Id: I09c386d45e219e472d2da1449f678b9d17fe3e2e
ROCm · Jun 16, 2023 · be02e1e · be02e1e
1 parent a6e72de
commit be02e1e
Showing 1 changed file with 84 additions and 0 deletions.
diff --git a/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/hipamd/include/hip/amd_detail/amd_hip_atomic.h
@@ -576,6 +576,34 @@ unsigned long long atomicMin_system(unsigned long long* address, unsigned long l
 #endif // __gfx941__
 }
 
+__device__
+inline
+long long atomicMin(long long* address, long long val) {
+#if defined(__gfx941__)
+  return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+      address, val, [](long long x, long long y) { return x < y; },
+      [=]() {
+        return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+      });
+#else
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif  // __gfx941__
+}
+
+__device__
+inline
+long long atomicMin_system(long long* address, long long val) {
+#if defined(__gfx941__)
+  return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+      address, val, [](long long x, long long y) { return x < y; },
+      [=]() {
+        return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+      });
+#else
+  return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif  // __gfx941__
+}
+
 __device__
 inline
 float atomicMin(float* addr, float val) {
@@ -792,6 +820,34 @@ unsigned long long atomicMax_system(unsigned long long* address, unsigned long l
 #endif // __gfx941__
 }
 
+__device__
+inline
+long long atomicMax(long long* address, long long val) {
+  #if defined(__gfx941__)
+  return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+      address, val, [](long long x, long long y) { return y < x; },
+      [=]() {
+        return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+      });
+#else
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+long long atomicMax_system(long long* address, long long val) {
+#if defined(__gfx941__)
+  return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+      address, val, [](long long x, long long y) { return y < x; },
+      [=]() {
+        return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+      });
+#else
+  return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif  // __gfx941__
+}
+
 __device__
 inline
 float atomicMax(float* addr, float val) {
@@ -1437,6 +1493,20 @@ unsigned long long atomicMin(
 
     return tmp;
 }
+__device__ inline long long atomicMin(long long* address, long long val) {
+    long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+    while (val < tmp) {
+        const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
+
+        if (tmp1 != tmp) {
+          tmp = tmp1;
+          continue;
+        }
+
+        tmp = atomicCAS(address, tmp, val);
+    }
+    return tmp;
+}
 
 __device__
 inline
@@ -1466,6 +1536,20 @@ unsigned long long atomicMax(
 
     return tmp;
 }
+__device__ inline long long atomicMax(long long* address, long long val) {
+    long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+    while (tmp < val) {
+        const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
+
+        if (tmp1 != tmp) {
+          tmp = tmp1;
+          continue;
+        }
+
+        tmp = atomicCAS(address, tmp, val);
+    }
+    return tmp;
+}
 
 __device__
 inline