AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

EXSWHTEC-265 - Implement Tests for Atomic Exchange (#196)

Przeglądaj źródła 
Change-Id: I4d332dcaad2d0a788b6711f5be05ca6d61a9e34b
This commit is contained in:
Mirza Halilcevic
2023-11-26 19:44:41 +00:00
zatwierdzone przez Maneesh Gupta
rodzic f05f97758e
commit f24d2dd934
15 zmienionych plików z 1566 dodań i 23 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
catch/hipTestMain/main.cc
+3 -3
Wyświetl plik
@@ -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
catch/include/cmd_options.hh
+2 -2
Wyświetl plik
@@ -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;
catch/include/hip_test_common.hh
+13
Wyświetl plik
@@ -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)); }
catch/include/hip_test_defgroups.hh
+41 -1
Wyświetl plik
@@ -20,6 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
// 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.
* @}
*/
catch/include/resource_guards.hh
+81 -16
Wyświetl plik
@@ -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 <typename T> 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 <typename T> class LinearAllocGuard {
case LinearAllocs::hipMallocManaged:
HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&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 <typename T> 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<void>(hipEventDestroy(e));
for (auto& e : events_) static_cast<void>(hipEventDestroy(e));
}
hipEvent_t& operator[](int index) { return events_[index]; }
@@ -255,21 +320,21 @@ public:
std::vector<hipEvent_t>& event_list() { return events_; }
private:
private:
std::vector<hipEvent_t> 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<void>(hipStreamDestroy(s));
for (auto& s : streams_) static_cast<void>(hipStreamDestroy(s));
}
hipStream_t& operator[](int index) { return streams_[index]; }
@@ -278,6 +343,6 @@ public:
std::vector<hipStream_t>& stream_list() { return streams_; }
private:
private:
std::vector<hipStream_t> streams_;
};
catch/unit/CMakeLists.txt
+2 -1
Wyświetl plik
@@ -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)
catch/unit/atomics/CMakeLists.txt
+47
Wyświetl plik
@@ -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)
catch/unit/atomics/atomicExch.cc
+199
Wyświetl plik
@@ -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<TestType, AtomicScopes::device>();
}
}
/**
* 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<TestType, AtomicScopes::device>(1, sizeof(TestType));
}
DYNAMIC_SECTION("Adjacent addresses " << current) {
AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::device>(warp_size,
sizeof(TestType));
}
DYNAMIC_SECTION("Scattered addresses " << current) {
AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::device>(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<TestType, AtomicScopes::device>(2, 1,
sizeof(TestType));
}
DYNAMIC_SECTION("Adjacent addresses " << current) {
AtomicExchSingleDeviceMultipleKernelTest<TestType, AtomicScopes::device>(2, warp_size,
sizeof(TestType));
}
DYNAMIC_SECTION("Scattered addresses " << current) {
AtomicExchSingleDeviceMultipleKernelTest<TestType, AtomicScopes::device>(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);
}
catch/unit/atomics/atomicExch_common.hh
+379
Wyświetl plik
@@ -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 <hip_test_common.hh>
#include <resource_guards.hh>
#include <hip/hip_cooperative_groups.h>
#include <cmd_options.hh>
enum class AtomicScopes { device, system };
template <typename T, AtomicScopes scope> __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 <typename T, bool use_shared_mem, AtomicScopes scope>
__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<T, scope>(mem, static_cast<T>(tid + 1));
if constexpr (use_shared_mem) {
__syncthreads();
if (tid == 0) global_mem[0] = mem[0];
}
}
template <typename T>
__host__ __device__ T* pitched_offset(T* const ptr, const unsigned int pitch,
const unsigned int idx) {
const auto byte_ptr = reinterpret_cast<uint8_t*>(ptr);
return reinterpret_cast<T*>(byte_ptr + idx * pitch);
}
template <typename T, bool use_shared_mem, AtomicScopes scope>
__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<T*>(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<T, scope>(pitched_offset(mem, pitch, tid % width),
base_val + static_cast<T>(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 <typename TestType, bool use_shared_mem, AtomicScopes scope>
void AtomicExchSameAddress(const dim3 blocks, const dim3 threads, const LinearAllocs alloc_type) {
LinearAllocGuard<TestType> 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<TestType> old_vals_dev(LinearAllocs::hipMalloc, old_vals_alloc_size);
std::vector<TestType> old_vals(thread_count + 1);
HIP_CHECK(hipMemset(mem_dev.ptr(), 0, sizeof(TestType)));
atomic_exch_kernel_compile_time<TestType, use_shared_mem, scope>
<<<blocks, threads>>>(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 <typename TestType, AtomicScopes scope> 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<TestType, false, AtomicScopes::device>(blocks, threads, allocation_type);
}
SECTION("Shared memory") {
const auto blocks = dim3(1);
AtomicExchSameAddress<TestType, true, AtomicScopes::device>(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 <typename Derived, typename T, bool use_shared_mem, AtomicScopes scope>
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<LinearAllocGuard<T>> old_vals_devs;
std::vector<StreamGuard> 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<T> 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<T> 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<const Derived&>(*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<T>(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<std::thread> 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 <typename T, bool use_shared_mem, AtomicScopes scope>
class AtomicExch
: public AtomicExchCRTP<AtomicExch<T, use_shared_mem, scope>, 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<T, use_shared_mem, scope><<<p.blocks, p.threads, shared_mem_size, stream>>>(
mem, old_vals, p.width, p.pitch, base_val);
}
void ValidateResults(std::vector<T>& 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 <typename TestType, AtomicScopes scope>
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<TestType, false, scope>().run(params);
}
}
}
SECTION("Shared memory") {
params.blocks = dim3(1);
params.alloc_type = LinearAllocs::hipMalloc;
AtomicExch<TestType, true, scope>().run(params);
}
}
template <typename TestType, AtomicScopes scope>
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<TestType, false, scope>().run(params);
}
}
}
template <typename TestType>
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<TestType, false, AtomicScopes::system>().run(params);
}
}
}
catch/unit/atomics/atomicExch_negative_kernels.cc
+94
Wyświetl plik
@@ -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 <hip_test_common.hh>
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); }
catch/unit/atomics/atomicExch_negative_kernels_rtc.hh
+124
Wyświetl plik
@@ -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); }
)"};
catch/unit/atomics/atomicExch_system.cc
+220
Wyświetl plik
@@ -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<TestType>(2, 2, 1, sizeof(TestType));
}
DYNAMIC_SECTION("Adjacent addresses " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(2, 2, warp_size,
sizeof(TestType));
}
DYNAMIC_SECTION("Scattered addresses " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(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<TestType>(1, 1, 1, sizeof(TestType), 4);
}
DYNAMIC_SECTION("Adjacent addresses " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(1, 1, warp_size, sizeof(TestType),
4);
}
DYNAMIC_SECTION("Scattered addresses " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(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<TestType>(2, 2, 1, sizeof(TestType), 4);
}
DYNAMIC_SECTION("Adjacent addresses " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(2, 2, warp_size, sizeof(TestType),
4);
}
DYNAMIC_SECTION("Scattered addresses " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(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);
}
catch/unit/atomics/atomicExch_system_negative_kernels.cc
+112
Wyświetl plik
@@ -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 <hip_test_common.hh>
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); }
catch/unit/atomics/atomicExch_system_negative_kernels_rtc.hh
+142
Wyświetl plik
@@ -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); }
)"};
catch/unit/compileAndCaptureOutput.py
+107
Wyświetl plik
@@ -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]])