diff --git a/catch/unit/atomics/CMakeLists.txt b/catch/unit/atomics/CMakeLists.txt index 5974902732..ab155fae4a 100644 --- a/catch/unit/atomics/CMakeLists.txt +++ b/catch/unit/atomics/CMakeLists.txt @@ -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 diff --git a/catch/unit/atomics/__hip_atomic_compare_exchange_strong.cc b/catch/unit/atomics/__hip_atomic_compare_exchange_strong.cc new file mode 100644 index 0000000000..69fd72ec51 --- /dev/null +++ b/catch/unit/atomics/__hip_atomic_compare_exchange_strong.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 + +/** + * @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(1, sizeof(TestType)); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + SingleDeviceSingleKernelTest(warp_size, sizeof(TestType)); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + SingleDeviceSingleKernelTest(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(1, sizeof(TestType)); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + SingleDeviceSingleKernelTest(warp_size, sizeof(TestType)); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + SingleDeviceSingleKernelTest(warp_size, cache_line_size); + } + } +} \ No newline at end of file diff --git a/catch/unit/atomics/arithmetic_common.hh b/catch/unit/atomics/arithmetic_common.hh index 11a5d771ff..d8f31e5a6e 100644 --- a/catch/unit/atomics/arithmetic_common.hh +++ b/catch/unit/atomics/arithmetic_common.hh @@ -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(mem, val); } } @@ -266,7 +269,8 @@ std::tuple, std::vector> 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)) {