From c8a720088f3a657e433c562c31d29fea8685819c Mon Sep 17 00:00:00 2001 From: "Belton-Schure, Aidan" Date: Thu, 29 May 2025 17:08:20 +0100 Subject: [PATCH] SWDEV-527851 SWDEV-527890 SWDEV-529456 - Update vector types test (#172) * 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 [ROCm/hip-tests commit: ef6c95f6ce96e0eacc0eac438e57461b7c01a8e3] --- .../hip-tests/catch/unit/complex/complex.cc | 24 ++++ .../unit/deviceLib/hipVectorTypesDevice.cc | 72 ++++++------ .../unit/deviceLib/hipVectorTypesHost.cc | 111 +++++++++--------- projects/hip-tests/catch/unit/module/log | 1 + .../vector_types/vector_operations_common.hh | 17 ++- .../catch/unit/vector_types/vector_types.cc | 61 +++++++++- .../unit/vector_types/vector_types_common.hh | 6 + 7 files changed, 197 insertions(+), 95 deletions(-) create mode 100644 projects/hip-tests/catch/unit/module/log diff --git a/projects/hip-tests/catch/unit/complex/complex.cc b/projects/hip-tests/catch/unit/complex/complex.cc index 2d483db577..8511a4c457 100644 --- a/projects/hip-tests/catch/unit/complex/complex.cc +++ b/projects/hip-tests/catch/unit/complex/complex.cc @@ -388,6 +388,30 @@ TEMPLATE_TEST_CASE("Unit_Device_Complex_Cast_Host_Sanity_Positive", "", hipFloat REQUIRE(result.y == static_cast().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 result = CastComplexType>(input); + + REQUIRE(result.x == static_cast().x)>(input_r)); + REQUIRE(result.y != static_cast().x)>(input_r)); +} + /** * End doxygen group ComplexTest. * @} diff --git a/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesDevice.cc b/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesDevice.cc index af869fcbd9..800d1fc870 100644 --- a/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesDevice.cc +++ b/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesDevice.cc @@ -26,6 +26,8 @@ THE SOFTWARE. #include #include +#include "../vector_types/vector_types_common.hh" + using namespace std; template __device__ typename std::add_rvalue_reference::type _declval() noexcept; @@ -38,21 +40,21 @@ __device__ constexpr bool integer_unary_tests(const V&, const V&) { template ().x)>{}>* = nullptr> __device__ bool integer_unary_tests(V& f1, V& f2) { f1 %= f2; - if (f1 != V{0}) return false; + if (f1 != MakeVector(0)) return false; f1 &= f2; - if (f1 != V{0}) return false; + if (f1 != MakeVector(0)) return false; f1 |= f2; - if (f1 != V{1}) return false; + if (f1 != MakeVector(1)) return false; f1 ^= f2; - if (f1 != V{0}) return false; - f1 = V{1}; + if (f1 != MakeVector(0)) return false; + f1 = MakeVector(1); f1 <<= f2; - if (f1 != V{2}) return false; + if (f1 != MakeVector(2)) return false; f1 >>= f2; - if (f1 != V{1}) return false; + if (f1 != MakeVector(1)) return false; f2 = ~f1; - return f2 == V{~1}; + return f2 == MakeVector(~1); return true; } @@ -65,35 +67,35 @@ __device__ constexpr bool integer_binary_tests(const V&, const V&, const V&) { template ().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(0)) return false; f1 = f3 & f2; - if (f1 != V{0}) return false; + if (f1 != MakeVector(0)) return false; f2 = f1 ^ f3; - if (f2 != V{0}) return false; - f1 = V{1}; - f2 = V{2}; + if (f2 != MakeVector(0)) return false; + f1 = MakeVector(1); + f2 = MakeVector(2); f3 = f1 << f2; - if (f3 != V{4}) return false; + if (f3 != MakeVector(4)) return false; f2 = f3 >> f1; - return f2 == V{2}; + return f2 == MakeVector(2); } template __device__ bool TestVectorType() { - constexpr V v1{1}; - constexpr V v2{2}; - constexpr V v3{3}; - constexpr V v4{4}; + const V v1 = MakeVector(1); + const V v2 = MakeVector(2); + const V v3 = MakeVector(3); + const V v4 = MakeVector(4); - V f1{1}; - V f2{1}; + V f1 = MakeVector(1); + V f2 = MakeVector(1); V f3 = f1 + f2; - if (f3 != V{2}) return false; + if (f3 != MakeVector(2)) return false; f2 = f3 - f1; - if (f2 != V{1}) return false; + if (f2 != MakeVector(1)) return false; f1 = f2 * f3; - if (f1 != V{2}) return false; + if (f1 != MakeVector(2)) return false; f2 = f1 / f3; - if (f2 != V{1}) return false; + if (f2 != MakeVector(1)) return false; if (!integer_binary_tests(f1, f2, f3)) return false; f1 = v2; @@ -166,23 +168,23 @@ __global__ void CheckVectorTypes(bool* ptr) { template __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(1); + const V v2 = MakeVector(2); + const V v3 = MakeVector(3); + const V v4 = MakeVector(4); __shared__ V f1, f2, f3; *ptr = true; - f1 = V{1}; - f2 = V{1}; + f1 = MakeVector(1); + f2 = MakeVector(1); f3 = f1 + f2; - *ptr = *ptr && f3 == V{2}; + *ptr = *ptr && f3 == MakeVector(2); f2 = f3 - f1; - *ptr = *ptr && f2 == V{1}; + *ptr = *ptr && f2 == MakeVector(1); f1 = f2 * f3; - *ptr = *ptr && f1 == V{2}; + *ptr = *ptr && f1 == MakeVector(2); f2 = f1 / f3; - *ptr = *ptr && f2 == V{1}; + *ptr = *ptr && f2 == MakeVector(1); *ptr = *ptr && integer_binary_tests(f1, f2, f3); f1 = v2; diff --git a/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesHost.cc b/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesHost.cc index f275230916..354aae252f 100644 --- a/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesHost.cc +++ b/projects/hip-tests/catch/unit/deviceLib/hipVectorTypesHost.cc @@ -19,6 +19,8 @@ THE SOFTWARE. #include #include +#include "../vector_types/vector_types_common.hh" + #include #include #include @@ -37,41 +39,39 @@ bool integer_binary_tests(V&, V&, V&...) { return true; } -template().x)>{}>* = nullptr> -bool integer_unary_tests(V f1, V f2) { +template ().x)>{}>* = nullptr> +void integer_unary_tests(V f1, V f2) { f1 %= f2; - if (f1 != V{0}) return false; + REQUIRE(f1 == MakeVector(0)); f1 &= f2; - if (f1 != V{0}) return false; + REQUIRE(f1 == MakeVector(0)); f1 |= f2; - if (f1 != V{1}) return false; + REQUIRE(f1 == MakeVector(1)); f1 ^= f2; - if (f1 != V{0}) return false; - f1 = V{1}; + REQUIRE(f1 == MakeVector(0)); + f1 = MakeVector(1); f1 <<= f2; - if (f1 != V{2}) return false; + REQUIRE(f1 == MakeVector(2)); f1 >>= f2; - if (f1 != V{1}) return false; + REQUIRE(f1 == MakeVector(1)); f2 = ~f1; - return f2 == V{~1}; + REQUIRE(f2 == MakeVector(~1)); } -template().x)>{}>* = nullptr> -bool integer_binary_tests(V f1, V f2, V f3) { +template ().x)>{}>* = nullptr> +void integer_binary_tests(V f1, V f2, V f3) { f3 = f1 % f2; - if (f3 != V{0}) return false; + REQUIRE(f3 == MakeVector(0)); f1 = f3 & f2; - if (f1 != V{0}) return false; + REQUIRE(f1 == MakeVector(0)); f2 = f1 ^ f3; - if (f2 != V{0}) return false; - f1 = V{1}; - f2 = V{2}; + REQUIRE(f2 == MakeVector(0)); + f1 = MakeVector(1); + f2 = MakeVector(2); f3 = f1 << f2; - if (f3 != V{4}) return false; + REQUIRE(f3 == MakeVector(4)); f2 = f3 >> f1; - return f2 == V{2}; + REQUIRE(f2 == MakeVector(2)); } template @@ -94,56 +94,57 @@ bool constructor_tests() { template bool TestVectorType() { - constexpr V v1{1}; - constexpr V v2{2}; - constexpr V v3{3}; - constexpr V v4{4}; + const V v1 = MakeVector(1); + const V v2 = MakeVector(2); + const V v3 = MakeVector(3); + const V v4 = MakeVector(4); - V f1{1}; - V f2{1}; + V f1 = MakeVector(1); + V f2 = MakeVector(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(2); + f2 = MakeVector(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()) return false; + REQUIRE(constructor_tests() == 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(3)); + REQUIRE(y == MakeVector(4)); + REQUIRE(z == MakeVector(3)); + REQUIRE(w == MakeVector(4)); stringstream str; str << f1.x; str >> f2.x; - if (f1.x != f2.x) return false; + REQUIRE(f1.x == f2.x); return true; } diff --git a/projects/hip-tests/catch/unit/module/log b/projects/hip-tests/catch/unit/module/log new file mode 100644 index 0000000000..355ffc4545 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/log @@ -0,0 +1 @@ +bash: ./ModuleTest: No such file or directory diff --git a/projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh b/projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh index 9684e0ecb1..ede6b010ac 100644 --- a/projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh +++ b/projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh @@ -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) { 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) { 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) { if (operation == VectorOperation::kModulo) { diff --git a/projects/hip-tests/catch/unit/vector_types/vector_types.cc b/projects/hip-tests/catch/unit/vector_types/vector_types.cc index 9bc221d498..7e73a45567 100644 --- a/projects/hip-tests/catch/unit/vector_types/vector_types.cc +++ b/projects/hip-tests/catch/unit/vector_types/vector_types.cc @@ -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(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(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(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(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(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 /** diff --git a/projects/hip-tests/catch/unit/vector_types/vector_types_common.hh b/projects/hip-tests/catch/unit/vector_types/vector_types_common.hh index cbaf3f24a6..65530618d5 100644 --- a/projects/hip-tests/catch/unit/vector_types/vector_types_common.hh +++ b/projects/hip-tests/catch/unit/vector_types/vector_types_common.hh @@ -165,6 +165,12 @@ __host__ __device__ void MakeVectorType(T* vector_ptr, decltype(T().x) value) { } } +template __host__ __device__ T MakeVector(decltype(T().x) value) { + T vector{}; + MakeVectorType(&vector, value); + return vector; +} + template T MakeVectorTypeHost(decltype(T().x) value) { T vector{}; MakeVectorType(&vector, value);