diff --git a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt index bfe6e6bf59..bb3b8db915 100644 --- a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt @@ -33,6 +33,8 @@ set(TEST_SRC unsafeAtomicMin.cc safeAtomicMax.cc unsafeAtomicMax.cc + __hip_atomic_fetch_min.cc + __hip_atomic_fetch_max.cc atomicExch.cc atomicExch_system.cc __hip_atomic_fetch_and.cc diff --git a/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_max.cc b/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_max.cc new file mode 100644 index 0000000000..cc42309333 --- /dev/null +++ b/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_max.cc @@ -0,0 +1,187 @@ +/* +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 "min_max_common.hh" + +#include + +/** + * @addtogroup __hip_atomic_fetch_max __hip_atomic_fetch_max + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MAX with memory scope WAVEFRONT from multiple threads on the same + * address. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_max.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Same address " << current) { + MinMax::SingleDeviceSingleKernelTest(1, sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MAX with memory scope WAVEFRONT from multiple threads on adjacent + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_max.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Adjacent address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MAX with memory scope WAVEFRONT from multiple threads on scattered + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_max.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses", "", 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("Scattered address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + cache_line_size); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MAX with memory scope WORKGROUP from multiple threads on the same + * address. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_max.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Same address " << current) { + MinMax::SingleDeviceSingleKernelTest(1, sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MAX with memory scope WORKGROUP from multiple threads on adjacent + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_max.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Adjacent address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MAX with memory scope WORKGROUP from multiple threads on scattered + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_max.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses", "", 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("Scattered address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + cache_line_size); + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_min.cc b/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_min.cc new file mode 100644 index 0000000000..f09a3732f9 --- /dev/null +++ b/projects/hip-tests/catch/unit/atomics/__hip_atomic_fetch_min.cc @@ -0,0 +1,187 @@ +/* +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 "min_max_common.hh" + +#include + +/** + * @addtogroup __hip_atomic_fetch_min __hip_atomic_fetch_min + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MIN with memory scope WAVEFRONT from multiple threads on the same + * address. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_min.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Same address " << current) { + MinMax::SingleDeviceSingleKernelTest(1, sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MIN with memory scope WAVEFRONT from multiple threads on adjacent + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_min.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Adjacent address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MIN with memory scope WAVEFRONT from multiple threads on scattered + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_min.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses", "", 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("Scattered address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + cache_line_size); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MIN with memory scope WORKGROUP from multiple threads on the same + * address. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_min.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Same address " << current) { + MinMax::SingleDeviceSingleKernelTest(1, sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MIN with memory scope WORKGROUP from multiple threads on adjacent + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_min.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses", "", int, + unsigned int, unsigned long, unsigned long long, float, double) { + int warp_size = 0; + HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); + + for (auto current = 0; current < cmd_options.iterations; ++current) { + DYNAMIC_SECTION("Adjacent address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + sizeof(TestType)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs a builtin atomic MIN with memory scope WORKGROUP from multiple threads on scattered + * addresses. + * - Uses only one device and launches one kernel. + * Test source + * ------------------------ + * - unit/atomics/__hip_atomic_fetch_min.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses", "", 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("Scattered address " << current) { + MinMax::SingleDeviceSingleKernelTest(warp_size, + cache_line_size); + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/atomics/min_max_common.hh b/projects/hip-tests/catch/unit/atomics/min_max_common.hh index c171c6f3c6..bd2934f4ee 100644 --- a/projects/hip-tests/catch/unit/atomics/min_max_common.hh +++ b/projects/hip-tests/catch/unit/atomics/min_max_common.hh @@ -39,6 +39,9 @@ enum class AtomicOperation { kUnsafeMin, kSafeMax, kUnsafeMax + kUnsafeMax, + kBuiltinMin, + kBuiltinMax }; constexpr auto kIntegerTestValue = 5; @@ -58,7 +61,7 @@ __host__ __device__ TestType GetTestValue() { return test_value + 2; } -template +template __device__ TestType PerformAtomicOperation(TestType* const mem) { const auto val = GetTestValue(); @@ -78,10 +81,15 @@ __device__ TestType PerformAtomicOperation(TestType* const mem) { return unsafeAtomicMax(mem, val); } else if constexpr (operation == AtomicOperation::kSafeMax) { return safeAtomicMax(mem, val); + } else if constexpr (operation == AtomicOperation::kBuiltinMin) { + return __hip_atomic_fetch_min(mem, val, __ATOMIC_RELAXED, memory_scope); + } else if constexpr (operation == AtomicOperation::kBuiltinMax) { + return __hip_atomic_fetch_max(mem, val, __ATOMIC_RELAXED, memory_scope); } } -template +template __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals) { __shared__ TestType shared_mem; @@ -94,7 +102,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(); @@ -109,7 +117,16 @@ __host__ __device__ TestType* PitchedOffset(TestType* const ptr, const unsigned return reinterpret_cast(byte_ptr + idx * pitch); } -template +__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; + } +} + +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[]; @@ -126,8 +143,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(); @@ -172,12 +199,14 @@ std::tuple, std::vector> TestKernelHostRef(const if constexpr (operation == AtomicOperation::kMin || operation == AtomicOperation::kMinSystem || operation == AtomicOperation::kUnsafeMin || - operation == AtomicOperation::kSafeMin) { + operation == AtomicOperation::kSafeMin || + operation == AtomicOperation::kBuiltinMin) { res = std::min(res, val); } else if constexpr (operation == AtomicOperation::kMax || operation == AtomicOperation::kMaxSystem || operation == AtomicOperation::kUnsafeMax || - operation == AtomicOperation::kSafeMax) { + operation == AtomicOperation::kSafeMax || + operation == AtomicOperation::kBuiltinMax) { res = std::max(res, val); } } @@ -202,19 +231,21 @@ 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); } -template +template void TestCore(const TestParams& p) { const auto old_vals_alloc_size = p.kernel_count * p.ThreadCount() * sizeof(TestType); std::vector> old_vals_devs; @@ -247,7 +278,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); } } @@ -262,17 +294,45 @@ void TestCore(const TestParams& p) { Verify(p, res_vals, old_vals); } -template +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)); +} + +template void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) { TestParams params; params.num_devices = 1; params.kernel_count = 1; - params.threads = GENERATE(dim3(1023)); + if constexpr ((operation == AtomicOperation::kBuiltinMin || + operation == AtomicOperation::kBuiltinMax) && + memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD) { + params.threads = 1; + } else if constexpr ((operation == AtomicOperation::kBuiltinMin || + operation == AtomicOperation::kBuiltinMax) && + 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 = GENERATE(dim3(3)); + if constexpr ((operation == AtomicOperation::kBuiltinMin || + operation == AtomicOperation::kBuiltinMax) && + (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}) { @@ -303,8 +363,8 @@ void SingleDeviceMultipleKernelTest(const unsigned int kernel_count, const unsig TestParams params; params.num_devices = 1; params.kernel_count = kernel_count; - params.blocks = GENERATE(dim3(3)); - params.threads = GENERATE(dim3(1023)); + params.blocks = GenerateThreadDimensions(); + params.threads = GenerateBlockDimensions(); params.width = width; params.pitch = pitch; @@ -344,8 +404,8 @@ void MultipleDeviceMultipleKernelTest(const unsigned int num_devices, TestParams params; params.num_devices = num_devices; params.kernel_count = kernel_count; - params.blocks = GENERATE(dim3(3)); - params.threads = GENERATE(dim3(1023)); + params.blocks = GenerateThreadDimensions(); + params.threads = GenerateBlockDimensions(); params.width = width; params.pitch = pitch; @@ -353,8 +413,9 @@ void MultipleDeviceMultipleKernelTest(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); } } } + } // namespace MinMax