diff --git a/catch/hipTestMain/main.cc b/catch/hipTestMain/main.cc index 030267f04c..74096c4af4 100644 --- a/catch/hipTestMain/main.cc +++ b/catch/hipTestMain/main.cc @@ -30,9 +30,9 @@ int main(int argc, char** argv) { | Opt(cmd_options.progress) ["-P"]["--progress"] ("Show progress bar when running performance tests") - | Opt(cmd_options.extended_run) - ["-E"]["--extended-run"] - ("TODO: Description goes here") + | Opt(cmd_options.cg_extended_run, "cg_extened_run") + ["-E"]["--cg-extended-run"] + ("Number of iterations used for cooperative groups sync tests (default: 5)") ; // clang-format on diff --git a/catch/include/cmd_options.hh b/catch/include/cmd_options.hh index 5dbd2f300c..99e565491b 100644 --- a/catch/include/cmd_options.hh +++ b/catch/include/cmd_options.hh @@ -23,11 +23,11 @@ THE SOFTWARE. #pragma once struct CmdOptions { - int iterations = 1000; + int iterations = 10; int warmups = 100; + int cg_extended_run = 5; bool no_display = false; bool progress = false; - bool extended_run = false; }; extern CmdOptions cmd_options; \ No newline at end of file diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 9274bbf6db..147abe0941 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -102,6 +102,19 @@ THE SOFTWARE. } \ } +// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError. +#define HIPRTC_CHECK_ERROR(errorExpr, expectedError) \ + { \ + auto localError = errorExpr; \ + INFO("Matching Errors: " \ + << "\n Expected Error: " << hiprtcGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << '\n' \ + << " Actual Error: " << hiprtcGetErrorString(localError) \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ + REQUIRE(localError == expectedError); \ + } + // Although its assert, it will be evaluated at runtime #define HIP_ASSERT(x) \ { REQUIRE((x)); } diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 7ade70efd7..5294e11a5c 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -20,6 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include + // Test groups are named based on the group names from hip_api_runtime.h, with adding "Test" suffix /** @@ -95,8 +97,46 @@ THE SOFTWARE. /** * @defgroup KernelTest Kernel Functions Management +* @{ +* This section describes the various kernel functions invocation. +* @} +*/ + +/** + * @defgroup AtomicsTest Device Atomics * @{ - * This section describes the various kernel functions invocation. + * This section describes tests for the Device Atomic APIs. + * @} + */ + + /** + * @addtogroup atomicExch atomicExch + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Compiles atomicExch with invalid parameters. + * - Compiles the source with specialized Python tool. + * -# Utilizes sub-process to invoke compilation of faulty source. + * -# Performs post-processing of compiler output and counts errors. + * Test source + * ------------------------ + * - unit/atomics/CMakeLists.txt + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_atomicExch_Negative_Parameters") {} +/** + * End doxygen group atomicExch. + * @} + */ + +/** + * End doxygen group AtomicsTest. * @} */ diff --git a/catch/include/resource_guards.hh b/catch/include/resource_guards.hh index 0f8697e121..3ec6aabf8d 100644 --- a/catch/include/resource_guards.hh +++ b/catch/include/resource_guards.hh @@ -29,10 +29,30 @@ enum class LinearAllocs { hipHostMalloc, hipMalloc, hipMallocManaged, + noAlloc }; +inline std::string to_string(const LinearAllocs allocation_type) { + switch (allocation_type) { + case LinearAllocs::malloc: + return "host pageable"; + case LinearAllocs::mallocAndRegister: + return "registered"; + case LinearAllocs::hipHostMalloc: + return "host pinned"; + case LinearAllocs::hipMalloc: + return "device malloc"; + case LinearAllocs::hipMallocManaged: + return "managed"; + default: + return "unknown alloc type"; + } +} + template class LinearAllocGuard { public: + LinearAllocGuard() = default; + LinearAllocGuard(const LinearAllocs allocation_type, const size_t size, const unsigned int flags = 0u) : allocation_type_{allocation_type} { @@ -55,15 +75,36 @@ template class LinearAllocGuard { case LinearAllocs::hipMallocManaged: HIP_CHECK(hipMallocManaged(reinterpret_cast(&ptr_), size, flags ? flags : 1u)); host_ptr_ = ptr_; + break; + case LinearAllocs::noAlloc: + break; } } LinearAllocGuard(const LinearAllocGuard&) = delete; - LinearAllocGuard(LinearAllocGuard&&) = delete; + + LinearAllocGuard(LinearAllocGuard&& o) + : allocation_type_{o.allocation_type_}, ptr_{o.ptr_}, host_ptr_{o.host_ptr_} { + o.allocation_type_ = LinearAllocs::noAlloc; + o.ptr_ = nullptr; + o.host_ptr_ = nullptr; + } + + LinearAllocGuard& operator=(LinearAllocGuard&& o) { + allocation_type_ = o.allocation_type_; + ptr_ = o.ptr_; + host_ptr_ = o.host_ptr_; + + o.allocation_type_ = LinearAllocs::noAlloc; + o.ptr_ = nullptr; + o.host_ptr_ = nullptr; + } ~LinearAllocGuard() { // No Catch macros, don't want to possibly throw in the destructor switch (allocation_type_) { + case LinearAllocs::noAlloc: + break; case LinearAllocs::malloc: free(ptr_); break; @@ -85,7 +126,7 @@ template class LinearAllocGuard { T* host_ptr() const { return host_ptr_; } private: - const LinearAllocs allocation_type_; + LinearAllocs allocation_type_ = LinearAllocs::noAlloc; T* ptr_ = nullptr; T* host_ptr_ = nullptr; }; @@ -200,7 +241,10 @@ enum class Streams { nullstream, perThread, created, withFlags, withPriority }; class StreamGuard { public: - StreamGuard(const Streams stream_type, unsigned int flags = hipStreamDefault, int priority = 0) : stream_type_{stream_type}, flags_{flags}, priority_{priority} { + StreamGuard() = default; + + StreamGuard(const Streams stream_type, unsigned int flags = hipStreamDefault, int priority = 0) + : stream_type_{stream_type}, flags_{flags}, priority_{priority} { switch (stream_type_) { case Streams::nullstream: stream_ = nullptr; @@ -219,7 +263,28 @@ class StreamGuard { } StreamGuard(const StreamGuard&) = delete; - StreamGuard(StreamGuard&&) = delete; + + StreamGuard(StreamGuard&& o) + : stream_type_{o.stream_type_}, flags_{o.flags_}, priority_{o.priority_}, stream_{o.stream_} { + o.stream_type_ = Streams::nullstream; + o.flags_ = 0u; + o.priority_ = 0; + o.stream_ = nullptr; + } + + StreamGuard& operator=(StreamGuard&& o) { + stream_type_ = o.stream_type_; + flags_ = o.flags_; + priority_ = o.priority_; + stream_ = o.stream_; + + o.stream_type_ = Streams::nullstream; + o.flags_ = 0u; + o.priority_ = 0; + o.stream_ = nullptr; + + return *this; + } ~StreamGuard() { if (stream_type_ == Streams::created) { @@ -230,23 +295,23 @@ class StreamGuard { hipStream_t stream() const { return stream_; } private: - const Streams stream_type_; - unsigned int flags_; - int priority_; - hipStream_t stream_; + Streams stream_type_ = Streams::nullstream; + unsigned int flags_ = 0u; + int priority_ = 0; + hipStream_t stream_ = nullptr; }; class EventsGuard { -public: + public: EventsGuard(size_t N) : events_(N) { - for (auto &e : events_) HIP_CHECK(hipEventCreate(&e)); + for (auto& e : events_) HIP_CHECK(hipEventCreate(&e)); } EventsGuard(const EventsGuard&) = delete; EventsGuard(EventsGuard&&) = delete; ~EventsGuard() { - for (auto &e : events_) static_cast(hipEventDestroy(e)); + for (auto& e : events_) static_cast(hipEventDestroy(e)); } hipEvent_t& operator[](int index) { return events_[index]; } @@ -255,21 +320,21 @@ public: std::vector& event_list() { return events_; } -private: + private: std::vector events_; }; class StreamsGuard { -public: + public: StreamsGuard(size_t N) : streams_(N) { - for (auto &s : streams_) HIP_CHECK(hipStreamCreate(&s)); + for (auto& s : streams_) HIP_CHECK(hipStreamCreate(&s)); } StreamsGuard(const StreamsGuard&) = delete; StreamsGuard(StreamsGuard&&) = delete; ~StreamsGuard() { - for (auto &s : streams_) static_cast(hipStreamDestroy(s)); + for (auto& s : streams_) static_cast(hipStreamDestroy(s)); } hipStream_t& operator[](int index) { return streams_[index]; } @@ -278,6 +343,6 @@ public: std::vector& stream_list() { return streams_; } -private: + private: std::vector streams_; }; diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 4f54b262c8..5543643f96 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -42,6 +42,8 @@ add_subdirectory(g++) add_subdirectory(module) add_subdirectory(channelDescriptor) add_subdirectory(executionControl) +add_subdirectory(vector_types) +add_subdirectory(atomics) add_subdirectory(p2p) add_subdirectory(gcc) @@ -52,4 +54,3 @@ add_subdirectory(clock) add_subdirectory(vulkan_interop) add_subdirectory(gl_interop) # Disabled on NVIDIA due to defect - EXSWHTEC-246 endif() -add_subdirectory(vector_types) diff --git a/catch/unit/atomics/CMakeLists.txt b/catch/unit/atomics/CMakeLists.txt new file mode 100644 index 0000000000..e9712c53a5 --- /dev/null +++ b/catch/unit/atomics/CMakeLists.txt @@ -0,0 +1,47 @@ +# 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. + +set(TEST_SRC + atomicExch.cc + atomicExch_system.cc +) + +if(HIP_PLATFORM MATCHES "nvidia") + set_source_files_properties(atomicExch_system.cc PROPERTIES COMPILE_FLAGS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") + hip_add_exe_to_target(NAME AtomicsTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS "nvrtc -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") +elseif(HIP_PLATFORM MATCHES "amd") + hip_add_exe_to_target(NAME AtomicsTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS hiprtc) +endif() + +add_test(NAME Unit_atomicExch_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + atomicExch_negative_kernels.cc 40) + +add_test(NAME Unit_atomicExch_system_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + atomicExch_system_negative_kernels.cc 40) diff --git a/catch/unit/atomics/atomicExch.cc b/catch/unit/atomics/atomicExch.cc new file mode 100644 index 0000000000..5cc8bf19f4 --- /dev/null +++ b/catch/unit/atomics/atomicExch.cc @@ -0,0 +1,199 @@ +/* +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" +#include "atomicExch_negative_kernels_rtc.hh" + +/** + * @addtogroup atomicExch atomicExch + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Executes a kernel wherein all threads will perform an atomic exchange in the same(compile + * time deducible) memory location. Each thread will exchange its own grid wide linear index + 1 + * 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]. + * + * - The test is run for: + * - All overloads of atomicExch + * - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory + * - Exchange memory located in shared memory + * - Several grid and block dimension combinations(only one block is used for shared memory) + * Test source + * ------------------------ + * - unit/atomics/atomicExch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_atomicExch_Positive_Same_Address_Compile_Time", "", int, unsigned int, + unsigned long, unsigned long long, float, double) { + for (auto current = 0; current < cmd_options.iterations; ++current) { + AtomicExchSameAddressTest(); + } +} + + +/** + * 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 + * - Several grid and block dimension combinations(only one block is used for shared memory) + * Test source + * ------------------------ + * - unit/atomics/atomicExch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_atomicExch_Positive", "", 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 kernel two times concurrently 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 + * - Several grid and block dimension combinations + * Test source + * ------------------------ + * - unit/atomics/atomicExch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_atomicExch_Positive_Multi_Kernel", "", 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) { + AtomicExchSingleDeviceMultipleKernelTest(2, 1, + sizeof(TestType)); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + AtomicExchSingleDeviceMultipleKernelTest(2, warp_size, + sizeof(TestType)); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + AtomicExchSingleDeviceMultipleKernelTest(2, warp_size, + cache_line_size); + } + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass combinations of arguments of invalid types for all overloads of + * atomicExch + * Test source + * ------------------------ + * - unit/atomics/atomicExch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_atomicExch_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + const auto program_source = GENERATE(kAtomicExchInt, kAtomicExchUnsignedInt, kAtomicExchULL, + kAtomicExchFloat, kAtomicExchDouble); + HIPRTC_CHECK( + hiprtcCreateProgram(&program, program_source, "atomicExch_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + 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 expected_error_count{8}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} diff --git a/catch/unit/atomics/atomicExch_common.hh b/catch/unit/atomics/atomicExch_common.hh new file mode 100644 index 0000000000..0888b17955 --- /dev/null +++ b/catch/unit/atomics/atomicExch_common.hh @@ -0,0 +1,379 @@ +/* +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 + +enum class AtomicScopes { device, system }; + +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); + } +} + +template +__global__ void atomic_exch_kernel_compile_time(T* const global_mem, T* const old_vals) { + __shared__ T shared_mem; + + const auto tid = cooperative_groups::this_grid().thread_rank(); + + T* 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] = perform_atomic_exch(mem, static_cast(tid + 1)); + + if constexpr (use_shared_mem) { + __syncthreads(); + if (tid == 0) global_mem[0] = mem[0]; + } +} + +template +__host__ __device__ T* pitched_offset(T* const ptr, const unsigned int pitch, + const unsigned int idx) { + const auto byte_ptr = reinterpret_cast(ptr); + return reinterpret_cast(byte_ptr + idx * pitch); +} + +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[]; + + const auto tid = cooperative_groups::this_grid().thread_rank(); + + T* const mem = use_shared_mem ? reinterpret_cast(shared_mem) : global_mem; + + if constexpr (use_shared_mem) { + if (tid < width) { + const auto target = pitched_offset(mem, pitch, tid); + *target = *pitched_offset(global_mem, pitch, tid); + }; + __syncthreads(); + } + + old_vals[tid] = perform_atomic_exch(pitched_offset(mem, pitch, tid % width), + base_val + static_cast(tid + width)); + + if constexpr (use_shared_mem) { + __syncthreads(); + if (tid < width) { + const auto target = pitched_offset(global_mem, pitch, tid); + *target = *pitched_offset(mem, pitch, tid); + }; + } +} + + +template +void AtomicExchSameAddress(const dim3 blocks, const dim3 threads, const LinearAllocs alloc_type) { + LinearAllocGuard mem_dev(alloc_type, sizeof(TestType)); + + const auto thread_count = blocks.x * blocks.y * blocks.z * threads.x * threads.y * threads.z; + const auto old_vals_alloc_size = thread_count * sizeof(TestType); + LinearAllocGuard old_vals_dev(LinearAllocs::hipMalloc, old_vals_alloc_size); + std::vector old_vals(thread_count + 1); + + + HIP_CHECK(hipMemset(mem_dev.ptr(), 0, sizeof(TestType))); + atomic_exch_kernel_compile_time + <<>>(mem_dev.ptr(), old_vals_dev.ptr()); + HIP_CHECK( + hipMemcpy(old_vals.data(), old_vals_dev.ptr(), old_vals_alloc_size, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(old_vals.data() + thread_count, mem_dev.ptr(), sizeof(TestType), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + // Every thread will exchange its grid-wide linear id into a target location within mem_dev, + // receiving back the value previously present therein. This previous value is written to + // old_vals_dev. + // old_vals_dev will not contain values that the final scheduled warp exchanged into mem_dev, but + // mem_dev obviously will. + // Given that mem_dev initially contains values in the range [0, width) and that the maximum value + // the final thread shall write is thread_count + width - 1, presuming correct operation of + // atomicExch, the union of mem_dev and old_vals_dev shall contain values in the range + //[0, thread_count + width) + std::sort(old_vals.begin(), old_vals.end()); + for (auto i = 0u; i < old_vals.size(); ++i) { + REQUIRE(i == old_vals[i]); + } +} + +template void AtomicExchSameAddressTest() { + const auto threads = GENERATE(dim3(1024), dim3(1023), dim3(511), dim3(17), dim3(31)); + + SECTION("Global memory") { + const auto blocks = GENERATE(dim3(20)); + using LA = LinearAllocs; + const auto allocation_type = + GENERATE(LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister); + AtomicExchSameAddress(blocks, threads, allocation_type); + } + + SECTION("Shared memory") { + const auto blocks = dim3(1); + AtomicExchSameAddress(blocks, threads, + LinearAllocs::hipMalloc); + } +} + +struct AtomicExchParams { + dim3 blocks; + dim3 threads; + unsigned int num_devices = 1u; + unsigned int kernel_count = 1u; + unsigned int width = 1u; + unsigned int pitch = 0u; + unsigned int host_thread_count = 0u; + LinearAllocs alloc_type; +}; + + +template +class AtomicExchCRTP { + public: + void run(const AtomicExchParams& p) const { + const auto thread_count = + p.blocks.x * p.blocks.y * p.blocks.z * p.threads.x * p.threads.y * p.threads.z; + + const auto old_vals_alloc_size = p.kernel_count * thread_count * sizeof(T); + 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); + + const auto host_iters_per_thread = + std::max(p.num_devices * p.kernel_count * thread_count / 20, p.width); + + std::vector old_vals(p.num_devices * p.kernel_count * thread_count + p.width + + p.host_thread_count * host_iters_per_thread); + std::iota(old_vals.begin(), old_vals.begin() + p.width, 0); + + HIP_CHECK(hipMemcpy2D(mem_dev.ptr(), p.pitch, old_vals.data(), sizeof(T), sizeof(T), p.width, + hipMemcpyHostToDevice)); + + const auto shared_mem_size = use_shared_mem ? mem_alloc_size : 0u; + for (auto i = 0u; i < p.num_devices; ++i) { + const auto device_offset = i * p.kernel_count * thread_count; + for (auto j = 0u; j < p.kernel_count; ++j) { + const auto& stream = streams[i * p.kernel_count + j].stream(); + const auto kern_offset = j * thread_count; + const auto old_vals = old_vals_devs[i].ptr() + kern_offset; + CastToDerived().LaunchKernel(shared_mem_size, stream, mem_dev.ptr(), old_vals, + device_offset + kern_offset, p); + } + } + + PerformHostAtomicExchange(p.host_thread_count, host_iters_per_thread, mem_dev.host_ptr(), + old_vals.data(), p); + + for (auto i = 0u; i < p.num_devices; ++i) { + const auto device_offset = i * p.kernel_count * thread_count; + HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(), + old_vals_alloc_size, hipMemcpyDeviceToHost)); + } + HIP_CHECK(hipMemcpy2D(old_vals.data() + p.num_devices * p.kernel_count * thread_count, + sizeof(T), mem_dev.ptr(), p.pitch, sizeof(T), p.width, + hipMemcpyDeviceToHost)); + + CastToDerived().ValidateResults(old_vals); + } + + private: + const Derived& CastToDerived() const { return static_cast(*this); } + + static void HostAtomicExchange(const unsigned int iterations, T* mem, T* const old_vals, + const unsigned int width, const unsigned pitch, T base_val) { + for (auto i = 0u; i < iterations; ++i) { + T new_val = base_val + static_cast(i); + T old_val; + __atomic_exchange(pitched_offset(mem, pitch, i % width), &new_val, &old_val, + __ATOMIC_RELAXED); + old_vals[i] = old_val; + } + } + + void PerformHostAtomicExchange(const unsigned int thread_count, const unsigned int iterations, + T* mem, T* const old_vals, const AtomicExchParams& p) const { + if (thread_count == 0) { + return; + } + const auto dev_threads = + p.blocks.x * p.blocks.y * p.blocks.z * p.threads.x * p.threads.y * p.threads.z; + const auto host_base_val = p.num_devices * p.kernel_count * dev_threads + p.width; + + std::vector threads; + for (auto i = 0u; i < thread_count; ++i) { + const auto thread_base_val = host_base_val + i * iterations; + threads.push_back(std::thread(HostAtomicExchange, iterations, mem, old_vals + thread_base_val, + p.width, p.pitch, thread_base_val)); + } + + for (auto& th : threads) { + th.join(); + } + } +}; + +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); + } + + void ValidateResults(std::vector& old_vals) const { + std::sort(old_vals.begin(), old_vals.end()); + for (auto i = 0u; i < old_vals.size(); ++i) { + REQUIRE(i == old_vals[i]); + } + } +}; + +inline dim3 GenerateAtomicExchThreadDimensions() { return GENERATE(dim3(16), dim3(1024)); } + +inline dim3 GenerateAtomicExchBlockDimensions() { + 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 AtomicExchSingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) { + AtomicExchParams params; + params.num_devices = 1; + params.kernel_count = 1; + params.threads = GenerateAtomicExchThreadDimensions(); + params.width = width; + params.pitch = pitch; + + SECTION("Global memory") { + 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); + } + } + } + + SECTION("Shared memory") { + params.blocks = dim3(1); + params.alloc_type = LinearAllocs::hipMalloc; + AtomicExch().run(params); + } +} + +template +void AtomicExchSingleDeviceMultipleKernelTest(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; + } + + AtomicExchParams params; + params.num_devices = 1; + params.kernel_count = kernel_count; + params.blocks = GenerateAtomicExchBlockDimensions(); + params.threads = GenerateAtomicExchThreadDimensions(); + 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)) { + AtomicExch().run(params); + } + } +} + +template +void AtomicExchMultipleDeviceMultipleKernelAndHostTest(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; + } + } + } + + AtomicExchParams params; + params.num_devices = num_devices; + params.kernel_count = kernel_count; + params.blocks = GenerateAtomicExchBlockDimensions(); + params.threads = GenerateAtomicExchThreadDimensions(); + 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)) { + AtomicExch().run(params); + } + } +} \ No newline at end of file diff --git a/catch/unit/atomics/atomicExch_negative_kernels.cc b/catch/unit/atomics/atomicExch_negative_kernels.cc new file mode 100644 index 0000000000..083d27747f --- /dev/null +++ b/catch/unit/atomics/atomicExch_negative_kernels.cc @@ -0,0 +1,94 @@ +/* +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 + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +/*int atomicExch(int*, int)*/ +__global__ void atomicExch_int_n1(int* p, int v) { atomicExch(p, p); } +__global__ void atomicExch_int_n2(int* p, int v) { atomicExch(&p, v); } +__global__ void atomicExch_int_n3(char* p, int v) { atomicExch(p, v); } +__global__ void atomicExch_int_n4(short* p, int v) { atomicExch(p, v); } +__global__ void atomicExch_int_n5(long* p, int v) { atomicExch(p, v); } +__global__ void atomicExch_int_n6(long long* p, int v) { atomicExch(p, v); } +__global__ void atomicExch_int_n7(Dummy* p, int v) { atomicExch(p, v); } +__global__ void atomicExch_int_n8(int* p, Dummy v) { atomicExch(p, v); } + +/*unsigned int atomicExch(unsigned int*, unsigned int)*/ +__global__ void atomicExch_unsigned_int_n1(unsigned int* p, unsigned int v) { atomicExch(p, p); } +__global__ void atomicExch_unsigned_int_n2(unsigned int* p, unsigned int v) { atomicExch(&p, v); } +__global__ void atomicExch_unsigned_int_n3(char* p, unsigned int v) { atomicExch(p, v); } +__global__ void atomicExch_unsigned_int_n4(short* p, unsigned int v) { atomicExch(p, v); } +__global__ void atomicExch_unsigned_int_n5(long* p, unsigned int v) { atomicExch(p, v); } +__global__ void atomicExch_unsigned_int_n6(long long* p, unsigned int v) { atomicExch(p, v); } +__global__ void atomicExch_unsigned_int_n7(Dummy* p, unsigned int v) { atomicExch(p, v); } +__global__ void atomicExch_unsigned_int_n8(unsigned int* p, Dummy v) { atomicExch(p, v); } + +// /*unsigned long long atomicExch(unsigned long long*, unsigned long long)*/ +__global__ void atomicExch_unsigned_long_long_n1(unsigned long long* p, unsigned long long v) { + atomicExch(p, p); +} +__global__ void atomicExch_unsigned_long_long_n2(unsigned long long* p, unsigned long long v) { + atomicExch(&p, v); +} +__global__ void atomicExch_unsigned_long_long_n3(char* p, unsigned long long v) { + atomicExch(p, v); +} +__global__ void atomicExch_unsigned_long_long_n4(short* p, unsigned long long v) { + atomicExch(p, v); +} +__global__ void atomicExch_unsigned_long_long_n5(long* p, unsigned long long v) { + atomicExch(p, v); +} +__global__ void atomicExch_unsigned_long_long_n6(long long* p, unsigned long long v) { + atomicExch(p, v); +} +__global__ void atomicExch_unsigned_long_long_n7(Dummy* p, unsigned long long v) { + atomicExch(p, v); +} +__global__ void atomicExch_unsigned_long_long_n8(unsigned long long* p, Dummy v) { + atomicExch(p, v); +} + +// /*float atomicExch(float*, float)*/ +__global__ void atomicExch_float_n1(float* p, float v) { atomicExch(p, p); } +__global__ void atomicExch_float_n2(float* p, float v) { atomicExch(&p, v); } +__global__ void atomicExch_float_n3(char* p, float v) { atomicExch(p, v); } +__global__ void atomicExch_float_n4(short* p, float v) { atomicExch(p, v); } +__global__ void atomicExch_float_n5(long* p, float v) { atomicExch(p, v); } +__global__ void atomicExch_float_n6(long long* p, float v) { atomicExch(p, v); } +__global__ void atomicExch_float_n7(Dummy* p, float v) { atomicExch(p, v); } +__global__ void atomicExch_float_n8(float* p, Dummy v) { atomicExch(p, v); } + +// /*double atomicExch(double*, double)*/ +__global__ void atomicExch_double_n1(double* p, double v) { atomicExch(p, p); } +__global__ void atomicExch_double_n2(double* p, double v) { atomicExch(&p, v); } +__global__ void atomicExch_double_n3(char* p, double v) { atomicExch(p, v); } +__global__ void atomicExch_double_n4(short* p, double v) { atomicExch(p, v); } +__global__ void atomicExch_double_n5(long* p, double v) { atomicExch(p, v); } +__global__ void atomicExch_double_n6(long long* p, double v) { atomicExch(p, v); } +__global__ void atomicExch_double_n7(Dummy* p, double v) { atomicExch(p, v); } +__global__ void atomicExch_double_n8(double* p, Dummy v) { atomicExch(p, v); } \ No newline at end of file diff --git a/catch/unit/atomics/atomicExch_negative_kernels_rtc.hh b/catch/unit/atomics/atomicExch_negative_kernels_rtc.hh new file mode 100644 index 0000000000..01387bba61 --- /dev/null +++ b/catch/unit/atomics/atomicExch_negative_kernels_rtc.hh @@ -0,0 +1,124 @@ +/* +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 + +static constexpr auto kAtomicExchInt{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_int_n1(int* p, int v) { atomicExch(p, p); } + __global__ void atomicExch_int_n2(int* p, int v) { atomicExch(&p, v); } + __global__ void atomicExch_int_n3(char* p, int v) { atomicExch(p, v); } + __global__ void atomicExch_int_n4(short* p, int v) { atomicExch(p, v); } + __global__ void atomicExch_int_n5(long* p, int v) { atomicExch(p, v); } + __global__ void atomicExch_int_n6(long long* p, int v) { atomicExch(p, v); } + __global__ void atomicExch_int_n7(Dummy* p, int v) { atomicExch(p, v); } + __global__ void atomicExch_int_n8(int* p, Dummy v) { atomicExch(p, v); } + )"}; + +static constexpr auto kAtomicExchUnsignedInt{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_unsigned_int_n1(unsigned int* p, unsigned int v) { atomicExch(p, p); } + __global__ void atomicExch_unsigned_int_n2(unsigned int* p, unsigned int v) { atomicExch(&p, v); } + __global__ void atomicExch_unsigned_int_n3(char* p, unsigned int v) { atomicExch(p, v); } + __global__ void atomicExch_unsigned_int_n4(short* p, unsigned int v) { atomicExch(p, v); } + __global__ void atomicExch_unsigned_int_n5(long* p, unsigned int v) { atomicExch(p, v); } + __global__ void atomicExch_unsigned_int_n6(long long* p, unsigned int v) { atomicExch(p, v); } + __global__ void atomicExch_unsigned_int_n7(Dummy* p, unsigned int v) { atomicExch(p, v); } + __global__ void atomicExch_unsigned_int_n8(unsigned int* p, Dummy v) { atomicExch(p, v); } + )"}; + +static constexpr auto kAtomicExchULL{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_unsigned_long_long_n1(unsigned long long* p, unsigned long long v) { + atomicExch(p, p); + } + __global__ void atomicExch_unsigned_long_long_n2(unsigned long long* p, unsigned long long v) { + atomicExch(&p, v); + } + __global__ void atomicExch_unsigned_long_long_n3(char* p, unsigned long long v) { + atomicExch(p, v); + } + __global__ void atomicExch_unsigned_long_long_n4(short* p, unsigned long long v) { + atomicExch(p, v); + } + __global__ void atomicExch_unsigned_long_long_n5(long* p, unsigned long long v) { + atomicExch(p, v); + } + __global__ void atomicExch_unsigned_long_long_n6(long long* p, unsigned long long v) { + atomicExch(p, v); + } + __global__ void atomicExch_unsigned_long_long_n7(Dummy* p, unsigned long long v) { + atomicExch(p, v); + } + __global__ void atomicExch_unsigned_long_long_n8(unsigned long long* p, Dummy v) { + atomicExch(p, v); + } + )"}; + +static constexpr auto kAtomicExchFloat{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_float_n1(float* p, float v) { atomicExch(p, p); } + __global__ void atomicExch_float_n2(float* p, float v) { atomicExch(&p, v); } + __global__ void atomicExch_float_n3(char* p, float v) { atomicExch(p, v); } + __global__ void atomicExch_float_n4(short* p, float v) { atomicExch(p, v); } + __global__ void atomicExch_float_n5(long* p, float v) { atomicExch(p, v); } + __global__ void atomicExch_float_n6(long long* p, float v) { atomicExch(p, v); } + __global__ void atomicExch_float_n7(Dummy* p, float v) { atomicExch(p, v); } + __global__ void atomicExch_float_n8(float* p, Dummy v) { atomicExch(p, v); } + )"}; + +static constexpr auto kAtomicExchDouble{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_double_n1(double* p, double v) { atomicExch(p, p); } + __global__ void atomicExch_double_n2(double* p, double v) { atomicExch(&p, v); } + __global__ void atomicExch_double_n3(char* p, double v) { atomicExch(p, v); } + __global__ void atomicExch_double_n4(short* p, double v) { atomicExch(p, v); } + __global__ void atomicExch_double_n5(long* p, double v) { atomicExch(p, v); } + __global__ void atomicExch_double_n6(long long* p, double v) { atomicExch(p, v); } + __global__ void atomicExch_double_n7(Dummy* p, double v) { atomicExch(p, v); } + __global__ void atomicExch_double_n8(double* p, Dummy v) { atomicExch(p, v); } + )"}; \ No newline at end of file diff --git a/catch/unit/atomics/atomicExch_system.cc b/catch/unit/atomics/atomicExch_system.cc new file mode 100644 index 0000000000..5b929b056e --- /dev/null +++ b/catch/unit/atomics/atomicExch_system.cc @@ -0,0 +1,220 @@ +/* +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" +#include "atomicExch_system_negative_kernels_rtc.hh" + +/** + * @addtogroup atomicExch_system atomicExch_system + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Executes a kernel two times concurrently on two devices 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_system + * - hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory + * - Several grid and block dimension combinations + * Test source + * ------------------------ + * - unit/atomics/atomicExch_system.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_atomicExch_system_Positive_Peer_GPUs", "", 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) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(2, 2, 1, sizeof(TestType)); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(2, 2, warp_size, + sizeof(TestType)); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(2, 2, warp_size, cache_line_size); + } + } +} + +/** + * Test Description + * ------------------------ + * - Executes a 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. While the kernel is running, the host performs atomic exchanges, in 4 + * threads, into the same memory location(s). Once complete, the union of output array, exchange + * memory, and host output is validated to contain all values in the range [0, number_of_threads + + * number_of_exchange_memory_slots + number_of_host_iterations). 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_system + * - hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory + * - Several grid and block dimension combinations + * Test source + * ------------------------ + * - unit/atomics/atomicExch_system.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_atomicExch_system_Positive_Host_And_GPU", "", 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) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(1, 1, 1, sizeof(TestType), 4); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(1, 1, warp_size, sizeof(TestType), + 4); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(1, 1, warp_size, cache_line_size, + 4); + } + } +} + +/** + * Test Description + * ------------------------ + * - Executes a kernel two times concurrently on two devices 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. While the kernels are running, the + * host performs atomic exchanges, in 4 threads, into the same memory location(s). Once complete, + * the union of output array, exchange memory, and host output is validated to contain all values in + * the range [0, number_of_threads + number_of_exchange_memory_slots + number_of_host_iterations). + * 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_system + * - hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory + * - Several grid and block dimension combinations + * Test source + * ------------------------ + * - unit/atomics/atomicExch_system.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_atomicExch_system_Positive_Host_And_Peer_GPUs", "", 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) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(2, 2, 1, sizeof(TestType), 4); + } + + DYNAMIC_SECTION("Adjacent addresses " << current) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(2, 2, warp_size, sizeof(TestType), + 4); + } + + DYNAMIC_SECTION("Scattered addresses " << current) { + AtomicExchMultipleDeviceMultipleKernelAndHostTest(2, 2, warp_size, cache_line_size, + 4); + } + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass combinations of arguments of invalid types for all overloads of + * atomicExch_system + * Test source + * ------------------------ + * - unit/atomics/atomicExch_system.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_atomicExch_system_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + const auto program_source = + GENERATE(kAtomicExchSystemInt, kAtomicExchSystemUnsignedInt, kAtomicExchSystemULL, + kAtomicExchSystemFloat, kAtomicExchSystemDouble); + HIPRTC_CHECK( + hiprtcCreateProgram(&program, program_source, "atomicExch_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + 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 expected_error_count{8}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} diff --git a/catch/unit/atomics/atomicExch_system_negative_kernels.cc b/catch/unit/atomics/atomicExch_system_negative_kernels.cc new file mode 100644 index 0000000000..15085f7da3 --- /dev/null +++ b/catch/unit/atomics/atomicExch_system_negative_kernels.cc @@ -0,0 +1,112 @@ +/* +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 + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +/*int atomicExch_system(int*, int)*/ +__global__ void atomicExch_system_int_n1(int* p, int v) { atomicExch_system(p, p); } +__global__ void atomicExch_system_int_n2(int* p, int v) { atomicExch_system(&p, v); } +__global__ void atomicExch_system_int_n3(char* p, int v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_int_n4(short* p, int v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_int_n5(long* p, int v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_int_n6(long long* p, int v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_int_n7(Dummy* p, int v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_int_n8(int* p, Dummy v) { atomicExch_system(p, v); } + +/*unsigned int atomicExch_system(unsigned int*, unsigned int)*/ +__global__ void atomicExch_system_unsigned_int_n1(unsigned int* p, unsigned int v) { + atomicExch_system(p, p); +} +__global__ void atomicExch_system_unsigned_int_n2(unsigned int* p, unsigned int v) { + atomicExch_system(&p, v); +} +__global__ void atomicExch_system_unsigned_int_n3(char* p, unsigned int v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_int_n4(short* p, unsigned int v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_int_n5(long* p, unsigned int v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_int_n6(long long* p, unsigned int v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_int_n7(Dummy* p, unsigned int v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_int_n8(unsigned int* p, Dummy v) { + atomicExch_system(p, v); +} + +// /*unsigned long long atomicExch_system(unsigned long long*, unsigned long long)*/ +__global__ void atomicExch_system_unsigned_long_long_n1(unsigned long long* p, + unsigned long long v) { + atomicExch_system(p, p); +} +__global__ void atomicExch_system_unsigned_long_long_n2(unsigned long long* p, + unsigned long long v) { + atomicExch_system(&p, v); +} +__global__ void atomicExch_system_unsigned_long_long_n3(char* p, unsigned long long v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_long_long_n4(short* p, unsigned long long v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_long_long_n5(long* p, unsigned long long v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_long_long_n6(long long* p, unsigned long long v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_long_long_n7(Dummy* p, unsigned long long v) { + atomicExch_system(p, v); +} +__global__ void atomicExch_system_unsigned_long_long_n8(unsigned long long* p, Dummy v) { + atomicExch_system(p, v); +} + +// /*float atomicExch_system(float*, float)*/ +__global__ void atomicExch_system_float_n1(float* p, float v) { atomicExch_system(p, p); } +__global__ void atomicExch_system_float_n2(float* p, float v) { atomicExch_system(&p, v); } +__global__ void atomicExch_system_float_n3(char* p, float v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_float_n4(short* p, float v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_float_n5(long* p, float v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_float_n6(long long* p, float v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_float_n7(Dummy* p, float v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_float_n8(float* p, Dummy v) { atomicExch_system(p, v); } + +// /*double atomicExch_system(double*, double)*/ +__global__ void atomicExch_system_double_n1(double* p, double v) { atomicExch_system(p, p); } +__global__ void atomicExch_system_double_n2(double* p, double v) { atomicExch_system(&p, v); } +__global__ void atomicExch_system_double_n3(char* p, double v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_double_n4(short* p, double v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_double_n5(long* p, double v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_double_n6(long long* p, double v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_double_n7(Dummy* p, double v) { atomicExch_system(p, v); } +__global__ void atomicExch_system_double_n8(double* p, Dummy v) { atomicExch_system(p, v); } \ No newline at end of file diff --git a/catch/unit/atomics/atomicExch_system_negative_kernels_rtc.hh b/catch/unit/atomics/atomicExch_system_negative_kernels_rtc.hh new file mode 100644 index 0000000000..69fc31fb35 --- /dev/null +++ b/catch/unit/atomics/atomicExch_system_negative_kernels_rtc.hh @@ -0,0 +1,142 @@ +/* +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 + +static constexpr auto kAtomicExchSystemInt{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_system_int_n1(int* p, int v) { atomicExch_system(p, p); } + __global__ void atomicExch_system_int_n2(int* p, int v) { atomicExch_system(&p, v); } + __global__ void atomicExch_system_int_n3(char* p, int v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_int_n4(short* p, int v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_int_n5(long* p, int v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_int_n6(long long* p, int v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_int_n7(Dummy* p, int v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_int_n8(int* p, Dummy v) { atomicExch_system(p, v); } + )"}; + +static constexpr auto kAtomicExchSystemUnsignedInt{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_system_unsigned_int_n1(unsigned int* p, unsigned int v) { + atomicExch_system(p, p); + } + __global__ void atomicExch_system_unsigned_int_n2(unsigned int* p, unsigned int v) { + atomicExch_system(&p, v); + } + __global__ void atomicExch_system_unsigned_int_n3(char* p, unsigned int v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_int_n4(short* p, unsigned int v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_int_n5(long* p, unsigned int v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_int_n6(long long* p, unsigned int v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_int_n7(Dummy* p, unsigned int v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_int_n8(unsigned int* p, Dummy v) { + atomicExch_system(p, v); + } + )"}; + +static constexpr auto kAtomicExchSystemULL{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_system_unsigned_long_long_n1(unsigned long long* p, + unsigned long long v) { + atomicExch_system(p, p); + } + __global__ void atomicExch_system_unsigned_long_long_n2(unsigned long long* p, + unsigned long long v) { + atomicExch_system(&p, v); + } + __global__ void atomicExch_system_unsigned_long_long_n3(char* p, unsigned long long v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_long_long_n4(short* p, unsigned long long v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_long_long_n5(long* p, unsigned long long v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_long_long_n6(long long* p, unsigned long long v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_long_long_n7(Dummy* p, unsigned long long v) { + atomicExch_system(p, v); + } + __global__ void atomicExch_system_unsigned_long_long_n8(unsigned long long* p, Dummy v) { + atomicExch_system(p, v); + } + )"}; + +static constexpr auto kAtomicExchSystemFloat{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_system_float_n1(float* p, float v) { atomicExch_system(p, p); } + __global__ void atomicExch_system_float_n2(float* p, float v) { atomicExch_system(&p, v); } + __global__ void atomicExch_system_float_n3(char* p, float v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_float_n4(short* p, float v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_float_n5(long* p, float v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_float_n6(long long* p, float v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_float_n7(Dummy* p, float v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_float_n8(float* p, Dummy v) { atomicExch_system(p, v); } + )"}; + +static constexpr auto kAtomicExchSystemDouble{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void atomicExch_system_double_n1(double* p, double v) { atomicExch_system(p, p); } + __global__ void atomicExch_system_double_n2(double* p, double v) { atomicExch_system(&p, v); } + __global__ void atomicExch_system_double_n3(char* p, double v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_double_n4(short* p, double v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_double_n5(long* p, double v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_double_n6(long long* p, double v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_double_n7(Dummy* p, double v) { atomicExch_system(p, v); } + __global__ void atomicExch_system_double_n8(double* p, Dummy v) { atomicExch_system(p, v); } + )"}; \ No newline at end of file diff --git a/catch/unit/compileAndCaptureOutput.py b/catch/unit/compileAndCaptureOutput.py new file mode 100644 index 0000000000..c06feb5858 --- /dev/null +++ b/catch/unit/compileAndCaptureOutput.py @@ -0,0 +1,107 @@ +# Copyright (c) 2022 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. + +import subprocess +import sys +import unittest + +class CompileAndCapture(unittest.TestCase): + path = None + expected_error_count = 0 + expected_warning_count = 0 + hip_path = None + file = None + error_string = None + warning_string = None + platform = None + + def setUp(self): + self.error_string = 'error:' + self.warning_string = 'warning:' + self.assertFalse(self.hip_path == None) + self.assertFalse(self.path == None) + self.assertFalse(self.file == None) + self.assertTrue(self.platform == 'amd' or self.platform == 'nvidia') + + def test(self): + compiler_args = [ + self.hip_path + '/bin/hipcc', + '-I' + self.path + '/../../external/Catch2', + '-I' + self.path + '/../../include', + '-I' + self.path + '/../../external/picojson', + '-c', + self.path + '/' + self.file, + ] + # HIP compiler on AMD platforms has limit of 20 errors, and some negative + # test cases expect that more errors are detected. + if (self.platform == 'amd'): + compiler_args.append('-ferror-limit=100') + compiler_output = subprocess.run(compiler_args, stderr=subprocess.PIPE) + # Get the compiler output in the stdout if -V flag is raised during ctest invocation. + compiler_stderr = compiler_output.stderr.decode('UTF-8') + print(compiler_stderr) + + error_count = compiler_stderr.count(self.error_string) + if self.expected_error_count < 0: + self.assertGreater(error_count, 0) + else: + self.assertEqual(error_count, self.expected_error_count) + + warning_count = compiler_stderr.count(self.warning_string) + if self.expected_warning_count < 0: + self.assertGreater(warning_count, 0) + else: + self.assertEqual(warning_count, self.expected_warning_count) + +if __name__ == '__main__': + try: + CompileAndCapture.path = sys.argv[1] + except IndexError: + CompileAndCapture.path = None + + try: + CompileAndCapture.platform = sys.argv[2] + except IndexError: + CompileAndCapture.platform = None + + try: + CompileAndCapture.hip_path = sys.argv[3] + except IndexError: + CompileAndCapture.hip_path = None + + try: + CompileAndCapture.file = sys.argv[4] + except IndexError: + CompileAndCapture.file = None + + try: + CompileAndCapture.expected_error_count = int(sys.argv[5]) + except IndexError: + CompileAndCapture.expected_error_count = 0 + + try: + CompileAndCapture.expected_warning_count = int(sys.argv[6]) + except IndexError: + CompileAndCapture.expected_warning_count = 0 + + # Unittest looks at the same argv's as the __main__ and doesn't know how + # to handle arguments other than the executable (0). Therefore passing only + # executable as the argv for unittest module. + unittest.main(argv=[sys.argv[0]])