From 73e8f26cf464f14eeb478de300775454d01020d1 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:11:02 +0100 Subject: [PATCH] EXSWHTEC-298 - Extend tests for atomic exchange operations #290 Change-Id: I19b4cb6177f3caa74c3f889de107e349ebb1d37c --- catch/unit/atomics/CMakeLists.txt | 1 + catch/unit/atomics/__hip_atomic_exchange.cc | 136 ++++++++++++++++++++ catch/unit/atomics/atomicExch_common.hh | 75 ++++++++--- 3 files changed, 194 insertions(+), 18 deletions(-) create mode 100644 catch/unit/atomics/__hip_atomic_exchange.cc diff --git a/catch/unit/atomics/CMakeLists.txt b/catch/unit/atomics/CMakeLists.txt index bb3b8db915..ed622195a9 100644 --- a/catch/unit/atomics/CMakeLists.txt +++ b/catch/unit/atomics/CMakeLists.txt @@ -40,6 +40,7 @@ set(TEST_SRC __hip_atomic_fetch_and.cc __hip_atomic_fetch_or.cc __hip_atomic_fetch_xor.cc + __hip_atomic_exchange.cc ) if(HIP_PLATFORM MATCHES "nvidia") diff --git a/catch/unit/atomics/__hip_atomic_exchange.cc b/catch/unit/atomics/__hip_atomic_exchange.cc new file mode 100644 index 0000000000..a518aaafbe --- /dev/null +++ b/catch/unit/atomics/__hip_atomic_exchange.cc @@ -0,0 +1,136 @@ +/* +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 "atomicExch_common.hh" + +/** + * @addtogroup __hip_atomic_exchange __hip_atomic_exchange + * @{ + * @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 + * exchange into a runtime determined memory location. Each thread will exchange its own grid wide + * linear index + offset into the memory location, storing the return value into a separate output + * array slot corresponding to it. Once complete, the union of output array and exchange memory is + * validated to contain all values in the range [0, number_of_threads + + * number_of_exchange_memory_slots). Several memory access patterns are tested: + * -# All threads exchange to a single memory location + * -# Each thread exchanges into an array containing warp_size elements, using tid % warp_size + * for indexing + * -# Same as the above, but the exchange elements are spread out by L1 cache line size bytes. + * + * - The test is run for: + * - All overloads of atomicExch + * - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory + * - Exchange memory located in shared memory + * - WAVEFRONT memory scope + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_exchange.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_exchange_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) { + AtomicExchSingleDeviceSingleKernelTest(1, sizeof(TestType)); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + AtomicExchSingleDeviceSingleKernelTest(warp_size, + sizeof(TestType)); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + AtomicExchSingleDeviceSingleKernelTest(warp_size, + cache_line_size); + } + } +} + +/** + * Test Description + * ------------------------ + * - Executes a single kernel on a single device wherein all threads will perform an atomic + * exchange into a runtime determined memory location. Each thread will exchange its own grid wide + * linear index + offset into the memory location, storing the return value into a separate output + * array slot corresponding to it. Once complete, the union of output array and exchange memory is + * validated to contain all values in the range [0, number_of_threads + + * number_of_exchange_memory_slots). Several memory access patterns are tested: + * -# All threads exchange to a single memory location + * -# Each thread exchanges into an array containing warp_size elements, using tid % warp_size + * for indexing + * -# Same as the above, but the exchange elements are spread out by L1 cache line size bytes. + * + * - The test is run for: + * - All overloads of atomicExch + * - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory + * - Exchange memory located in shared memory + * - WORKGROUP memory scope + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_exchange.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_exchange_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) { + AtomicExchSingleDeviceSingleKernelTest(1, sizeof(TestType)); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + AtomicExchSingleDeviceSingleKernelTest(warp_size, + sizeof(TestType)); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + AtomicExchSingleDeviceSingleKernelTest(warp_size, + cache_line_size); + } + } +} \ No newline at end of file diff --git a/catch/unit/atomics/atomicExch_common.hh b/catch/unit/atomics/atomicExch_common.hh index 1b4add5253..90e78edcc2 100644 --- a/catch/unit/atomics/atomicExch_common.hh +++ b/catch/unit/atomics/atomicExch_common.hh @@ -22,24 +22,26 @@ THE SOFTWARE. #pragma once -#include - +#include #include #include #include -#include -enum class AtomicScopes { device, system }; +enum class AtomicScopes { device, system, builtin }; -template __device__ T perform_atomic_exch(T* address, T val) { +template +__device__ T perform_atomic_exch(T* address, T val) { if constexpr (scope == AtomicScopes::device) { return atomicExch(address, val); } else if (scope == AtomicScopes::system) { return atomicExch_system(address, val); + } else if (scope == AtomicScopes::builtin) { + return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, memory_scope); } } -template +template __global__ void atomic_exch_kernel_compile_time(T* const global_mem, T* const old_vals) { __shared__ T shared_mem; @@ -52,7 +54,7 @@ __global__ void atomic_exch_kernel_compile_time(T* const global_mem, T* const ol __syncthreads(); } - old_vals[tid] = perform_atomic_exch(mem, static_cast(tid + 1)); + old_vals[tid] = perform_atomic_exch(mem, static_cast(tid + 1)); if constexpr (use_shared_mem) { __syncthreads(); @@ -67,7 +69,16 @@ __host__ __device__ T* pitched_offset(T* const ptr, const unsigned int pitch, return reinterpret_cast(byte_ptr + idx * pitch); } -template +__device__ void generate_memory_traffic(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; + } +} + +template __global__ void atomic_exch_kernel(T* const global_mem, T* const old_vals, const unsigned int width, const unsigned pitch, const T base_val = 0) { extern __shared__ uint8_t shared_mem[]; @@ -84,8 +95,18 @@ __global__ void atomic_exch_kernel(T* const global_mem, T* const old_vals, const __syncthreads(); } - old_vals[tid] = perform_atomic_exch(pitched_offset(mem, pitch, tid % width), - base_val + static_cast(tid + width)); + const auto n = cooperative_groups::this_grid().size() - width; + + T* atomic_addr = pitched_offset(mem, pitch, tid % width); + + if (tid < n) { + old_vals[tid] = perform_atomic_exch( + pitched_offset(mem, pitch, tid % width), base_val + static_cast(tid + width)); + } else { + uint8_t* const begin_addr = reinterpret_cast(atomic_addr + 1); + uint8_t* const end_addr = reinterpret_cast(atomic_addr) + pitch; + generate_memory_traffic(begin_addr, end_addr); + } if constexpr (use_shared_mem) { __syncthreads(); @@ -255,14 +276,16 @@ class AtomicExchCRTP { } }; -template +template class AtomicExch : public AtomicExchCRTP, T, use_shared_mem, scope> { public: void LaunchKernel(const unsigned int shared_mem_size, const hipStream_t stream, T* const mem, T* const old_vals, const T base_val, const AtomicExchParams& p) const { - atomic_exch_kernel<<>>( - mem, old_vals, p.width, p.pitch, base_val); + atomic_exch_kernel + <<>>(mem, old_vals, p.width, p.pitch, + base_val); } void ValidateResults(std::vector& old_vals) const { @@ -281,23 +304,39 @@ inline dim3 GenerateAtomicExchBlockDimensions() { return GENERATE_COPY(dim3(sm_count), dim3(sm_count + sm_count / 2)); } -template +template void AtomicExchSingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) { AtomicExchParams params; params.num_devices = 1; params.kernel_count = 1; - params.threads = GenerateAtomicExchThreadDimensions(); + if constexpr (scope == AtomicScopes::builtin && memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD) { + params.threads = 1; + } else if constexpr (scope == AtomicScopes::builtin && + 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 = GenerateAtomicExchThreadDimensions(); + } params.width = width; params.pitch = pitch; SECTION("Global memory") { - params.blocks = GenerateAtomicExchBlockDimensions(); + if constexpr (scope == AtomicScopes::builtin && + (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 = GenerateAtomicExchBlockDimensions(); + } 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)) { - AtomicExch().run(params); + AtomicExch().run(params); } } } @@ -305,7 +344,7 @@ void AtomicExchSingleDeviceSingleKernelTest(const unsigned int width, const unsi SECTION("Shared memory") { params.blocks = dim3(1); params.alloc_type = LinearAllocs::hipMalloc; - AtomicExch().run(params); + AtomicExch().run(params); } }