From d6dd4fd05b9135644bfa6c7904d9c7d6050a962b 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 16:51:22 +0100 Subject: [PATCH] EXSWHTEC-297 - Introduce common utilities for memory ordering tests for builtin atomic operations #389 Change-Id: Iae1db918eab6a722c85ff00183c973dd8dd54e9b --- catch/unit/atomics/memory_order_common.hh | 436 ++++++++++++++++++++++ 1 file changed, 436 insertions(+) create mode 100644 catch/unit/atomics/memory_order_common.hh diff --git a/catch/unit/atomics/memory_order_common.hh b/catch/unit/atomics/memory_order_common.hh new file mode 100644 index 0000000000..5e37575c07 --- /dev/null +++ b/catch/unit/atomics/memory_order_common.hh @@ -0,0 +1,436 @@ +/* +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 + +enum class BuiltinAtomicOperation { + kLoadStore = 0, + kExchange, + kCompareExchangeStrong, + kCompareExchangeWeak, + kAdd, + kAnd, + kOr, + kXor, + kMin, + kMax +}; + +template +__host__ __device__ void SetFlag(int* const flag) { +#ifdef __HIP_DEVICE_COMPILE__ + if constexpr (operation == BuiltinAtomicOperation::kLoadStore) { + static_assert(memory_order != __ATOMIC_ACQ_REL); + __hip_atomic_store(flag, 1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kExchange) { + __hip_atomic_exchange(flag, 1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kCompareExchangeStrong) { + int compare = 0; + __hip_atomic_compare_exchange_strong(flag, &compare, 1, memory_order, __ATOMIC_RELAXED, + memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kCompareExchangeWeak) { + int compare = 0; + while (!__hip_atomic_compare_exchange_weak(flag, &compare, 1, memory_order, __ATOMIC_RELAXED, + memory_scope)) + compare = 0; + } else if constexpr (operation == BuiltinAtomicOperation::kAdd) { + __hip_atomic_fetch_add(flag, 1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kAnd) { + __hip_atomic_fetch_and(flag, 0x0, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kOr) { + __hip_atomic_fetch_or(flag, 0x1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kXor) { + __hip_atomic_fetch_xor(flag, 0x1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kMin) { + __hip_atomic_fetch_min(flag, -1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kMax) { + __hip_atomic_fetch_max(flag, 1, memory_order, memory_scope); + } +#else + if constexpr (operation == BuiltinAtomicOperation::kAnd) { + __atomic_store_n(flag, 0, __ATOMIC_RELEASE); + } else { + __atomic_store_n(flag, 1, __ATOMIC_RELEASE); + } +#endif +} + +template +__host__ __device__ int FetchFlag(int* const flag) { +#ifdef __HIP_DEVICE_COMPILE__ + if constexpr (operation == BuiltinAtomicOperation::kLoadStore) { + static_assert(memory_order != __ATOMIC_ACQ_REL); + return __hip_atomic_load(flag, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kExchange) { + return __hip_atomic_exchange(flag, 0, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kCompareExchangeStrong) { + int compare = 1; + __hip_atomic_compare_exchange_strong( + flag, &compare, 1, memory_order, + memory_order == __ATOMIC_ACQ_REL ? __ATOMIC_ACQUIRE : memory_order, memory_scope); + return compare; + } else if constexpr (operation == BuiltinAtomicOperation::kCompareExchangeWeak) { + int compare = 1; + __hip_atomic_compare_exchange_weak( + flag, &compare, 1, memory_order, + memory_order == __ATOMIC_ACQ_REL ? __ATOMIC_ACQUIRE : memory_order, memory_scope); + return compare; + } else if constexpr (operation == BuiltinAtomicOperation::kAdd) { + return __hip_atomic_fetch_add(flag, 0, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kAnd) { + return !__hip_atomic_fetch_and(flag, 0x1, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kOr) { + return __hip_atomic_fetch_or(flag, 0x0, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kXor) { + return __hip_atomic_fetch_xor(flag, 0x0, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kMin) { + return __hip_atomic_fetch_min(flag, 0, memory_order, memory_scope); + } else if constexpr (operation == BuiltinAtomicOperation::kMax) { + return __hip_atomic_fetch_max(flag, 0, memory_order, memory_scope); + } +#else + if constexpr (operation == BuiltinAtomicOperation::kAnd) { + return !__atomic_load_n(flag, __ATOMIC_ACQUIRE); + } else { + return __atomic_load_n(flag, __ATOMIC_ACQUIRE); + } +#endif +} + +namespace AcquireRelease { + +constexpr auto kTestValue = 42; + +template +__host__ __device__ void Producer(int* const flag, int* const data) { + constexpr int actual_memory_order = + memory_order == __ATOMIC_ACQUIRE ? __ATOMIC_RELEASE : memory_order; + + data[0] = kTestValue; + + SetFlag(flag); +} + +template +__host__ __device__ void Consumer(int* const flag, int* const data, int* const ret) { + constexpr int actual_memory_order = + memory_order == __ATOMIC_RELEASE ? __ATOMIC_ACQUIRE : memory_order; + + while (!FetchFlag(flag)) + ; + + ret[0] = data[0]; +} + +template +__global__ void TestKernel(int* const flag, int* data, int* const ret) { + __shared__ int shared_mem; + + if (data == nullptr) data = &shared_mem; + + if (blockIdx.x == 0 && threadIdx.x == 0) { + if constexpr (operation == BuiltinAtomicOperation::kAnd) + *flag = 1; + else + *flag = 0; + } + __syncthreads(); + + bool producer = false, consumer = false; + + if constexpr (memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) { + producer = blockIdx.x == 0 && threadIdx.x == 0; + consumer = blockIdx.x == 0 && threadIdx.x == 1; + } else if constexpr (memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP) { + producer = blockIdx.x == 0 && threadIdx.x == 0; + consumer = blockIdx.x == 0 && threadIdx.x == warpSize; + } else if constexpr (memory_scope == __HIP_MEMORY_SCOPE_AGENT) { + producer = blockIdx.x == 0 && threadIdx.x == 0; + consumer = blockIdx.x == 1 && threadIdx.x == 0; + } + + if (producer) { + Producer(flag, data); + return; + } + + if (consumer) { + Consumer(flag, data, ret); + return; + } +} + +template +__global__ void ProducerKernel(int* const flag, int* const data) { + if (!(blockIdx.x == 0 && threadIdx.x == 0)) { + return; + } + + Producer(flag, data); +} + +template +__global__ void ConsumerKernel(int* const flag, int* const data, int* const ret) { + if (!(blockIdx.x == 0 && threadIdx.x == 0)) { + return; + } + + Consumer(flag, data, ret); +} + +template void Test() { + int blocks = 1, threads = 1; + if (memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) { + blocks = 1; + threads = 2; + } else if (memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP) { + blocks = 1; + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + threads = warp_size * 2; + } else if (memory_scope == __HIP_MEMORY_SCOPE_AGENT) { + blocks = 2; + threads = 1; + } + + LinearAllocGuard flag(LinearAllocs::hipMalloc, sizeof(int)); + LinearAllocGuard ret(LinearAllocs::hipMallocManaged, sizeof(int)); + + SECTION("Global memory") { + const auto alloc_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipMallocManaged); + LinearAllocGuard data(alloc_type, sizeof(int)); + TestKernel + <<>>(flag.ptr(), data.ptr(), ret.ptr()); + } + + if (memory_scope != __HIP_MEMORY_SCOPE_AGENT && memory_scope != __HIP_MEMORY_SCOPE_SYSTEM) { + SECTION("Shared memory") { + TestKernel + <<>>(flag.ptr(), nullptr, ret.ptr()); + } + } + + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(ret.ptr()[0] == kTestValue); +} + +template void SystemTest() { + std::thread host_thread; + + LinearAllocGuard flag(LinearAllocs::hipMallocManaged, sizeof(int)); + LinearAllocGuard ret(LinearAllocs::hipMallocManaged, sizeof(int)); + + SECTION("Global memory") { + const auto alloc_type = GENERATE(LinearAllocs::hipHostMalloc, LinearAllocs::hipMallocManaged); + LinearAllocGuard data(alloc_type, sizeof(int)); + + SECTION("Host producer - Device consumer") { + ConsumerKernel + <<<1, 1>>>(flag.ptr(), data.ptr(), ret.ptr()); + host_thread = std::thread([&] { + Producer(flag.ptr(), data.ptr()); + }); + } + + SECTION("Device producer - Host consumer") { + host_thread = std::thread([&] { + Consumer(flag.ptr(), data.ptr(), + ret.ptr()); + }); + ProducerKernel + <<<1, 1>>>(flag.ptr(), data.ptr()); + } + } + + HIP_CHECK(hipDeviceSynchronize()); + host_thread.join(); + + REQUIRE(ret.ptr()[0] == kTestValue); +} + +} /* namespace AcquireRelease */ + +namespace SequentialConsistency { + +template +__host__ __device__ void Producer(int* const flag) { + __atomic_store_n(flag, 1, __ATOMIC_SEQ_CST); +} + +template +__host__ __device__ void Consumer(int* const flag1, int* const flag2, int* const counter) { + while (!FetchFlag(flag1)) + ; + if (FetchFlag(flag2)) { +#ifdef __HIP_DEVICE_COMPILE__ + __hip_atomic_fetch_add(counter, 1, __ATOMIC_SEQ_CST, memory_scope); +#else + __atomic_fetch_add(counter, 1, __ATOMIC_SEQ_CST); +#endif + } +} + +template +__global__ void TestKernel(int* flag1, int* flag2, int* const counter) { + __shared__ int shared_mem[2]; + + if (flag1 == nullptr) flag1 = &shared_mem[0]; + if (flag2 == nullptr) flag2 = &shared_mem[1]; + + if (blockIdx.x == 0 && threadIdx.x == 0) { + if constexpr (operation == BuiltinAtomicOperation::kAnd) { + *flag1 = 1; + *flag2 = 1; + } else { + *flag1 = 0; + *flag2 = 0; + } + } + __syncthreads(); + + bool producer1 = false, producer2 = false, consumer1 = false, consumer2 = false; + + if constexpr (memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) { + producer1 = blockIdx.x == 0 && threadIdx.x == 0; + consumer1 = blockIdx.x == 0 && threadIdx.x == 1; + producer2 = blockIdx.x == 0 && threadIdx.x == 2; + consumer2 = blockIdx.x == 0 && threadIdx.x == 3; + } else if constexpr (memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP) { + producer1 = blockIdx.x == 0 && threadIdx.x == 0; + consumer1 = blockIdx.x == 0 && threadIdx.x == warpSize; + producer2 = blockIdx.x == 0 && threadIdx.x == warpSize * 2; + consumer2 = blockIdx.x == 0 && threadIdx.x == warpSize * 3; + } else if constexpr (memory_scope == __HIP_MEMORY_SCOPE_AGENT) { + producer1 = blockIdx.x == 0 && threadIdx.x == 0; + consumer1 = blockIdx.x == 1 && threadIdx.x == 0; + producer2 = blockIdx.x == 2 && threadIdx.x == 0; + consumer2 = blockIdx.x == 3 && threadIdx.x == 0; + } + + if (producer1) { + Producer(flag1); + return; + } + + if (consumer1) { + Consumer(flag1, flag2, counter); + return; + } + + if (producer2) { + Producer(flag2); + return; + } + + if (consumer2) { + Consumer(flag2, flag1, counter); + return; + } +} + +template +__global__ void ProducerKernel(int* const flag) { + if (!(blockIdx.x == 0 && threadIdx.x == 0)) { + return; + } + + Producer(flag); +} + +template +__global__ void ConsumerKernel(int* const flag1, int* const flag2, int* const counter) { + if (!(blockIdx.x == 0 && threadIdx.x == 0)) { + return; + } + + Consumer(flag1, flag2, counter); +} + +template void Test() { + int blocks = 1, threads = 1; + if (memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) { + blocks = 1; + threads = 4; + } else if (memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP) { + blocks = 1; + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + threads = warp_size * 4; + } else if (memory_scope == __HIP_MEMORY_SCOPE_AGENT) { + blocks = 4; + threads = 1; + } + + LinearAllocGuard counter(LinearAllocs::hipMallocManaged, sizeof(int)); + + SECTION("Global memory") { + const auto alloc_type = GENERATE(LinearAllocs::hipMalloc); + LinearAllocGuard flag1(alloc_type, sizeof(int)); + LinearAllocGuard flag2(alloc_type, sizeof(int)); + TestKernel + <<>>(flag1.ptr(), flag2.ptr(), counter.ptr()); + } + + if (memory_scope != __HIP_MEMORY_SCOPE_AGENT && memory_scope != __HIP_MEMORY_SCOPE_SYSTEM) { + SECTION("Shared memory") { + TestKernel<<>>(nullptr, nullptr, counter.ptr()); + } + } + + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(counter.ptr()[0] != 0); +} + +template void SystemTest() { + std::thread host_producer, host_consumer; + + LinearAllocGuard counter(LinearAllocs::hipMallocManaged, sizeof(int)); + + SECTION("Global memory") { + const auto alloc_type = GENERATE(LinearAllocs::hipMallocManaged); + LinearAllocGuard flag1(alloc_type, sizeof(int)); + LinearAllocGuard flag2(alloc_type, sizeof(int)); + + ConsumerKernel + <<<1, 1>>>(flag1.ptr(), flag2.ptr(), counter.ptr()); + host_consumer = std::thread([&] { + Consumer(flag2.ptr(), flag1.ptr(), counter.ptr()); + }); + + ProducerKernel<<<1, 1>>>(flag1.ptr()); + host_producer = + std::thread([&] { Producer(flag2.ptr()); }); + } + + HIP_CHECK(hipDeviceSynchronize()); + host_producer.join(); + host_consumer.join(); + + REQUIRE(counter.ptr()[0] != 0); +} + +} // namespace SequentialConsistency \ No newline at end of file