From 8ee015d1ef00869536b2c14b7f312912753273f1 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 15:51:36 +0100 Subject: [PATCH] EXSWHTEC-262 - Introduce common utilities for atomic arithmetic operations #392 Change-Id: I016e571d15b6b6c3a1b91a4bb4b5410e338efe49 --- catch/unit/atomics/arithmetic_common.hh | 466 ++++++++++++++++++++++++ 1 file changed, 466 insertions(+) create mode 100644 catch/unit/atomics/arithmetic_common.hh diff --git a/catch/unit/atomics/arithmetic_common.hh b/catch/unit/atomics/arithmetic_common.hh new file mode 100644 index 0000000000..0be866390e --- /dev/null +++ b/catch/unit/atomics/arithmetic_common.hh @@ -0,0 +1,466 @@ +/* +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. +*/ + +#pragma once + +#include +#include +#include +#include + +namespace cg = cooperative_groups; + +// Atomic operations for which the tests in this file apply for +enum class AtomicOperation { + kAdd = 0, + kAddSystem, + kSub, + kSubSystem, + kInc, + kDec, + kUnsafeAdd, + kSafeAdd +}; + +// Constants that are passed as operands to the atomic operations +constexpr auto kIntegerTestValue = 7; +constexpr auto kFloatingPointTestValue = 3.125; +constexpr auto kIncDecWraparoundValue = 1023; + +// Retrieves test value constant based on the atomic operation and test type: +// - kIncDecWraparoundValue for increment and decrement operations +// - kFloatingPointTestValue for floating point test type +// - kIntegerTestValue for integer test type +template +__host__ __device__ TestType GetTestValue() { + if constexpr (operation == AtomicOperation::kInc || operation == AtomicOperation::kDec) { + return kIncDecWraparoundValue; + } + + return std::is_floating_point_v ? kFloatingPointTestValue : kIntegerTestValue; +} + +// Performs an atomic operation on parameter `mem` based on the `operation` enumerator. +template +__device__ TestType PerformAtomicOperation(TestType* const mem) { + const auto val = GetTestValue(); + + if constexpr (operation == AtomicOperation::kAdd) { + return atomicAdd(mem, val); + } else if constexpr (operation == AtomicOperation::kAddSystem) { + return atomicAdd_system(mem, val); + } else if constexpr (operation == AtomicOperation::kSub) { + return atomicSub(mem, val); + } else if constexpr (operation == AtomicOperation::kSubSystem) { + return atomicSub_system(mem, val); + } else if constexpr (operation == AtomicOperation::kInc) { + return atomicInc(mem, val); + } else if constexpr (operation == AtomicOperation::kDec) { + return atomicDec(mem, val); + } else if constexpr (operation == AtomicOperation::kUnsafeAdd) { + return unsafeAtomicAdd(mem, val); + } else if constexpr (operation == AtomicOperation::kSafeAdd) { + return safeAtomicAdd(mem, val); + } +} + +// This kernel executes the atomic operation specified by the enumerator `operation`. Results of +// the atomic operations are stored in `old_vals`. Each thread executes the atomic operation on the +// 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 +__global__ void TestKernel(TestType* const global_mem, TestType* const old_vals) { + __shared__ TestType shared_mem; + + const auto tid = cg::this_grid().thread_rank(); + + TestType* const mem = use_shared_mem ? &shared_mem : global_mem; + + if constexpr (use_shared_mem) { + if (tid == 0) mem[0] = global_mem[0]; + __syncthreads(); + } + + old_vals[tid] = PerformAtomicOperation(mem); + + if constexpr (use_shared_mem) { + __syncthreads(); + if (tid == 0) global_mem[0] = mem[0]; + } +} + +// Indexes array `ptr`, with the size in bytes of each element specified by `pitch` +template +__host__ __device__ TestType* PitchedOffset(TestType* const ptr, const unsigned int pitch, + const unsigned int idx) { + const auto byte_ptr = reinterpret_cast(ptr); + return reinterpret_cast(byte_ptr + idx * pitch); +} + +// 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 +// id (tid % width). +// The elements of `global_mem` can be larger than sizeof(TestType) with the actual size in bytes +// specified by `pitch`. This is done so we can test scenarios where threads target memory locations +// 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`. +// +// For example, given that sizeof(TestType) is 1, `width` is 3, and `pitch` is 4: +// +// 0 1 2 3 4 5 6 7 8 9 10 11 +// global_mem -> | x | | | | x | | | | x | | | | +// | pitch | pitch | pitch | +// +// In this scenario, the atomic operations will target the elements denoted with `x` (addresses 0, +// 4, 8). +template +__global__ void TestKernel(TestType* const global_mem, TestType* const old_vals, + const unsigned int width, const unsigned pitch) { + extern __shared__ uint8_t shared_mem[]; + + const auto tid = cg::this_grid().thread_rank(); + + TestType* const mem = use_shared_mem ? reinterpret_cast(shared_mem) : global_mem; + + if constexpr (use_shared_mem) { + if (tid < width) { + const auto target = PitchedOffset(mem, pitch, tid); + *target = *PitchedOffset(global_mem, pitch, tid); + }; + __syncthreads(); + } + + old_vals[tid] = + PerformAtomicOperation(PitchedOffset(mem, pitch, tid % width)); + + if constexpr (use_shared_mem) { + __syncthreads(); + if (tid < width) { + const auto target = PitchedOffset(global_mem, pitch, tid); + *target = *PitchedOffset(mem, pitch, tid); + }; + } +} + +// Used to configure test run +struct TestParams { + auto ThreadCount() const { + return blocks.x * blocks.y * blocks.z * threads.x * threads.y * threads.z; + } + + auto HostIterationsPerThread() const { + return std::max(num_devices * kernel_count * ThreadCount() / 20, width); + } + + dim3 blocks; // number of blocks per kernel launch + dim3 threads; // number of threads per kernel launch + unsigned int num_devices = 1u; // number of devices used + unsigned int kernel_count = 1u; // number of kernels launched per device + unsigned int width = 1u; // number of memory locations targeted + unsigned int pitch = 0u; // defines spacing between memory locations + unsigned int host_thread_count = 0u; // number of host threads launched + LinearAllocs alloc_type; // type of allocation used +}; + +// Reference implementation used to verify results +template +std::tuple, std::vector> TestKernelHostRef(const TestParams& p) { + const auto val = GetTestValue(); + + const auto total_thread_count = p.num_devices * p.kernel_count * p.ThreadCount() + + p.host_thread_count * p.HostIterationsPerThread(); + + std::vector res_vals(p.width); + std::vector old_vals; + old_vals.reserve(total_thread_count); + + auto perform_op = [&](unsigned id) { + auto& res = res_vals[id % p.width]; + old_vals.push_back(res); + + if constexpr (operation == AtomicOperation::kAdd || operation == AtomicOperation::kAddSystem || + operation == AtomicOperation::kUnsafeAdd || + operation == AtomicOperation::kSafeAdd) { + res = res + val; + } else if constexpr (operation == AtomicOperation::kSub || + operation == AtomicOperation::kSubSystem) { + res = res - val; + } else if constexpr (operation == AtomicOperation::kInc) { + res = (res >= val) ? 0 : res + 1; + } else if constexpr (operation == AtomicOperation::kDec) { + res = ((res == 0) || (res > val)) ? val : res - 1; + } + }; + + 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) { + perform_op(tid); + } + } + } + + for (auto i = 0u; i < p.host_thread_count; ++i) { + for (auto j = 0u; j < p.HostIterationsPerThread(); ++j) { + perform_op(j); + } + } + + return {res_vals, old_vals}; +} + +// Compares the results of the test kernel stored in `res_vals` with results generated by the +// reference implementation +template +void Verify(const TestParams& p, std::vector& res_vals, std::vector& old_vals) { + auto [expected_res_vals, expected_old_vals] = TestKernelHostRef(p); + + for (auto i = 0u; i < res_vals.size(); ++i) { + INFO("Results index: " << i); + REQUIRE(expected_res_vals[i] == res_vals[i]); + } + + std::sort(begin(old_vals), end(old_vals)); + std::sort(begin(expected_old_vals), end(expected_old_vals)); + for (auto i = 0u; i < old_vals.size(); ++i) { + INFO("Old values index: " << i); + REQUIRE(expected_old_vals[i] == old_vals[i]); + } +} + +// Launches the test kernel +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 + <<>>(mem_ptr, old_vals); + else + TestKernel + <<>>(mem_ptr, old_vals, p.width, p.pitch); +} + +// Performs a host atomic operation on parameter `mem` based on the `operation` enumerator. +template +void HostAtomicOperation(const unsigned int iterations, TestType* mem, TestType* const old_vals, + const unsigned int width, const unsigned pitch, TestType base_val) { + const auto val = GetTestValue(); + + for (auto i = 0u; i < iterations; ++i) { + if constexpr (operation == AtomicOperation::kAddSystem) { + 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); + } + } +} + +// Launches host threads based on TestParams::host_thread_count that compete with the test kernel +// for the same resources +template +void PerformHostAtomicOperation(const TestParams& p, TestType* mem, TestType* const old_vals) { + if (p.host_thread_count == 0) { + return; + } + + const auto host_base_val = p.num_devices * p.kernel_count * p.ThreadCount(); + + std::vector threads; + for (auto i = 0u; i < p.host_thread_count; ++i) { + const auto iterations = p.HostIterationsPerThread(); + const auto thread_base_val = host_base_val + i * iterations; + threads.push_back(std::thread(HostAtomicOperation, iterations, mem, + old_vals + thread_base_val, p.width, p.pitch, thread_base_val)); + } + + for (auto& th : threads) { + th.join(); + } +} + +// This is the main body of the test: +// 1. Allocate memory based on TestParams::alloc_type +// 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 +void TestCore(const TestParams& p) { + const unsigned int flags = + p.alloc_type == LinearAllocs::mallocAndRegister ? hipHostRegisterMapped : 0u; + + const auto old_vals_alloc_size = p.kernel_count * p.ThreadCount() * sizeof(TestType); + std::vector> old_vals_devs; + std::vector streams; + for (auto i = 0; i < p.num_devices; ++i) { + HIP_CHECK(hipSetDevice(i)); + old_vals_devs.emplace_back(LinearAllocs::hipMalloc, old_vals_alloc_size); + for (auto j = 0; j < p.kernel_count; ++j) { + streams.emplace_back(Streams::created); + } + } + + const auto mem_alloc_size = p.width * p.pitch; + LinearAllocGuard mem_dev(p.alloc_type, mem_alloc_size, flags); + + std::vector old_vals(p.num_devices * p.kernel_count * p.ThreadCount() + + p.host_thread_count * p.HostIterationsPerThread()); + std::vector res_vals(p.width); + + TestType* const mem_ptr = + p.alloc_type == LinearAllocs::hipMalloc ? mem_dev.ptr() : mem_dev.host_ptr(); + + HIP_CHECK(hipMemset(mem_ptr, 0, mem_alloc_size)); + + for (auto i = 0u; i < p.num_devices; ++i) { + 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); + } + } + + PerformHostAtomicOperation(p, mem_dev.host_ptr(), old_vals.data()); + + for (auto i = 0u; i < p.num_devices; ++i) { + const auto device_offset = i * p.kernel_count * p.ThreadCount(); + HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(), + old_vals_alloc_size, hipMemcpyDeviceToHost)); + } + HIP_CHECK(hipMemcpy2D(res_vals.data(), sizeof(TestType), mem_ptr, p.pitch, sizeof(TestType), + p.width, hipMemcpyDeviceToHost)); + + Verify(p, res_vals, old_vals); +} + +inline dim3 GenerateThreadDimensions() { return GENERATE(dim3(16), dim3(1024)); } + +inline dim3 GenerateBlockDimensions() { + int sm_count = 0; + HIP_CHECK(hipDeviceGetAttribute(&sm_count, hipDeviceAttributeMultiprocessorCount, 0)); + return GENERATE_COPY(dim3(sm_count), dim3(sm_count + sm_count / 2)); +} + +// Configures and creates the TestCore for a single device, and a single kernel launch +template +void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) { + TestParams params; + params.num_devices = 1; + params.kernel_count = 1; + params.threads = GenerateThreadDimensions(); + params.width = width; + params.pitch = pitch; + + SECTION("Global memory") { + 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); + } + } + } + + SECTION("Shared memory") { + params.blocks = dim3(1); + params.alloc_type = LinearAllocs::hipMalloc; + TestCore(params); + } +} + +// Configures and creates the TestCore for a single device, and multiple kernel launches +template +void SingleDeviceMultipleKernelTest(const unsigned int kernel_count, const unsigned int width, + const unsigned int pitch) { + int concurrent_kernels = 0; + HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, 0)); + if (!concurrent_kernels) { + HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution"); + return; + } + + TestParams params; + params.num_devices = 1; + params.kernel_count = kernel_count; + params.blocks = GenerateBlockDimensions(); + params.threads = GenerateThreadDimensions(); + params.width = width; + params.pitch = pitch; + + 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); + } + } +} + +// Configures and creates the TestCore for a multiple devices (and host), and multiple kernel +// launches +template +void MultipleDeviceMultipleKernelAndHostTest(const unsigned int num_devices, + const unsigned int kernel_count, + const unsigned int width, const unsigned int pitch, + const unsigned int host_thread_count = 0u) { + if (num_devices > 1) { + if (HipTest::getDeviceCount() < num_devices) { + std::string msg = std::to_string(num_devices) + " devices are required"; + HipTest::HIP_SKIP_TEST(msg.c_str()); + return; + } + } + + if (kernel_count > 1) { + for (auto i = 0u; i < num_devices; ++i) { + int concurrent_kernels = 0; + HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i)); + if (!concurrent_kernels) { + HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution"); + return; + } + } + } + + TestParams params; + params.num_devices = num_devices; + params.kernel_count = kernel_count; + params.blocks = GenerateBlockDimensions(); + params.threads = GenerateThreadDimensions(); + params.width = width; + params.pitch = pitch; + params.host_thread_count = host_thread_count; + + using LA = LinearAllocs; + 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); + } + } +} \ No newline at end of file