EXSWHTEC-298 - Extend tests for atomic CAS operations #287
Change-Id: Ieb3e7effc1d3f767c77c0cd7c0b20c391c4665af
This commit is contained in:
Mirza Halilčević
کامیت شده توسط
Rakesh Roy
Rakesh Roy
والد
abf39d2dcd
کامیت
8f8e30e1c6
@@ -49,6 +49,7 @@ set(TEST_SRC
|
||||
atomicCAS.cc
|
||||
atomicCAS_system.cc
|
||||
__hip_atomic_fetch_add.cc
|
||||
__hip_atomic_compare_exchange_strong.cc
|
||||
atomicExch.cc
|
||||
atomicExch_system.cc
|
||||
__hip_atomic_fetch_and.cc
|
||||
|
||||
@@ -0,0 +1,129 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "arithmetic_common.hh"
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup __hip_atomic_compare_exchange_strong __hip_atomic_compare_exchange_strong
|
||||
* @{
|
||||
* @ingroup AtomicsTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes a single kernel on a single device wherein all threads will perform an atomic
|
||||
* addition on a target memory location. Each thread will add the same value to the memory location,
|
||||
* storing the return value into a separate output array slot corresponding to it. Once complete,
|
||||
* the output array and target memory is validated to contain all the expected values. Several
|
||||
* memory access patterns are tested:
|
||||
* -# All threads add to a single, compile time deducible, memory location
|
||||
* -# Each thread targets an array containing warp_size elements, using tid % warp_size
|
||||
* for indexing
|
||||
* -# Same as the above, but the elements are spread out by L1 cache line size bytes.
|
||||
*
|
||||
* - The test is run for:
|
||||
* - All overloads of __hip_atomic_compare_exchange_strong
|
||||
* - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated memory
|
||||
* - Shared memory
|
||||
* - WAVEFRONT memory scope.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/atomics/__hip_atomic_compare_exchange_strong.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit___hip_atomic_compare_exchange_strong_Positive_Wavefront", "", int,
|
||||
unsigned int, unsigned long, unsigned long long, float, double) {
|
||||
int warp_size = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
|
||||
const auto cache_line_size = 128u;
|
||||
|
||||
for (auto current = 0; current < cmd_options.iterations; ++current) {
|
||||
DYNAMIC_SECTION("Same address " << current) {
|
||||
SingleDeviceSingleKernelTest<TestType, AtomicOperation::kBuiltinCAS,
|
||||
__HIP_MEMORY_SCOPE_WAVEFRONT>(1, sizeof(TestType));
|
||||
}
|
||||
|
||||
DYNAMIC_SECTION("Adjacent addresses " << current) {
|
||||
SingleDeviceSingleKernelTest<TestType, AtomicOperation::kBuiltinCAS,
|
||||
__HIP_MEMORY_SCOPE_WAVEFRONT>(warp_size, sizeof(TestType));
|
||||
}
|
||||
|
||||
DYNAMIC_SECTION("Scattered addresses " << current) {
|
||||
SingleDeviceSingleKernelTest<TestType, AtomicOperation::kBuiltinCAS,
|
||||
__HIP_MEMORY_SCOPE_WAVEFRONT>(warp_size, cache_line_size);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes a single kernel on a single device wherein all threads will perform an atomic
|
||||
* addition on a target memory location. Each thread will add the same value to the memory location,
|
||||
* storing the return value into a separate output array slot corresponding to it. Once complete,
|
||||
* the output array and target memory is validated to contain all the expected values. Several
|
||||
* memory access patterns are tested:
|
||||
* -# All threads add to a single, compile time deducible, memory location
|
||||
* -# Each thread targets an array containing warp_size elements, using tid % warp_size
|
||||
* for indexing
|
||||
* -# Same as the above, but the elements are spread out by L1 cache line size bytes.
|
||||
*
|
||||
* - The test is run for:
|
||||
* - All overloads of __hip_atomic_compare_exchange_strong
|
||||
* - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated memory
|
||||
* - Shared memory
|
||||
* - WORKGROUP memory scope.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/atomics/__hip_atomic_compare_exchange_strong.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit___hip_atomic_compare_exchange_strong_Positive_Workgroup", "", int,
|
||||
unsigned int, unsigned long, unsigned long long, float, double) {
|
||||
int warp_size = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
|
||||
const auto cache_line_size = 128u;
|
||||
|
||||
for (auto current = 0; current < cmd_options.iterations; ++current) {
|
||||
DYNAMIC_SECTION("Same address " << current) {
|
||||
SingleDeviceSingleKernelTest<TestType, AtomicOperation::kBuiltinCAS,
|
||||
__HIP_MEMORY_SCOPE_WORKGROUP>(1, sizeof(TestType));
|
||||
}
|
||||
|
||||
DYNAMIC_SECTION("Adjacent addresses " << current) {
|
||||
SingleDeviceSingleKernelTest<TestType, AtomicOperation::kBuiltinCAS,
|
||||
__HIP_MEMORY_SCOPE_WORKGROUP>(warp_size, sizeof(TestType));
|
||||
}
|
||||
|
||||
DYNAMIC_SECTION("Scattered addresses " << current) {
|
||||
SingleDeviceSingleKernelTest<TestType, AtomicOperation::kBuiltinCAS,
|
||||
__HIP_MEMORY_SCOPE_WORKGROUP>(warp_size, cache_line_size);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -41,7 +41,8 @@ enum class AtomicOperation {
|
||||
kSafeAdd,
|
||||
kCASAdd,
|
||||
kCASAddSystem,
|
||||
kBuiltinAdd
|
||||
kBuiltinAdd,
|
||||
kBuiltinCAS
|
||||
};
|
||||
|
||||
// Constants that are passed as operands to the atomic operations
|
||||
@@ -117,6 +118,8 @@ __device__ TestType PerformAtomicOperation(TestType* const mem) {
|
||||
return CASAtomicAddSystem(mem, val);
|
||||
} else if constexpr (operation == AtomicOperation::kBuiltinAdd) {
|
||||
return __hip_atomic_fetch_add(mem, val, __ATOMIC_RELAXED, memory_scope);
|
||||
} else if constexpr (operation == AtomicOperation::kBuiltinCAS) {
|
||||
return BuiltinCASAtomicAdd<TestType, memory_scope>(mem, val);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -266,7 +269,8 @@ std::tuple<std::vector<TestType>, std::vector<TestType>> TestKernelHostRef(const
|
||||
operation == AtomicOperation::kUnsafeAdd ||
|
||||
operation == AtomicOperation::kSafeAdd || operation == AtomicOperation::kCASAdd ||
|
||||
operation == AtomicOperation::kCASAddSystem ||
|
||||
operation == AtomicOperation::kBuiltinAdd) {
|
||||
operation == AtomicOperation::kBuiltinAdd ||
|
||||
operation == AtomicOperation::kBuiltinCAS) {
|
||||
res = res + val;
|
||||
} else if constexpr (operation == AtomicOperation::kSub ||
|
||||
operation == AtomicOperation::kSubSystem) {
|
||||
@@ -337,7 +341,8 @@ void HostAtomicOperation(const unsigned int iterations, TestType* mem, TestType*
|
||||
for (auto i = 0u; i < iterations; ++i) {
|
||||
if constexpr (operation == AtomicOperation::kAddSystem ||
|
||||
operation == AtomicOperation::kCASAddSystem ||
|
||||
operation == AtomicOperation::kBuiltinAdd) {
|
||||
operation == AtomicOperation::kBuiltinAdd ||
|
||||
operation == AtomicOperation::kBuiltinCAS) {
|
||||
old_vals[i] = __atomic_fetch_add(PitchedOffset(mem, pitch, i % width), val, __ATOMIC_RELAXED);
|
||||
} else if constexpr (operation == AtomicOperation::kSubSystem) {
|
||||
old_vals[i] = __atomic_fetch_sub(PitchedOffset(mem, pitch, i % width), val, __ATOMIC_RELAXED);
|
||||
@@ -438,10 +443,12 @@ void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int p
|
||||
TestParams params;
|
||||
params.num_devices = 1;
|
||||
params.kernel_count = 1;
|
||||
if constexpr (operation == AtomicOperation::kBuiltinAdd &&
|
||||
if constexpr ((operation == AtomicOperation::kBuiltinAdd ||
|
||||
operation == AtomicOperation::kBuiltinCAS) &&
|
||||
memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD) {
|
||||
params.threads = 1;
|
||||
} else if constexpr (operation == AtomicOperation::kBuiltinAdd &&
|
||||
} else if constexpr ((operation == AtomicOperation::kBuiltinAdd ||
|
||||
operation == AtomicOperation::kBuiltinCAS) &&
|
||||
memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) {
|
||||
int warp_size = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
|
||||
@@ -453,7 +460,8 @@ void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int p
|
||||
params.pitch = pitch;
|
||||
|
||||
SECTION("Global memory") {
|
||||
if constexpr (operation == AtomicOperation::kBuiltinAdd &&
|
||||
if constexpr ((operation == AtomicOperation::kBuiltinAdd ||
|
||||
operation == AtomicOperation::kBuiltinCAS) &&
|
||||
(memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD ||
|
||||
memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT ||
|
||||
memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP)) {
|
||||
|
||||
مرجع در شماره جدید
Block a user