AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-523483 - Updated hip-tests for atomics add-operation (#448)

Browse Source 
[ROCm/hip-tests commit: a7344b87c2]
This commit is contained in:
Jiang, Julia
2025-07-18 21:19:03 -04:00
committed by GitHub
parent b8ef4025e2
commit 4b41394fe6
1 changed files with 23 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
+23 -8
View File
@@ -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.