diff --git a/projects/hip-tests/catch/unit/math/CMakeLists.txt b/projects/hip-tests/catch/unit/math/CMakeLists.txt index a194675b4b..1a9cd98f89 100644 --- a/projects/hip-tests/catch/unit/math/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/math/CMakeLists.txt @@ -28,6 +28,7 @@ set(TEST_SRC root_funcs.cc pow_funcs.cc log_funcs.cc + special_funcs.cc ) if(HIP_PLATFORM MATCHES "nvidia") @@ -96,3 +97,7 @@ 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) +add_test(NAME Unit_Device_special_funcs_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + math_special_func_kernels.cc 76) diff --git a/projects/hip-tests/catch/unit/math/math_special_func_kernels.cc b/projects/hip-tests/catch/unit/math/math_special_func_kernels.cc new file mode 100644 index 0000000000..7dec752ddc --- /dev/null +++ b/projects/hip-tests/catch/unit/math/math_special_func_kernels.cc @@ -0,0 +1,60 @@ +/* +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) \ + __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); } + +#define NEGATIVE_KERNELS_SHELL_TWO_ARGS(func_name) \ + __global__ void func_name##_kernel_v1(int* x, double y) { double result = func_name(x, y); } \ + __global__ void func_name##_kernel_v2(int x, double* y) { double result = func_name(x, y); } \ + __global__ void func_name##_kernel_v3(Dummy x, double y) { double result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(int x, Dummy y) { double result = func_name(x, y); } \ + __global__ void func_name##f_kernel_v1(int* x, float y) { float result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v2(int x, float* y) { float result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v3(Dummy x, float y) { float result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v4(int x, Dummy y) { float result = func_name##f(x, y); } + +NEGATIVE_KERNELS_SHELL_ONE_ARG(erf) +NEGATIVE_KERNELS_SHELL_ONE_ARG(erfc) +NEGATIVE_KERNELS_SHELL_ONE_ARG(erfinv) +NEGATIVE_KERNELS_SHELL_ONE_ARG(erfcinv) +NEGATIVE_KERNELS_SHELL_ONE_ARG(erfcx) +NEGATIVE_KERNELS_SHELL_ONE_ARG(normcdf) +NEGATIVE_KERNELS_SHELL_ONE_ARG(normcdfinv) +NEGATIVE_KERNELS_SHELL_ONE_ARG(lgamma) +NEGATIVE_KERNELS_SHELL_ONE_ARG(tgamma) +NEGATIVE_KERNELS_SHELL_ONE_ARG(j0) +NEGATIVE_KERNELS_SHELL_ONE_ARG(j1) +NEGATIVE_KERNELS_SHELL_TWO_ARGS(jn) +NEGATIVE_KERNELS_SHELL_ONE_ARG(y0) +NEGATIVE_KERNELS_SHELL_ONE_ARG(y1) +NEGATIVE_KERNELS_SHELL_TWO_ARGS(yn) +NEGATIVE_KERNELS_SHELL_ONE_ARG(cyl_bessel_i0) +NEGATIVE_KERNELS_SHELL_ONE_ARG(cyl_bessel_i1) diff --git a/projects/hip-tests/catch/unit/math/math_special_func_kernels_rtc.hh b/projects/hip-tests/catch/unit/math/math_special_func_kernels_rtc.hh new file mode 100644 index 0000000000..e829db2f3a --- /dev/null +++ b/projects/hip-tests/catch/unit/math/math_special_func_kernels_rtc.hh @@ -0,0 +1,236 @@ +/* +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 special function negative Test Cases that are using RTC. +*/ + +static constexpr auto kErf{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void erf_kernel_v1(double* x) { double result = erf(x); } + __global__ void erf_kernel_v2(Dummy x) { double result = erf(x); } + __global__ void erff_kernel_v1(float* x) { float result = erff(x); } + __global__ void erff_kernel_v2(Dummy x) { float result = erff(x); } +)"}; + +static constexpr auto kErfc{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void erfc_kernel_v1(double* x) { double result = erfc(x); } + __global__ void erfc_kernel_v2(Dummy x) { double result = erfc(x); } + __global__ void erfcf_kernel_v1(float* x) { float result = erfcf(x); } + __global__ void erfcf_kernel_v2(Dummy x) { float result = erfcf(x); } +)"}; + +static constexpr auto kErfinv{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void erfinv_kernel_v1(double* x) { double result = erfinv(x); } + __global__ void erfinv_kernel_v2(Dummy x) { double result = erfinv(x); } + __global__ void erfinvf_kernel_v1(float* x) { float result = erfinvf(x); } + __global__ void erfinvf_kernel_v2(Dummy x) { float result = erfinvf(x); } +)"}; + +static constexpr auto kErfcinv{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void erfcinv_kernel_v1(double* x) { double result = erfcinv(x); } + __global__ void erfcinv_kernel_v2(Dummy x) { double result = erfcinv(x); } + __global__ void erfcinvf_kernel_v1(float* x) { float result = erfcinvf(x); } + __global__ void erfcinvf_kernel_v2(Dummy x) { float result = erfcinvf(x); } +)"}; + +static constexpr auto kErfcx{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void erfcx_kernel_v1(double* x) { double result = erfcx(x); } + __global__ void erfcx_kernel_v2(Dummy x) { double result = erfcx(x); } + __global__ void erfcxf_kernel_v1(float* x) { float result = erfcxf(x); } + __global__ void erfcxf_kernel_v2(Dummy x) { float result = erfcxf(x); } +)"}; + +static constexpr auto kNormcdf{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void normcdf_kernel_v1(double* x) { double result = normcdf(x); } + __global__ void normcdf_kernel_v2(Dummy x) { double result = normcdf(x); } + __global__ void normcdff_kernel_v1(float* x) { float result = normcdff(x); } + __global__ void normcdff_kernel_v2(Dummy x) { float result = normcdff(x); } +)"}; + +static constexpr auto kNormcdfinv{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void normcdfinv_kernel_v1(double* x) { double result = normcdfinv(x); } + __global__ void normcdfinv_kernel_v2(Dummy x) { double result = normcdfinv(x); } + __global__ void normcdfinvf_kernel_v1(float* x) { float result = normcdfinvf(x); } + __global__ void normcdfinvf_kernel_v2(Dummy x) { float result = normcdfinvf(x); } +)"}; + +static constexpr auto kLgamma{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void lgamma_kernel_v1(double* x) { double result = lgamma(x); } + __global__ void lgamma_kernel_v2(Dummy x) { double result = lgamma(x); } + __global__ void lgammaf_kernel_v1(float* x) { float result = lgammaf(x); } + __global__ void lgammaf_kernel_v2(Dummy x) { float result = lgammaf(x); } +)"}; + +static constexpr auto kTgamma{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void tgamma_kernel_v1(double* x) { double result = tgamma(x); } + __global__ void tgamma_kernel_v2(Dummy x) { double result = tgamma(x); } + __global__ void tgammaf_kernel_v1(float* x) { float result = tgammaf(x); } + __global__ void tgammaf_kernel_v2(Dummy x) { float result = tgammaf(x); } +)"}; + +static constexpr auto kJ0{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void j0_kernel_v1(double* x) { double result = j0(x); } + __global__ void j0_kernel_v2(Dummy x) { double result = j0(x); } + __global__ void j0f_kernel_v1(float* x) { float result = j0f(x); } + __global__ void j0f_kernel_v2(Dummy x) { float result = j0f(x); } +)"}; + +static constexpr auto kJ1{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void j1_kernel_v1(double* x) { double result = j1(x); } + __global__ void j1_kernel_v2(Dummy x) { double result = j1(x); } + __global__ void j1f_kernel_v1(float* x) { float result = j1f(x); } + __global__ void j1f_kernel_v2(Dummy x) { float result = j1f(x); } +)"}; + +static constexpr auto kJn{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void jn_kernel_v1(int* x, double y) { double result = jn(x, y); } + __global__ void jn_kernel_v2(int x, double* y) { double result = jn(x, y); } + __global__ void jn_kernel_v3(Dummy x, double y) { double result = jn(x, y); } + __global__ void jn_kernel_v4(int x, Dummy y) { double result = jn(x, y); } + __global__ void jnf_kernel_v1(int* x, float y) { float result = jnf(x, y); } + __global__ void jnf_kernel_v2(int x, float* y) { float result = jnf(x, y); } + __global__ void jnf_kernel_v3(Dummy x, float y) { float result = jnf(x, y); } + __global__ void jnf_kernel_v4(int x, Dummy y) { float result = jnf(x, y); } +)"}; + +static constexpr auto kY0{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void y0_kernel_v1(double* x) { double result = y0(x); } + __global__ void y0_kernel_v2(Dummy x) { double result = y0(x); } + __global__ void y0f_kernel_v1(float* x) { float result = y0f(x); } + __global__ void y0f_kernel_v2(Dummy x) { float result = y0f(x); } +)"}; + +static constexpr auto kY1{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void y1_kernel_v1(double* x) { double result = y1(x); } + __global__ void y1_kernel_v2(Dummy x) { double result = y1(x); } + __global__ void y1f_kernel_v1(float* x) { float result = y1f(x); } + __global__ void y1f_kernel_v2(Dummy x) { float result = y1f(x); } +)"}; + +static constexpr auto kYn{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void yn_kernel_v1(int* x, double y) { double result = yn(x, y); } + __global__ void yn_kernel_v2(int x, double* y) { double result = yn(x, y); } + __global__ void yn_kernel_v3(Dummy x, double y) { double result = yn(x, y); } + __global__ void yn_kernel_v4(int x, Dummy y) { double result = yn(x, y); } + __global__ void ynf_kernel_v1(int* x, float y) { float result = ynf(x, y); } + __global__ void ynf_kernel_v2(int x, float* y) { float result = ynf(x, y); } + __global__ void ynf_kernel_v3(Dummy x, float y) { float result = ynf(x, y); } + __global__ void ynf_kernel_v4(int x, Dummy y) { float result = ynf(x, y); } +)"}; + +static constexpr auto kCylBesselI0{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void cyl_bessel_i0_kernel_v1(double* x) { double result = cyl_bessel_i0(x); } + __global__ void cyl_bessel_i0_kernel_v2(Dummy x) { double result = cyl_bessel_i0(x); } + __global__ void cyl_bessel_i0f_kernel_v1(float* x) { float result = cyl_bessel_i0f(x); } + __global__ void cyl_bessel_i0f_kernel_v2(Dummy x) { float result = cyl_bessel_i0f(x); } +)"}; + +static constexpr auto kCylBesselI1{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void cyl_bessel_i1_kernel_v1(double* x) { double result = cyl_bessel_i1(x); } + __global__ void cyl_bessel_i1_kernel_v2(Dummy x) { double result = cyl_bessel_i1(x); } + __global__ void cyl_bessel_i1f_kernel_v1(float* x) { float result = cyl_bessel_i1f(x); } + __global__ void cyl_bessel_i1f_kernel_v2(Dummy x) { float result = cyl_bessel_i1f(x); } +)"}; diff --git a/projects/hip-tests/catch/unit/math/special_common.hh b/projects/hip-tests/catch/unit/math/special_common.hh new file mode 100644 index 0000000000..4b55a88fee --- /dev/null +++ b/projects/hip-tests/catch/unit/math/special_common.hh @@ -0,0 +1,145 @@ +/* +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. +*/ + +#pragma once + +#include "math_common.hh" +#include "math_special_values.hh" + +#include + +namespace cg = cooperative_groups; + +#define MATH_BESSEL_N_KERNEL_DEF(func_name) \ + template \ + __global__ void func_name##_kernel(T* const ys, const size_t num_xs, int* n, 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] = func_name##f(n[i], xs[i]); \ + } else if constexpr (std::is_same_v) { \ + ys[i] = func_name(n[i], xs[i]); \ + } \ + } \ + } + +template using kernel_bessel_n_sig = void (*)(T*, const size_t, int*, T*); + +template using ref_bessel_n_sig = T (*)(int, T); + +template +void BesselDoublePrecisionBruteForceTest(kernel_bessel_n_sig kernel, + ref_bessel_n_sig ref_func, + const ValidatorBuilder& validator_builder, int n_input = 0, + const double a = std::numeric_limits::lowest(), + const double b = std::numeric_limits::max()) { + const auto [grid_size, block_size] = GetOccupancyMaxPotentialBlockSize(kernel); + const uint64_t num_iterations = GetTestIterationCount(); + const auto max_batch_size = std::min( + GetMaxAllowedDeviceMemoryUsage() / (sizeof(double) * 2 + sizeof(int)), num_iterations); + LinearAllocGuard x1s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(int)}; + LinearAllocGuard x2s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(double)}; + + MathTest math_test(kernel, max_batch_size); + std::fill_n(x1s.ptr(), max_batch_size, n_input); + + auto batch_size = max_batch_size; + const auto num_threads = thread_pool.thread_count(); + for (uint64_t i = 0ul; i < num_iterations; i += batch_size) { + batch_size = std::min(max_batch_size, num_iterations - i); + + const auto min_sub_batch_size = batch_size / num_threads; + const auto tail = batch_size % num_threads; + + auto base_idx = 0u; + for (auto i = 0u; i < num_threads; ++i) { + const auto sub_batch_size = min_sub_batch_size + (i < tail); + thread_pool.Post([=, &x2s] { + const auto generator = [=] { + static thread_local std::mt19937 rng(std::random_device{}()); + std::uniform_real_distribution> unif_dist(a, b); + return static_cast(unif_dist(rng)); + }; + std::generate(x2s.ptr() + base_idx, x2s.ptr() + base_idx + sub_batch_size, generator); + }); + base_idx += sub_batch_size; + } + + thread_pool.Wait(); + + math_test.Run(validator_builder, grid_size, block_size, ref_func, batch_size, x1s.ptr(), + x2s.ptr()); + } +} + +template +void BesselSinglePrecisionRangeTest(kernel_bessel_n_sig kernel, + ref_bessel_n_sig ref_func, + const ValidatorBuilder& validator_builder, int n_input, + const float a, const float b) { + const auto [grid_size, block_size] = GetOccupancyMaxPotentialBlockSize(kernel); + const auto max_batch_size = GetMaxAllowedDeviceMemoryUsage() / (sizeof(float) * 2 + sizeof(int)); + LinearAllocGuard x1s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(int)}; + LinearAllocGuard x2s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(float)}; + + MathTest math_test(kernel, max_batch_size); + std::fill_n(x1s.ptr(), max_batch_size, n_input); + + size_t inserted = 0u; + for (float v = a; v != b; v = std::nextafter(v, b)) { + x2s.ptr()[inserted++] = v; + if (inserted < max_batch_size) continue; + + math_test.Run(validator_builder, grid_size, block_size, ref_func, inserted, x1s.ptr(), + x2s.ptr()); + inserted = 0u; + } +} + +template +void SpecialSimpleTest(F kernel, const ValidatorBuilder& validator_builder, const T* x, + const T* ref, size_t num_args) { + LinearAllocGuard x_dev{LinearAllocs::hipMalloc, num_args * sizeof(T)}; + LinearAllocGuard y{LinearAllocs::hipHostMalloc, num_args * sizeof(T)}; + LinearAllocGuard y_dev{LinearAllocs::hipMalloc, num_args * sizeof(T)}; + + HIP_CHECK(hipMemcpy(x_dev.ptr(), x, num_args * sizeof(T), hipMemcpyHostToDevice)); + + kernel<<<1, num_args>>>(y_dev.ptr(), num_args, x_dev.ptr()); + HIP_CHECK(hipGetLastError()); + + HIP_CHECK(hipMemcpy(y.ptr(), y_dev.ptr(), num_args * sizeof(T), hipMemcpyDeviceToHost)); + + for (auto i = 0u; i < num_args; ++i) { + const auto actual_val = y.ptr()[i]; + const auto ref_val = ref[i]; + const auto validator = validator_builder(ref_val); + + if (!validator->match(actual_val)) { + std::stringstream ss; + ss << "Input value(s): " << std::scientific + << std::setprecision(std::numeric_limits::max_digits10 - 1); + ss << x[i] << " " << actual_val << " " << ref_val << "\n"; + INFO(ss.str()); + REQUIRE(false); + } + } +} diff --git a/projects/hip-tests/catch/unit/math/special_funcs.cc b/projects/hip-tests/catch/unit/math/special_funcs.cc new file mode 100644 index 0000000000..5461afadd1 --- /dev/null +++ b/projects/hip-tests/catch/unit/math/special_funcs.cc @@ -0,0 +1,1117 @@ +/* +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 "binary_common.hh" +#include "special_common.hh" +#include "math_special_func_kernels_rtc.hh" + +#include + + +/** + * @addtogroup SpecialMathFuncs SpecialMathFuncs + * @{ + * @ingroup MathTest + */ + +/********** Unary Functions **********/ + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `erff(x)` for all possible inputs and `erf(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::erf(T)`. The maximum ulp error is 2. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(erf, 2, 2) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for erff and erf. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erf_erff_Negative_RTC") { NegativeTestRTCWrapper<4>(kErf); } + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `erfcf(x)` for all possible inputs and `erfc(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::erfc(T)`. The maximum ulp error for single + * precision is 4 and for double precision is 5. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(erfc, 4, 5) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for erfcf and erfc. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfc_erfcf_Negative_RTC") { NegativeTestRTCWrapper<4>(kErfc); } + +MATH_UNARY_KERNEL_DEF(erfinv) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `erfinvf(x)` for all possible inputs. The results are + * compared against reference function `double boost::math::erf_inv(double)`. The maximum ulp error + * is 2. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfinvf_Accuracy_Positive") { + auto erfinv_ref = [](double arg) -> double { + if (arg == 0) return 0; + if (arg == 1) + return std::numeric_limits::infinity(); + else if (arg == -1) + return -std::numeric_limits::infinity(); + else if (arg < -1 || arg > 1) + return std::numeric_limits::quiet_NaN(); + return boost::math::erf_inv(arg); + }; + double (*ref)(double) = erfinv_ref; + UnarySinglePrecisionTest(erfinv_kernel, ref, ULPValidatorBuilderFactory(2)); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `erfinv(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double boost::math::erf_inv(long double)`. The maximum + * ulp error is 5. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfinv_Accuracy_Positive") { + auto erfinv_ref = [](long double arg) -> long double { + if (arg == 0) return 0; + if (arg == 1) + return std::numeric_limits::infinity(); + else if (arg == -1) + return -std::numeric_limits::infinity(); + else if (arg < -1 || arg > 1) + return std::numeric_limits::quiet_NaN(); + return boost::math::erf_inv(arg); + }; + long double (*ref)(long double) = erfinv_ref; + UnaryDoublePrecisionTest(erfinv_kernel, ref, ULPValidatorBuilderFactory(5)); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for erfinvf and erfinv. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfinv_erfinvf_Negative_RTC") { NegativeTestRTCWrapper<4>(kErfinv); } + +MATH_UNARY_KERNEL_DEF(erfcinv) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `erfcinvf(x)` for all possible inputs. The results are + * compared against reference function `double boost::math::erfc_inv(double)`. The maximum ulp error + * is 4. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfcinvf_Accuracy_Positive") { + auto erfcinv_ref = [](double arg) -> double { + if (arg == 0) + return std::numeric_limits::infinity(); + else if (arg == 2) + return -std::numeric_limits::infinity(); + else if (arg < 0 || arg > 2) + return std::numeric_limits::quiet_NaN(); + return boost::math::erfc_inv(arg); + }; + double (*ref)(double) = erfcinv_ref; + UnarySinglePrecisionTest(erfcinv_kernel, ref, ULPValidatorBuilderFactory(4)); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `erfcinv(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double boost::math::erfc_inv(long double)`. The maximum + * ulp error is 6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfcinv_Accuracy_Positive") { + auto erfcinv_ref = [](long double arg) -> long double { + if (arg == 0) + return std::numeric_limits::infinity(); + else if (arg == 2) + return -std::numeric_limits::infinity(); + else if (arg < 0 || arg > 2) + return std::numeric_limits::quiet_NaN(); + return boost::math::erfc_inv(arg); + }; + long double (*ref)(long double) = erfcinv_ref; + UnaryDoublePrecisionTest(erfcinv_kernel, ref, ULPValidatorBuilderFactory(6)); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for erfcinvf and erfcinv. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfcinv_erfcinvf_Negative_RTC") { NegativeTestRTCWrapper<4>(kErfcinv); } + +MATH_UNARY_KERNEL_DEF(erfcx) + +/** + * Test Description + * ------------------------ + * - Sanity test for `erfcxf(x)`. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfcxf_Sanity_Positive") { + constexpr std::array input{-std::numeric_limits::infinity(), + -1000.f, + -100.f, + -5.f, + -0.5f, + 0., + 0.75f, + 15.f, + 200.f, + 500.f, + std::numeric_limits::infinity()}; + constexpr std::array reference{std::numeric_limits::infinity(), + std::numeric_limits::infinity(), + std::numeric_limits::infinity(), + 1.44009806e11f, + 1.95236027f, + 1.0f, + 5.06937683e-1f, + 3.75296101e-2f, + 2.82091252e-3f, + 1.12837693e-3f, + 0.f}; + SpecialSimpleTest(erfcx_kernel, ULPValidatorBuilderFactory(4), input.data(), + reference.data(), input.size()); +} + +/** + * Test Description + * ------------------------ + * - Sanity test for `erfcx(x)`. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfcx_Sanity_Positive") { + constexpr std::array input{ + -std::numeric_limits::infinity(), -1000., -100., -5., -0.5, 0., 0.75, 15., 200., 500., + std::numeric_limits::infinity()}; + constexpr std::array reference{std::numeric_limits::infinity(), + std::numeric_limits::infinity(), + std::numeric_limits::infinity(), + 1.4400979867466104e11, + 1.9523604891825568, + 1.0, + 5.0693765029314475e-1, + 3.7529606388505762e-2, + 2.8209126572120466e-3, + 1.1283769103507188e-3, + 0.}; + SpecialSimpleTest(erfcx_kernel, ULPValidatorBuilderFactory(4), + input.data(), reference.data(), input.size()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for erfcxf and erfcx. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_erfcx_erfcxf_Negative_RTC") { NegativeTestRTCWrapper<4>(kErfcx); } + +MATH_UNARY_KERNEL_DEF(normcdf) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `normcdff(x)` for all possible inputs. The maximum ulp error + * is 5. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_normcdff_Accuracy_Positive") { + auto normcdf_ref = [](double arg) -> double { return std::erfc(-arg / std::sqrt(2)) / 2; }; + double (*ref)(double) = normcdf_ref; + UnarySinglePrecisionTest(normcdf_kernel, ref, ULPValidatorBuilderFactory(5)); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `normcdf(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The maximum ulp error is 5. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_normcdf_Accuracy_Positive") { + auto normcdf_ref = [](long double arg) -> long double { + return std::erfc(-arg / std::sqrt(2.L)) / 2; + }; + long double (*ref)(long double) = normcdf_ref; + UnaryDoublePrecisionTest(normcdf_kernel, ref, ULPValidatorBuilderFactory(5)); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for normcdff and normcdf. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_normcdf_normcdff_Negative_RTC") { NegativeTestRTCWrapper<4>(kNormcdf); } + +MATH_UNARY_KERNEL_DEF(normcdfinv) + +/** + * Test Description + * ------------------------ + * - Sanity test for `normcdfinvf(x)`. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_normcdfinvf_Sanity_Positive") { + constexpr std::array input{0.f, 0.1f, 0.25f, 0.4f, 0.5f, 0.6f, 0.75f, 0.9f, 1.f}; + constexpr std::array reference{-std::numeric_limits::infinity(), + -1.28155160f, + -0.674489737f, + -0.253347069f, + 0, + 0.253347158f, + 0.674489737f, + 1.28155148f, + std::numeric_limits::infinity()}; + SpecialSimpleTest(normcdfinv_kernel, ULPValidatorBuilderFactory(5), + input.data(), reference.data(), input.size()); +} + +/** + * Test Description + * ------------------------ + * - Sanity test for `normcdfinv(x)`. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_normcdfinv_Sanity_Positive") { + constexpr std::array input{0., 0.1, 0.25, 0.4, 0.5, 0.6, 0.75, 0.9, 1.}; + constexpr std::array reference{-std::numeric_limits::infinity(), + -1.2815515655446004, + -0.67448975019608159, + -0.25334710313579972, + 0, + 0.25334710313579972, + 0.67448975019608159, + 1.2815515655446006, + std::numeric_limits::infinity()}; + SpecialSimpleTest(normcdfinv_kernel, ULPValidatorBuilderFactory(5), + input.data(), reference.data(), input.size()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for normcdfinvf and normcdfinv. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_normcdfinv_normcdfinvf_Negative_RTC") { + NegativeTestRTCWrapper<4>(kNormcdfinv); +} + +MATH_UNARY_KERNEL_DEF(tgamma) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `tgammaf(x)` for all possible inputs below 171.7 and that + * are not very small negative numbers, as they lead to overflow for IEEE compatible double. The + * results are compared against reference function `double std::tgamma(double)`. The maximum ulp + * error is 5. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_tgammaf_Accuracy_Limited_Positive") { + double (*ref)(double) = std::tgamma; + UnarySinglePrecisionRangeTest(tgamma_kernel, ref, ULPValidatorBuilderFactory(5), + std::numeric_limits::lowest(), -0.001f); + UnarySinglePrecisionRangeTest(tgamma_kernel, ref, ULPValidatorBuilderFactory(5), 0, + 171.7); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `tgamma(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double std::tgamma(long double)`. The maximum ulp error + * is 10. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_tgamma_Accuracy_Limited_Positive") { + long double (*ref)(long double) = std::tgamma; + UnaryDoublePrecisionTest(tgamma_kernel, ref, ULPValidatorBuilderFactory(10)); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for tgammaf and tgamma. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_tgamma_tgammaf_Negative_RTC") { NegativeTestRTCWrapper<4>(kTgamma); } + +MATH_UNARY_KERNEL_DEF(lgamma) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `lgammaf(x)` for all possible inputs. The results are + * compared against reference function `double std::lgamma(double)`. For `x` outside interval + * -11.0001 … -2.2637, the maximum ulp error is 4, and larger otherwise. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_lgammaf_Accuracy_Limited_Positive") { + double (*ref)(double) = std::lgamma; + UnarySinglePrecisionRangeTest(lgamma_kernel, ref, ULPValidatorBuilderFactory(6), + std::numeric_limits::lowest(), -11.0001f); + UnarySinglePrecisionRangeTest(lgamma_kernel, ref, ULPValidatorBuilderFactory(6), + -2.2636f, std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `lgamma(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are compared against + * reference function `long double std::lgamma(long double)`. For `x` outside interval -11.0001 … + * -2.2637, the maximum ulp error is 4, and larger otherwise. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_lgamma_Accuracy_Limited_Positive") { + long double (*ref)(long double) = std::lgamma; + UnaryDoublePrecisionBruteForceTest(lgamma_kernel, ref, + ULPValidatorBuilderFactory(4), + std::numeric_limits::lowest(), -11.0001); + UnaryDoublePrecisionBruteForceTest(lgamma_kernel, ref, + ULPValidatorBuilderFactory(4), -2.2636, + std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for lgammaf and lgamma. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_lgamma_lgammaf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLgamma); } + +MATH_UNARY_KERNEL_DEF(cyl_bessel_i0) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `cyl_bessel_i0f(x)` for all possible inputs in range [0, + * 10000). The results are compared against reference function `double std::cyl_bessel_i(0, + * double)`. The maximum ulp error is 6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_cyl_bessel_i0f_Accuracy_Limited_Positive") { + auto cyl_bessel_i0_ref = [](double arg) -> double { return std::cyl_bessel_i(0, arg); }; + double (*ref)(double) = cyl_bessel_i0_ref; + UnarySinglePrecisionRangeTest(cyl_bessel_i0_kernel, ref, + ULPValidatorBuilderFactory(6), 0, 10000); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `cyl_bessel_i0(x)` against a table of difficult values, + * followed by a large number of randomly generated values from range [0, 10000). The results are + * compared against reference function `long double std::cyl_bessel_i(0, long double)`. The maximum + * ulp error is 6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_cyl_bessel_i0_Accuracy_Limited_Positive") { + auto cyl_bessel_i0_ref = [](long double arg) -> long double { return std::cyl_bessel_i(0, arg); }; + long double (*ref)(long double) = cyl_bessel_i0_ref; + UnaryDoublePrecisionBruteForceTest(cyl_bessel_i0_kernel, ref, + ULPValidatorBuilderFactory(6), 0, 10000); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for cyl_bessel_i0f and cyl_bessel_i0. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_cyl_bessel_i0_cyl_bessel_i0f_Negative_RTC") { + NegativeTestRTCWrapper<4>(kCylBesselI0); +} + +MATH_UNARY_KERNEL_DEF(cyl_bessel_i1) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `cyl_bessel_i1f(x)` for all possible inputs in range [0, + * 10000). The results are compared against reference function `double std::cyl_bessel_i(1, + * double)`. The maximum ulp error is 6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_cyl_bessel_i1f_Accuracy_Limited_Positive") { + auto cyl_bessel_i1_ref = [](double arg) -> double { return std::cyl_bessel_i(1, arg); }; + double (*ref)(double) = cyl_bessel_i1_ref; + UnarySinglePrecisionRangeTest(cyl_bessel_i1_kernel, ref, + ULPValidatorBuilderFactory(6), 0, 10000); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `cyl_bessel_i1(x)` against a table of difficult values, + * followed by a large number of randomly generated values from range [0, 10000). The results are + * compared against reference function `long double std::cyl_bessel_i(1, long double)`. The maximum + * ulp error is 6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_cyl_bessel_i1_Accuracy_Limited_Positive") { + auto cyl_bessel_i1_ref = [](long double arg) -> long double { return std::cyl_bessel_i(1, arg); }; + long double (*ref)(long double) = cyl_bessel_i1_ref; + UnaryDoublePrecisionBruteForceTest(cyl_bessel_i1_kernel, ref, + ULPValidatorBuilderFactory(6), 0, 10000); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for cyl_bessel_i1f and cyl_bessel_i1. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_cyl_bessel_i1_cyl_bessel_i1f_Negative_RTC") { + NegativeTestRTCWrapper<4>(kCylBesselI1); +} + +/********** Bessel Functions **********/ + +MATH_UNARY_KERNEL_DEF(y0) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `y0f(x)` for all possible inputs. The results are + * compared against reference function `double y0(double)`. For `x` outside [-8, 8], the maximum + * absolute error is 2.2x10^-6, otherwise, the maximum ulp error is 9. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_y0f_Accuracy_Limited_Positive") { +#ifdef __unix__ + double (*ref)(double) = y0; +#elif _WIN64 + double (*ref)(double) = _y0; +#endif + UnarySinglePrecisionRangeTest(y0_kernel, ref, ULPValidatorBuilderFactory(9), -8.f, + 8.f); + UnarySinglePrecisionRangeTest(y0_kernel, ref, AbsValidatorBuilderFactory(0.0000022), + 8.f, std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `y0(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double y0l(long double)`. The maximum absolute error is + * 5x10^-12. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_y0_Accuracy_Limited_Positive") { +#ifdef __unix__ + long double (*ref)(long double) = y0l; +#elif _WIN64 + long double (*ref)(long double) = _y0l; +#endif + UnaryDoublePrecisionBruteForceTest(y0_kernel, ref, + AbsValidatorBuilderFactory(5.e-12), -8., + std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for y0f and y0. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_y0_y0f_Negative_RTC") { NegativeTestRTCWrapper<4>(kY0); } + +MATH_UNARY_KERNEL_DEF(y1) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `y1f(x)` for all possible inputs. The results are + * compared against reference function `double y1(double)`. For `x` outside [-8, 8], the maximum + * absolute error is 2.2x10^-6, otherwise, the maximum ulp error is 9. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_y1f_Accuracy_Limited_Positive") { +#ifdef __unix__ + double (*ref)(double) = y1; +#elif _WIN64 + double (*ref)(double) = _y1; +#endif + UnarySinglePrecisionRangeTest(y1_kernel, ref, ULPValidatorBuilderFactory(9), -8.f, + 8.f); + UnarySinglePrecisionRangeTest(y1_kernel, ref, AbsValidatorBuilderFactory(0.0000022), + 8.f, std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `y1(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double y1l(long double)`. The maximum absolute error is + * 5x10^-12. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_y1_Accuracy_Limited_Positive") { +#ifdef __unix__ + long double (*ref)(long double) = y1l; +#elif _WIN64 + long double (*ref)(long double) = _y1l; +#endif + UnaryDoublePrecisionBruteForceTest(y1_kernel, ref, + AbsValidatorBuilderFactory(5.e-12), -8., + std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for y1f and y1. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_y1_y1f_Negative_RTC") { NegativeTestRTCWrapper<4>(kY1); } + +MATH_BESSEL_N_KERNEL_DEF(yn) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `ynf(x)` for all possible inputs and n equal to 5, 25 or + * 120. The results are compared against reference function `double yn(int, double)`. For `x` larger + * than n, the maximum absolute error is 2.2x10^-6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_ynf_Accuracy_Limited_Positive") { +#ifdef __unix__ + double (*ref)(int, double) = yn; +#elif _WIN64 + double (*ref)(int, double) = _yn; +#endif + int n = GENERATE(5, 25, 120); + BesselSinglePrecisionRangeTest(yn_kernel, ref, AbsValidatorBuilderFactory(0.0000022), n, n, + std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `yn(x)` against a table of difficult values, + * followed by a large number of randomly generated values from range and n equal to 5, 25, or 120. + * The results are compared against reference function `long double ynl(int, long double)`. For `x` + * larger than 1.5n, the maximum absolute error is 5x10^-12. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_yn_Accuracy_Limited_Positive") { +#ifdef __unix__ + long double (*ref)(int, long double) = ynl; +#elif _WIN64 + long double (*ref)(int, long double) = _ynl; +#endif + int n = GENERATE(5, 25, 120); + BesselDoublePrecisionBruteForceTest(yn_kernel, ref, + AbsValidatorBuilderFactory(5.e-12), n, 1.5 * n, + std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for ynf and yn. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_yn_ynf_Negative_RTC") { NegativeTestRTCWrapper<8>(kYn); } + +MATH_UNARY_KERNEL_DEF(j0) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `j0f(x)` for all possible inputs. The results are + * compared against reference function `double j0(double)`. For `x` outside [-8, 8], the maximum + * absolute error is 2.2x10^-6, otherwise, the maximum ulp error is 9. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_j0f_Accuracy_Limited_Positive") { +#ifdef __unix__ + double (*ref)(double) = j0; +#elif _WIN64 + double (*ref)(double) = _j0; +#endif + UnarySinglePrecisionRangeTest(j0_kernel, ref, AbsValidatorBuilderFactory(0.0000022), + std::numeric_limits::lowest(), -8.f); + UnarySinglePrecisionRangeTest(j0_kernel, ref, ULPValidatorBuilderFactory(9), -8.f, + 8.f); + UnarySinglePrecisionRangeTest(j0_kernel, ref, AbsValidatorBuilderFactory(0.0000022), + 8.f, std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `j0(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double j0l(long double)`. The maximum absolute error is + * 5x10^-12. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_j0_Accuracy_Limited_Positive") { +#ifdef __unix__ + long double (*ref)(long double) = j0l; +#elif _WIN64 + long double (*ref)(long double) = _j0l; +#endif + UnaryDoublePrecisionBruteForceTest( + j0_kernel, ref, AbsValidatorBuilderFactory(5.e-12), + std::numeric_limits::lowest(), std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for j0f and j0. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_j0_j0f_Negative_RTC") { NegativeTestRTCWrapper<4>(kJ0); } + +MATH_UNARY_KERNEL_DEF(j1) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `j1f(x)` for all possible inputs. The results are + * compared against reference function `double j1(double)`. For `x` outside [-8, 8], the maximum + * absolute error is 2.2x10^-6, otherwise, the maximum ulp error is 9. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_j1f_Accuracy_Limited_Positive") { +#ifdef __unix__ + double (*ref)(double) = j1; +#elif _WIN64 + double (*ref)(double) = _j1; +#endif + UnarySinglePrecisionRangeTest(j1_kernel, ref, AbsValidatorBuilderFactory(0.0000022), + std::numeric_limits::lowest(), -8.f); + UnarySinglePrecisionRangeTest(j1_kernel, ref, ULPValidatorBuilderFactory(9), -8.f, + 8.f); + UnarySinglePrecisionRangeTest(j1_kernel, ref, AbsValidatorBuilderFactory(0.0000022), + 8.f, std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `j1(x)` against a table of difficult values, + * followed by a large number of randomly generated values. The results are + * compared against reference function `long double j1l(long double)`. The maximum absolute error is + * 5x10^-12. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_j1_Accuracy_Limited_Positive") { +#ifdef __unix__ + long double (*ref)(long double) = j1l; +#elif _WIN64 + long double (*ref)(long double) = _j1l; +#endif + UnaryDoublePrecisionBruteForceTest( + j1_kernel, ref, AbsValidatorBuilderFactory(5.e-12), + std::numeric_limits::lowest(), std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for j1f and j1. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_j1_j1f_Negative_RTC") { NegativeTestRTCWrapper<4>(kJ1); } + +MATH_BESSEL_N_KERNEL_DEF(jn) + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `jnf(x)` for all possible inputs and n equal to 5, 25 or + * 120. The results are compared against reference function `double jn(int, double)`. For `x` larger + * than n, the maximum absolute error is 2.2x10^-6. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_jnf_Accuracy_Limited_Positive") { +#ifdef __unix__ + double (*ref)(int, double) = jn; +#elif _WIN64 + double (*ref)(int, double) = _jn; +#endif + int n = GENERATE(5, 25, 120); + BesselSinglePrecisionRangeTest(jn_kernel, ref, AbsValidatorBuilderFactory(0.0000022), n, n, + std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - Tests the numerical accuracy of `jn(x)` against a table of difficult values, + * followed by a large number of randomly generated values from range and n equal to 5, 25, or 120. + * The results are compared against reference function `long double jnl(int, long double)`. The + * maximum absolute error is 5x10^-12. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_jn_Accuracy_Limited_Positive") { +#ifdef __unix__ + long double (*ref)(int, long double) = jnl; +#elif _WIN64 + long double (*ref)(int, long double) = _jnl; +#endif + int n = GENERATE(5, 25, 120); + BesselDoublePrecisionBruteForceTest( + jn_kernel, ref, AbsValidatorBuilderFactory(5.e-12), n, + std::numeric_limits::lowest(), std::numeric_limits::max()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for jnf and jn. + * + * Test source + * ------------------------ + * - unit/math/special_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device_jn_jnf_Negative_RTC") { NegativeTestRTCWrapper<8>(kJn); }