AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

EXSWHTEC-262 - Introduce build dependencies for atomic arithmetic operations #180

Procházet zdrojové kódy 
Change-Id: Ic2293f273158613312dc39914da087fc0a462242
Tento commit je obsažen v:
Mirza Halilčević
2023-12-28 15:46:50 +01:00
odevzdal Rakesh Roy
rodič c6a6475743
revize 2eb17f0696
3 změnil soubory, kde provedl 161 přidání a 43 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
catch/include/hip_test_defgroups.hh
+115
Zobrazit soubor
@@ -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
* @{
catch/include/resource_guards.hh
+44 -41
Zobrazit soubor
@@ -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 <typename T> 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 <typename T> 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 <typename T> 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<void>(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;
}
catch/unit/atomics/CMakeLists.txt
+2 -2
Zobrazit soubor
@@ -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}