From 2eb17f0696c3ff9f45d15d1d74a3c40f84fbeb18 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 15:46:50 +0100 Subject: [PATCH] EXSWHTEC-262 - Introduce build dependencies for atomic arithmetic operations #180 Change-Id: Ic2293f273158613312dc39914da087fc0a462242 --- catch/include/hip_test_defgroups.hh | 115 ++++++++++++++++++++++++++++ catch/include/resource_guards.hh | 85 ++++++++++---------- catch/unit/atomics/CMakeLists.txt | 4 +- 3 files changed, 161 insertions(+), 43 deletions(-) diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 680dfa8a04..e108f296fd 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -172,6 +172,121 @@ THE SOFTWARE. * @} */ +/** + * @defgroup AtomicsTest Device Atomics + * @{ + * This section describes tests for the Device Atomic APIs. + */ + +/** + * @addtogroup atomicAdd atomicAdd + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Compiles atomicAdd 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_atomicAdd_Negative_Parameters") {} +/** + * End doxygen group atomicAdd. + * @} + */ + +/** + * @addtogroup atomicSub atomicSub + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Compiles atomicSub 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_atomicSub_Negative_Parameters") {} +/** + * End doxygen group atomicSub. + * @} + */ + +/** + * @addtogroup atomicInc atomicInc + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Compiles atomicInc 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_atomicInc_Negative_Parameters") {} +/** + * End doxygen group atomicInc. + * @} + */ + +/** + * @addtogroup atomicDec atomicDec + * @{ + * @ingroup AtomicsTest + */ + +/** + * Test Description + * ------------------------ + * - Compiles atomicDec 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_atomicDec_Negative_Parameters") {} +/** + * End doxygen group atomicDec. + * @} + */ + +/** + * End doxygen group AtomicsTest. + * @} + */ + /** * @defgroup PrintfTest Printf API Management * @{ diff --git a/catch/include/resource_guards.hh b/catch/include/resource_guards.hh index 5f8f2cbee4..262b7b4437 100644 --- a/catch/include/resource_guards.hh +++ b/catch/include/resource_guards.hh @@ -35,15 +35,15 @@ enum class LinearAllocs { inline std::string to_string(const LinearAllocs allocation_type) { switch (allocation_type) { case LinearAllocs::malloc: - return "host pageable"; + return "malloc"; case LinearAllocs::mallocAndRegister: - return "registered"; + return "malloc + hipHostRegister"; case LinearAllocs::hipHostMalloc: - return "host pinned"; + return "hipHostMalloc"; case LinearAllocs::hipMalloc: - return "device malloc"; + return "hipMalloc"; case LinearAllocs::hipMallocManaged: - return "managed"; + return "hipMallocManaged"; default: return "unknown alloc type"; } @@ -83,24 +83,35 @@ template class LinearAllocGuard { LinearAllocGuard(const 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(LinearAllocGuard&& o) { *this = std::move(o); } LinearAllocGuard& operator=(LinearAllocGuard&& o) { - allocation_type_ = o.allocation_type_; - ptr_ = o.ptr_; - host_ptr_ = o.host_ptr_; + if (this != &o) { + dealloc(); - o.allocation_type_ = LinearAllocs::noAlloc; - o.ptr_ = nullptr; - o.host_ptr_ = nullptr; + 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; + } + + return *this; } - ~LinearAllocGuard() { + ~LinearAllocGuard() { dealloc(); } + + T* ptr() const { return ptr_; }; + T* host_ptr() const { return host_ptr_; } + + private: + LinearAllocs allocation_type_ = LinearAllocs::noAlloc; + T* ptr_ = nullptr; + T* host_ptr_ = nullptr; + + void dealloc() { // No Catch macros, don't want to possibly throw in the destructor if (ptr_ != nullptr) { switch (allocation_type_) { @@ -123,14 +134,6 @@ template class LinearAllocGuard { } } } - - T* ptr() const { return ptr_; }; - T* host_ptr() const { return host_ptr_; } - - private: - LinearAllocs allocation_type_ = LinearAllocs::noAlloc; - T* ptr_ = nullptr; - T* host_ptr_ = nullptr; }; template class LinearAllocGuardMultiDim { @@ -266,24 +269,24 @@ class StreamGuard { StreamGuard(const 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(StreamGuard&& o) { *this = std::move(o); } StreamGuard& operator=(StreamGuard&& o) { - stream_type_ = o.stream_type_; - flags_ = o.flags_; - priority_ = o.priority_; - stream_ = o.stream_; + if (this != &o) { + if (stream_type_ == Streams::created) { + static_cast(hipStreamDestroy(stream_)); + } - o.stream_type_ = Streams::nullstream; - o.flags_ = 0u; - o.priority_ = 0; - o.stream_ = nullptr; + 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; } diff --git a/catch/unit/atomics/CMakeLists.txt b/catch/unit/atomics/CMakeLists.txt index d8066a2f1a..2aef527440 100644 --- a/catch/unit/atomics/CMakeLists.txt +++ b/catch/unit/atomics/CMakeLists.txt @@ -24,11 +24,11 @@ set(TEST_SRC ) 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") + set_source_files_properties(atomicExch_system.cc PROPERTIES COMPILE_FLAGS "-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") + LINKER_LIBS "nvrtc -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}