AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-527851 SWDEV-527890 SWDEV-529456 - Update vector types test (#172)

Browse Source 
* Update vector types test

* update MakeVector to be host+device function

* Add regression testing

* Add tests for subscript operator

Change-Id: Ifac85aaa6cd6d6bd96c4239414e38f284e2a2d54

* Fix unused vars

Change-Id: I6bc2da76dbf962db5d75ea5a84bf16b66f8ba6ba

---------

Co-authored-by: Aidan Belton <abeltons@amd.com>

[ROCm/hip-tests commit: ef6c95f6ce]
This commit is contained in:
Belton-Schure, Aidan
2025-05-29 17:08:20 +01:00
committed by GitHub
parent 98c34321f1
commit c8a720088f
7 changed files with 197 additions and 95 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/complex/complex.cc
+24
View File
@@ -388,6 +388,30 @@ TEMPLATE_TEST_CASE("Unit_Device_Complex_Cast_Host_Sanity_Positive", "", hipFloat
REQUIRE(result.y == static_cast<decltype(CastType_t<TestType>().x)>(input_i));
}
/**
* Test Description
* ------------------------
* - Test that checks hipComplexDoubleToFloat/hipComplexFloatToDouble single value
* constructor only initializes the real component.
*
* Test source
* ------------------------
* - unit/complex/complex.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEMPLATE_TEST_CASE("Unit_Device_Complex_Constructor_Host", "", hipFloatComplex,
hipDoubleComplex) {
decltype(TestType().x) input_r = GENERATE(-0.25, 0.25);
TestType input{input_r};
CastType_t<TestType> result = CastComplexType<CastType_t<TestType>>(input);
REQUIRE(result.x == static_cast<decltype(CastType_t<TestType>().x)>(input_r));
REQUIRE(result.y != static_cast<decltype(CastType_t<TestType>().x)>(input_r));
}
/**
* End doxygen group ComplexTest.
* @}
projects/hip-tests/catch/unit/deviceLib/hipVectorTypesDevice.cc
+37 -35
View File
@@ -26,6 +26,8 @@ THE SOFTWARE.
#include <type_traits>
#include <utility>
#include "../vector_types/vector_types_common.hh"
using namespace std;
template <class T> __device__ typename std::add_rvalue_reference<T>::type _declval() noexcept;
@@ -38,21 +40,21 @@ __device__ constexpr bool integer_unary_tests(const V&, const V&) {
template <typename V, enable_if_t<is_integral<decltype(_declval<V>().x)>{}>* = nullptr>
__device__ bool integer_unary_tests(V& f1, V& f2) {
f1 %= f2;
if (f1 != V{0}) return false;
if (f1 != MakeVector<V>(0)) return false;
f1 &= f2;
if (f1 != V{0}) return false;
if (f1 != MakeVector<V>(0)) return false;
f1 |= f2;
if (f1 != V{1}) return false;
if (f1 != MakeVector<V>(1)) return false;
f1 ^= f2;
if (f1 != V{0}) return false;
f1 = V{1};
if (f1 != MakeVector<V>(0)) return false;
f1 = MakeVector<V>(1);
f1 <<= f2;
if (f1 != V{2}) return false;
if (f1 != MakeVector<V>(2)) return false;
f1 >>= f2;
if (f1 != V{1}) return false;
if (f1 != MakeVector<V>(1)) return false;
f2 = ~f1;
return f2 == V{~1};
return f2 == MakeVector<V>(~1);
return true;
}
@@ -65,35 +67,35 @@ __device__ constexpr bool integer_binary_tests(const V&, const V&, const V&) {
template <typename V, enable_if_t<is_integral<decltype(_declval<V>().x)>{}>* = nullptr>
__device__ bool integer_binary_tests(V& f1, V& f2, V& f3) {
f3 = f1 % f2;
if (f3 != V{0}) return false;
if (f3 != MakeVector<V>(0)) return false;
f1 = f3 & f2;
if (f1 != V{0}) return false;
if (f1 != MakeVector<V>(0)) return false;
f2 = f1 ^ f3;
if (f2 != V{0}) return false;
f1 = V{1};
f2 = V{2};
if (f2 != MakeVector<V>(0)) return false;
f1 = MakeVector<V>(1);
f2 = MakeVector<V>(2);
f3 = f1 << f2;
if (f3 != V{4}) return false;
if (f3 != MakeVector<V>(4)) return false;
f2 = f3 >> f1;
return f2 == V{2};
return f2 == MakeVector<V>(2);
}
template <typename V> __device__ bool TestVectorType() {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
const V v1 = MakeVector<V>(1);
const V v2 = MakeVector<V>(2);
const V v3 = MakeVector<V>(3);
const V v4 = MakeVector<V>(4);
V f1{1};
V f2{1};
V f1 = MakeVector<V>(1);
V f2 = MakeVector<V>(1);
V f3 = f1 + f2;
if (f3 != V{2}) return false;
if (f3 != MakeVector<V>(2)) return false;
f2 = f3 - f1;
if (f2 != V{1}) return false;
if (f2 != MakeVector<V>(1)) return false;
f1 = f2 * f3;
if (f1 != V{2}) return false;
if (f1 != MakeVector<V>(2)) return false;
f2 = f1 / f3;
if (f2 != V{1}) return false;
if (f2 != MakeVector<V>(1)) return false;
if (!integer_binary_tests(f1, f2, f3)) return false;
f1 = v2;
@@ -166,23 +168,23 @@ __global__ void CheckVectorTypes(bool* ptr) {
template <typename V> __global__ void CheckSharedVectorType(bool* ptr) {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
const V v1 = MakeVector<V>(1);
const V v2 = MakeVector<V>(2);
const V v3 = MakeVector<V>(3);
const V v4 = MakeVector<V>(4);
__shared__ V f1, f2, f3;
*ptr = true;
f1 = V{1};
f2 = V{1};
f1 = MakeVector<V>(1);
f2 = MakeVector<V>(1);
f3 = f1 + f2;
*ptr = *ptr && f3 == V{2};
*ptr = *ptr && f3 == MakeVector<V>(2);
f2 = f3 - f1;
*ptr = *ptr && f2 == V{1};
*ptr = *ptr && f2 == MakeVector<V>(1);
f1 = f2 * f3;
*ptr = *ptr && f1 == V{2};
*ptr = *ptr && f1 == MakeVector<V>(2);
f2 = f1 / f3;
*ptr = *ptr && f2 == V{1};
*ptr = *ptr && f2 == MakeVector<V>(1);
*ptr = *ptr && integer_binary_tests(f1, f2, f3);
f1 = v2;
projects/hip-tests/catch/unit/deviceLib/hipVectorTypesHost.cc
+56 -55
View File
@@ -19,6 +19,8 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <hip/hip_vector_types.h>
#include "../vector_types/vector_types_common.hh"
#include <string>
#include <sstream>
#include <type_traits>
@@ -37,41 +39,39 @@ bool integer_binary_tests(V&, V&, V&...) {
return true;
}
template<typename V,
enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
bool integer_unary_tests(V f1, V f2) {
template <typename V, enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
void integer_unary_tests(V f1, V f2) {
f1 %= f2;
if (f1 != V{0}) return false;
REQUIRE(f1 == MakeVector<V>(0));
f1 &= f2;
if (f1 != V{0}) return false;
REQUIRE(f1 == MakeVector<V>(0));
f1 |= f2;
if (f1 != V{1}) return false;
REQUIRE(f1 == MakeVector<V>(1));
f1 ^= f2;
if (f1 != V{0}) return false;
f1 = V{1};
REQUIRE(f1 == MakeVector<V>(0));
f1 = MakeVector<V>(1);
f1 <<= f2;
if (f1 != V{2}) return false;
REQUIRE(f1 == MakeVector<V>(2));
f1 >>= f2;
if (f1 != V{1}) return false;
REQUIRE(f1 == MakeVector<V>(1));
f2 = ~f1;
return f2 == V{~1};
REQUIRE(f2 == MakeVector<V>(~1));
}
template<typename V,
enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
bool integer_binary_tests(V f1, V f2, V f3) {
template <typename V, enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
void integer_binary_tests(V f1, V f2, V f3) {
f3 = f1 % f2;
if (f3 != V{0}) return false;
REQUIRE(f3 == MakeVector<V>(0));
f1 = f3 & f2;
if (f1 != V{0}) return false;
REQUIRE(f1 == MakeVector<V>(0));
f2 = f1 ^ f3;
if (f2 != V{0}) return false;
f1 = V{1};
f2 = V{2};
REQUIRE(f2 == MakeVector<V>(0));
f1 = MakeVector<V>(1);
f2 = MakeVector<V>(2);
f3 = f1 << f2;
if (f3 != V{4}) return false;
REQUIRE(f3 == MakeVector<V>(4));
f2 = f3 >> f1;
return f2 == V{2};
REQUIRE(f2 == MakeVector<V>(2));
}
template<typename V>
@@ -94,56 +94,57 @@ bool constructor_tests() {
template<typename V>
bool TestVectorType() {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
const V v1 = MakeVector<V>(1);
const V v2 = MakeVector<V>(2);
const V v3 = MakeVector<V>(3);
const V v4 = MakeVector<V>(4);
V f1{1};
V f2{1};
V f1 = MakeVector<V>(1);
V f2 = MakeVector<V>(1);
V f3 = f1 + f2;
if (f3 != v2) return false;
REQUIRE(f3 == v2);
f2 = f3 - f1;
if (f2 != v1) return false;
REQUIRE(f2 == v1);
f1 = f2 * f3;
if (f1 != v2) return false;
REQUIRE(f1 == v2);
f2 = f1 / f3;
if (f2 != v1) return false;
if (!integer_binary_tests(f1, f2, f3)) return false;
REQUIRE(f2 == v1);
integer_binary_tests(f1, f2, f3);
f1 = V{2};
f2 = V{1};
f1 = MakeVector<V>(2);
f2 = MakeVector<V>(1);
f1 += f2;
if (f1 != v3) return false;
REQUIRE(f1 == v3);
f1 -= f2;
if (f1 != v2) return false;
REQUIRE(f1 == v2);
f1 *= f2;
if (f1 != v2) return false;
REQUIRE(f1 == v2);
f1 /= f2;
if (f1 != v2) return false;
if (!integer_unary_tests(f1, f2)) return false;
REQUIRE(f1 == v2);
integer_unary_tests(f1, f2);
f1 = v2;
f2 = f1++;
if (f1 != v3) return false;
if (f2 != v2) return false;
REQUIRE(f1 == v3);
REQUIRE(f2 == v2);
f2 = f1--;
if (f2 != v3) return false;
if (f1 != v2) return false;
REQUIRE(f2 == v3);
REQUIRE(f1 == v2);
f2 = ++f1;
if (f1 != v3) return false;
if (f2 != v3) return false;
REQUIRE(f1 == v3);
REQUIRE(f2 == v3);
f2 = --f1;
if (f1 != v2) return false;
if (f2 != v2) return false;
REQUIRE(f1 == v2);
REQUIRE(f2 == v2);
if (!constructor_tests<V>()) return false;
REQUIRE(constructor_tests<V>() == true);
f1 = v3;
f2 = v4;
f3 = v3;
if (f1 == f2) return false;
if (!(f1 != f2)) return false;
REQUIRE(!(f1 == f2));
REQUIRE(f1 != f2);
using T = typename V::value_type;
@@ -152,16 +153,16 @@ bool TestVectorType() {
const volatile T& z = f3.x;
volatile T& w = f2.x;
if (x != T{3}) return false;
if (y != T{4}) return false;
if (z != T{3}) return false;
if (w != T{4}) return false;
REQUIRE(x == MakeVector<V>(3));
REQUIRE(y == MakeVector<V>(4));
REQUIRE(z == MakeVector<V>(3));
REQUIRE(w == MakeVector<V>(4));
stringstream str;
str << f1.x;
str >> f2.x;
if (f1.x != f2.x) return false;
REQUIRE(f1.x == f2.x);
return true;
}
projects/hip-tests/catch/unit/module/log
+1
View File
@@ -0,0 +1 @@
bash: ./ModuleTest: No such file or directory
projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh
+16 -1
View File
@@ -47,7 +47,8 @@ enum class VectorOperation {
kBitwiseOr,
kBitwiseAnd,
kRightShift,
kLeftShift
kLeftShift,
kSubscript,
};
inline std::string to_string(VectorOperation operation) {
@@ -108,6 +109,8 @@ inline std::string to_string(VectorOperation operation) {
return "right shift";
case VectorOperation::kLeftShift:
return "left shift";
case VectorOperation::kSubscript:
return "subscript";
default:
return "Unknown";
}
@@ -144,6 +147,8 @@ void SanityCheck(VectorOperation operation, T vector, decltype(T().x) value1,
value1 = (value1 == value2) ? 2 * value1 : 3 * value1;
} else if (operation == VectorOperation::kNotEqual) {
value1 = (value1 != value2) ? 2 * value1 : 3 * value1;
} else if (operation == VectorOperation::kSubscript) {
value1 = value2;
} else {
if constexpr (std::is_signed_v<decltype(T().x)>) {
if (operation == VectorOperation::kNegate) {
@@ -214,6 +219,11 @@ __device__ __host__ void PerformVectorOperation(VectorOperation operation, T* ve
*vector1 = (*vector1 == vector2) ? 2 * *vector1 : 3 * *vector1;
} else if (operation == VectorOperation::kNotEqual) {
*vector1 = (*vector1 != vector2) ? 2 * *vector1 : 3 * *vector1;
} else if (operation == VectorOperation::kSubscript) {
constexpr size_t n_elements = sizeof(T) / sizeof(decltype(T().x));
for (int i = 0; i < n_elements; ++i) {
(*vector1)[i] = vector2[i];
}
} else {
if constexpr (std::is_signed_v<decltype(T().x)>) {
if (operation == VectorOperation::kNegate) {
@@ -275,6 +285,11 @@ __device__ __host__ void PerformVectorOperation(VectorOperation operation, T* ve
*vector = (*vector == value) ? 2 * *vector : 3 * *vector;
} else if (operation == VectorOperation::kNotEqual) {
*vector = (*vector != value) ? 2 * *vector : 3 * *vector;
} else if (operation == VectorOperation::kSubscript) {
constexpr size_t n_elements = sizeof(T) / sizeof(decltype(T().x));
for (int i = 0; i < n_elements; ++i) {
(*vector)[i] = value;
}
} else {
if constexpr (std::is_integral_v<decltype(T().x)>) {
if (operation == VectorOperation::kModulo) {
projects/hip-tests/catch/unit/vector_types/vector_types.cc
+57 -4
View File
@@ -153,7 +153,8 @@ TEMPLATE_TEST_CASE("Unit_VectorAndVectorOperations_SanityCheck_Basic_Host", "",
VectorOperation::kBitwiseOr,
VectorOperation::kBitwiseAnd,
VectorOperation::kRightShift,
VectorOperation::kLeftShift}) {
VectorOperation::kLeftShift,
VectorOperation::kSubscript}) {
DYNAMIC_SECTION("operation: " << to_string(operation)) {
TestType vector = PerformVectorOperationHost<TestType>(operation, value1, value2);
SanityCheck(operation, vector, value1, value2);
@@ -190,7 +191,7 @@ TEMPLATE_TEST_CASE("Unit_VectorAndValueTypeOperations_SanityCheck_Basic_Host", "
VectorOperation::kSubtract, VectorOperation::kMultiply, VectorOperation::kDivide,
VectorOperation::kEqual, VectorOperation::kNotEqual, VectorOperation::kModulo,
VectorOperation::kBitwiseXor, VectorOperation::kBitwiseOr, VectorOperation::kBitwiseAnd,
VectorOperation::kRightShift, VectorOperation::kLeftShift}) {
VectorOperation::kRightShift, VectorOperation::kLeftShift, VectorOperation::kSubscript}) {
DYNAMIC_SECTION("operation: " << to_string(operation)) {
TestType vector = PerformVectorOperationHost<TestType, false>(operation, value1, value2);
SanityCheck(operation, vector, value1, value2);
@@ -248,7 +249,8 @@ TEMPLATE_TEST_CASE("Unit_VectorAndVectorOperations_SanityCheck_Basic_Device", ""
VectorOperation::kBitwiseOr,
VectorOperation::kBitwiseAnd,
VectorOperation::kRightShift,
VectorOperation::kLeftShift}) {
VectorOperation::kLeftShift,
VectorOperation::kSubscript}) {
DYNAMIC_SECTION("operation: " << to_string(operation)) {
TestType vector = PerformVectorOperationDevice<TestType>(operation, value1, value2);
SanityCheck(operation, vector, value1, value2);
@@ -285,13 +287,64 @@ TEMPLATE_TEST_CASE("Unit_VectorAndValueTypeOperations_SanityCheck_Basic_Device",
VectorOperation::kSubtract, VectorOperation::kMultiply, VectorOperation::kDivide,
VectorOperation::kEqual, VectorOperation::kNotEqual, VectorOperation::kModulo,
VectorOperation::kBitwiseXor, VectorOperation::kBitwiseOr, VectorOperation::kBitwiseAnd,
VectorOperation::kRightShift, VectorOperation::kLeftShift}) {
VectorOperation::kRightShift, VectorOperation::kLeftShift, VectorOperation::kSubscript}) {
DYNAMIC_SECTION("operation: " << to_string(operation)) {
TestType vector = PerformVectorOperationDevice<TestType, false>(operation, value1, value2);
SanityCheck(operation, vector, value1, value2);
}
}
}
/**
* Test Description
* ------------------------
* - Checks that vectors can be used with structured bindigns
* - Tests from the host side
* Test source
* ------------------------
* - unit/vector_types/vector_types.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEMPLATE_TEST_CASE("Unit_VectorStructuredBindings_SanityCheck_Basic_host", "", float3, double3) {
auto value = GetTestValue<decltype(TestType().x)>(0);
TestType vec3 = {value, value, value};
auto&& [ret1, ret2, ret3] = vec3;
REQUIRE(ret1 == value);
}
__host__ __device__ constexpr bool func()
{
int3 vec3 = int3{0};
int exp = int{0};
return vec3.x == exp;
}
__global__ void generate_my_kernel()
{
static_assert(func());
}
/**
* Test Description
* ------------------------
* - Checks that vectors can be used with structured bindigns
* - Tests from the host and device side
* Test source
* ------------------------
* - unit/vector_types/vector_types.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_VectorConstexpr_SanityCheck_Basic_host_device", "") {
generate_my_kernel<<<1, 1>>>();
static_assert(func());
}
#endif // HT_AMD
/**
projects/hip-tests/catch/unit/vector_types/vector_types_common.hh
+6
View File
@@ -165,6 +165,12 @@ __host__ __device__ void MakeVectorType(T* vector_ptr, decltype(T().x) value) {
}
}
template <typename T> __host__ __device__ T MakeVector(decltype(T().x) value) {
T vector{};
MakeVectorType(&vector, value);
return vector;
}
template <typename T> T MakeVectorTypeHost(decltype(T().x) value) {
T vector{};
MakeVectorType(&vector, value);