diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 441b3dc9ed..0f6f626367 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -130,6 +130,17 @@ "Unit_deviceAllocation_InOneThread_AccessInAllThreads", "=== Below test is disabled due to defect EXSWHTEC-347 ===", "Unit_hipPointerSetAttribute_Positive_SyncMemops", + "=== Below 2 tests are disable due to defect EXSWHTEC-356 ===", + "Unit_Device___hisinf2_Accuracy_Positive", + "Unit_Device___hisnan2_Accuracy_Positive", + "Unit_Device___hbequ2_Accuracy_Positive", + "Unit_Device___hne_Accuracy_Positive", + "Unit_Device___hne2_Accuracy_Positive", + "Unit_Device___hbne2_Accuracy_Positive", + "Unit_Device___hbgeu2_Accuracy_Positive", + "Unit_Device___hbgtu2_Accuracy_Positive", + "Unit_Device___hbleu2_Accuracy_Positive", + "Unit_Device___hbltu2_Accuracy_Positive", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index dacd9280c5..1522525fea 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -455,6 +455,17 @@ "Unit_Assert_Positive_Basic_KernelFail", "Unit_StaticAssert_Positive_Basic", "Unit_StaticAssert_Negative_Basic", + "=== Below tests are disabled due to defect EXSWHTEC-356 ===", + "Unit_Device___hisinf2_Accuracy_Positive", + "Unit_Device___hisnan2_Accuracy_Positive", + "Unit_Device___hbequ2_Accuracy_Positive", + "Unit_Device___hne_Accuracy_Positive", + "Unit_Device___hne2_Accuracy_Positive", + "Unit_Device___hbne2_Accuracy_Positive", + "Unit_Device___hbgeu2_Accuracy_Positive", + "Unit_Device___hbgtu2_Accuracy_Positive", + "Unit_Device___hbleu2_Accuracy_Positive", + "Unit_Device___hbltu2_Accuracy_Positive", #endif "End of json" ] diff --git a/projects/hip-tests/catch/unit/math/CMakeLists.txt b/projects/hip-tests/catch/unit/math/CMakeLists.txt index e3490dbb51..b67c251392 100644 --- a/projects/hip-tests/catch/unit/math/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/math/CMakeLists.txt @@ -34,7 +34,8 @@ set(TEST_SRC casting_int_funcs.cc casting_half2_funcs.cc half_precision_math.cc - half_precision_arithmetic.cc + half_precision_arithmetic.cc + half_precision_comparison.cc ) if(HIP_PLATFORM MATCHES "nvidia") @@ -132,3 +133,7 @@ add_test(NAME Unit_Half_Precision_Arithmetic_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} half_precision_arithmetic_negative_kernels.cc 88) +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) diff --git a/projects/hip-tests/catch/unit/math/half_precision_comparison.cc b/projects/hip-tests/catch/unit/math/half_precision_comparison.cc new file mode 100644 index 0000000000..c736054e6d --- /dev/null +++ b/projects/hip-tests/catch/unit/math/half_precision_comparison.cc @@ -0,0 +1,847 @@ +/* +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" + +/** + * @addtogroup HalfPrecisionComparison HalfPrecisionComparison + * @{ + * @ingroup MathTest + */ + +/********** Unary Functions **********/ + +#define MATH_BOOL_UNARY_HP_TEST_DEF(func_name, ref_func) \ + __global__ void func_name##_kernel(bool* const ys, const size_t num_xs, Float16* const xs) { \ + 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] = func_name(xs[i]); \ + } \ + } \ + \ + TEST_CASE("Unit_Device_" #func_name "_Accuracy_Positive") { \ + UnaryHalfPrecisionTest(func_name##_kernel, ref_func, EqValidatorBuilderFactory()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hisinf(x)` for all possible inputs. The results are + * compared against reference function `bool std::isinf(float)`. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_BOOL_UNARY_HP_TEST_DEF(__hisinf, static_cast(std::isinf)) + +static float __hisinf2_ref(float x) { return static_cast(std::isinf(x)); } + +MATH_UNARY_HP_KERNEL_DEF(__hisinf2) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hisinf2(x)` for all possible inputs. The results are + * compared against reference function `float std::isinf(float)`. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_HP_TEST_DEF_IMPL(__hisinf2, __hisinf2_ref, EqValidatorBuilderFactory()); + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hisnan(x)` for all possible inputs. The results are + * compared against reference function `bool std::isnan(float)`. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_BOOL_UNARY_HP_TEST_DEF(__hisnan, static_cast(std::isnan)) + +static float __hisnan2_ref(float x) { return static_cast(std::isnan(x)); } + +MATH_UNARY_HP_KERNEL_DEF(__hisnan2) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hisnan2(x)` for all possible inputs. The results are + * compared against reference function `float std::isnan(float)`. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_HP_TEST_DEF_IMPL(__hisnan2, __hisnan2_ref, EqValidatorBuilderFactory()); + +/********** Binary Functions **********/ + +#define MATH_COMPARISON_HP_TEST_DEF(func_name, ref_func, T, RT, nan_value) \ + __global__ void func_name##_kernel(T* const ys, const size_t num_xs, Float16* const x1s, \ + Float16* 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] = func_name(x1s[i], x2s[i]); \ + } \ + } \ + \ + TEST_CASE("Unit_Device_" #func_name "_Accuracy_Positive") { \ + BinaryFloatingPointTest(func_name##_kernel, ref_func, \ + EqValidatorBuilderFactory()); \ + } + + +template static T __heq_ref(float x1, float x2) { + if (std::isnan(x1) || std::isnan(x2)) { + return static_cast(nan_value); + } + return x1 == x2; +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__heq(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'equal + * to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__heq, __heq_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbeq2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'equal + * to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbeq2, __heq_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hequ(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'equal + * to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hequ, __heq_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbequ2(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against result + * of 'equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbequ2, __heq_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__heq2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'equal + * to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__heq2, __heq_ref, Float16, float, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hequ2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'equal + * to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hequ2, __heq_ref, Float16, float, true) + + +template static T __hne_ref(float x1, float x2) { + if (std::isnan(x1) || std::isnan(x2)) { + return static_cast(nan_value); + } + return x1 != x2; +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hne(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'not + * equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hne, __hne_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbne2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'not + * equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbne2, __hne_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hneu(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'not + * equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hneu, __hne_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbneu2(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against result + * of 'not equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbneu2, __hne_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hne2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'not + * equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hne2, __hne_ref, Float16, float, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hneu2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'not + * equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hneu2, __hne_ref, Float16, float, true) + + +template static T __hge_ref(float x1, float x2) { + if (std::isnan(x1) || std::isnan(x2)) { + return static_cast(nan_value); + } + return x1 >= x2; +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hge(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hge, __hge_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbge2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbge2, __hge_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hgeu(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hgeu, __hge_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbgeu2(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against result + * of 'greater than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbgeu2, __hge_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hge2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hge2, __hge_ref, Float16, float, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hgeu2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hgeu2, __hge_ref, Float16, float, true) + + +template static T __hgt_ref(float x1, float x2) { + if (std::isnan(x1) || std::isnan(x2)) { + return static_cast(nan_value); + } + return x1 > x2; +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hgt(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hgt, __hgt_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbgt2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbgt2, __hgt_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hgtu(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hgtu, __hgt_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbgtu2(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against result + * of 'greater than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbgtu2, __hgt_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hgt2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hgt2, __hgt_ref, Float16, float, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hgtu2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of + * 'greater than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hgtu2, __hgt_ref, Float16, float, true) + + +template static T __hle_ref(float x1, float x2) { + if (std::isnan(x1) || std::isnan(x2)) { + return static_cast(nan_value); + } + return x1 <= x2; +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hle(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hle, __hle_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hble2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hble2, __hle_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hleu(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hleu, __hle_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbleu2(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against result + * of 'less than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbleu2, __hle_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hle2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hle2, __hle_ref, Float16, float, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hleu2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than equal to' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hleu2, __hle_ref, Float16, float, true) + + +template static T __hlt_ref(float x1, float x2) { + if (std::isnan(x1) || std::isnan(x2)) { + return static_cast(nan_value); + } + return x1 < x2; +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hlt(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hlt, __hlt_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hblt2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hblt2, __hlt_ref, bool, bool, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hltu(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hltu, __hlt_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hbltu2(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against result + * of 'less than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hbltu2, __hlt_ref, bool, bool, true) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hlt2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hlt2, __hlt_ref, Float16, float, false) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hltu2(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against result of 'less + * than' relational operator for float operands. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_COMPARISON_HP_TEST_DEF(__hltu2, __hlt_ref, Float16, float, true) + +MATH_BINARY_HP_KERNEL_DEF(__hmax) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hmax(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against reference + * function `float std::fmax(float, float)` + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_BINARY_HP_TEST_DEF_IMPL(__hmax, static_cast(std::fmax), + EqValidatorBuilderFactory()) + +MATH_BINARY_HP_KERNEL_DEF(__hmin) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hmin(x,y)` against a table of difficult values, followed + * by a large number of randomly generated values. The results are compared against reference + * function `float std::fmin(float, float)` + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_BINARY_HP_TEST_DEF_IMPL(__hmin, static_cast(std::fmin), + EqValidatorBuilderFactory()) + +static float __hmax_nan_ref(float x1, float x2) { + if (std::isnan(x1)) + return x1; + else if (std::isnan(x2)) + return x2; + else + return std::fmax(x1, x2); +} + +MATH_BINARY_HP_KERNEL_DEF(__hmax_nan) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hmax_nan(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against + * reference function `float std::fmax(float, float)` with modified result when an operand is nan. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_BINARY_HP_TEST_DEF_IMPL(__hmax_nan, __hmax_nan_ref, EqValidatorBuilderFactory()) + +static float __hmin_nan_ref(float x1, float x2) { + if (std::isnan(x1)) + return x1; + else if (std::isnan(x2)) + return x2; + else + return std::fmin(x1, x2); +} + +MATH_BINARY_HP_KERNEL_DEF(__hmin_nan) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `__hmin_nan(x,y)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against + * reference function `float std::fmin(float, float)` with modified result when an operand is nan. + * + * Test source + * ------------------------ + * - unit/math/half_precision_comparison.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_BINARY_HP_TEST_DEF_IMPL(__hmin_nan, __hmin_nan_ref, EqValidatorBuilderFactory()) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/math/half_precision_comparison_negative_kernels.cc b/projects/hip-tests/catch/unit/math/half_precision_comparison_negative_kernels.cc new file mode 100644 index 0000000000..a045af211c --- /dev/null +++ b/projects/hip-tests/catch/unit/math/half_precision_comparison_negative_kernels.cc @@ -0,0 +1,120 @@ +/* +Copyright (c) 2022 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 UNARY_BOOL_HALF_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(__half* x) { bool result = func_name(x); } \ + __global__ void func_name##_kernel_v2(Dummy x) { bool result = func_name(x); } + +#define BINARY_BOOL_HALF_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(__half* x, __half y) { bool result = func_name(x, y); } \ + __global__ void func_name##_kernel_v2(__half x, __half* y) { bool result = func_name(x, y); } \ + __global__ void func_name##_kernel_v3(Dummy x, __half y) { bool result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(__half x, Dummy y) { bool result = func_name(x, y); } + + +#define BINARY_HALF_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(__half* x, __half y) { __half result = func_name(x, y); } \ + __global__ void func_name##_kernel_v2(__half x, __half* y) { __half result = func_name(x, y); } \ + __global__ void func_name##_kernel_v3(Dummy x, __half y) { __half result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(__half x, Dummy y) { __half result = func_name(x, y); } + + +UNARY_BOOL_HALF_NEGATIVE_KERNELS(__hisinf) +UNARY_BOOL_HALF_NEGATIVE_KERNELS(__hisnan) + +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__heq) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hequ) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hne) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hneu) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hge) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hgeu) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hgt) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hgtu) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hle) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hleu) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hlt) +BINARY_BOOL_HALF_NEGATIVE_KERNELS(__hltu) + +BINARY_HALF_NEGATIVE_KERNELS(__hmax) +BINARY_HALF_NEGATIVE_KERNELS(__hmax_nan) +BINARY_HALF_NEGATIVE_KERNELS(__hmin) +BINARY_HALF_NEGATIVE_KERNELS(__hmin_nan) + + +#define UNARY_HALF2_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(__half2* x) { __half2 result = func_name(x); } \ + __global__ void func_name##_kernel_v2(Dummy x) { __half2 result = func_name(x); } + +#define BINARY_HALF2_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(__half2* x, __half2 y) { \ + __half2 result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v2(__half2 x, __half2* y) { \ + __half2 result = func_name(x, y); \ + } \ + __global__ void func_name##_kernel_v3(Dummy x, __half2 y) { __half2 result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(__half2 x, Dummy y) { __half2 result = func_name(x, y); } + +#define BINARY_BOOL_HALF2_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(__half2* x, __half2 y) { bool result = func_name(x, y); } \ + __global__ void func_name##_kernel_v2(__half2 x, __half2* y) { bool result = func_name(x, y); } \ + __global__ void func_name##_kernel_v3(Dummy x, __half2 y) { bool result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(__half2 x, Dummy y) { bool result = func_name(x, y); } + +UNARY_HALF2_NEGATIVE_KERNELS(__hisinf2) +UNARY_HALF2_NEGATIVE_KERNELS(__hisnan2) + +BINARY_HALF2_NEGATIVE_KERNELS(__heq2) +BINARY_HALF2_NEGATIVE_KERNELS(__hequ2) +BINARY_HALF2_NEGATIVE_KERNELS(__hne2) +BINARY_HALF2_NEGATIVE_KERNELS(__hneu2) +BINARY_HALF2_NEGATIVE_KERNELS(__hge2) +BINARY_HALF2_NEGATIVE_KERNELS(__hgeu2) +BINARY_HALF2_NEGATIVE_KERNELS(__hgt2) +BINARY_HALF2_NEGATIVE_KERNELS(__hgtu2) +BINARY_HALF2_NEGATIVE_KERNELS(__hle2) +BINARY_HALF2_NEGATIVE_KERNELS(__hleu2) +BINARY_HALF2_NEGATIVE_KERNELS(__hlt2) +BINARY_HALF2_NEGATIVE_KERNELS(__hltu2) + +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbeq2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbequ2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbne2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbneu2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbge2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbgeu2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbgt2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbgtu2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hble2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbleu2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hblt2) +BINARY_BOOL_HALF2_NEGATIVE_KERNELS(__hbltu2) \ No newline at end of file