diff --git a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt index ed622195a9..1c40746dff 100644 --- a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt @@ -35,6 +35,7 @@ set(TEST_SRC unsafeAtomicMax.cc __hip_atomic_fetch_min.cc __hip_atomic_fetch_max.cc + atomic_builtins.cc atomicExch.cc atomicExch_system.cc __hip_atomic_fetch_and.cc diff --git a/projects/hip-tests/catch/unit/atomics/atomic_builtins.cc b/projects/hip-tests/catch/unit/atomics/atomic_builtins.cc new file mode 100644 index 0000000000..c5ade6b30a --- /dev/null +++ b/projects/hip-tests/catch/unit/atomics/atomic_builtins.cc @@ -0,0 +1,97 @@ +/* +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 +#include + +#include "atomic_builtins_kernels_rtc.hh" + +/** + * @addtogroup __hip_atomic_fetch_add __hip_atomic_fetch_add + * @{ + * @ingroup AtomicsTest + */ + +void AtomicBuiltinsRTCWrapper(const char* program_source, int expected_errors_num, + int expected_warnings_num) { + hiprtcProgram program{}; + HIPRTC_CHECK(hiprtcCreateProgram(&program, program_source, "atomics_builtins_kernels.cc", 0, + nullptr, nullptr)); + + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + int warning_count{0}; + + std::string error_message{"error:"}; + std::string warning_message{"warning:"}; + + size_t npos_e = log.find(error_message, 0); + while (npos_e != std::string::npos) { + ++error_count; + npos_e = log.find(error_message, npos_e + 1); + } + + size_t npos_w = log.find(warning_message, 0); + while (npos_w != std::string::npos) { + ++warning_count; + npos_w = log.find(warning_message, npos_w + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_errors_num); + REQUIRE(warning_count == expected_warnings_num); +} + +/** + * Test Description + * ------------------------ + * - Compiles atomic builtins while passing parameters that shall cause: + * -# Compiler warnings + * -# Compiler errors + * Test source + * ------------------------ + * - unit/atomics/atomic_builtins.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_AtomicBuiltins_Negative_Parameters_RTC") { + AtomicBuiltinsRTCWrapper(kBuiltinStore, 5, 5); + AtomicBuiltinsRTCWrapper(kBuiltinLoad, 4, 4); + /* Begin: Should be 5 errors, 6 warnings for both. See EXSWHTEC-309*/ + AtomicBuiltinsRTCWrapper(kBuiltinCompExWeak, 5, 2); + AtomicBuiltinsRTCWrapper(kBuiltinCompExStrong, 5, 2); + /* End. */ + AtomicBuiltinsRTCWrapper(kBuiltinExchange, 5, 2); + AtomicBuiltinsRTCWrapper(kBuiltinFetchAdd, 5, 2); + AtomicBuiltinsRTCWrapper(kBuiltinFetchAnd, 7, 2); + AtomicBuiltinsRTCWrapper(kBuiltinFetchOr, 7, 2); + AtomicBuiltinsRTCWrapper(kBuiltinFetchXor, 7, 2); + AtomicBuiltinsRTCWrapper(kBuiltinFetchMax, 5, 2); + AtomicBuiltinsRTCWrapper(kBuiltinFetchMin, 5, 2); +} diff --git a/projects/hip-tests/catch/unit/atomics/atomic_builtins_kernels_rtc.hh b/projects/hip-tests/catch/unit/atomics/atomic_builtins_kernels_rtc.hh new file mode 100644 index 0000000000..1339eaaa45 --- /dev/null +++ b/projects/hip-tests/catch/unit/atomics/atomic_builtins_kernels_rtc.hh @@ -0,0 +1,590 @@ +/* +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 + +/* +Positive and negative kernels used for the builtin atomic Test Cases that are using RTC. +*/ + +static constexpr auto kBuiltinStore{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void StoreCompileKernel(int* x) { + __hip_atomic_store(x, 1, __ATOMIC_RELAXED, kMemScope); + __hip_atomic_store(x, 1, __ATOMIC_RELEASE, kMemScope); + __hip_atomic_store(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + __hip_atomic_store(reinterpret_cast(x), 1, kMemOrder, kMemScope); + __hip_atomic_store(*x, 1, kMemOrder, kMemScope); + __hip_atomic_store(x, 1, __ATOMIC_CONSUME, kMemScope); + __hip_atomic_store(x, 1, __ATOMIC_ACQUIRE, kMemScope); + __hip_atomic_store(x, 1, __ATOMIC_ACQ_REL, kMemScope); + __hip_atomic_store(x, 1, -1, kMemScope); + __hip_atomic_store(x, 1, 10, kMemScope); + __hip_atomic_store(x, 1, kMemOrder, -1); + __hip_atomic_store(x, 1, kMemOrder, 10); + + Dummy dummy_a{}; + Dummy dummy_b{}; + __hip_atomic_store(&dummy_a, dummy_b, kMemOrder, kMemScope); + + DummyTC dummytc_a{}; + DummyTC dummytc_b{}; + __hip_atomic_store(&dummytc_a, dummytc_b, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinLoad{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void LoadCompileKernel(int* x, int* y) { + *y = __hip_atomic_load(x, __ATOMIC_RELAXED, kMemScope); + *y = __hip_atomic_load(x, __ATOMIC_CONSUME, kMemScope); + *y = __hip_atomic_load(x, __ATOMIC_ACQUIRE, kMemScope); + *y = __hip_atomic_load(x, __ATOMIC_SEQ_CST, kMemScope); + + *y = __hip_atomic_load(*x, kMemOrder, kMemScope); + *y = __hip_atomic_load(x, __ATOMIC_RELEASE, kMemScope); + *y = __hip_atomic_load(x, __ATOMIC_ACQ_REL, kMemScope); + *y = __hip_atomic_load(x, -1, kMemScope); + *y = __hip_atomic_load(x, 10, kMemScope); + *y = __hip_atomic_load(x, kMemOrder, -1); + *y = __hip_atomic_load(x, kMemOrder, 10); + + Dummy dummy_a{}; + Dummy dummy_b{}; + dummy_a = __hip_atomic_load(&dummy_b, kMemOrder, kMemScope); + + DummyTC dummytc_a{}; + DummyTC dummytc_b{}; + dummytc_a = __hip_atomic_load(&dummytc_b, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinCompExWeak{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void CompareWeakCompileKernel(int* x, int* expected) { + bool res{false}; + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_RELAXED, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_CONSUME, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_CONSUME, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_ACQUIRE, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_ACQUIRE, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_ACQUIRE, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_RELEASE, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_RELEASE, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_RELEASE, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_ACQ_REL, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_ACQ_REL, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_ACQ_REL, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_ACQ_REL, + kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST, + kMemScope); + + res = __hip_atomic_compare_exchange_weak(x, expected, 1, kMemOrder, __ATOMIC_RELEASE, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, kMemOrder, __ATOMIC_ACQ_REL, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, __ATOMIC_RELAXED, __ATOMIC_SEQ_CST, + kMemScope); + res = __hip_atomic_compare_exchange_weak(reinterpret_cast(x), expected, 1, kMemOrder, + kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_weak(*x, expected, 1, kMemOrder, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, -1, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, 10, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, kMemOrder, -1, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, kMemOrder, 10, kMemScope); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, kMemOrder, kMemOrder, -1); + res = __hip_atomic_compare_exchange_weak(x, expected, 1, kMemOrder, kMemOrder, 10); + + Dummy dummy_a{}; + Dummy dummy_b{}; + Dummy dummy_c{}; + res = __hip_atomic_compare_exchange_weak(&dummy_a, &dummy_b, dummy_c, kMemOrder, kMemOrder, + kMemScope); + DummyTC dummytc_a{}; + DummyTC dummytc_b{}; + DummyTC dummytc_c{}; + res = __hip_atomic_compare_exchange_weak(&dummytc_a, &dummytc_b, dummytc_c, kMemOrder, kMemOrder, + kMemScope); + } +)"}; + +static constexpr auto kBuiltinCompExStrong{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void CompareStrongCompileKernel(int* x, int* expected) { + bool res{false}; + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_RELAXED, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_CONSUME, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_CONSUME, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_ACQUIRE, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_ACQUIRE, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_ACQUIRE, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_RELEASE, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_RELEASE, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_RELEASE, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_ACQ_REL, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_ACQ_REL, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_ACQ_REL, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_RELAXED, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_CONSUME, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_ACQUIRE, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_ACQ_REL, + kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST, + kMemScope); + + res = + __hip_atomic_compare_exchange_strong(x, expected, 1, kMemOrder, __ATOMIC_RELEASE, kMemScope); + res = + __hip_atomic_compare_exchange_strong(x, expected, 1, kMemOrder, __ATOMIC_ACQ_REL, kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, __ATOMIC_RELAXED, __ATOMIC_SEQ_CST, + kMemScope); + res = __hip_atomic_compare_exchange_strong(reinterpret_cast(x), expected, 1, + kMemOrder, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_strong(*x, expected, 1, kMemOrder, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, -1, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, 10, kMemOrder, kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, kMemOrder, -1, kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, kMemOrder, 10, kMemScope); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, kMemOrder, kMemOrder, -1); + res = __hip_atomic_compare_exchange_strong(x, expected, 1, kMemOrder, kMemOrder, 10); + + Dummy dummy_a{}; + Dummy dummy_b{}; + Dummy dummy_c{}; + res = __hip_atomic_compare_exchange_strong(&dummy_a, &dummy_b, dummy_c, kMemOrder, kMemOrder, + kMemScope); + DummyTC dummytc_a{}; + DummyTC dummytc_b{}; + DummyTC dummytc_c{}; + res = __hip_atomic_compare_exchange_strong(&dummytc_a, &dummytc_b, dummytc_c, kMemOrder, + kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinExchange{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void ExchangeCompileKernel(int* x) { + int old{}; + old = __hip_atomic_exchange(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_exchange(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_exchange(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_exchange(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_exchange(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_exchange(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_exchange(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_exchange(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_exchange(x, 1, -1, kMemScope); + old = __hip_atomic_exchange(x, 1, 10, kMemScope); + old = __hip_atomic_exchange(x, 1, kMemOrder, -1); + old = __hip_atomic_exchange(x, 1, kMemOrder, 10); + + Dummy dummy_a{}; + Dummy dummy_b{}; + dummy_b = __hip_atomic_exchange(&dummy_a, dummy_b, kMemOrder, kMemScope); + + DummyTC dummytc_a{}; + DummyTC dummytc_b{}; + dummytc_b = __hip_atomic_exchange(&dummytc_a, dummytc_b, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinFetchAdd{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void FetchAddCompileKernel(int* x) { + int old{}; + old = __hip_atomic_fetch_add(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_fetch_add(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_fetch_add(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_fetch_add(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_fetch_add(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_fetch_add(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_fetch_add(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_add(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_add(x, 1, -1, kMemScope); + old = __hip_atomic_fetch_add(x, 1, 10, kMemScope); + old = __hip_atomic_fetch_add(x, 1, kMemOrder, -1); + old = __hip_atomic_fetch_add(x, 1, kMemOrder, 10); + + Dummy dummy{}; + old = __hip_atomic_fetch_add(&dummy, 1, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinFetchAnd{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void FetchAndCompileKernel(int* x) { + int old{}; + old = __hip_atomic_fetch_and(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_fetch_and(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_fetch_and(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_fetch_and(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_fetch_and(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_fetch_and(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_fetch_and(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_and(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_and(x, 1, -1, kMemScope); + old = __hip_atomic_fetch_and(x, 1, 10, kMemScope); + old = __hip_atomic_fetch_and(x, 1, kMemOrder, -1); + old = __hip_atomic_fetch_and(x, 1, kMemOrder, 10); + + Dummy dummy{}; + old = __hip_atomic_fetch_and(&dummy, 1, kMemOrder, kMemScope); + float float_var{1.5f}; + old = __hip_atomic_fetch_and(&float_var, 1, kMemOrder, kMemScope); + double double_var{1.5}; + old = __hip_atomic_fetch_and(&double_var, 1, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinFetchOr{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void FetchOrCompileKernel(int* x) { + int old{}; + old = __hip_atomic_fetch_or(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_fetch_or(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_fetch_or(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_fetch_or(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_fetch_or(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_fetch_or(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_fetch_or(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_or(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_or(x, 1, -1, kMemScope); + old = __hip_atomic_fetch_or(x, 1, 10, kMemScope); + old = __hip_atomic_fetch_or(x, 1, kMemOrder, -1); + old = __hip_atomic_fetch_or(x, 1, kMemOrder, 10); + + Dummy dummy{}; + old = __hip_atomic_fetch_or(&dummy, 1, kMemOrder, kMemScope); + float float_var{1.5f}; + old = __hip_atomic_fetch_or(&float_var, 1, kMemOrder, kMemScope); + double double_var{1.5}; + old = __hip_atomic_fetch_or(&double_var, 1, kMemOrder, kMemScope); + } +)"}; + +static auto constexpr kBuiltinFetchXor{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void FetchXorCompileKernel(int* x) { + int old{}; + old = __hip_atomic_fetch_xor(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_fetch_xor(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_xor(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, -1, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, 10, kMemScope); + old = __hip_atomic_fetch_xor(x, 1, kMemOrder, -1); + old = __hip_atomic_fetch_xor(x, 1, kMemOrder, 10); + + Dummy dummy{}; + old = __hip_atomic_fetch_xor(&dummy, 1, kMemOrder, kMemScope); + float float_var{1.5f}; + old = __hip_atomic_fetch_xor(&float_var, 1, kMemOrder, kMemScope); + double double_var{1.5}; + old = __hip_atomic_fetch_xor(&double_var, 1, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinFetchMax{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void FetchMaxCompileKernel(int* x) { + int old{}; + old = __hip_atomic_fetch_max(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_fetch_max(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_fetch_max(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_fetch_max(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_fetch_max(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_fetch_max(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_fetch_max(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_max(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_max(x, 1, -1, kMemScope); + old = __hip_atomic_fetch_max(x, 1, 10, kMemScope); + old = __hip_atomic_fetch_max(x, 1, kMemOrder, -1); + old = __hip_atomic_fetch_max(x, 1, kMemOrder, 10); + + Dummy dummy{}; + old = __hip_atomic_fetch_max(&dummy, 1, kMemOrder, kMemScope); + } +)"}; + +static constexpr auto kBuiltinFetchMin{R"( + constexpr int kMemOrder = __ATOMIC_RELAXED; + constexpr int kMemScope = __HIP_MEMORY_SCOPE_SYSTEM; + + class DummyTC { + public: + __device__ DummyTC() {} + __device__ ~DummyTC() = default; + __device__ DummyTC(const DummyTC&) = default; + __device__ DummyTC& operator=(const DummyTC&) = default; + __device__ DummyTC(DummyTC&&) = default; + __device__ DummyTC& operator=(DummyTC&&) = default; + }; + + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void FetchMinCompileKernel(int* x) { + int old{}; + old = __hip_atomic_fetch_min(x, 1, __ATOMIC_RELAXED, kMemScope); + old = __hip_atomic_fetch_min(x, 1, __ATOMIC_CONSUME, kMemScope); + old = __hip_atomic_fetch_min(x, 1, __ATOMIC_ACQUIRE, kMemScope); + old = __hip_atomic_fetch_min(x, 1, __ATOMIC_RELEASE, kMemScope); + old = __hip_atomic_fetch_min(x, 1, __ATOMIC_ACQ_REL, kMemScope); + old = __hip_atomic_fetch_min(x, 1, __ATOMIC_SEQ_CST, kMemScope); + + old = __hip_atomic_fetch_min(reinterpret_cast(x), 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_min(*x, 1, kMemOrder, kMemScope); + old = __hip_atomic_fetch_min(x, 1, -1, kMemScope); + old = __hip_atomic_fetch_min(x, 1, 10, kMemScope); + old = __hip_atomic_fetch_min(x, 1, kMemOrder, -1); + old = __hip_atomic_fetch_min(x, 1, kMemOrder, 10); + + Dummy dummy{}; + old = __hip_atomic_fetch_min(&dummy, 1, kMemOrder, kMemScope); + } +)"};