From 4b41394fe649eaef245c465d8e0b2c1dc9d15226 Mon Sep 17 00:00:00 2001
From: "Jiang, Julia" <Julia.Jiang@amd.com>
Date: Fri, 18 Jul 2025 21:19:03 -0400
Subject: [PATCH] SWDEV-523483 - Updated hip-tests for atomics add-operation
 (#448)

[ROCm/hip-tests commit: a7344b87c219925059a2b6d4e54a1c7849136f1b]
---
 .../catch/unit/atomics/arithmetic_common.hh   | 31 ++++++++++++++-----
 1 file changed, 23 insertions(+), 8 deletions(-)

diff --git a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
index add775a279..12ae1a431b 100644
--- a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
+++ b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
@@ -29,6 +29,12 @@ THE SOFTWARE.
 
 namespace cg = cooperative_groups;
 
+#if defined(__has_extension) && __has_extension(clang_atomic_attributes)
+#define HIP_TEST_FINE_GRAINED_MEMORY [[clang::atomic(fine_grained_memory)]]
+#else
+#define HIP_TEST_FINE_GRAINED_MEMORY
+#endif
+
 // Atomic operations for which the tests in this file apply for
 enum class AtomicOperation {
   kAdd = 0,
@@ -110,7 +116,7 @@ __device__ TestType BuiltinCASAtomicAdd(TestType* address, TestType val) {
 // Performs an atomic operation on parameter `mem` based on the `operation` enumerator.
 // `memory_scope` is forwarded to the builtin operations and is by default device-wide.
 template <typename TestType, AtomicOperation operation, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
-__device__ TestType PerformAtomicOperation(TestType* const mem) {
+__device__ TestType PerformAtomicOperation(TestType* const mem, const LinearAllocs allocType) {
   const auto val = GetTestValue<TestType, operation>();
 
   if constexpr (operation == AtomicOperation::kAdd) {
@@ -134,7 +140,13 @@ __device__ TestType PerformAtomicOperation(TestType* const mem) {
   } else if constexpr (operation == AtomicOperation::kCASAddSystem) {
     return CASAtomicAddSystem(mem, val);
   } else if constexpr (operation == AtomicOperation::kBuiltinAdd) {
-    return __hip_atomic_fetch_add(mem, val, __ATOMIC_RELAXED, memory_scope);
+    if (std::is_floating_point_v<TestType> && allocType == LinearAllocs::hipHostMalloc) {
+      HIP_TEST_FINE_GRAINED_MEMORY {
+        return __hip_atomic_fetch_add(mem, val, __ATOMIC_RELAXED, memory_scope);
+      }
+    } else {
+      return __hip_atomic_fetch_add(mem, val, __ATOMIC_RELAXED, memory_scope);
+    }
   } else if constexpr (operation == AtomicOperation::kBuiltinCAS) {
     return BuiltinCASAtomicAdd<TestType, memory_scope>(mem, val);
   }
@@ -147,7 +159,8 @@ __device__ TestType PerformAtomicOperation(TestType* const mem) {
 // operations are executed on shared memory, and the result is copied back to `global_mem`.
 template <typename TestType, AtomicOperation operation, bool use_shared_mem,
           int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
-__global__ void TestKernel(TestType* const global_mem, TestType* const old_vals) {
+__global__ void TestKernel(TestType* const global_mem, TestType* const old_vals,
+                           const LinearAllocs allocType) {
   __shared__ TestType shared_mem;
 
   const auto tid = cg::this_grid().thread_rank();
@@ -159,7 +172,7 @@ __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals)
     __syncthreads();
   }
 
-  old_vals[tid]  = PerformAtomicOperation<TestType, operation, memory_scope>(mem);
+  old_vals[tid]  = PerformAtomicOperation<TestType, operation, memory_scope>(mem, allocType);
 
 
   if constexpr (use_shared_mem) {
@@ -210,7 +223,8 @@ __device__ void GenerateMemoryTraffic(uint8_t* const begin_addr, uint8_t* const
 template <typename TestType, AtomicOperation operation, bool use_shared_mem,
           int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
 __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals,
-                           const unsigned int width, const unsigned int pitch) {
+                           const unsigned int width, const unsigned pitch,
+                           const LinearAllocs allocType) {
   extern __shared__ uint8_t shared_mem[];
 
   const auto tid = cg::this_grid().thread_rank();
@@ -231,7 +245,7 @@ __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals,
 
   if (tid < n) {
     old_vals[tid] = PerformAtomicOperation<TestType, operation, memory_scope>(
-        PitchedOffset(mem, pitch, tid % width));
+        PitchedOffset(mem, pitch, tid % width), allocType);
   } else {
     uint8_t* const begin_addr = reinterpret_cast<uint8_t*>(atomic_addr + 1);
     uint8_t* const end_addr = reinterpret_cast<uint8_t*>(atomic_addr) + pitch;
@@ -344,10 +358,11 @@ void LaunchKernel(const TestParams& p, hipStream_t stream, TestType* const mem_p
   const auto shared_mem_size = use_shared_mem ? p.width * p.pitch : 0u;
   if (p.width == 1 && p.pitch == sizeof(TestType))
     TestKernel<TestType, operation, use_shared_mem, memory_scope>
-        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem_ptr, old_vals);
+        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem_ptr, old_vals, p.alloc_type);
   else
     TestKernel<TestType, operation, use_shared_mem, memory_scope>
-        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem_ptr, old_vals, p.width, p.pitch);
+        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem_ptr, old_vals, p.width, p.pitch,
+            p.alloc_type);
 }
 
 // Performs a host atomic operation on parameter `mem` based on the `operation` enumerator.