diff --git a/projects/hip-tests/catch/CMakeLists.txt b/projects/hip-tests/catch/CMakeLists.txt index f31db5f10f..69d9d9aa15 100644 --- a/projects/hip-tests/catch/CMakeLists.txt +++ b/projects/hip-tests/catch/CMakeLists.txt @@ -173,7 +173,7 @@ if (WIN32) endif() if(HIP_PLATFORM STREQUAL "amd") - add_compile_options(-Wall -Wextra -Wvla -Werror -Wno-deprecated -Wno-option-ignored) + add_compile_options(-Wall -Wextra -Wvla -Werror -Wno-deprecated -Wno-option-ignored -Wno-unused-parameter -Wunused-variable) endif() cmake_policy(PUSH) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 1f38291d20..cf72d409d4 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -131,6 +131,9 @@ "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", + "=== Below 2 tests are disable due to defect EXSWHTEC-369 ===", + "Unit_Device_ilogbf_Accuracy_Positive", + "Unit_Device_ilogb_Accuracy_Positive", "Unit_hipMemAddressFree_negative", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipMemAddressReserve_Negative", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 74c5bca32a..b3b396f6ba 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -222,6 +222,9 @@ "NOTE: The following test is disabled due to defect - EXSWHTEC-244", "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters", "Unit_hipMemAddressFree_negative", + "=== Below 2 tests are disable due to defect EXSWHTEC-369 ===", + "Unit_Device_ilogbf_Accuracy_Positive", + "Unit_Device_ilogb_Accuracy_Positive", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", "Unit_hipMemCreate_ChkWithKerLaunch", diff --git a/projects/hip-tests/catch/unit/math/CMakeLists.txt b/projects/hip-tests/catch/unit/math/CMakeLists.txt index dba9476a0e..a194675b4b 100644 --- a/projects/hip-tests/catch/unit/math/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/math/CMakeLists.txt @@ -27,6 +27,7 @@ set(TEST_SRC integer_intrinsics.cc root_funcs.cc pow_funcs.cc + log_funcs.cc ) if(HIP_PLATFORM MATCHES "nvidia") @@ -91,3 +92,7 @@ add_test(NAME Unit_Device_pow_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} math_pow_negative_kernels.cc 76) +add_test(NAME Unit_Device_log_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + math_log_negative_kernels.cc 24) diff --git a/projects/hip-tests/catch/unit/math/log_funcs.cc b/projects/hip-tests/catch/unit/math/log_funcs.cc new file mode 100644 index 0000000000..83ec1806f3 --- /dev/null +++ b/projects/hip-tests/catch/unit/math/log_funcs.cc @@ -0,0 +1,260 @@ +/* +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 "unary_common.hh" +#include "math_log_negative_kernels_rtc.hh" + +/** + * @addtogroup LogMathFuncs LogMathFuncs + * @{ + * @ingroup MathTest + */ + +/********** Unary Functions **********/ + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `logf(x)` for all possible inputs and `log(x)` against a + * table of difficult values, followed by a large number of randomly generated values. The results + * are compared against reference function `T std::log(T)`. The maximum ulp error is 1. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log, 1, 1) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for logf and log. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_log_logf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog); } + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `log2f(x)` for all possible inputs and `log2(x)` against a + * table of difficult values, followed by a large number of randomly generated values. The results + * are compared against reference function `T std::log2(T)`. The maximum ulp error is 1. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log2, 1, 1) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for log2f and log2. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_log2_log2f_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog2); } + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `log10f(x)` for all possible inputs and `log10(x)` against a + * table of difficult values, followed by a large number of randomly generated values. The results + * are compared against reference function `T std::log10(T)`. The maximum ulp error for single + * precision is 2 and for double precision is 1. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log10, 2, 1) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for log10f and log10. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_log10_log10f_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog10); } + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `log1pf(x)` for all possible inputs and `log1p(x)` against a + * table of difficult values, followed by a large number of randomly generated values. The results + * are compared against reference function `T std::log1p(T)`. The maximum ulp error is 1. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log1p, 1, 1) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for log1pf and log1p. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_log1p_log1pf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog1p); } + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `logb(x)` for all possible inputs and `logb(x)` against a + * table of difficult values, followed by a large number of randomly generated values. The results + * are compared against reference function `T std::logb(T)`. The maximum ulp error is 0. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(logb, 0, 0) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for logbf and logb. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_logb_logbf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLogb); } + + +template +__global__ void ilogb_kernel(int* const ys, const size_t num_xs, T* 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) { + if constexpr (std::is_same_v) { + ys[i] = ilogbf(xs[i]); + } else if constexpr (std::is_same_v) { + ys[i] = ilogb(xs[i]); + } + } +} + +template int ilogb_ref(T arg) { + if (arg == 0) { + return std::numeric_limits::min(); + } else if (std::isnan(arg)) { + return std::numeric_limits::min(); + } else if (std::isinf(arg)) { + return std::numeric_limits::max(); + } else { + return std::ilogb(arg); + } +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `ilogbf(x)` for all possible inputs. The results are + * compared against reference function `int std::ilogb(double)`. The maximum ulp error is 0. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_ilogbf_Accuracy_Positive") { + UnarySinglePrecisionTest(ilogb_kernel, ilogb_ref, + EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `ilogb(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `int std::ilogb(long double)`. The maximum ulp error is 0. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_ilogb_Accuracy_Positive") { + UnaryDoublePrecisionTest(ilogb_kernel, ilogb_ref, + EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for ilogbf and ilogb. + * + * Test source + * ------------------------ + * - unit/math/log_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_ilogb_ilogbf_Negative_RTC") { NegativeTestRTCWrapper<4>(kIlogb); } diff --git a/projects/hip-tests/catch/unit/math/math_log_negative_kernels.cc b/projects/hip-tests/catch/unit/math/math_log_negative_kernels.cc new file mode 100644 index 0000000000..732fc62a08 --- /dev/null +++ b/projects/hip-tests/catch/unit/math/math_log_negative_kernels.cc @@ -0,0 +1,39 @@ +/* +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(func_name) \ + __global__ void func_name##_kernel_v1(double* x) { double result = func_name(x); } \ + __global__ void func_name##_kernel_v2(Dummy x) { double result = func_name(x); } \ + __global__ void func_name##f_kernel_v1(float* x) { float result = func_name##f(x); } \ + __global__ void func_name##f_kernel_v2(Dummy x) { float result = func_name##f(x); } + +NEGATIVE_KERNELS_SHELL(log) +NEGATIVE_KERNELS_SHELL(log2) +NEGATIVE_KERNELS_SHELL(log10) +NEGATIVE_KERNELS_SHELL(log1p) +NEGATIVE_KERNELS_SHELL(logb) +NEGATIVE_KERNELS_SHELL(ilogb) diff --git a/projects/hip-tests/catch/unit/math/math_log_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/math/math_log_negative_kernels_rtc.hh new file mode 100644 index 0000000000..fd1cbdfcaf --- /dev/null +++ b/projects/hip-tests/catch/unit/math/math_log_negative_kernels_rtc.hh @@ -0,0 +1,96 @@ +/* +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 math log negative Test Cases that are using RTC. +*/ + +static constexpr auto kLog{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void log_kernel_v1(double* x) { double result = log(x); } + __global__ void log_kernel_v2(Dummy x) { double result = log(x); } + __global__ void logf_kernel_v1(float* x) { float result = logf(x); } + __global__ void logf_kernel_v2(Dummy x) { float result = logf(x); } +)"}; + +static constexpr auto kLog2{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void log2_kernel_v1(double* x) { double result = log2(x); } + __global__ void log2_kernel_v2(Dummy x) { double result = log2(x); } + __global__ void log2f_kernel_v1(float* x) { float result = log2f(x); } + __global__ void log2f_kernel_v2(Dummy x) { float result = log2f(x); } +)"}; + +static constexpr auto kLog10{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void log10_kernel_v1(double* x) { double result = log10(x); } + __global__ void log10_kernel_v2(Dummy x) { double result = log10(x); } + __global__ void log10f_kernel_v1(float* x) { float result = log10f(x); } + __global__ void log10f_kernel_v2(Dummy x) { float result = log10f(x); } +)"}; + +static constexpr auto kLog1p{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void log1p_kernel_v1(double* x) { double result = log1p(x); } + __global__ void log1p_kernel_v2(Dummy x) { double result = log1p(x); } + __global__ void log1pf_kernel_v1(float* x) { float result = log1pf(x); } + __global__ void log1pf_kernel_v2(Dummy x) { float result = log1pf(x); } +)"}; + +static constexpr auto kLogb{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void logb_kernel_v1(double* x) { double result = logb(x); } + __global__ void logb_kernel_v2(Dummy x) { double result = logb(x); } + __global__ void logbf_kernel_v1(float* x) { float result = logbf(x); } + __global__ void logbf_kernel_v2(Dummy x) { float result = logbf(x); } +)"}; + +static constexpr auto kIlogb{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void ilogb_kernel_v1(double* x) { double result = ilogb(x); } + __global__ void ilogb_kernel_v2(Dummy x) { double result = ilogb(x); } + __global__ void ilogbf_kernel_v1(float* x) { float result = ilogbf(x); } + __global__ void ilogbf_kernel_v2(Dummy x) { float result = ilogbf(x); } +)"}; diff --git a/projects/hip-tests/catch/unit/math/ternary_common.hh b/projects/hip-tests/catch/unit/math/ternary_common.hh index a335073916..4bc7fe26cc 100644 --- a/projects/hip-tests/catch/unit/math/ternary_common.hh +++ b/projects/hip-tests/catch/unit/math/ternary_common.hh @@ -139,4 +139,6 @@ void TernaryFloatingPointTest(kernel_sig kernel, ref_sig, ref, \ ULPValidatorBuilderFactory(ulp)); \ + \ } +