diff --git a/catch/unit/math/CMakeLists.txt b/catch/unit/math/CMakeLists.txt index b67c251392..d3d06e09f4 100644 --- a/catch/unit/math/CMakeLists.txt +++ b/catch/unit/math/CMakeLists.txt @@ -36,6 +36,8 @@ set(TEST_SRC half_precision_math.cc half_precision_arithmetic.cc half_precision_comparison.cc + casting_half2int_funcs.cc + casting_int2half_funcs.cc ) if(HIP_PLATFORM MATCHES "nvidia") @@ -137,3 +139,12 @@ add_test(NAME Unit_Half_Precision_Comparison_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} half_precision_comparison_negative_kernels.cc 168) +add_test(NAME Unit_Device_casting_half2int_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + casting_half2int_negative_kernels.cc 78) + +add_test(NAME Unit_Device_casting_int2half_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + casting_int2half_negative_kernels.cc 78) diff --git a/catch/unit/math/casting_half2int_funcs.cc b/catch/unit/math/casting_half2int_funcs.cc new file mode 100644 index 0000000000..77c32fcb9c --- /dev/null +++ b/catch/unit/math/casting_half2int_funcs.cc @@ -0,0 +1,440 @@ +/* +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 "half_precision_common.hh" +#include "casting_common.hh" + +/** + * @addtogroup HalfPrecisionCastingIntTypes HalfPrecisionCastingIntTypes + * @{ + * @ingroup MathTest + */ + +#define CAST_HALF2INT_RN_TEST_DEF(kern_name, T) \ + CAST_KERNEL_DEF(kern_name, T, Float16) \ + CAST_F2I_RZ_REF_DEF(kern_name, T, Float16) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive") { \ + T (*ref)(Float16) = kern_name##_ref; \ + CastUnaryHalfPrecisionTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2int_rn` for all possible inputs. The results are compared against + * reference function which performs __half cast to int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2int_rn, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2int_rz` for all possible inputs. The results are compared against + * reference function which performs __half cast to int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2int_rz, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2int_rd` for all possible inputs. The results are compared against + * reference function which performs __half cast to int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2int_rd, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2int_ru` for all possible inputs. The results are compared against + * reference function which performs __half cast to int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2int_ru, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2uint_rn` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2uint_rn, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2uint_rz` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2uint_rz, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2uint_rd` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2uint_rd, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2uint_ru` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned int. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2uint_ru, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2short_rn` for all possible inputs. The results are compared + * against reference function which performs __half cast to short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2short_rn, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2short_rz` for all possible inputs. The results are compared + * against reference function which performs __half cast to short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2short_rz, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2short_rd` for all possible inputs. The results are compared + * against reference function which performs __half cast to short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2short_rd, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2short_ru` for all possible inputs. The results are compared + * against reference function which performs __half cast to short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2short_ru, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ushort_rn` for all possible inputs. The results are compared + * against reference function which performs __half cast to unsigned short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ushort_rn, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ushort_rz` for all possible inputs. The results are compared + * against reference function which performs __half cast to unsigned short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ushort_rz, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ushort_rd` for all possible inputs. The results are compared + * against reference function which performs __half cast to unsigned short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ushort_rd, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ushort_ru` for all possible inputs. The results are compared + * against reference function which performs __half cast to unsigned short. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ushort_ru, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ll_rn` for all possible inputs. The results are compared against + * reference function which performs __half cast to long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ll_rn, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ll_rz` for all possible inputs. The results are compared against + * reference function which performs __half cast to long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ll_rz, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ll_rd` for all possible inputs. The results are compared against + * reference function which performs __half cast to long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ll_rd, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ll_ru` for all possible inputs. The results are compared against + * reference function which performs __half cast to long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ll_ru, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ull_rn` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ull_rn, unsigned long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ull_rz` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ull_rz, unsigned long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ull_rd` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ull_rd, unsigned long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half2ull_ru` for all possible inputs. The results are compared against + * reference function which performs __half cast to unsigned long long. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_HALF2INT_RN_TEST_DEF(__half2ull_ru, unsigned long long) + +CAST_KERNEL_DEF(__half_as_short, short, Float16) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half_as_short` for all possible inputs. The results are compared + * against reference function which performs copy of __half value to short variable. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___half_as_short_Accuracy_Positive") { + short (*ref)(Float16) = type2_as_type1_ref; + CastUnaryHalfPrecisionTest(__half_as_short_kernel, ref, EqValidatorBuilderFactory()); +} + +CAST_KERNEL_DEF(__half_as_ushort, unsigned short, Float16) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__half_as_ushort` for all possible inputs. The results are compared + * against reference function which performs copy of __half value to unsigned short variable. + * + * Test source + * ------------------------ + * - unit/math/casting_half2int_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___half_as_ushort_Accuracy_Positive") { + unsigned short (*ref)(Float16) = type2_as_type1_ref; + CastUnaryHalfPrecisionTest(__half_as_ushort_kernel, ref, + EqValidatorBuilderFactory()); +} \ No newline at end of file diff --git a/catch/unit/math/casting_half2int_negative_kernels.cc b/catch/unit/math/casting_half2int_negative_kernels.cc new file mode 100644 index 0000000000..6b7d75040f --- /dev/null +++ b/catch/unit/math/casting_half2int_negative_kernels.cc @@ -0,0 +1,59 @@ +/* +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 +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_KERNELS_SHELL(func_name, T) \ + __global__ void func_name##_kernel_v1(T* result, __half* x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v2(T* result, Dummy x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v3(Dummy* result, __half x) { *result = unc_name(x); } + +NEGATIVE_KERNELS_SHELL(__half2int_rn, int) +NEGATIVE_KERNELS_SHELL(__half2int_rz, int) +NEGATIVE_KERNELS_SHELL(__half2int_rd, int) +NEGATIVE_KERNELS_SHELL(__half2int_ru, int) +NEGATIVE_KERNELS_SHELL(__half2uint_rn, unsigned int) +NEGATIVE_KERNELS_SHELL(__half2uint_rz, unsigned int) +NEGATIVE_KERNELS_SHELL(__half2uint_rd, unsigned int) +NEGATIVE_KERNELS_SHELL(__half2uint_ru, unsigned int) +NEGATIVE_KERNELS_SHELL(__half2short_rn, short) +NEGATIVE_KERNELS_SHELL(__half2short_rz, short) +NEGATIVE_KERNELS_SHELL(__half2short_rd, short) +NEGATIVE_KERNELS_SHELL(__half2short_ru, short) +NEGATIVE_KERNELS_SHELL(__half_as_short, short) +NEGATIVE_KERNELS_SHELL(__half2ushort_rn, unsigned short) +NEGATIVE_KERNELS_SHELL(__half2ushort_rz, unsigned short) +NEGATIVE_KERNELS_SHELL(__half2ushort_rd, unsigned short) +NEGATIVE_KERNELS_SHELL(__half2ushort_ru, unsigned short) +NEGATIVE_KERNELS_SHELL(__half_as_ushort, unsigned short) +NEGATIVE_KERNELS_SHELL(__half2ll_rn, long long) +NEGATIVE_KERNELS_SHELL(__half2ll_rz, long long) +NEGATIVE_KERNELS_SHELL(__half2ll_rd, long long) +NEGATIVE_KERNELS_SHELL(__half2ll_ru, long long) +NEGATIVE_KERNELS_SHELL(__half2ull_rn, unsigned long long) +NEGATIVE_KERNELS_SHELL(__half2ull_rz, unsigned long long) +NEGATIVE_KERNELS_SHELL(__half2ull_rd, unsigned long long) +NEGATIVE_KERNELS_SHELL(__half2ull_ru, unsigned long long) \ No newline at end of file diff --git a/catch/unit/math/casting_int2half_funcs.cc b/catch/unit/math/casting_int2half_funcs.cc new file mode 100644 index 0000000000..d0d404ebc9 --- /dev/null +++ b/catch/unit/math/casting_int2half_funcs.cc @@ -0,0 +1,448 @@ +/* +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 "half_precision_common.hh" +#include "casting_common.hh" + +/** + * @addtogroup HalfPrecisionCastingIntTypes HalfPrecisionCastingIntTypes + * @{ + * @ingroup MathTest + */ + +#define CAST_INT2HALF_RN_TEST_DEF(kern_name, T) \ + CAST_KERNEL_DEF(kern_name, Float16, T) \ + CAST_REF_DEF(kern_name, Float16, T) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive") { \ + Float16 (*ref)(T) = kern_name##_ref; \ + CastIntRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2half_rn` for all possible inputs. The results are compared against + * reference function which performs int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__int2half_rn, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2half_rz` for all possible inputs. The results are compared against + * reference function which performs int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__int2half_rz, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2half_rd` for all possible inputs. The results are compared against + * reference function which performs int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__int2half_rd, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__int2half_ru` for all possible inputs. The results are compared against + * reference function which performs int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__int2half_ru, int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2half_rn` for all possible inputs. The results are compared against + * reference function which performs unsigned int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__uint2half_rn, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2half_rz` for all possible inputs. The results are compared against + * reference function which performs unsigned int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__uint2half_rz, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2half_rd` for all possible inputs. The results are compared against + * reference function which performs unsigned int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__uint2half_rd, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__uint2half_ru` for all possible inputs. The results are compared against + * reference function which performs unsigned int cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__uint2half_ru, unsigned int) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__short2half_rn` for all possible inputs. The results are compared + * against reference function which performs short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__short2half_rn, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__short2half_rz` for all possible inputs. The results are compared + * against reference function which performs short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__short2half_rz, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__short2half_rd` for all possible inputs. The results are compared + * against reference function which performs short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__short2half_rd, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__short2half_ru` for all possible inputs. The results are compared + * against reference function which performs short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__short2half_ru, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ushort2half_rn` for all possible inputs. The results are compared + * against reference function which performs unsigned short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__ushort2half_rn, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ushort2half_rz` for all possible inputs. The results are compared + * against reference function which performs unsigned short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__ushort2half_rz, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ushort2half_rd` for all possible inputs. The results are compared + * against reference function which performs unsigned short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__ushort2half_rd, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ushort2half_ru` for all possible inputs. The results are compared + * against reference function which performs unsigned short cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_INT2HALF_RN_TEST_DEF(__ushort2half_ru, unsigned short) + +#define CAST_LL2HALF_TEST_DEF(kern_name, T) \ + CAST_KERNEL_DEF(kern_name, Float16, T) \ + CAST_REF_DEF(kern_name, Float16, T) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive") { \ + Float16 (*ref)(T) = kern_name##_ref; \ + CastIntBruteForceTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2half_rn` against a large number of randomly generated values. The + * results are compared against reference function which performs long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ll2half_rn, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2half_rz` against a large number of randomly generated values. The + * results are compared against reference function which performs long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ll2half_rz, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2half_rd` against a large number of randomly generated values. The + * results are compared against reference function which performs long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ll2half_rd, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ll2half_ru` against a large number of randomly generated values. The + * results are compared against reference function which performs long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ll2half_ru, long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2half_rn` against a large number of randomly generated values. The + * results are compared against reference function which performs unsigned long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ull2half_rn, unsigned long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2half_rz` against a large number of randomly generated values. The + * results are compared against reference function which performs unsigned long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ull2half_rz, unsigned long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2half_rd` against a large number of randomly generated values. The + * results are compared against reference function which performs unsigned long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ull2half_rd, unsigned long long) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ull2half_ru` against a large number of randomly generated values. The + * results are compared against reference function which performs unsigned long long cast to __half. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_LL2HALF_TEST_DEF(__ull2half_ru, unsigned long long) + +CAST_KERNEL_DEF(__short_as_half, Float16, short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__short_as_half` for all possible inputs. The results are compared + * against reference function which performs copy of short value to __half variable. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___short_as_half_Accuracy_Positive") { + Float16 (*ref)(short) = type2_as_type1_ref; + CastIntBruteForceTest(__short_as_half_kernel, ref, EqValidatorBuilderFactory()); +} + +CAST_KERNEL_DEF(__ushort_as_half, Float16, unsigned short) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__ushort_as_half` for all possible inputs. The results are compared + * against reference function which performs copy of unsigned short value to __half variable. + * + * Test source + * ------------------------ + * - unit/math/casting_int2half_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___ushort_as_half_Accuracy_Positive") { + Float16 (*ref)(unsigned short) = type2_as_type1_ref; + CastIntBruteForceTest(__ushort_as_half_kernel, ref, EqValidatorBuilderFactory()); +} \ No newline at end of file diff --git a/catch/unit/math/casting_int2half_negative_kernels.cc b/catch/unit/math/casting_int2half_negative_kernels.cc new file mode 100644 index 0000000000..a23da476e6 --- /dev/null +++ b/catch/unit/math/casting_int2half_negative_kernels.cc @@ -0,0 +1,59 @@ +/* +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 +#include + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_KERNELS_SHELL(func_name, T) \ + __global__ void func_name##_kernel_v1(__half* result, T* x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v2(__half* result, Dummy x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v3(Dummy* result, T x) { *result = func_name(x); } + +NEGATIVE_KERNELS_SHELL(__int2half_rn, int) +NEGATIVE_KERNELS_SHELL(__int2half_rz, int) +NEGATIVE_KERNELS_SHELL(__int2half_rd, int) +NEGATIVE_KERNELS_SHELL(__int2half_ru, int) +NEGATIVE_KERNELS_SHELL(__uint2half_rn, unsigned int) +NEGATIVE_KERNELS_SHELL(__uint2half_rz, unsigned int) +NEGATIVE_KERNELS_SHELL(__uint2half_rd, unsigned int) +NEGATIVE_KERNELS_SHELL(__uint2half_ru, unsigned int) +NEGATIVE_KERNELS_SHELL(__short2half_rn, short) +NEGATIVE_KERNELS_SHELL(__short2half_rz, short) +NEGATIVE_KERNELS_SHELL(__short2half_rd, short) +NEGATIVE_KERNELS_SHELL(__short2half_ru, short) +NEGATIVE_KERNELS_SHELL(__short_as_half, short) +NEGATIVE_KERNELS_SHELL(__ushort2half_rn, unsigned short) +NEGATIVE_KERNELS_SHELL(__ushort2half_rz, unsigned short) +NEGATIVE_KERNELS_SHELL(__ushort2half_rd, unsigned short) +NEGATIVE_KERNELS_SHELL(__ushort2half_ru, unsigned short) +NEGATIVE_KERNELS_SHELL(__ushort_as_half, unsigned short) +NEGATIVE_KERNELS_SHELL(__ll2half_rn, long long) +NEGATIVE_KERNELS_SHELL(__ll2half_rz, long long) +NEGATIVE_KERNELS_SHELL(__ll2half_rd, long long) +NEGATIVE_KERNELS_SHELL(__ll2half_ru, long long) +NEGATIVE_KERNELS_SHELL(__ull2half_rn, unsigned long long) +NEGATIVE_KERNELS_SHELL(__ull2half_rz, unsigned long long) +NEGATIVE_KERNELS_SHELL(__ull2half_rd, unsigned long long) +NEGATIVE_KERNELS_SHELL(__ull2half_ru, unsigned long long) \ No newline at end of file