From 2520691374af73ee7575210248366577439c653a Mon Sep 17 00:00:00 2001 From: Milos Mozetic Date: Thu, 28 Dec 2023 18:35:58 +0100 Subject: [PATCH] EXSWHTEC-314 - Implement unit tests for short vector types #297 Change-Id: I609548b79ae45530b0e82c532f93cf1f411343e4 [ROCm/hip-tests commit: 536c20b62d2366271f9344396ed591aee01abf21] --- .../catch/include/hip_test_defgroups.hh | 3 + .../catch/unit/vector_types/CMakeLists.txt | 58 ++ .../vector_types/negative_bitwise_double.cc | 25 + .../vector_types/negative_bitwise_float.cc | 25 + .../negative_bitwise_float_double_rtc.hh | 272 +++++++ ...calculate_assign_with_unsigned_value_1D.cc | 26 + ...calculate_assign_with_unsigned_value_2D.cc | 26 + ...calculate_assign_with_unsigned_value_3D.cc | 26 + ...calculate_assign_with_unsigned_value_4D.cc | 26 + ...negative_calculate_assign_with_value_1D.cc | 26 + ...negative_calculate_assign_with_value_2D.cc | 26 + ...negative_calculate_assign_with_value_3D.cc | 26 + ...negative_calculate_assign_with_value_4D.cc | 26 + ...egative_calculate_assign_with_value_rtc.hh | 750 ++++++++++++++++++ .../vector_types/negative_macros_common.hh | 78 ++ .../vector_types/negative_negate_unsigned.cc | 41 + .../negative_negate_unsigned_rtc.hh | 150 ++++ .../vector_types/vector_operations_common.hh | 338 ++++++++ .../catch/unit/vector_types/vector_types.cc | 357 +++++++++ .../unit/vector_types/vector_types_common.hh | 187 +++++ 20 files changed, 2492 insertions(+) create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_bitwise_double.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_bitwise_float.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_bitwise_float_double_rtc.hh create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_1D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_2D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_3D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_4D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_1D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_2D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_3D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_4D.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_rtc.hh create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_macros_common.hh create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned_rtc.hh create mode 100644 projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh create mode 100644 projects/hip-tests/catch/unit/vector_types/vector_types.cc create mode 100644 projects/hip-tests/catch/unit/vector_types/vector_types_common.hh diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 1ab256ec10..58ecdae020 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -388,4 +388,7 @@ TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_ParseError") {} /** * End doxygen group DeviceLanguageTest. * @} + * @defgroup VectorTypeTest Vector types + * @{ + * This section describes tests for the Vector type functions and operators. */ diff --git a/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt b/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt index 49619275f3..f3aa0fb1af 100644 --- a/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/vector_types/CMakeLists.txt @@ -21,9 +21,67 @@ # Common Tests - Test independent of all platforms set(TEST_SRC + vector_types.cc dim3.cc ) hip_add_exe_to_target(NAME VectorTypesTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) + +if(HIP_PLATFORM MATCHES "amd") + add_test(NAME Unit_NegateUnsigned_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_negate_unsigned.cc 40) + + add_test(NAME Unit_BitwiseFloat_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_bitwise_float.cc 96) + + add_test(NAME Unit_BitwiseDouble_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_bitwise_double.cc 96) + + add_test(NAME Unit_CalculateAssign1D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_value_1D.cc 60) + + add_test(NAME Unit_CalculateAssign2D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_value_2D.cc 60) + + add_test(NAME Unit_CalculateAssign3D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_value_3D.cc 60) + + add_test(NAME Unit_CalculateAssign4D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_value_4D.cc 60) + + add_test(NAME Unit_CalculateAssignUnsigned1D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_unsigned_value_1D.cc 60) + + add_test(NAME Unit_CalculateAssignUnsigned2D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_unsigned_value_2D.cc 60) + + add_test(NAME Unit_CalculateAssignUnsigned3D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_unsigned_value_3D.cc 60) + + add_test(NAME Unit_CalculateAssignUnsigned4D_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + negative_calculate_assign_with_unsigned_value_4D.cc 60) +endif() diff --git a/projects/hip-tests/catch/unit/vector_types/negative_bitwise_double.cc b/projects/hip-tests/catch/unit/vector_types/negative_bitwise_double.cc new file mode 100644 index 0000000000..e3a88bfe28 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_bitwise_double.cc @@ -0,0 +1,25 @@ +/* +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 "negative_macros_common.hh" + +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(double1) +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(double2) +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(double3) +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(double4) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_bitwise_float.cc b/projects/hip-tests/catch/unit/vector_types/negative_bitwise_float.cc new file mode 100644 index 0000000000..cc7185f098 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_bitwise_float.cc @@ -0,0 +1,25 @@ +/* +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 "negative_macros_common.hh" + +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(float1) +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(float2) +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(float3) +BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(float4) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_bitwise_float_double_rtc.hh b/projects/hip-tests/catch/unit/vector_types/negative_bitwise_float_double_rtc.hh new file mode 100644 index 0000000000..bf5eed0424 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_bitwise_float_double_rtc.hh @@ -0,0 +1,272 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +static constexpr auto kBitwiseFloat{R"( + __global__ void BitwiseDevice(float1* vector1_dev_ptr, float1* vector2_dev_ptr) { \ + float1 vector1_dev = *vector1_dev_ptr; \ + float1 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(float1& vector1_host, float1& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + + __global__ void BitwiseDevice(float2* vector1_dev_ptr, float2* vector2_dev_ptr) { \ + float2 vector1_dev = *vector1_dev_ptr; \ + float2 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(float2& vector1_host, float2& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + + __global__ void BitwiseDevice(float3* vector1_dev_ptr, float3* vector2_dev_ptr) { \ + float3 vector1_dev = *vector1_dev_ptr; \ + float3 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(float3& vector1_host, float3& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + + __global__ void BitwiseDevice(float4* vector1_dev_ptr, float4* vector2_dev_ptr) { \ + float4 vector1_dev = *vector1_dev_ptr; \ + float4 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(float4& vector1_host, float4& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } +)"}; + +static constexpr auto kBitwiseDouble{R"( + __global__ void BitwiseDevice(double1* vector1_dev_ptr, double1* vector2_dev_ptr) { \ + double1 vector1_dev = *vector1_dev_ptr; \ + double1 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(double1& vector1_host, double1& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + + __global__ void BitwiseDevice(double2* vector1_dev_ptr, double2* vector2_dev_ptr) { \ + double2 vector1_dev = *vector1_dev_ptr; \ + double2 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(double2& vector1_host, double2& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + + __global__ void BitwiseDevice(double3* vector1_dev_ptr, double3* vector2_dev_ptr) { \ + double3 vector1_dev = *vector1_dev_ptr; \ + double3 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(double3& vector1_host, double3& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + + __global__ void BitwiseDevice(double4* vector1_dev_ptr, double4* vector2_dev_ptr) { \ + double4 vector1_dev = *vector1_dev_ptr; \ + double4 vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(double4& vector1_host, double4& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } +)"}; diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_1D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_1D.cc new file mode 100644 index 0000000000..f95f2e5f0d --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_1D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uchar1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ushort1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uint1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulong1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulonglong1) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_2D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_2D.cc new file mode 100644 index 0000000000..9f69fc18e7 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_2D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uchar2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ushort2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uint2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulong2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulonglong2) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_3D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_3D.cc new file mode 100644 index 0000000000..99f9d8130e --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_3D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uchar3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ushort3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uint3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulong3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulonglong3) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_4D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_4D.cc new file mode 100644 index 0000000000..301ae99c61 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_unsigned_value_4D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uchar4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ushort4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(uint4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulong4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(ulonglong4) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_1D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_1D.cc new file mode 100644 index 0000000000..09d7267186 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_1D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(char1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(short1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(int1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(long1) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(longlong1) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_2D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_2D.cc new file mode 100644 index 0000000000..173d5fcb4e --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_2D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(char2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(short2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(int2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(long2) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(longlong2) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_3D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_3D.cc new file mode 100644 index 0000000000..8cd3145aa8 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_3D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(char3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(short3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(int3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(long3) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(longlong3) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_4D.cc b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_4D.cc new file mode 100644 index 0000000000..9679d17548 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_4D.cc @@ -0,0 +1,26 @@ +/* +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 "negative_macros_common.hh" + +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(char4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(short4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(int4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(long4) +CALCULATE_ASSIGN_VECTOR_FUNCTIONS(longlong4) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_rtc.hh b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_rtc.hh new file mode 100644 index 0000000000..ee921a016c --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_calculate_assign_with_value_rtc.hh @@ -0,0 +1,750 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of longge, 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +static constexpr auto kCalculateAssignChar{R"( + __global__ void CalculateAssignDevice(char1* vector_dev_ptr, decltype(char1().x) value) { + char1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(char1& vector_host, decltype(char1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uchar1* vector_dev_ptr, decltype(uchar1().x) value) { + uchar1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uchar1& vector_host, decltype(uchar1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(char2* vector_dev_ptr, decltype(char2().x) value) { + char2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(char2& vector_host, decltype(char2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uchar2* vector_dev_ptr, decltype(uchar2().x) value) { + uchar2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uchar2& vector_host, decltype(uchar2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(char3* vector_dev_ptr, decltype(char3().x) value) { + char3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(char3& vector_host, decltype(char3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uchar3* vector_dev_ptr, decltype(uchar3().x) value) { + uchar3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uchar3& vector_host, decltype(uchar3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(char4* vector_dev_ptr, decltype(char4().x) value) { + char4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(char4& vector_host, decltype(char4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uchar4* vector_dev_ptr, decltype(uchar4().x) value) { + uchar4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uchar4& vector_host, decltype(uchar4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } +)"}; + +static constexpr auto kCalculateAssignShort{R"( + __global__ void CalculateAssignDevice(short1* vector_dev_ptr, decltype(short1().x) value) { + short1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(short1& vector_host, decltype(short1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ushort1* vector_dev_ptr, decltype(ushort1().x) value) { + ushort1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ushort1& vector_host, decltype(ushort1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(short2* vector_dev_ptr, decltype(short2().x) value) { + short2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(short2& vector_host, decltype(short2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ushort2* vector_dev_ptr, decltype(ushort2().x) value) { + ushort2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ushort2& vector_host, decltype(ushort2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(short3* vector_dev_ptr, decltype(short3().x) value) { + short3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(short3& vector_host, decltype(short3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ushort3* vector_dev_ptr, decltype(ushort3().x) value) { + ushort3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ushort3& vector_host, decltype(ushort3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(short4* vector_dev_ptr, decltype(short4().x) value) { + short4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(short4& vector_host, decltype(short4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ushort4* vector_dev_ptr, decltype(ushort4().x) value) { + ushort4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ushort4& vector_host, decltype(ushort4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } +)"}; + +static constexpr auto kCalculateAssignInt{R"( + __global__ void CalculateAssignDevice(int1* vector_dev_ptr, decltype(int1().x) value) { + int1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(int1& vector_host, decltype(int1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uint1* vector_dev_ptr, decltype(uint1().x) value) { + uint1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uint1& vector_host, decltype(uint1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(int2* vector_dev_ptr, decltype(int2().x) value) { + int2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(int2& vector_host, decltype(int2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uint2* vector_dev_ptr, decltype(uint2().x) value) { + uint2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uint2& vector_host, decltype(uint2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(int3* vector_dev_ptr, decltype(int3().x) value) { + int3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(int3& vector_host, decltype(int3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uint3* vector_dev_ptr, decltype(uint3().x) value) { + uint3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uint3& vector_host, decltype(uint3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(int4* vector_dev_ptr, decltype(int4().x) value) { + int4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(int4& vector_host, decltype(int4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(uint4* vector_dev_ptr, decltype(uint4().x) value) { + uint4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(uint4& vector_host, decltype(uint4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } +)"}; + +static constexpr auto kCalculateAssignLong{R"( + __global__ void CalculateAssignDevice(long1* vector_dev_ptr, decltype(long1().x) value) { + long1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(long1& vector_host, decltype(long1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulong1* vector_dev_ptr, decltype(ulong1().x) value) { + ulong1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulong1& vector_host, decltype(ulong1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(long2* vector_dev_ptr, decltype(long2().x) value) { + long2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(long2& vector_host, decltype(long2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulong2* vector_dev_ptr, decltype(ulong2().x) value) { + ulong2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulong2& vector_host, decltype(ulong2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(long3* vector_dev_ptr, decltype(long3().x) value) { + long3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(long3& vector_host, decltype(long3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulong3* vector_dev_ptr, decltype(ulong3().x) value) { + ulong3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulong3& vector_host, decltype(ulong3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(long4* vector_dev_ptr, decltype(long4().x) value) { + long4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(long4& vector_host, decltype(long4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulong4* vector_dev_ptr, decltype(ulong4().x) value) { + ulong4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulong4& vector_host, decltype(ulong4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } +)"}; + +static constexpr auto kCalculateAssignLongLong{R"( + __global__ void CalculateAssignDevice(longlong1* vector_dev_ptr, decltype(longlong1().x) value) { + longlong1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(longlong1& vector_host, decltype(longlong1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulonglong1* vector_dev_ptr, decltype(ulonglong1().x) value) { + ulonglong1 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulonglong1& vector_host, decltype(ulonglong1().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(longlong2* vector_dev_ptr, decltype(longlong2().x) value) { + longlong2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(longlong2& vector_host, decltype(longlong2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulonglong2* vector_dev_ptr, decltype(ulonglong2().x) value) { + ulonglong2 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulonglong2& vector_host, decltype(ulonglong2().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(longlong3* vector_dev_ptr, decltype(longlong3().x) value) { + longlong3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(longlong3& vector_host, decltype(longlong3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulonglong3* vector_dev_ptr, decltype(ulonglong3().x) value) { + ulonglong3 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulonglong3& vector_host, decltype(ulonglong3().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(longlong4* vector_dev_ptr, decltype(longlong4().x) value) { + longlong4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(longlong4& vector_host, decltype(longlong4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } + + __global__ void CalculateAssignDevice(ulonglong4* vector_dev_ptr, decltype(ulonglong4().x) value) { + ulonglong4 vector_dev = *vector_dev_ptr; + vector_dev %= value; + vector_dev ^= value; + vector_dev |= value; + vector_dev &= value; + vector_dev >>= value; + vector_dev <<= value; + } + void CalculateAssignHost(ulonglong4& vector_host, decltype(ulonglong4().x) value) { + vector_host %= value; + vector_host ^= value; + vector_host |= value; + vector_host &= value; + vector_host >>= value; + vector_host <<= value; + } +)"}; diff --git a/projects/hip-tests/catch/unit/vector_types/negative_macros_common.hh b/projects/hip-tests/catch/unit/vector_types/negative_macros_common.hh new file mode 100644 index 0000000000..2967d44800 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_macros_common.hh @@ -0,0 +1,78 @@ +/* +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 + +#define NEGATE_UNSIGNED_VECTOR_FUNCTIONS(type) \ + __global__ void NegateDevice(type* vector_dev_ptr) { \ + type vector_dev = *vector_dev_ptr; \ + vector_dev = -vector_dev; \ + } \ + void NegateHost(type& vector_host) { vector_host = -vector_host; } + +#define BITWISE_FLOATING_POINT_VECTOR_FUNCTIONS(type) \ + __global__ void BitwiseDevice(type* vector1_dev_ptr, type* vector2_dev_ptr) { \ + type vector1_dev = *vector1_dev_ptr; \ + type vector2_dev = *vector2_dev_ptr; \ + vector1_dev = ~vector1_dev; \ + vector1_dev %= vector2_dev; \ + vector1_dev ^= vector2_dev; \ + vector1_dev |= vector2_dev; \ + vector1_dev &= vector2_dev; \ + vector1_dev >>= vector2_dev; \ + vector1_dev <<= vector2_dev; \ + vector1_dev = vector1_dev ^ vector2_dev; \ + vector1_dev = vector1_dev | vector2_dev; \ + vector1_dev = vector1_dev & vector2_dev; \ + vector1_dev = vector1_dev >> vector2_dev; \ + vector1_dev = vector1_dev << vector2_dev; \ + } \ + void BitwiseHost(type& vector1_host, type& vector2_host) { \ + vector1_host = ~vector1_host; \ + vector1_host %= vector2_host; \ + vector1_host ^= vector2_host; \ + vector1_host |= vector2_host; \ + vector1_host &= vector2_host; \ + vector1_host >>= vector2_host; \ + vector1_host <<= vector2_host; \ + vector1_host = vector1_host ^ vector2_host; \ + vector1_host = vector1_host | vector2_host; \ + vector1_host = vector1_host & vector2_host; \ + vector1_host = vector1_host >> vector2_host; \ + vector1_host = vector1_host << vector2_host; \ + } + +#define CALCULATE_ASSIGN_VECTOR_FUNCTIONS(type) \ + __global__ void CalculateAssignDevice(type* vector_dev_ptr, decltype(type().x) value) { \ + type vector_dev = *vector_dev_ptr; \ + vector_dev %= value; \ + vector_dev ^= value; \ + vector_dev |= value; \ + vector_dev &= value; \ + vector_dev >>= value; \ + vector_dev <<= value; \ + } \ + void CalculateAssignHost(type& vector_host, decltype(type().x) value) { \ + vector_host %= value; \ + vector_host ^= value; \ + vector_host |= value; \ + vector_host &= value; \ + vector_host >>= value; \ + vector_host <<= value; \ + } diff --git a/projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned.cc b/projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned.cc new file mode 100644 index 0000000000..c8bd39c7a8 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned.cc @@ -0,0 +1,41 @@ +/* +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 "negative_macros_common.hh" + +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uchar1) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uchar2) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uchar3) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uchar4) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ushort1) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ushort2) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ushort3) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ushort4) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uint1) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uint2) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uint3) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(uint4) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulong1) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulong2) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulong3) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulong4) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulonglong1) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulonglong2) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulonglong3) +NEGATE_UNSIGNED_VECTOR_FUNCTIONS(ulonglong4) diff --git a/projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned_rtc.hh b/projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned_rtc.hh new file mode 100644 index 0000000000..c393570a09 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/negative_negate_unsigned_rtc.hh @@ -0,0 +1,150 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +static constexpr auto kNegateUnsignedChar{R"( + __global__ void NegateDevice(uchar1* vector_dev_ptr) { + uchar1 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uchar1& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(uchar2* vector_dev_ptr) { + uchar2 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uchar2& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(uchar3* vector_dev_ptr) { + uchar3 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uchar3& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(uchar4* vector_dev_ptr) { + uchar4 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uchar4& vector_host) { vector_host = -vector_host; } +)"}; + +static constexpr auto kNegateUnsignedShort{R"( + __global__ void NegateDevice(ushort1* vector_dev_ptr) { + ushort1 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ushort1& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ushort2* vector_dev_ptr) { + ushort2 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ushort2& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ushort3* vector_dev_ptr) { + ushort3 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ushort3& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ushort4* vector_dev_ptr) { + ushort4 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ushort4& vector_host) { vector_host = -vector_host; } +)"}; + +static constexpr auto kNegateUnsignedInt{R"( + __global__ void NegateDevice(uint1* vector_dev_ptr) { + uint1 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uint1& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(uint2* vector_dev_ptr) { + uint2 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uint2& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(uint3* vector_dev_ptr) { + uint3 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uint3& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(uint4* vector_dev_ptr) { + uint4 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(uint4& vector_host) { vector_host = -vector_host; } +)"}; + +static constexpr auto kNegateUnsignedLong{R"( + __global__ void NegateDevice(ulong1* vector_dev_ptr) { + ulong1 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulong1& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ulong2* vector_dev_ptr) { + ulong2 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulong2& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ulong3* vector_dev_ptr) { + ulong3 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulong3& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ulong4* vector_dev_ptr) { + ulong4 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulong4& vector_host) { vector_host = -vector_host; } +)"}; + +static constexpr auto kNegateUnsignedLongLong{R"( + __global__ void NegateDevice(ulonglong1* vector_dev_ptr) { + ulonglong1 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulonglong1& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ulonglong2* vector_dev_ptr) { + ulonglong2 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulonglong2& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ulonglong3* vector_dev_ptr) { + ulonglong3 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulonglong3& vector_host) { vector_host = -vector_host; } + + __global__ void NegateDevice(ulonglong4* vector_dev_ptr) { + ulonglong4 vector_dev = *vector_dev_ptr; + vector_dev = -vector_dev; + } + void NegateHost(ulonglong4& vector_host) { vector_host = -vector_host; } +)"}; 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 new file mode 100644 index 0000000000..9684e0ecb1 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/vector_operations_common.hh @@ -0,0 +1,338 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vector_types_common.hh" + +enum class VectorOperation { + kIncrementPrefix, + kIncrementPostfix, + kDecrementPrefix, + kDecrementPostfix, + kAddAssign, + kSubtractAssign, + kMultiplyAssign, + kDivideAssign, + kNegate, + kBitwiseNot, + kModuloAssign, + kBitwiseXorAssign, + kBitwiseOrAssign, + kBitwiseAndAssign, + kRightShiftAssign, + kLeftShiftAssign, + kAdd, + kSubtract, + kMultiply, + kDivide, + kEqual, + kNotEqual, + kModulo, + kBitwiseXor, + kBitwiseOr, + kBitwiseAnd, + kRightShift, + kLeftShift +}; + +inline std::string to_string(VectorOperation operation) { + switch (operation) { + case VectorOperation::kIncrementPrefix: + return "increment (prefix)"; + case VectorOperation::kIncrementPostfix: + return "increment (postfix)"; + case VectorOperation::kDecrementPrefix: + return "decrement (prefix)"; + case VectorOperation::kDecrementPostfix: + return "decrement (postfix)"; + case VectorOperation::kAddAssign: + return "add and assign"; + case VectorOperation::kSubtractAssign: + return "subtract and assign"; + case VectorOperation::kMultiplyAssign: + return "multiply and assign"; + case VectorOperation::kDivideAssign: + return "divide and assign"; + case VectorOperation::kNegate: + return "negate"; + case VectorOperation::kBitwiseNot: + return "bitwise not"; + case VectorOperation::kModuloAssign: + return "modulo and assign"; + case VectorOperation::kBitwiseXorAssign: + return "bitwise XOR and assign"; + case VectorOperation::kBitwiseOrAssign: + return "bitwise OR and assign"; + case VectorOperation::kBitwiseAndAssign: + return "bitwise AND and assign"; + case VectorOperation::kRightShiftAssign: + return "right shift and assign"; + case VectorOperation::kLeftShiftAssign: + return "left shift and assign"; + case VectorOperation::kAdd: + return "add"; + case VectorOperation::kSubtract: + return "subtract"; + case VectorOperation::kMultiply: + return "multiply"; + case VectorOperation::kDivide: + return "divide"; + case VectorOperation::kEqual: + return "equal"; + case VectorOperation::kNotEqual: + return "not equal"; + case VectorOperation::kModulo: + return "modulo"; + case VectorOperation::kBitwiseXor: + return "bitwise XOR"; + case VectorOperation::kBitwiseOr: + return "bitwise OR"; + case VectorOperation::kBitwiseAnd: + return "bitwise AND"; + case VectorOperation::kRightShift: + return "right shift"; + case VectorOperation::kLeftShift: + return "left shift"; + default: + return "Unknown"; + } +} + +template +void SanityCheck(VectorOperation operation, T vector, decltype(T().x) value1, + decltype(T().x) value2) { + if (operation == VectorOperation::kIncrementPrefix) { + ++value1; + } else if (operation == VectorOperation::kIncrementPostfix) { + value1++; + } else if (operation == VectorOperation::kDecrementPrefix) { + --value1; + } else if (operation == VectorOperation::kDecrementPostfix) { + value1--; + } else if (operation == VectorOperation::kAddAssign) { + value1 += value2; + } else if (operation == VectorOperation::kSubtractAssign) { + value1 -= value2; + } else if (operation == VectorOperation::kMultiplyAssign) { + value1 *= value2; + } else if (operation == VectorOperation::kDivideAssign) { + value1 /= value2; + } else if (operation == VectorOperation::kAdd) { + value1 = value1 + value2; + } else if (operation == VectorOperation::kSubtract) { + value1 = value1 - value2; + } else if (operation == VectorOperation::kMultiply) { + value1 = value1 * value2; + } else if (operation == VectorOperation::kDivide) { + value1 = value1 / value2; + } else if (operation == VectorOperation::kEqual) { + value1 = (value1 == value2) ? 2 * value1 : 3 * value1; + } else if (operation == VectorOperation::kNotEqual) { + value1 = (value1 != value2) ? 2 * value1 : 3 * value1; + } else { + if constexpr (std::is_signed_v) { + if (operation == VectorOperation::kNegate) { + value1 = -value1; + } + } + if constexpr (std::is_integral_v) { + if (operation == VectorOperation::kBitwiseNot) { + value1 = ~value1; + } else if (operation == VectorOperation::kModuloAssign) { + value1 %= value2; + } else if (operation == VectorOperation::kBitwiseXorAssign) { + value1 ^= value2; + } else if (operation == VectorOperation::kBitwiseOrAssign) { + value1 |= value2; + } else if (operation == VectorOperation::kBitwiseAndAssign) { + value1 &= value2; + } else if (operation == VectorOperation::kRightShiftAssign) { + value1 >>= value2; + } else if (operation == VectorOperation::kLeftShiftAssign) { + value1 <<= value2; + } else if (operation == VectorOperation::kModulo) { + value1 = value1 % value2; + } else if (operation == VectorOperation::kBitwiseXor) { + value1 = value1 ^ value2; + } else if (operation == VectorOperation::kBitwiseOr) { + value1 = value1 | value2; + } else if (operation == VectorOperation::kBitwiseAnd) { + value1 = value1 & value2; + } else if (operation == VectorOperation::kRightShift) { + value1 = value1 >> value2; + } else if (operation == VectorOperation::kLeftShift) { + value1 = value1 << value2; + } + } + } + SanityCheck(vector, value1); +} + +template +__device__ __host__ void PerformVectorOperation(VectorOperation operation, T* vector1, + const T& vector2) { + if (operation == VectorOperation::kIncrementPrefix) { + ++(*vector1); + } else if (operation == VectorOperation::kIncrementPostfix) { + (*vector1)++; + } else if (operation == VectorOperation::kDecrementPrefix) { + --(*vector1); + } else if (operation == VectorOperation::kDecrementPostfix) { + (*vector1)--; + } else if (operation == VectorOperation::kAddAssign) { + *vector1 += vector2; + } else if (operation == VectorOperation::kSubtractAssign) { + *vector1 -= vector2; + } else if (operation == VectorOperation::kMultiplyAssign) { + *vector1 *= vector2; + } else if (operation == VectorOperation::kDivideAssign) { + *vector1 /= vector2; + } else if (operation == VectorOperation::kAdd) { + *vector1 = *vector1 + vector2; + } else if (operation == VectorOperation::kSubtract) { + *vector1 = *vector1 - vector2; + } else if (operation == VectorOperation::kMultiply) { + *vector1 = *vector1 * vector2; + } else if (operation == VectorOperation::kDivide) { + *vector1 = *vector1 / vector2; + } else if (operation == VectorOperation::kEqual) { + *vector1 = (*vector1 == vector2) ? 2 * *vector1 : 3 * *vector1; + } else if (operation == VectorOperation::kNotEqual) { + *vector1 = (*vector1 != vector2) ? 2 * *vector1 : 3 * *vector1; + } else { + if constexpr (std::is_signed_v) { + if (operation == VectorOperation::kNegate) { + *vector1 = -(*vector1); + } + } + if constexpr (std::is_integral_v) { + if (operation == VectorOperation::kBitwiseNot) { + *vector1 = ~(*vector1); + } else if (operation == VectorOperation::kModuloAssign) { + *vector1 %= vector2; + } else if (operation == VectorOperation::kBitwiseXorAssign) { + *vector1 ^= vector2; + } else if (operation == VectorOperation::kBitwiseOrAssign) { + *vector1 |= vector2; + } else if (operation == VectorOperation::kBitwiseAndAssign) { + *vector1 &= vector2; + } else if (operation == VectorOperation::kRightShiftAssign) { + *vector1 >>= vector2; + } else if (operation == VectorOperation::kLeftShiftAssign) { + *vector1 <<= vector2; + } else if (operation == VectorOperation::kModulo) { + *vector1 = *vector1 % vector2; + } else if (operation == VectorOperation::kBitwiseXor) { + *vector1 = *vector1 ^ vector2; + } else if (operation == VectorOperation::kBitwiseOr) { + *vector1 = *vector1 | vector2; + } else if (operation == VectorOperation::kBitwiseAnd) { + *vector1 = *vector1 & vector2; + } else if (operation == VectorOperation::kRightShift) { + *vector1 = *vector1 >> vector2; + } else if (operation == VectorOperation::kLeftShift) { + *vector1 = *vector1 << vector2; + } + } + } +} + +template +__device__ __host__ void PerformVectorOperation(VectorOperation operation, T* vector, + decltype(T().x) value) { + if (operation == VectorOperation::kAddAssign) { + *vector += value; + } else if (operation == VectorOperation::kSubtractAssign) { + *vector -= value; + } else if (operation == VectorOperation::kMultiplyAssign) { + *vector *= value; + } else if (operation == VectorOperation::kDivideAssign) { + *vector /= value; + } else if (operation == VectorOperation::kAdd) { + *vector = *vector + value; + } else if (operation == VectorOperation::kSubtract) { + *vector = *vector - value; + } else if (operation == VectorOperation::kMultiply) { + *vector = *vector * value; + } else if (operation == VectorOperation::kDivide) { + *vector = *vector / value; + } else if (operation == VectorOperation::kEqual) { + *vector = (*vector == value) ? 2 * *vector : 3 * *vector; + } else if (operation == VectorOperation::kNotEqual) { + *vector = (*vector != value) ? 2 * *vector : 3 * *vector; + } else { + if constexpr (std::is_integral_v) { + if (operation == VectorOperation::kModulo) { + *vector = *vector % value; + } else if (operation == VectorOperation::kBitwiseXor) { + *vector = *vector ^ value; + } else if (operation == VectorOperation::kBitwiseOr) { + *vector = *vector | value; + } else if (operation == VectorOperation::kBitwiseAnd) { + *vector = *vector & value; + } else if (operation == VectorOperation::kRightShift) { + *vector = *vector >> value; + } else if (operation == VectorOperation::kLeftShift) { + *vector = *vector << value; + } + } + } +} + +template +T PerformVectorOperationHost(VectorOperation operation, decltype(T().x) value1, + decltype(T().x) value2) { + T vector1{}; + MakeVectorType(&vector1, value1); + + if constexpr (two_vectors) { + T vector2{}; + MakeVectorType(&vector2, value2); + PerformVectorOperation(operation, &vector1, vector2); + } else { + PerformVectorOperation(operation, &vector1, value2); + } + + return vector1; +} + +template +__global__ void VectorOperationKernel(VectorOperation operation, T* vector1, decltype(T().x) value1, + decltype(T().x) value2) { + MakeVectorType(vector1, value1); + if constexpr (two_vectors) { + T vector2{}; + MakeVectorType(&vector2, value2); + PerformVectorOperation(operation, vector1, vector2); + } else { + PerformVectorOperation(operation, vector1, value2); + } +} + +template +T PerformVectorOperationDevice(VectorOperation operation, decltype(T().x) value1, + decltype(T().x) value2) { + T vector_h{}; + T* vector_d; + HIP_CHECK(hipMalloc(&vector_d, sizeof(T))); + HIP_CHECK(hipMemcpy(vector_d, &vector_h, sizeof(T), hipMemcpyHostToDevice)); + VectorOperationKernel<<<1, 1, 0, 0>>>(operation, vector_d, value1, value2); + HIP_CHECK(hipMemcpy(&vector_h, vector_d, sizeof(T), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(vector_d)); + return vector_h; +} diff --git a/projects/hip-tests/catch/unit/vector_types/vector_types.cc b/projects/hip-tests/catch/unit/vector_types/vector_types.cc new file mode 100644 index 0000000000..2a4b5eda0c --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/vector_types.cc @@ -0,0 +1,357 @@ +/* +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 "vector_operations_common.hh" +#if HT_AMD +#include "negative_negate_unsigned_rtc.hh" +#include "negative_bitwise_float_double_rtc.hh" +#include "negative_calculate_assign_with_value_rtc.hh" +#endif + +/** + * @addtogroup make_vector make_vector + * @{ + * @ingroup VectorTypeTest + */ + +/** + * Test Description + * ------------------------ + * - Creates vectors for all supported types: + * -# make_char1, make_char2, make_char3, make_char4 + * -# make_uchar1, make_uchar2, make_uchar3, make_uchar4 + * -# make_short1, make_short2, make_short3, make_short4 + * -# make_ushort1, make_ushort2, make_ushort3, make_ushort4 + * -# make_int1, make_int2, make_int3, make_int4 + * -# make_uint1, make_uint2, make_uint4, make_uint4 + * -# make_long1, make_long2, make_long3, make_long4 + * -# make_ulong1, make_ulong2, make_ulong3, make_ulong4 + * -# make_longlong1, make_longlong2, make_longlong3, make_longlong4 + * -# make_ulonglong1, make_ulonglong2, make_ulonglong3, make_ulonglong4 + * -# make_float1, make_float2, make_float3, make_float4 + * -# make_double1, make_double2, make_double3, make_double4 + * - Checks that each vector type is created as expected + * - Calls make function from the host side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_make_vector_SanityCheck_Basic_Host", "", char1, uchar1, char2, uchar2, + char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, short3, ushort3, + short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, uint4, long1, + ulong1, long2, ulong2, long3, ulong3, long4, ulong4, longlong1, ulonglong1, + longlong2, ulonglong2, longlong3, ulonglong3, longlong4, ulonglong4, float1, + float2, float3, float4, double1, double2, double3, double4) { + auto value = GetTestValue(0); + TestType vector = MakeVectorTypeHost(value); + SanityCheck(vector, value); +} + +/** + * Test Description + * ------------------------ + * - Creates vectors for all supported types: + * -# make_char1, make_char2, make_char3, make_char4 + * -# make_uchar1, make_uchar2, make_uchar3, make_uchar4 + * -# make_short1, make_short2, make_short3, make_short4 + * -# make_ushort1, make_ushort2, make_ushort3, make_ushort4 + * -# make_int1, make_int2, make_int3, make_int4 + * -# make_uint1, make_uint2, make_uint4, make_uint4 + * -# make_long1, make_long2, make_long3, make_long4 + * -# make_ulong1, make_ulong2, make_ulong3, make_ulong4 + * -# make_longlong1, make_longlong2, make_longlong3, make_longlong4 + * -# make_ulonglong1, make_ulonglong2, make_ulonglong3, make_ulonglong4 + * -# make_float1, make_float2, make_float3, make_float4 + * -# make_double1, make_double2, make_double3, make_double4 + * - Checks that each vector type is created as expected + * - Calls make function from the device side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_make_vector_SanityCheck_Basic_Device", "", char1, uchar1, char2, uchar2, + char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, short3, ushort3, + short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, uint4, long1, + ulong1, long2, ulong2, long3, ulong3, long4, ulong4, longlong1, ulonglong1, + longlong2, ulonglong2, longlong3, ulonglong3, longlong4, ulonglong4, float1, + float2, float3, float4, double1, double2, double3, double4) { + auto value = GetTestValue(0); + TestType vector = MakeVectorTypeDevice(value); + SanityCheck(vector, value); +} + +#if HT_AMD +/** + * Test Description + * ------------------------ + * - Performs supported operations between all supported vector types + * - Checks that the operators are overloaded as expected by comparing results to the manually + * calculated ones + * - Calls operations from the host side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_VectorAndVectorOperations_SanityCheck_Basic_Host", "", char1, uchar1, + char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, + short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, + uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, longlong1, + ulonglong1, longlong2, ulonglong2, longlong3, ulonglong3, longlong4, ulonglong4, + float1, float2, float3, float4, double1, double2, double3, double4) { + auto value1 = GetTestValue(0); + auto value2 = GetTestValue(1); + + for (const auto operation : {VectorOperation::kIncrementPrefix, + VectorOperation::kIncrementPostfix, + VectorOperation::kDecrementPrefix, + VectorOperation::kDecrementPostfix, + VectorOperation::kAddAssign, + VectorOperation::kSubtractAssign, + VectorOperation::kMultiplyAssign, + VectorOperation::kDivideAssign, + VectorOperation::kNegate, + VectorOperation::kBitwiseNot, + VectorOperation::kModuloAssign, + VectorOperation::kBitwiseXorAssign, + VectorOperation::kBitwiseOrAssign, + VectorOperation::kBitwiseAndAssign, + VectorOperation::kRightShiftAssign, + VectorOperation::kLeftShiftAssign, + VectorOperation::kAdd, + VectorOperation::kSubtract, + VectorOperation::kMultiply, + VectorOperation::kDivide, + VectorOperation::kEqual, + VectorOperation::kNotEqual, + VectorOperation::kModulo, + VectorOperation::kBitwiseXor, + VectorOperation::kBitwiseOr, + VectorOperation::kBitwiseAnd, + VectorOperation::kRightShift, + VectorOperation::kLeftShift}) { + DYNAMIC_SECTION("operation: " << to_string(operation)) { + TestType vector = PerformVectorOperationHost(operation, value1, value2); + SanityCheck(operation, vector, value1, value2); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs supported operations between vector and underlying vector type (scalar) + * - Checks that the operators are overloaded as expected by comparing results to the manually + * calculated ones + * - Calls operations from the host side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_VectorAndValueTypeOperations_SanityCheck_Basic_Host", "", char1, uchar1, + char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, + short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, + uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, longlong1, + ulonglong1, longlong2, ulonglong2, longlong3, ulonglong3, longlong4, ulonglong4, + float1, float2, float3, float4, double1, double2, double3, double4) { + auto value1 = GetTestValue(0); + auto value2 = GetTestValue(1); + + for (const auto operation : + {VectorOperation::kAddAssign, VectorOperation::kSubtractAssign, + VectorOperation::kMultiplyAssign, VectorOperation::kDivideAssign, VectorOperation::kAdd, + VectorOperation::kSubtract, VectorOperation::kMultiply, VectorOperation::kDivide, + VectorOperation::kEqual, VectorOperation::kNotEqual, VectorOperation::kModulo, + VectorOperation::kBitwiseXor, VectorOperation::kBitwiseOr, VectorOperation::kBitwiseAnd, + VectorOperation::kRightShift, VectorOperation::kLeftShift}) { + DYNAMIC_SECTION("operation: " << to_string(operation)) { + TestType vector = PerformVectorOperationHost(operation, value1, value2); + SanityCheck(operation, vector, value1, value2); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs supported operations between all supported vector types + * - Checks that the operators are overloaded as expected by comparing results to the manually + * calculated ones + * - Calls operations from the device side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_VectorAndVectorOperations_SanityCheck_Basic_Device", "", char1, uchar1, + char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, + short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, + uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, longlong1, + ulonglong1, longlong2, ulonglong2, longlong3, ulonglong3, longlong4, ulonglong4, + float1, float2, float3, float4, double1, double2, double3, double4) { + auto value1 = GetTestValue(0); + auto value2 = GetTestValue(1); + + for (const auto operation : {VectorOperation::kIncrementPrefix, + VectorOperation::kIncrementPostfix, + VectorOperation::kDecrementPrefix, + VectorOperation::kDecrementPostfix, + VectorOperation::kAddAssign, + VectorOperation::kSubtractAssign, + VectorOperation::kMultiplyAssign, + VectorOperation::kDivideAssign, + VectorOperation::kNegate, + VectorOperation::kBitwiseNot, + VectorOperation::kModuloAssign, + VectorOperation::kBitwiseXorAssign, + VectorOperation::kBitwiseOrAssign, + VectorOperation::kBitwiseAndAssign, + VectorOperation::kRightShiftAssign, + VectorOperation::kLeftShiftAssign, + VectorOperation::kAdd, + VectorOperation::kSubtract, + VectorOperation::kMultiply, + VectorOperation::kDivide, + VectorOperation::kEqual, + VectorOperation::kNotEqual, + VectorOperation::kModulo, + VectorOperation::kBitwiseXor, + VectorOperation::kBitwiseOr, + VectorOperation::kBitwiseAnd, + VectorOperation::kRightShift, + VectorOperation::kLeftShift}) { + DYNAMIC_SECTION("operation: " << to_string(operation)) { + TestType vector = PerformVectorOperationDevice(operation, value1, value2); + SanityCheck(operation, vector, value1, value2); + } + } +} + +/** + * Test Description + * ------------------------ + * - Performs supported operations between vector and underlying vector type (scalar) + * - Checks that the operators are overloaded as expected by comparing results to the manually + * calculated ones + * - Calls operations from the device side + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_VectorAndValueTypeOperations_SanityCheck_Basic_Device", "", char1, uchar1, + char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, + short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, + uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, longlong1, + ulonglong1, longlong2, ulonglong2, longlong3, ulonglong3, longlong4, ulonglong4, + float1, float2, float3, float4, double1, double2, double3, double4) { + auto value1 = GetTestValue(0); + auto value2 = GetTestValue(1); + + for (const auto operation : + {VectorOperation::kAddAssign, VectorOperation::kSubtractAssign, + VectorOperation::kMultiplyAssign, VectorOperation::kDivideAssign, VectorOperation::kAdd, + VectorOperation::kSubtract, VectorOperation::kMultiply, VectorOperation::kDivide, + VectorOperation::kEqual, VectorOperation::kNotEqual, VectorOperation::kModulo, + VectorOperation::kBitwiseXor, VectorOperation::kBitwiseOr, VectorOperation::kBitwiseAnd, + VectorOperation::kRightShift, VectorOperation::kLeftShift}) { + DYNAMIC_SECTION("operation: " << to_string(operation)) { + TestType vector = PerformVectorOperationDevice(operation, value1, value2); + SanityCheck(operation, vector, value1, value2); + } + } +} + +template void VectorTypesRTCWrapper(const char* program_source) { + hiprtcProgram program{}; + HIPRTC_CHECK(hiprtcCreateProgram(&program, program_source, "vector_types_kernels.cc", 0, nullptr, + nullptr)); + +#if HT_AMD + std::string args = std::string("-ferror-limit=100"); + const char* options[] = {args.c_str()}; + hiprtcResult result{hiprtcCompileProgram(program, 1, options)}; +#else + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; +#endif + + 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}; + + std::string error_message{"error:"}; + + size_t npos_e = log.find(error_message, 0); + while (npos_e != std::string::npos) { + ++error_count; + npos_e = log.find(error_message, npos_e + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_errors_num); +} + +/** + * Test Description + * ------------------------ + * - Compiles kernels and host functions with negative scenarios: + * -# Negate (-) operator on the unsigned vectors + * -# Bitwise operators on the floating-point vectors + * -# Calculate-assign operators that are not supported between vector and scalar + * - Utilizes HIP RTC for compilation + * Test source + * ------------------------ + * - unit/vector_types/vector_types.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_VectorOperators_Negative_Parameters_RTC") { + VectorTypesRTCWrapper<8>(kNegateUnsignedChar); + VectorTypesRTCWrapper<8>(kNegateUnsignedShort); + VectorTypesRTCWrapper<8>(kNegateUnsignedInt); + VectorTypesRTCWrapper<8>(kNegateUnsignedLong); + VectorTypesRTCWrapper<8>(kNegateUnsignedLongLong); + VectorTypesRTCWrapper<96>(kBitwiseFloat); + VectorTypesRTCWrapper<96>(kBitwiseDouble); + VectorTypesRTCWrapper<96>(kCalculateAssignChar); + VectorTypesRTCWrapper<96>(kCalculateAssignShort); + VectorTypesRTCWrapper<96>(kCalculateAssignInt); + VectorTypesRTCWrapper<96>(kCalculateAssignLong); + VectorTypesRTCWrapper<96>(kCalculateAssignLongLong); +} +#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 new file mode 100644 index 0000000000..cbaf3f24a6 --- /dev/null +++ b/projects/hip-tests/catch/unit/vector_types/vector_types_common.hh @@ -0,0 +1,187 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +constexpr auto kIntegerTestValueFirst = 42; +constexpr auto kIntegerTestValueSecond = 4; +constexpr auto kFloatingPointTestValueFirst = 42.125; +constexpr auto kFloatingPointTestValueSecond = 4.875; + +template T GetTestValue(int index) { + if (index == 0) { + return std::is_floating_point_v ? static_cast(kIntegerTestValueFirst) + : static_cast(kFloatingPointTestValueFirst); + } else { + return std::is_floating_point_v ? static_cast(kIntegerTestValueSecond) + : static_cast(kFloatingPointTestValueSecond); + } +} + +template +typename std::enable_if::type SanityCheck( + T vector, decltype(T().x) expected_value) { + REQUIRE(vector.x == expected_value); +} + +template +typename std::enable_if::type SanityCheck( + T vector, decltype(T().x) expected_value) { + REQUIRE(vector.x == expected_value); + REQUIRE(vector.y == expected_value); +} + +template +typename std::enable_if::type SanityCheck( + T vector, decltype(T().x) expected_value) { + REQUIRE(vector.x == expected_value); + REQUIRE(vector.y == expected_value); + REQUIRE(vector.z == expected_value); +} + +template +typename std::enable_if::type SanityCheck( + T vector, decltype(T().x) expected_value) { + REQUIRE(vector.x == expected_value); + REQUIRE(vector.y == expected_value); + REQUIRE(vector.z == expected_value); + REQUIRE(vector.w == expected_value); +} + +template +__host__ __device__ void MakeVectorType(T* vector_ptr, decltype(T().x) value) { + if constexpr (std::is_same_v) { + *vector_ptr = make_char1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uchar1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_char2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uchar2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_char3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uchar3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_char4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uchar4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_short1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ushort1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_short2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ushort2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_short3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ushort3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_short4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ushort4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_int1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uint1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_int2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uint2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_int3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uint3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_int4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_uint4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_long1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulong1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_long2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulong2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_long3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulong3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_long4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulong4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_longlong1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulonglong1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_longlong2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulonglong2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_longlong3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulonglong3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_longlong4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_ulonglong4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_float1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_float2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_float3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_float4(value, value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_double1(value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_double2(value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_double3(value, value, value); + } else if constexpr (std::is_same_v) { + *vector_ptr = make_double4(value, value, value, value); + } +} + +template T MakeVectorTypeHost(decltype(T().x) value) { + T vector{}; + MakeVectorType(&vector, value); + return vector; +} + +template __global__ void VectorTypeKernel(T* vector, decltype(T().x) value) { + MakeVectorType(vector, value); +} + +template T MakeVectorTypeDevice(decltype(T().x) value) { + T vector_h{}; + T* vector_d; + HIP_CHECK(hipMalloc(&vector_d, sizeof(T))); + HIP_CHECK(hipMemcpy(vector_d, &vector_h, sizeof(T), hipMemcpyHostToDevice)); + VectorTypeKernel<<<1, 1, 0, 0>>>(vector_d, value); + HIP_CHECK(hipMemcpy(&vector_h, vector_d, sizeof(T), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(vector_d)); + return vector_h; +}