From ab6ba8c0a874fa73ca0a70d73c585fc43478800d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 17:13:17 +0100 Subject: [PATCH] EXSWHTEC-299 - Extend tests for atomic arithmetic operations #286 Change-Id: I221332c33be92ec152a2cd2ede34379aaa73d996 [ROCm/hip-tests commit: abf39d2dcda4f675465031f0ecb7eb4c9ad24821] --- .../catch/unit/atomics/CMakeLists.txt | 1 + .../unit/atomics/__hip_atomic_fetch_add.cc | 132 ++++++++++++++++++ .../catch/unit/atomics/arithmetic_common.hh | 99 ++++++++++--- 3 files changed, 209 insertions(+), 23 deletions(-) create mode 100644 projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_add.cc diff --git a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt index 1ec472bffc..5974902732 100644 --- a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt @@ -48,6 +48,7 @@ set(TEST_SRC atomicDec.cc atomicCAS.cc atomicCAS_system.cc + __hip_atomic_fetch_add.cc atomicExch.cc atomicExch_system.cc __hip_atomic_fetch_and.cc diff --git a/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_add.cc b/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_add.cc new file mode 100644 index 0000000000..075b2b858e --- /dev/null +++ b/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_add.cc @@ -0,0 +1,132 @@ +/* +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_fetch_add __hip_atomic_fetch_add + * @{ + * @ingroup AtomicsTest + * ________________________ + * Test cases from other modules: + * - @ref Unit_AtomicBuiltins_Negative_Parameters_RTC + */ + +/** + * 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_fetch_add + * - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated memory + * - Shared memory + * - WAVEFRONT memory scope. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_add.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_add_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_fetch_add + * - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated memory + * - Shared memory + * - WORKGROUP memory scope. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_add.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_add_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/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh index cc701a06a0..11a5d771ff 100644 --- a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh +++ b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh @@ -40,7 +40,8 @@ enum class AtomicOperation { kUnsafeAdd, kSafeAdd, kCASAdd, - kCASAddSystem + kCASAddSystem, + kBuiltinAdd }; // Constants that are passed as operands to the atomic operations @@ -88,6 +89,9 @@ __device__ TestType CASAtomicAddSystem(TestType* address, TestType val) { // Performs an atomic operation on parameter `mem` based on the `operation` enumerator. template +// 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 __device__ TestType PerformAtomicOperation(TestType* const mem) { const auto val = GetTestValue(); @@ -111,6 +115,8 @@ __device__ TestType PerformAtomicOperation(TestType* const mem) { return CASAtomicAdd(mem, val); } 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); } } @@ -119,7 +125,8 @@ __device__ TestType PerformAtomicOperation(TestType* const mem) { // same memory location `global_mem`. // If `use_shared_mem` is true, `global_mem` is copied to shared memory first, the atomic // operations are executed on shared memory, and the result is copied back to `global_mem`. -template +template __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals) { __shared__ TestType shared_mem; @@ -132,7 +139,7 @@ __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals) __syncthreads(); } - old_vals[tid] = PerformAtomicOperation(mem); + old_vals[tid] = PerformAtomicOperation(mem); if constexpr (use_shared_mem) { __syncthreads(); @@ -148,6 +155,15 @@ __host__ __device__ TestType* PitchedOffset(TestType* const ptr, const unsigned return reinterpret_cast(byte_ptr + idx * pitch); } +// Executes arbitrary load-store operations on the range specified by `begin_addr` and `end_addr` +__device__ void GenerateMemoryTraffic(uint8_t* const begin_addr, uint8_t* const end_addr) { + for (volatile uint8_t* addr = begin_addr; addr != end_addr; ++addr) { + uint8_t val = *addr; + val ^= 0xAB; + *addr = val; + } +} + // This kernel executes the atomic operation specified by the enumerator `operation`. Results of the // atomic operations are stored in `old_vals`. `global_mem` is an array with `width` number of // elements. Each thread performs the atomic operation on the element that corresponds to its thread @@ -157,6 +173,9 @@ __host__ __device__ TestType* PitchedOffset(TestType* const ptr, const unsigned // that are scattered over different cache lines. // If `use_shared_mem` is true, `global_mem` is copied to shared memory first, the atomic operations // are executed on shared memory, and the result is copied back to `global_mem`. +// If `pitch` is greater than sizeof(TestType), random memory operations are performed in the empty +// space between consecutive atomic operations so that we can test that the atomic operations +// behaves correctly even with some interference. // // For example, given that sizeof(TestType) is 1, `width` is 3, and `pitch` is 4: // @@ -165,10 +184,12 @@ __host__ __device__ TestType* PitchedOffset(TestType* const ptr, const unsigned // | pitch | pitch | pitch | // // In this scenario, the atomic operations will target the elements denoted with `x` (addresses 0, -// 4, 8). -template +// 4, 8). Random memory traffic will be generated on the addresses in between (1, 2, 3, 5, 6, 7, 9, +// 10, 11) +template __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals, - const unsigned int width, const unsigned pitch) { + const unsigned int width, const unsigned int pitch) { extern __shared__ uint8_t shared_mem[]; const auto tid = cg::this_grid().thread_rank(); @@ -183,8 +204,18 @@ __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals, __syncthreads(); } - old_vals[tid] = - PerformAtomicOperation(PitchedOffset(mem, pitch, tid % width)); + const auto n = cooperative_groups::this_grid().size() - width; + + TestType* atomic_addr = PitchedOffset(mem, pitch, tid % width); + + if (tid < n) { + old_vals[tid] = PerformAtomicOperation( + PitchedOffset(mem, pitch, tid % width)); + } else { + uint8_t* const begin_addr = reinterpret_cast(atomic_addr + 1); + uint8_t* const end_addr = reinterpret_cast(atomic_addr) + pitch; + GenerateMemoryTraffic(begin_addr, end_addr); + } if constexpr (use_shared_mem) { __syncthreads(); @@ -201,7 +232,7 @@ struct TestParams { return blocks.x * blocks.y * blocks.z * threads.x * threads.y * threads.z; } - auto HostIterationsPerThread() const { + auto HostIterationsPerThread() const { // number of iterations per host thread return std::max(num_devices * kernel_count * ThreadCount() / 20, width); } @@ -234,7 +265,8 @@ std::tuple, std::vector> TestKernelHostRef(const if constexpr (operation == AtomicOperation::kAdd || operation == AtomicOperation::kAddSystem || operation == AtomicOperation::kUnsafeAdd || operation == AtomicOperation::kSafeAdd || operation == AtomicOperation::kCASAdd || - operation == AtomicOperation::kCASAddSystem) { + operation == AtomicOperation::kCASAddSystem || + operation == AtomicOperation::kBuiltinAdd) { res = res + val; } else if constexpr (operation == AtomicOperation::kSub || operation == AtomicOperation::kSubSystem) { @@ -248,7 +280,7 @@ std::tuple, std::vector> TestKernelHostRef(const for (auto i = 0u; i < p.num_devices; ++i) { for (auto j = 0u; j < p.kernel_count; ++j) { - for (auto tid = 0u; tid < p.ThreadCount(); ++tid) { + for (auto tid = 0u; tid < p.ThreadCount() - p.width; ++tid) { perform_op(tid); } } @@ -283,15 +315,16 @@ void Verify(const TestParams& p, std::vector& res_vals, std::vector +template void LaunchKernel(const TestParams& p, hipStream_t stream, TestType* const mem_ptr, TestType* const old_vals) { const auto shared_mem_size = use_shared_mem ? p.width * p.pitch : 0u; if (p.width == 1 && p.pitch == sizeof(TestType)) - TestKernel + TestKernel <<>>(mem_ptr, old_vals); else - TestKernel + TestKernel <<>>(mem_ptr, old_vals, p.width, p.pitch); } @@ -303,7 +336,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::kCASAddSystem || + operation == AtomicOperation::kBuiltinAdd) { 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); @@ -339,7 +373,8 @@ void PerformHostAtomicOperation(const TestParams& p, TestType* mem, TestType* co // 2. Launch kernels based on TestParams::num_devices and TestParams::kernel_count // 3. Launch host threads based on TestParams::host_thread_count // 4. Verify the results -template +template void TestCore(const TestParams& p) { const unsigned int flags = p.alloc_type == LinearAllocs::mallocAndRegister ? hipHostRegisterMapped : 0u; @@ -371,7 +406,8 @@ void TestCore(const TestParams& p) { for (auto j = 0u; j < p.kernel_count; ++j) { const auto& stream = streams[i * p.kernel_count + j].stream(); const auto old_vals = old_vals_devs[i].ptr() + j * p.ThreadCount(); - LaunchKernel(p, stream, mem_dev.ptr(), old_vals); + LaunchKernel(p, stream, mem_dev.ptr(), + old_vals); } } @@ -397,23 +433,40 @@ inline dim3 GenerateBlockDimensions() { } // Configures and creates the TestCore for a single device, and a single kernel launch -template +template void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) { TestParams params; params.num_devices = 1; params.kernel_count = 1; - params.threads = GenerateThreadDimensions(); + if constexpr (operation == AtomicOperation::kBuiltinAdd && + memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD) { + params.threads = 1; + } else if constexpr (operation == AtomicOperation::kBuiltinAdd && + memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) { + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + params.threads = dim3(warp_size); + } else { + params.threads = GenerateThreadDimensions(); + } params.width = width; params.pitch = pitch; SECTION("Global memory") { - params.blocks = GenerateBlockDimensions(); + if constexpr (operation == AtomicOperation::kBuiltinAdd && + (memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD || + memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT || + memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP)) { + params.blocks = dim3(1); + } else { + params.blocks = GenerateBlockDimensions(); + } using LA = LinearAllocs; for (const auto alloc_type : {LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) { params.alloc_type = alloc_type; DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) { - TestCore(params); + TestCore(params); } } } @@ -421,7 +474,7 @@ void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int p SECTION("Shared memory") { params.blocks = dim3(1); params.alloc_type = LinearAllocs::hipMalloc; - TestCore(params); + TestCore(params); } } @@ -493,7 +546,7 @@ void MultipleDeviceMultipleKernelAndHostTest(const unsigned int num_devices, for (const auto alloc_type : {LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) { params.alloc_type = alloc_type; DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) { - TestCore(params); + TestCore(params); } } } \ No newline at end of file