From f9cf87fe602d3ff6ec6553823edd5d5484d0657e Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Wed, 24 Jan 2024 22:01:49 +0530 Subject: [PATCH] EXSWHTEC-311 - Implement tests for integer type casting intrinsics #285 Change-Id: I6e6bee38dad6948d46ba2ce0d5d2e3b27c150d35 --- catch/include/cmd_options.hh | 4 + catch/unit/math/CMakeLists.txt | 5 + catch/unit/math/casting_int_funcs.cc | 735 ++++++++++++++++++ .../unit/math/casting_int_negative_kernels.cc | 79 ++ .../math/casting_int_negative_kernels_rtc.hh | 215 +++++ catch/unit/math/math_common.hh | 11 + 6 files changed, 1049 insertions(+) create mode 100644 catch/unit/math/casting_int_funcs.cc create mode 100644 catch/unit/math/casting_int_negative_kernels.cc create mode 100644 catch/unit/math/casting_int_negative_kernels_rtc.hh diff --git a/catch/include/cmd_options.hh b/catch/include/cmd_options.hh index 666f34ea82..71f21006e0 100644 --- a/catch/include/cmd_options.hh +++ b/catch/include/cmd_options.hh @@ -37,3 +37,7 @@ struct CmdOptions { }; extern CmdOptions cmd_options; +<<<<<<< HEAD +======= + +>>>>>>> c08a2a5d (Merge branch 'develop' into casting_int_tests) diff --git a/catch/unit/math/CMakeLists.txt b/catch/unit/math/CMakeLists.txt index b3e5a937d1..973e875abe 100644 --- a/catch/unit/math/CMakeLists.txt +++ b/catch/unit/math/CMakeLists.txt @@ -31,6 +31,7 @@ set(TEST_SRC special_funcs.cc casting_double_funcs.cc casting_float_funcs.cc + casting_int_funcs.cc ) if(HIP_PLATFORM MATCHES "nvidia") @@ -111,3 +112,7 @@ add_test(NAME Unit_Device_casting_float_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} casting_float_negative_kernels.cc 54) +add_test(NAME Unit_Device_casting_int_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + casting_int_negative_kernels.cc 92) diff --git a/catch/unit/math/casting_int_funcs.cc b/catch/unit/math/casting_int_funcs.cc new file mode 100644 index 0000000000..49e8ae7463 --- /dev/null +++ b/catch/unit/math/casting_int_funcs.cc @@ -0,0 +1,735 @@ +/* +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 "casting_common.hh" +#include "casting_int_negative_kernels_rtc.hh" + +/** + * @addtogroup CastingIntTypes CastingIntTypes + * @{ + * @ingroup MathTest + */ + +#define CAST_INT2FLOAT_TEST_DEF(kern_name, T1, T2, round_dir) \ + CAST_KERNEL_DEF(kern_name, T1, T2) \ + CAST_RND_REF_DEF(kern_name, T1, T2, round_dir) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T1 (*ref)(T2) = kern_name##_ref; \ + CastIntRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +#define CAST_INT2FLOAT_RN_TEST_DEF(kern_name, T1, T2) \ + CAST_KERNEL_DEF(kern_name, T1, T2) \ + CAST_REF_DEF(kern_name, T1, T2) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T1 (*ref)(T2) = kern_name##_ref; \ + CastIntRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2float_rd` for all possible inputs. The results are compared against + * reference function which performs cast to float with FE_DOWNWARD rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_TEST_DEF(__int2float_rd, float, int, FE_DOWNWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2float_rn` for all possible inputs. The results are compared against + * reference function which performs cast to float. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_RN_TEST_DEF(__int2float_rn, float, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2float_ru` for all possible inputs. The results are compared against + * reference function which performs cast to float with FE_UPWARD rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_TEST_DEF(__int2float_ru, float, int, FE_UPWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2float_rz` for all possible inputs. The results are compared against + * reference function which performs cast to float with FE_TOWARDZERO rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_TEST_DEF(__int2float_rz, float, int, FE_TOWARDZERO) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __int2float_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_int2float___Negative_RTC") { NegativeTestRTCWrapper<12>(kInt2Float); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2float_rd` for all possible inputs. The results are compared + * against reference function which performs cast to float with FE_DOWNWARD rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_TEST_DEF(__uint2float_rd, float, unsigned int, FE_DOWNWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2float_rn` for all possible inputs. The results are compared + * against reference function which performs cast to float. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_RN_TEST_DEF(__uint2float_rn, float, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2float_ru` for all possible inputs. The results are compared + * against reference function which performs cast to float with FE_UPWARD rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_TEST_DEF(__uint2float_ru, float, unsigned int, FE_UPWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2float_rz` for all possible inputs. The results are compared + * against reference function which performs cast to float with FE_TOWARDZERO rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_TEST_DEF(__uint2float_rz, float, unsigned int, FE_TOWARDZERO) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __uint2float_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___uint2float_Negative_RTC") { NegativeTestRTCWrapper<12>(kUint2Float); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2double_rn` for all possible inputs. The results are compared + * against reference function which performs cast to double. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_RN_TEST_DEF(__int2double_rn, double, int) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __int2double_rn. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___int2double_Negative_RTC") { NegativeTestRTCWrapper<3>(kInt2Double); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2double_rn` for all possible inputs. The results are compared + * against reference function which performs cast to double. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2FLOAT_RN_TEST_DEF(__uint2double_rn, double, unsigned int) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __uint2double_rn. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___uint2double_Negative_RTC") { NegativeTestRTCWrapper<3>(kUint2Double); } + +#define CAST_LL2FLOAT_TEST_DEF(kern_name, T1, T2, round_dir) \ + CAST_KERNEL_DEF(kern_name, T1, T2) \ + CAST_RND_REF_DEF(kern_name, T1, T2, round_dir) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T1 (*ref)(T2) = kern_name##_ref; \ + CastIntBruteForceTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +#define CAST_LL2FLOAT_RN_TEST_DEF(kern_name, T1, T2) \ + CAST_KERNEL_DEF(kern_name, T1, T2) \ + CAST_REF_DEF(kern_name, T1, T2) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T1 (*ref)(T2) = kern_name##_ref; \ + CastIntBruteForceTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2float_rd` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float with FE_DOWNWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ll2float_rd, float, long long int, FE_DOWNWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2float_rn` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_RN_TEST_DEF(__ll2float_rn, float, long long int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2float_ru` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float with FE_UPWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ll2float_ru, float, long long int, FE_UPWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2float_rz` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float with FE_TOWARDZERO + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ll2float_rz, float, long long int, FE_TOWARDZERO) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __ll2float_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___ll2float_Negative_RTC") { NegativeTestRTCWrapper<12>(kLL2Float); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2float_rd` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float with FE_DOWNWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ull2float_rd, float, unsigned long long int, FE_DOWNWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2float_rn` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_RN_TEST_DEF(__ull2float_rn, float, unsigned long long int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2float_ru` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float with FE_UPWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ull2float_ru, float, unsigned long long int, FE_UPWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2float_rz` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to float with FE_TOWARDZERO + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ull2float_rz, float, unsigned long long int, FE_TOWARDZERO) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __ull2float_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___ull2float_Negative_RTC") { NegativeTestRTCWrapper<12>(kULL2Float); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2double_rd` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double with FE_DOWNWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ll2double_rd, double, long long int, FE_DOWNWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2double_rn` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_RN_TEST_DEF(__ll2double_rn, double, long long int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2double_ru` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double with FE_UPWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ll2double_ru, double, long long int, FE_UPWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2double_rz` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double with FE_TOWARDZERO + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ll2double_rz, double, long long int, FE_TOWARDZERO) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __ll2double_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___ll2double_Negative_RTC") { NegativeTestRTCWrapper<12>(kLL2Double); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2double_rd` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double with FE_DOWNWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ull2double_rd, double, unsigned long long int, FE_DOWNWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2double_rn` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_RN_TEST_DEF(__ull2double_rn, double, unsigned long long int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2double_ru` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double with FE_UPWARD + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ull2double_ru, double, unsigned long long int, FE_UPWARD) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2double_rz` against a large number of randomly generated values. The + * results are compared against reference function which performs cast to double with FE_TOWARDZERO + * rounding mode. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2FLOAT_TEST_DEF(__ull2double_rz, double, unsigned long long int, FE_TOWARDZERO) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __ull2double_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___ull2double_Negative_RTC") { NegativeTestRTCWrapper<12>(kULL2Double); } + +CAST_KERNEL_DEF(__int_as_float, float, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int_as_float` for all possible inputs. The results are compared against + * reference function which performs copy of int value to float variable. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___int_as_float_Positive") { + float (*ref)(int) = type2_as_type1_ref; + CastIntRangeTest(__int_as_float_kernel, ref, EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __int_as_float. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___int_as_float_Negative_RTC") { NegativeTestRTCWrapper<3>(kIntAsFloat); } + +CAST_KERNEL_DEF(__uint_as_float, float, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint_as_float` for all possible inputs. The results are compared + * against reference function which performs copy of unsigned int value to float variable. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___uint_as_float_Positive") { + float (*ref)(unsigned int) = type2_as_type1_ref; + CastIntRangeTest(__uint_as_float_kernel, ref, EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __uint_as_float. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___uint_as_float_Negative_RTC") { NegativeTestRTCWrapper<3>(kUintAsFloat); } + +CAST_KERNEL_DEF(__longlong_as_double, double, long long int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__longlong_as_double` against a large number of randomly generated + * values. The results are compared against reference function which performs copy of long long int + * value to double variable. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___longlong_as_double_Positive") { + double (*ref)(long long int) = type2_as_type1_ref; + CastIntBruteForceTest(__longlong_as_double_kernel, ref, EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __longlong_as_double. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___longlong_as_double_Negative_RTC") { + NegativeTestRTCWrapper<3>(kLonglongAsDouble); +} + +__global__ void __hiloint2double_kernel(double* const ys, const size_t num_xs, int* const x1s, + int* const x2s) { + const auto tid = cg::this_grid().thread_rank(); + const auto stride = cg::this_grid().size(); + + for (auto i = tid; i < num_xs; i += stride) { + ys[i] = __hiloint2double(x1s[i], x2s[i]); + } +} + +double __hiloint2double_ref(int hi, int lo) { + uint64_t tmp0 = (static_cast(hi) << 32ull) | static_cast(lo); + double tmp1; + memcpy(&tmp1, &tmp0, sizeof(tmp0)); + + return tmp1; +} + +/** + * Test Description + * ------------------------ + * - Tests that checks `__hiloint2double` for all possible inputs for hi value. The results are + * compared against reference function which performs copy of hi int value to higher part of double + * variable and copy of lo int value to lower part of double variable. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___hiloint2double_Positive") { + double (*ref)(int, int) = __hiloint2double_ref; + CastBinaryIntRangeTest(__hiloint2double_kernel, ref, EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __hiloint2double. + * + * Test source + * ------------------------ + * - unit/math/casting_int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___hiloint2double_Negative_RTC") { NegativeTestRTCWrapper<5>(kHilo2Double); } \ No newline at end of file diff --git a/catch/unit/math/casting_int_negative_kernels.cc b/catch/unit/math/casting_int_negative_kernels.cc new file mode 100644 index 0000000000..3f2586d738 --- /dev/null +++ b/catch/unit/math/casting_int_negative_kernels.cc @@ -0,0 +1,79 @@ +/* +Copyright (c) 2021 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 + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_KERNELS_SHELL_ONE_ARG(func_name, T1, T2) \ + __global__ void func_name##_kernel_v1(T1* result, T2* x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v2(T1* result, Dummy x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v3(Dummy* result, T2 x) { *result = func_name(x); } + +#define NEGATIVE_KERNELS_SHELL_TWO_ARGS(func_name, T1, T2) \ + __global__ void func_name##_kernel_v1(T1* result, T2* x, T2 y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v2(T1* result, T2 x, T2* y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v3(T1* result, Dummy x, T2 y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v4(T1* result, T2 x, Dummy y) { \ + *result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v5(Dummy* result, T2 x, T2 y) { \ + *result = func_name(x, y); \ + } + +NEGATIVE_KERNELS_SHELL_ONE_ARG(__int2float_rd, float, int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__int2float_rn, float, int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__int2float_ru, float, int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__int2float_rz, float, int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__uint2float_rd, float, unsigned int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__uint2float_rn, float, unsigned int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__uint2float_ru, float, unsigned int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__uint2float_rz, float, unsigned int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2float_rd, float, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2float_rn, float, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2float_ru, float, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2float_rz, float, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2float_rd, float, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2float_rn, float, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2float_ru, float, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2float_rz, float, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__int2double_rn, double, int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__uint2double_rn, double, unsigned int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2double_rd, double, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2double_rn, double, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2double_ru, double, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ll2double_rz, double, long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2double_rd, double, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2double_rn, double, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2double_ru, double, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__ull2double_rz, double, unsigned long long int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__int_as_float, float, int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__uint_as_float, float, unsigned int) +NEGATIVE_KERNELS_SHELL_ONE_ARG(__longlong_as_double, double, long long int) +NEGATIVE_KERNELS_SHELL_TWO_ARGS(__hiloint2double, double, int) \ No newline at end of file diff --git a/catch/unit/math/casting_int_negative_kernels_rtc.hh b/catch/unit/math/casting_int_negative_kernels_rtc.hh new file mode 100644 index 0000000000..acdc621f8a --- /dev/null +++ b/catch/unit/math/casting_int_negative_kernels_rtc.hh @@ -0,0 +1,215 @@ +/* +Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +/* +Negative kernels used for the int/long long type casting negative Test Cases that are using RTC. +*/ + +static constexpr auto kInt2Float{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void int2float_rd_kernel_v1(float* result, int* x) { *result = __int2float_rd(x); } + __global__ void int2float_rd_kernel_v2(float* result, Dummy x) { *result = __int2float_rd(x); } + __global__ void int2float_rd_kernel_v3(Dummy* result, int x) { *result = __int2float_rd(x); } + __global__ void int2float_rn_kernel_v1(float* result, int* x) { *result = __int2float_rn(x); } + __global__ void int2float_rn_kernel_v2(float* result, Dummy x) { *result = __int2float_rn(x); } + __global__ void int2float_rn_kernel_v3(Dummy* result, int x) { *result = __int2float_rn(x); } + __global__ void int2float_ru_kernel_v1(float* result, int* x) { *result = __int2float_ru(x); } + __global__ void int2float_ru_kernel_v2(float* result, Dummy x) { *result = __int2float_ru(x); } + __global__ void int2float_ru_kernel_v3(Dummy* result, int x) { *result = __int2float_ru(x); } + __global__ void int2float_rz_kernel_v1(float* result, int* x) { *result = __int2float_rz(x); } + __global__ void int2float_rz_kernel_v2(float* result, Dummy x) { *result = __int2float_rz(x); } + __global__ void int2float_rz_kernel_v3(Dummy* result, int x) { *result = __int2float_rz(x); } +)"}; + +static constexpr auto kUint2Float{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void uint2float_rd_kernel_v1(float* result, unsigned int* x) { *result = __uint2float_rd(x); } + __global__ void uint2float_rd_kernel_v2(float* result, Dummy x) { *result = __uint2float_rd(x); } + __global__ void uint2float_rd_kernel_v3(Dummy* result, unsigned int x) { *result = __uint2float_rd(x); } + __global__ void uint2float_rn_kernel_v1(float* result, unsigned int* x) { *result = __uint2float_rn(x); } + __global__ void uint2float_rn_kernel_v2(float* result, Dummy x) { *result = __uint2float_rn(x); } + __global__ void uint2float_rn_kernel_v3(Dummy* result, unsigned int x) { *result = __uint2float_rn(x); } + __global__ void uint2float_ru_kernel_v1(float* result, unsigned int* x) { *result = __uint2float_ru(x); } + __global__ void uint2float_ru_kernel_v2(float* result, Dummy x) { *result = __uint2float_ru(x); } + __global__ void uint2float_ru_kernel_v3(Dummy* result, unsigned int x) { *result = __uint2float_ru(x); } + __global__ void uint2float_rz_kernel_v1(float* result, unsigned int* x) { *result = __uint2float_rz(x); } + __global__ void uint2float_rz_kernel_v2(float* result, Dummy x) { *result = __uint2float_rz(x); } + __global__ void uint2float_rz_kernel_v3(Dummy* result, unsigned int x) { *result = __uint2float_rz(x); } +)"}; + +static constexpr auto kLL2Float{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void ll2float_rd_kernel_v1(float* result, long long int* x) { *result = __ll2float_rd(x); } + __global__ void ll2float_rd_kernel_v2(float* result, Dummy x) { *result = __ll2float_rd(x); } + __global__ void ll2float_rd_kernel_v3(Dummy* result, long long int x) { *result = __ll2float_rd(x); } + __global__ void ll2float_rn_kernel_v1(float* result, long long int* x) { *result = __ll2float_rn(x); } + __global__ void ll2float_rn_kernel_v2(float* result, Dummy x) { *result = __ll2float_rn(x); } + __global__ void ll2float_rn_kernel_v3(Dummy* result, long long int x) { *result = __ll2float_rn(x); } + __global__ void ll2float_ru_kernel_v1(float* result, long long int* x) { *result = __ll2float_ru(x); } + __global__ void ll2float_ru_kernel_v2(float* result, Dummy x) { *result = __ll2float_ru(x); } + __global__ void ll2float_ru_kernel_v3(Dummy* result, long long int x) { *result = __ll2float_ru(x); } + __global__ void ll2float_rz_kernel_v1(float* result, long long int* x) { *result = __ll2float_rz(x); } + __global__ void ll2float_rz_kernel_v2(float* result, Dummy x) { *result = __ll2float_rz(x); } + __global__ void ll2float_rz_kernel_v3(Dummy* result, long long int x) { *result = __ll2float_rz(x); } +)"}; + +static constexpr auto kULL2Float{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void ull2float_rd_kernel_v1(float* result, unsigned long long int* x) { *result = __ull2float_rd(x); } + __global__ void ull2float_rd_kernel_v2(float* result, Dummy x) { *result = __ull2float_rd(x); } + __global__ void ull2float_rd_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2float_rd(x); } + __global__ void ull2float_rn_kernel_v1(float* result, unsigned long long int* x) { *result = __ull2float_rn(x); } + __global__ void ull2float_rn_kernel_v2(float* result, Dummy x) { *result = __ull2float_rn(x); } + __global__ void ull2float_rn_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2float_rn(x); } + __global__ void ull2float_ru_kernel_v1(float* result, unsigned long long int* x) { *result = __ull2float_ru(x); } + __global__ void ull2float_ru_kernel_v2(float* result, Dummy x) { *result = __ull2float_ru(x); } + __global__ void ull2float_ru_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2float_ru(x); } + __global__ void ull2float_rz_kernel_v1(float* result, unsigned long long int* x) { *result = __ull2float_rz(x); } + __global__ void ull2float_rz_kernel_v2(float* result, Dummy x) { *result = __ull2float_rz(x); } + __global__ void ull2float_rz_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2float_rz(x); } +)"}; + +static constexpr auto kIntAsFloat{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void int_as_float_kernel_v1(float* result, int* x) { *result = __int_as_float(x); } + __global__ void int_as_float_kernel_v2(float* result, Dummy x) { *result = __int_as_float(x); } + __global__ void int_as_float_kernel_v3(Dummy* result, int x) { *result = __int_as_float(x); } +)"}; + +static constexpr auto kUintAsFloat{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void uint_as_float_kernel_v1(float* result, unsigned int* x) { *result = __uint_as_float(x); } + __global__ void uint_as_float_kernel_v2(float* result, Dummy x) { *result = __uint_as_float(x); } + __global__ void uint_as_float_kernel_v3(Dummy* result, unsigned int x) { *result = __uint_as_float(x); } +)"}; + +static constexpr auto kInt2Double{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void int2double_rn_kernel_v1(double* result, int* x) { *result = __int2double_rn(x); } + __global__ void int2double_rn_kernel_v2(double* result, Dummy x) { *result = __int2double_rn(x); } + __global__ void int2double_rn_kernel_v3(Dummy* result, int x) { *result = __int2double_rn(x); } +)"}; + +static constexpr auto kUint2Double{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void uint2double_rn_kernel_v1(double* result, unsigned int* x) { *result = __uint2double_rn(x); } + __global__ void uint2double_rn_kernel_v2(double* result, Dummy x) { *result = __uint2double_rn(x); } + __global__ void uint2double_rn_kernel_v3(Dummy* result, unsigned int x) { *result = __uint2double_rn(x); } +)"}; + + +static constexpr auto kLL2Double{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void ll2double_rd_kernel_v1(double* result, long long int* x) { *result = __ll2double_rd(x); } + __global__ void ll2double_rd_kernel_v2(double* result, Dummy x) { *result = __ll2double_rd(x); } + __global__ void ll2double_rd_kernel_v3(Dummy* result, long long int x) { *result = __ll2double_rd(x); } + __global__ void ll2double_rn_kernel_v1(double* result, long long int* x) { *result = __ll2double_rn(x); } + __global__ void ll2double_rn_kernel_v2(double* result, Dummy x) { *result = __ll2double_rn(x); } + __global__ void ll2double_rn_kernel_v3(Dummy* result, long long int x) { *result = __ll2double_rn(x); } + __global__ void ll2double_ru_kernel_v1(double* result, long long int* x) { *result = __ll2double_ru(x); } + __global__ void ll2double_ru_kernel_v2(double* result, Dummy x) { *result = __ll2double_ru(x); } + __global__ void ll2double_ru_kernel_v3(Dummy* result, long long int x) { *result = __ll2double_ru(x); } + __global__ void ll2double_rz_kernel_v1(double* result, long long int* x) { *result = __ll2double_rz(x); } + __global__ void ll2double_rz_kernel_v2(double* result, Dummy x) { *result = __ll2double_rz(x); } + __global__ void ll2double_rz_kernel_v3(Dummy* result, long long int x) { *result = __ll2double_rz(x); } +)"}; + +static constexpr auto kULL2Double{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void ull2double_rd_kernel_v1(double* result, unsigned long long int* x) { *result = __ull2double_rd(x); } + __global__ void ull2double_rd_kernel_v2(double* result, Dummy x) { *result = __ull2double_rd(x); } + __global__ void ull2double_rd_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2double_rd(x); } + __global__ void ull2double_rn_kernel_v1(double* result, unsigned long long int* x) { *result = __ull2double_rn(x); } + __global__ void ull2double_rn_kernel_v2(double* result, Dummy x) { *result = __ull2double_rn(x); } + __global__ void ull2double_rn_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2double_rn(x); } + __global__ void ull2double_ru_kernel_v1(double* result, unsigned long long int* x) { *result = __ull2double_ru(x); } + __global__ void ull2double_ru_kernel_v2(double* result, Dummy x) { *result = __ull2double_ru(x); } + __global__ void ull2double_ru_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2double_ru(x); } + __global__ void ull2double_rz_kernel_v1(double* result, unsigned long long int* x) { *result = __ull2double_rz(x); } + __global__ void ull2double_rz_kernel_v2(double* result, Dummy x) { *result = __ull2double_rz(x); } + __global__ void ull2double_rz_kernel_v3(Dummy* result, unsigned long long int x) { *result = __ull2double_rz(x); } +)"}; + +static constexpr auto kLonglongAsDouble{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void longlong_as_double_kernel_v1(double* result, long long int* x) { *result = __longlong_as_double(x); } + __global__ void longlong_as_double_kernel_v2(double* result, Dummy x) { *result = __longlong_as_double(x); } + __global__ void longlong_as_double_kernel_v3(Dummy* result, long long int x) { *result = __longlong_as_double(x); } +)"}; + +static constexpr auto kHilo2Double{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void hiloint2double_kernel_v1(double* result, int* x, int y) { *result = __hiloint2double(x, y); } + __global__ void hiloint2double_kernel_v2(double* result, int x, int* y) { *result = __hiloint2double(x, y); } + __global__ void hiloint2double_kernel_v3(double* result, Dummy x, int y) { *result = __hiloint2double(x, y); } + __global__ void hiloint2double_kernel_v4(double* result, int x, Dummy y) { *result = __hiloint2double(x, y); } + __global__ void hiloint2double_kernel_v5(Dummy* result, int x, int y) { *result = __hiloint2double(x, y); } +)"}; + + diff --git a/catch/unit/math/math_common.hh b/catch/unit/math/math_common.hh index 738dbf66c0..0cd30db404 100644 --- a/catch/unit/math/math_common.hh +++ b/catch/unit/math/math_common.hh @@ -7,8 +7,15 @@ 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: +<<<<<<< HEAD The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. +======= + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +>>>>>>> c08a2a5d (Merge branch 'develop' into casting_int_tests) 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 @@ -100,7 +107,11 @@ template class MathTest { template void RunImpl(const ValidatorBuilder& validator_builder, const size_t grid_dim, const size_t block_dim, RT (*const ref_func)(RTs...), const size_t num_args, +<<<<<<< HEAD std::index_sequence is, const Ts*... xss) { +======= + std::index_sequence, const Ts*... xss) { +>>>>>>> c08a2a5d (Merge branch 'develop' into casting_int_tests) const auto xss_tup = std::make_tuple(xss...); constexpr auto f = [](auto dst, auto src, size_t size) {