From 34d572ebdc687ddf9a4c06f8975237aefcc6fb36 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Fri, 25 Aug 2023 10:32:17 -0500 Subject: [PATCH 01/10] use dpctl for trigonometric functions in dpnp --- dpnp/backend/extensions/vm/acos.hpp | 78 ++++ dpnp/backend/extensions/vm/asin.hpp | 78 ++++ dpnp/backend/extensions/vm/atan.hpp | 78 ++++ dpnp/backend/extensions/vm/atan2.hpp | 81 ++++ dpnp/backend/extensions/vm/tan.hpp | 78 ++++ dpnp/backend/extensions/vm/types_matrix.hpp | 95 ++++- dpnp/backend/extensions/vm/vm_py.cpp | 151 +++++++ dpnp/backend/include/dpnp_iface_fptr.hpp | 9 - dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 49 --- dpnp/dpnp_algo/dpnp_algo.pxd | 6 - dpnp/dpnp_algo/dpnp_algo_mathematical.pxi | 9 - dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi | 20 - dpnp/dpnp_algo/dpnp_elementwise_common.py | 358 ++++++++++++++-- dpnp/dpnp_iface_trigonometric.py | 391 +++++++++++------- tests/skipped_tests_gpu_no_fp64.tbl | 9 +- tests/test_sycl_queue.py | 25 +- tests/test_usm_type.py | 17 + .../cupy/math_tests/test_trigonometric.py | 5 +- 18 files changed, 1240 insertions(+), 297 deletions(-) create mode 100644 dpnp/backend/extensions/vm/acos.hpp create mode 100644 dpnp/backend/extensions/vm/asin.hpp create mode 100644 dpnp/backend/extensions/vm/atan.hpp create mode 100644 dpnp/backend/extensions/vm/atan2.hpp create mode 100644 dpnp/backend/extensions/vm/tan.hpp diff --git a/dpnp/backend/extensions/vm/acos.hpp b/dpnp/backend/extensions/vm/acos.hpp new file mode 100644 index 000000000000..10bd68bf66e1 --- /dev/null +++ b/dpnp/backend/extensions/vm/acos.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event acos_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::acos(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AcosContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AcosOutputType::value_type, void>) + { + return nullptr; + } + else { + return acos_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/asin.hpp b/dpnp/backend/extensions/vm/asin.hpp new file mode 100644 index 000000000000..8df471f79c12 --- /dev/null +++ b/dpnp/backend/extensions/vm/asin.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event asin_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::asin(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AsinContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AsinOutputType::value_type, void>) + { + return nullptr; + } + else { + return asin_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/atan.hpp b/dpnp/backend/extensions/vm/atan.hpp new file mode 100644 index 000000000000..35967bb720a3 --- /dev/null +++ b/dpnp/backend/extensions/vm/atan.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event atan_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::atan(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AtanContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AtanOutputType::value_type, void>) + { + return nullptr; + } + else { + return atan_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/atan2.hpp b/dpnp/backend/extensions/vm/atan2.hpp new file mode 100644 index 000000000000..bf3d95a9c7db --- /dev/null +++ b/dpnp/backend/extensions/vm/atan2.hpp @@ -0,0 +1,81 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event atan2_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + const char *in_b, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + const T *b = reinterpret_cast(in_b); + T *y = reinterpret_cast(out_y); + + return mkl_vm::atan2(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing 1st input vector of size n + b, // pointer `b` containing 2nd input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct Atan2ContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::Atan2OutputType::value_type, void>) + { + return nullptr; + } + else { + return atan2_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/tan.hpp b/dpnp/backend/extensions/vm/tan.hpp new file mode 100644 index 000000000000..4a9d0cc5e2e9 --- /dev/null +++ b/dpnp/backend/extensions/vm/tan.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event tan_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::tan(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct TanContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::TanOutputType::value_type, void>) + { + return nullptr; + } + else { + return tan_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/types_matrix.hpp b/dpnp/backend/extensions/vm/types_matrix.hpp index 3fabd38ce388..155dfc54d776 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -43,6 +43,23 @@ namespace vm { namespace types { +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::acos function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AcosOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::add function. @@ -68,6 +85,55 @@ struct AddOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::asin function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AsinOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::atan function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AtanOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::atan2 function. + * + * @tparam T Type of input vectors `a` and `b` and of result vector `y`. + */ +template +struct Atan2OutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::BinaryTypeMapResultEntry, + dpctl_td_ns::BinaryTypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::ceil function. @@ -110,10 +176,8 @@ template struct CosOutputType { using value_type = typename std::disjunction< - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::DefaultResultEntry>::result_type; @@ -253,10 +317,8 @@ template struct SinOutputType { using value_type = typename std::disjunction< - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, - dpctl_td_ns:: - TypeMapResultEntry, std::complex>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::TypeMapResultEntry, dpctl_td_ns::DefaultResultEntry>::result_type; @@ -321,6 +383,23 @@ struct SubOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::tan function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct TanOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::trunc function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index d1cec34680fb..e769a40c2ba9 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -30,7 +30,11 @@ #include #include +#include "acos.hpp" #include "add.hpp" +#include "asin.hpp" +#include "atan.hpp" +#include "atan2.hpp" #include "ceil.hpp" #include "common.hpp" #include "conj.hpp" @@ -45,6 +49,7 @@ #include "sqr.hpp" #include "sqrt.hpp" #include "sub.hpp" +#include "tan.hpp" #include "trunc.hpp" #include "types_matrix.hpp" @@ -54,7 +59,11 @@ namespace vm_ext = dpnp::backend::ext::vm; using vm_ext::binary_impl_fn_ptr_t; using vm_ext::unary_impl_fn_ptr_t; +static unary_impl_fn_ptr_t acos_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t add_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t asin_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t atan_dispatch_vector[dpctl_td_ns::num_types]; +static binary_impl_fn_ptr_t atan2_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ceil_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t cos_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types]; @@ -68,6 +77,7 @@ static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqr_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t sub_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t tan_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t trunc_dispatch_vector[dpctl_td_ns::num_types]; PYBIND11_MODULE(_vm_impl, m) @@ -75,6 +85,34 @@ PYBIND11_MODULE(_vm_impl, m) using arrayT = dpctl::tensor::usm_ndarray; using event_vecT = std::vector; + // UnaryUfunc: ==== Acos(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + acos_dispatch_vector); + + auto acos_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + acos_dispatch_vector); + }; + m.def("_acos", acos_pyapi, + "Call `acos` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto acos_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + acos_dispatch_vector); + }; + m.def("_mkl_acos_to_call", acos_need_to_call_pyapi, + "Check input arguments to answer if `acos` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // BinaryUfunc: ==== Add(x1, x2) ==== { vm_ext::init_ufunc_dispatch_vector( + asin_dispatch_vector); + + auto asin_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + asin_dispatch_vector); + }; + m.def("_asin", asin_pyapi, + "Call `asin` function from OneMKL VM library to compute " + "inverse sine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto asin_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + asin_dispatch_vector); + }; + m.def("_mkl_asin_to_call", asin_need_to_call_pyapi, + "Check input arguments to answer if `asin` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // UnaryUfunc: ==== Atan(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + atan_dispatch_vector); + + auto atan_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + atan_dispatch_vector); + }; + m.def("_atan", atan_pyapi, + "Call `atan` function from OneMKL VM library to compute " + "inverse tangent of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto atan_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + atan_dispatch_vector); + }; + m.def("_mkl_atan_to_call", atan_need_to_call_pyapi, + "Check input arguments to answer if `atan` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + + // BinaryUfunc: ==== Atan2(x1, x2) ==== + { + vm_ext::init_ufunc_dispatch_vector( + atan2_dispatch_vector); + + auto atan2_pyapi = [&](sycl::queue exec_q, arrayT src1, arrayT src2, + arrayT dst, const event_vecT &depends = {}) { + return vm_ext::binary_ufunc(exec_q, src1, src2, dst, depends, + atan2_dispatch_vector); + }; + m.def("_atan2", atan2_pyapi, + "Call `atan2` function from OneMKL VM library to compute element " + "by element inverse tangent of `x1/x2`", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst"), py::arg("depends") = py::list()); + + auto atan2_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src1, + arrayT src2, arrayT dst) { + return vm_ext::need_to_call_binary_ufunc(exec_q, src1, src2, dst, + atan2_dispatch_vector); + }; + m.def("_mkl_atan2_to_call", atan2_need_to_call_pyapi, + "Check input arguments to answer if `atan2` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst")); + } + // UnaryUfunc: ==== Ceil(x) ==== { vm_ext::init_ufunc_dispatch_vector( + tan_dispatch_vector); + + auto tan_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + tan_dispatch_vector); + }; + m.def("_tan", tan_pyapi, + "Call `tan` function from OneMKL VM library to compute " + "tangent of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto tan_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + tan_dispatch_vector); + }; + m.def("_mkl_tan_to_call", tan_need_to_call_pyapi, + "Check input arguments to answer if `tan` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Trunc(x) ==== { vm_ext::init_ufunc_dispatch_vector}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOS_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_acos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_acosh_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCCOSH][eft_LNG][eft_LNG] = { @@ -269,15 +260,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_asin_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_asin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_asinh_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_LNG][eft_LNG] = { @@ -305,15 +287,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_atan_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_atan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_atanh_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_LNG][eft_LNG] = { @@ -693,15 +666,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_TAN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_tan_c_default}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TAN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tan_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_tanh_c_default}; fmap[DPNPFuncName::DPNP_FN_TANH][eft_LNG][eft_LNG] = { @@ -1577,19 +1541,6 @@ static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) template static void func_map_elemwise_2arg_3type_short_core(func_map_t &fmap) { - ((fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][FT1][FTs] = - {get_floating_res_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type()>, - func_type_map_t::find_type, - func_type_map_t::find_type>, - get_floating_res_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_floating_res_type()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), - ...); ((fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][FT1][FTs] = {get_floating_res_type(), (void *)dpnp_copysign_c_ext< diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index eab29517f730..2197333af5c9 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -408,8 +408,6 @@ cpdef dpnp_descriptor dpnp_copy(dpnp_descriptor x1) """ Mathematical functions """ -cpdef dpnp_descriptor dpnp_arctan2(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, - dpnp_descriptor out=*, object where=*) cpdef dpnp_descriptor dpnp_hypot(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, dpnp_descriptor out=*, object where=*) cpdef dpnp_descriptor dpnp_maximum(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, @@ -444,11 +442,8 @@ cpdef dpnp_descriptor dpnp_argmin(dpnp_descriptor array1) """ Trigonometric functions """ -cpdef dpnp_descriptor dpnp_arccos(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_arccosh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arcsin(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_arcsinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arctan(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_arctanh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_cbrt(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_cosh(dpnp_descriptor array1) @@ -462,5 +457,4 @@ cpdef dpnp_descriptor dpnp_log2(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_radians(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_recip(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_sinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_tan(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_tanh(dpnp_descriptor array1) diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index c0df1c290d62..a6794b0626ae 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -37,7 +37,6 @@ and the rest of the library __all__ += [ "dpnp_absolute", - "dpnp_arctan2", "dpnp_copysign", "dpnp_cross", "dpnp_cumprod", @@ -107,14 +106,6 @@ cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1): return result -cpdef utils.dpnp_descriptor dpnp_arctan2(utils.dpnp_descriptor x1_obj, - utils.dpnp_descriptor x2_obj, - object dtype=None, - utils.dpnp_descriptor out=None, - object where=True): - return call_fptr_2in_1out_strides(DPNP_FN_ARCTAN2_EXT, x1_obj, x2_obj, dtype, out, where, func_name="arctan2") - - cpdef utils.dpnp_descriptor dpnp_copysign(utils.dpnp_descriptor x1_obj, utils.dpnp_descriptor x2_obj, object dtype=None, diff --git a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi index ec129a5285e6..12995c346fde 100644 --- a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi @@ -36,11 +36,8 @@ and the rest of the library # NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file __all__ += [ - 'dpnp_arccos', 'dpnp_arccosh', - 'dpnp_arcsin', 'dpnp_arcsinh', - 'dpnp_arctan', 'dpnp_arctanh', 'dpnp_cbrt', 'dpnp_cosh', @@ -54,32 +51,19 @@ __all__ += [ 'dpnp_radians', 'dpnp_recip', 'dpnp_sinh', - 'dpnp_tan', 'dpnp_tanh', 'dpnp_unwrap' ] -cpdef utils.dpnp_descriptor dpnp_arccos(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCCOS_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_arccosh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_ARCCOSH_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_arcsin(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_ARCSIN_EXT, x1, dtype=None, out=out, where=True, func_name='arcsin') - - cpdef utils.dpnp_descriptor dpnp_arcsinh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_ARCSINH_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_arctan(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_ARCTAN_EXT, x1, dtype=None, out=out, where=True, func_name='arctan') - - cpdef utils.dpnp_descriptor dpnp_arctanh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_ARCTANH_EXT, x1) @@ -132,10 +116,6 @@ cpdef utils.dpnp_descriptor dpnp_sinh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_SINH_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_tan(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_TAN_EXT, x1, dtype=None, out=out, where=True, func_name='tan') - - cpdef utils.dpnp_descriptor dpnp_tanh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_TANH_EXT, x1) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index c4d263ba514d..3da994049deb 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -42,7 +42,11 @@ __all__ = [ "check_nd_call_func", + "dpnp_acos", "dpnp_add", + "dpnp_asin", + "dpnp_atan", + "dpnp_atan2", "dpnp_bitwise_and", "dpnp_bitwise_or", "dpnp_bitwise_xor", @@ -81,6 +85,7 @@ "dpnp_sqrt", "dpnp_square", "dpnp_subtract", + "dpnp_tan", "dpnp_trunc", ] @@ -151,6 +156,60 @@ def check_nd_call_func( ) +_acos_docstring = """ +acos(x, out=None, order='K') + +Computes inverse cosine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse cosine, in radians + and in the closed interval `[-pi/2, pi/2]`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_acos(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_acos_to_call(sycl_queue, src, dst): + # call pybind11 extension for acos() function from OneMKL VM + return vmi._acos(sycl_queue, src, dst, depends) + return ti._acos(src, dst, sycl_queue, depends) + + +acos_func = UnaryElementwiseFunc( + "acos", ti._acos_result_type, _call_acos, _acos_docstring +) + + +def dpnp_acos(x, out=None, order="K"): + """ + Invokes acos() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for acos() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = acos_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _add_docstring_ = """ add(x1, x2, out=None, order="K") @@ -210,6 +269,180 @@ def dpnp_add(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_asin_docstring = """ +asin(x, out=None, order='K') + +Computes inverse sine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse sine, in radians + and in the closed interval `[-pi/2, pi/2]`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_asin(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_asin_to_call(sycl_queue, src, dst): + # call pybind11 extension for asin() function from OneMKL VM + return vmi._asin(sycl_queue, src, dst, depends) + return ti._asin(src, dst, sycl_queue, depends) + + +asin_func = UnaryElementwiseFunc( + "asin", ti._asin_result_type, _call_asin, _asin_docstring +) + + +def dpnp_asin(x, out=None, order="K"): + """ + Invokes asin() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for asin() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = asin_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_atan_docstring = """ +atan(x, out=None, order='K') + +Computes inverse tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse tangent, in radians + and in the closed interval `[-pi/2, pi/2]`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_atan(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_atan_to_call(sycl_queue, src, dst): + # call pybind11 extension for atan() function from OneMKL VM + return vmi._atan(sycl_queue, src, dst, depends) + return ti._atan(src, dst, sycl_queue, depends) + + +atan_func = UnaryElementwiseFunc( + "atan", ti._atan_result_type, _call_atan, _atan_docstring +) + + +def dpnp_atan(x, out=None, order="K"): + """ + Invokes atan() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for atan() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = atan_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + +_atan2_docstring_ = """ +atan2(x1, x2, out=None, order="K") + +Calculates the inverse tangent of the quotient `x1_i/x2_i` for each element +`x1_i` of the input array `x1` with the respective element `x2_i` of the +input array `x2`. Each element-wise result is expressed in radians. + +Args: + x1 (dpnp.ndarray): + First input array, expected to have a real-valued floating-point + data type. + x2 (dpnp.ndarray): + Second input array, also expected to have a real-valued + floating-point data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", None, optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + dpnp.ndarray: + An array containing the inverse tangent of the quotient `x1`/`x2`. + The returned array must have a real-valued floating-point data type + determined by Type Promotion Rules. +""" + + +def _call_atan2(src1, src2, dst, sycl_queue, depends=None): + """A callback to register in BinaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_atan2_to_call(sycl_queue, src1, src2, dst): + # call pybind11 extension for atan2() function from OneMKL VM + return vmi._atan2(sycl_queue, src1, src2, dst, depends) + return ti._atan2(src1, src2, dst, sycl_queue, depends) + + +atan2_func = BinaryElementwiseFunc( + "atan2", + ti._atan2_result_type, + _call_atan2, + _atan2_docstring_, +) + + +def dpnp_atan2(x1, x2, out=None, order="K"): + """ + Invokes atan2() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for atan2() function. + """ + + # dpctl.tensor only works with usm_ndarray or scalar + x1_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x1) + x2_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x2) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = atan2_func( + x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order + ) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _bitwise_and_docstring_ = """ bitwise_and(x1, x2, out=None, order='K') @@ -409,7 +642,9 @@ def dpnp_ceil(x, out=None, order="K"): _cos_docstring = """ cos(x, out=None, order='K') + Computes cosine for each element `x_i` for input array `x`. + Args: x (dpnp.ndarray): Input array, expected to have numeric data type. @@ -426,6 +661,38 @@ def dpnp_ceil(x, out=None, order="K"): """ +def _call_cos(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_cos_to_call(sycl_queue, src, dst): + # call pybind11 extension for cos() function from OneMKL VM + return vmi._cos(sycl_queue, src, dst, depends) + return ti._cos(src, dst, sycl_queue, depends) + + +cos_func = UnaryElementwiseFunc( + "cos", ti._cos_result_type, _call_cos, _cos_docstring +) + + +def dpnp_cos(x, out=None, order="K"): + """ + Invokes cos() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for cos() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = cos_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _conj_docstring = """ conj(x, out=None, order='K') @@ -464,36 +731,6 @@ def _call_conj(src, dst, sycl_queue, depends=None): ) -def dpnp_cos(x, out=None, order="K"): - """ - Invokes cos() function from pybind11 extension of OneMKL VM if possible. - - Otherwise fully relies on dpctl.tensor implementation for cos() function. - - """ - - def _call_cos(src, dst, sycl_queue, depends=None): - """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" - - if depends is None: - depends = [] - - if vmi._mkl_cos_to_call(sycl_queue, src, dst): - # call pybind11 extension for cos() function from OneMKL VM - return vmi._cos(sycl_queue, src, dst, depends) - return ti._cos(src, dst, sycl_queue, depends) - - # dpctl.tensor only works with usm_ndarray - x1_usm = dpnp.get_usm_ndarray(x) - out_usm = None if out is None else dpnp.get_usm_ndarray(out) - - func = UnaryElementwiseFunc( - "cos", ti._cos_result_type, _call_cos, _cos_docstring - ) - res_usm = func(x1_usm, out=out_usm, order=order) - return dpnp_array._create_from_usm_ndarray(res_usm) - - def dpnp_conj(x, out=None, order="K"): """ Invokes conj() function from pybind11 extension of OneMKL VM if possible. @@ -980,8 +1217,8 @@ def dpnp_isnan(x, out=None, order="K"): """ -leftt_shift_func = BinaryElementwiseFunc( - "bitwise_leftt_shift", +left_shift_func = BinaryElementwiseFunc( + "bitwise_left_shift", ti._bitwise_left_shift_result_type, ti._bitwise_left_shift, _left_shift_docstring_, @@ -996,7 +1233,7 @@ def dpnp_left_shift(x1, x2, out=None, order="K"): x2_usm_or_scalar = dpnp.get_usm_ndarray_or_scalar(x2) out_usm = None if out is None else dpnp.get_usm_ndarray(out) - res_usm = leftt_shift_func( + res_usm = left_shift_func( x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order ) return dpnp_array._create_from_usm_ndarray(res_usm) @@ -1789,7 +2026,9 @@ def dpnp_signbit(x, out=None, order="K"): _sin_docstring = """ sin(x, out=None, order='K') + Computes sine for each element `x_i` of input array `x`. + Args: x (dpnp.ndarray): Input array, expected to have numeric data type. @@ -2018,6 +2257,59 @@ def dpnp_subtract(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_tan_docstring = """ +tan(x, out=None, order='K') + +Computes tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise tangent. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_tan(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_tan_to_call(sycl_queue, src, dst): + # call pybind11 extension for tan() function from OneMKL VM + return vmi._tan(sycl_queue, src, dst, depends) + return ti._tan(src, dst, sycl_queue, depends) + + +tan_func = UnaryElementwiseFunc( + "tan", ti._tan_result_type, _call_tan, _tan_docstring +) + + +def dpnp_tan(x, out=None, order="K"): + """ + Invokes tan() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for tan() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = tan_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _trunc_docstring = """ trunc(x, out=None, order='K') diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index f01ff69a4801..39dba2d118ae 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -48,11 +48,16 @@ from .dpnp_algo.dpnp_elementwise_common import ( check_nd_call_func, + dpnp_acos, + dpnp_asin, + dpnp_atan, + dpnp_atan2, dpnp_cos, dpnp_log, dpnp_sin, dpnp_sqrt, dpnp_square, + dpnp_tan, ) __all__ = [ @@ -89,45 +94,67 @@ ] -def arccos(x1): +def arccos( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Trigonometric inverse cosine, element-wise. For full documentation refer to :obj:`numpy.arccos`. + Returns + ------- + out : dpnp.ndarray + The inverse cosine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- - :obj:`dpnp.cos` : Cosine element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. - :obj:`dpnp.arcsin` : Inverse sine, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. Examples -------- >>> import dpnp as np >>> x = np.array([1, -1]) - >>> out = np.arccos(x) - >>> [i for i in out] - [0.0, 3.14159265] + >>> np.arccos(x) + array([0.0, 3.14159265]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arccos, + dpnp_acos, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arccos(x1_desc).get_pyobj() - - return call_origin(numpy.arccos, x1, **kwargs) def arccosh(x1): """ - Trigonometric inverse hyperbolic cosine, element-wise. + Inverse hyperbolic cosine, element-wise. For full documentation refer to :obj:`numpy.arccosh`. @@ -163,50 +190,62 @@ def arccosh(x1): return call_origin(numpy.arccosh, x1, **kwargs) -def arcsin(x1, out=None, **kwargs): +def arcsin( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Inverse sine, element-wise. For full documentation refer to :obj:`numpy.arcsin`. + Returns + ------- + out : dpnp.ndarray + The inverse sine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- :obj:`dpnp.sin` : Trigonometric sine, element-wise. - :obj:`dpnp.cos` : Cosine element-wise. - :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. - :obj:`dpnp.tan` : Compute tangent element-wise. :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. - :obj:`dpnp.arctan2` : Element-wise arc tangent of ``x1/x2`` - choosing the quadrant correctly. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. Examples -------- >>> import dpnp as np >>> x = np.array([0, 1, -1]) - >>> out = np.arcsin(x) - >>> [i for i in out] - [0.0, 1.5707963267948966, -1.5707963267948966] + >>> np.arcsin(x) + array([0.0, 1.5707963267948966, -1.5707963267948966]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arcsin, + dpnp_asin, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_arcsin(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.arcsin, x1, out=out, **kwargs) def arcsinh(x1): @@ -239,51 +278,143 @@ def arcsinh(x1): return call_origin(numpy.arcsinh, x1, **kwargs) -def arctan(x1, out=None, **kwargs): +def arctan( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Trigonometric inverse tangent, element-wise. For full documentation refer to :obj:`numpy.arctan`. + Returns + ------- + out : dpnp.ndarray + The inverse tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- - :obj:`dpnp.arctan2` : Element-wise arc tangent of ``x1/x2`` - choosing the quadrant correctly. + :obj:`dpnp.arctan2` : Element-wise arc tangent of `x1/x2` choosing the quadrant correctly. :obj:`dpnp.angle` : Argument of complex values. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.arctanh` : Inverse hyperbolic tangent, element-wise. Examples -------- >>> import dpnp as np >>> x = np.array([0, 1]) - >>> out = np.arctan(x) - >>> [i for i in out] - [0.0, 0.78539816] + >>> np.arctan(x) + array([0.0, 0.78539816]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arctan, + dpnp_atan, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_arctan(x1_desc, out_desc).get_pyobj() - return call_origin(numpy.arctan, x1, out=out, **kwargs) + +def arctan2( + x1, + x2, + /, + out=None, + *, + where=True, + order="K", + dtype=None, + subok=True, + **kwargs, +): + """ + Element-wise arc tangent of `x1/x2` choosing the quadrant correctly. + + For full documentation refer to :obj:`numpy.arctan2`. + + Returns + ------- + out : dpnp.ndarray + The inverse tangent of `x1/x2`, element-wise. + + Limitations + ----------- + Parameters `x1` and `x2` are supported as either scalar, :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`, but both `x1` and `x2` can not be scalars at the same time. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword arguments `kwargs` are currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. + + See Also + -------- + :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. + :obj:`dpnp.tan` : Compute tangent element-wise. + :obj:`dpnp.angle` : Return the angle of the complex argument. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.arctanh` : Inverse hyperbolic tangent, element-wise. + + Examples + -------- + >>> import dpnp as np + >>> x1 = np.array([1., -1.]) + >>> x2 = np.array([0., 0.]) + >>> np.arctan2(x1, x2) + array([1.57079633, -1.57079633]) + + >>> x1 = np.array([0., 0., np.inf]) + >>> x2 = np.array([+0., -0., np.inf]) + >>> np.arctan2(x1, x2) + array([0.0 , 3.14159265, 0.78539816]) + + >>> x1 = np.array([-1, +1, +1, -1]) + >>> x2 = np.array([-1, -1, +1, +1]) + >>> np.arctan2(x1, x2) * 180 / np.pi + array([-135., -45., 45., 135.]) + + """ + + return check_nd_call_func( + numpy.arctan2, + dpnp_atan2, + x1, + x2, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, + ) def arctanh(x1): """ - Trigonometric hyperbolic inverse tangent, element-wise. + Hyperbolic inverse tangent, element-wise. For full documentation refer to :obj:`numpy.arctanh`. @@ -341,76 +472,6 @@ def cbrt(x1): return call_origin(numpy.cbrt, x1, **kwargs) -def arctan2(x1, x2, dtype=None, out=None, where=True, **kwargs): - """ - Element-wise arc tangent of ``x1/x2`` choosing the quadrant correctly. - - For full documentation refer to :obj:`numpy.arctan2`. - - Limitations - ----------- - Parameters `x1` and `x2` are supported as either :obj:`dpnp.ndarray` or scalar. - Parameters `dtype`, `out` and `where` are supported with their default values. - Keyword argument `kwargs` is currently unsupported. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. - - See Also - -------- - :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. - :obj:`dpnp.tan` : Compute tangent element-wise. - :obj:`dpnp.angle` : Return the angle of the complex argument. - - Examples - -------- - >>> import dpnp as np - >>> x1 = np.array([1., -1.]) - >>> x2 = np.array([0., 0.]) - >>> out = np.arctan2(x1, x2) - >>> [i for i in out] - [1.57079633, -1.57079633] - - """ - - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False - ) - x2_desc = dpnp.get_dpnp_descriptor( - x2, copy_when_strides=False, copy_when_nondefault_queue=False - ) - - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif dtype is not None: - pass - elif not where: - pass - else: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_arctan2( - x1_desc, x2_desc, dtype, out_desc, where - ).get_pyobj() - - return call_origin( - numpy.arctan2, x1, x2, dtype=dtype, out=out, where=where, **kwargs - ) - - def cos( x, /, @@ -440,6 +501,13 @@ def cos( Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. + :obj:`dpnp.sin` : Trigonometric sine, element-wise. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. + :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. + Examples -------- >>> import dpnp as np @@ -464,7 +532,7 @@ def cos( def cosh(x1): """ - Trigonometric hyperbolic cosine, element-wise. + Hyperbolic cosine, element-wise. For full documentation refer to :obj:`numpy.cosh`. @@ -546,7 +614,7 @@ def degrees(x1): def exp(x1, out=None, **kwargs): """ - Trigonometric exponent, element-wise. + Calculate the exponential, element-wise. For full documentation refer to :obj:`numpy.exp`. @@ -586,7 +654,7 @@ def exp(x1, out=None, **kwargs): def exp2(x1): """ - Trigonometric exponent2, element-wise. + Calculate `2**p` for all `p` in the input array. For full documentation refer to :obj:`numpy.exp2`. @@ -621,7 +689,7 @@ def exp2(x1): def expm1(x1): """ - Trigonometric exponent minus 1, element-wise. + Return the exponential of the input array minus one, element-wise. For full documentation refer to :obj:`numpy.expm1`. @@ -810,7 +878,7 @@ def log( def log10(x1): """ - Trigonometric logarithm, element-wise. + Return the base 10 logarithm of the input array, element-wise. For full documentation refer to :obj:`numpy.log10`. @@ -840,7 +908,7 @@ def log10(x1): def log1p(x1): """ - Trigonometric logarithm, element-wise. + Return the natural logarithm of one plus the input array, element-wise. For full documentation refer to :obj:`numpy.log1p`. @@ -874,7 +942,7 @@ def log1p(x1): def log2(x1): """ - Trigonometric logarithm, element-wise. + Return the base 2 logarithm of the input array, element-wise. For full documentation refer to :obj:`numpy.log2`. @@ -1025,9 +1093,10 @@ def sin( See Also -------- - :obj:`dpnp.arcsin` : Inverse sine, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. - :obj:`dpnp.cos` : Cosine element-wise. Examples -------- @@ -1053,7 +1122,7 @@ def sin( def sinh(x1): """ - Trigonometric hyperbolic sine, element-wise. + Hyperbolic sine, element-wise. For full documentation refer to :obj:`numpy.sinh`. @@ -1202,40 +1271,62 @@ def square( ) -def tan(x1, out=None, **kwargs): +def tan( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Compute tangent element-wise. + Trigonometric tangent, element-wise. For full documentation refer to :obj:`numpy.tan`. + Returns + ------- + out : dpnp.ndarray + The tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. + :obj:`dpnp.sin` : Trigonometric sine, element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([-np.pi, np.pi/2, np.pi]) - >>> out = np.tan(x) - >>> [i for i in out] - [1.22460635e-16, 1.63317787e+16, -1.22460635e-16] + >>> np.tan(x) + array([1.22460635e-16, 1.63317787e+16, -1.22460635e-16]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.tan, + dpnp_tan, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_tan(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.tan, x1, out=out, **kwargs) def tanh(x1): diff --git a/tests/skipped_tests_gpu_no_fp64.tbl b/tests/skipped_tests_gpu_no_fp64.tbl index 1339c7cfba50..c92ef28d471d 100644 --- a/tests/skipped_tests_gpu_no_fp64.tbl +++ b/tests/skipped_tests_gpu_no_fp64.tbl @@ -78,8 +78,8 @@ tests/test_strides.py::test_strides_tan[(10,)-None] tests/test_sycl_queue.py::test_array_creation[opencl:gpu:0-arange-arg0-kwargs0] tests/test_sycl_queue.py::test_array_creation[level_zero:gpu:0-arange-arg0-kwargs0] -tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-gradient-data10] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data10] +tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-gradient-data14] +tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data14] tests/test_sycl_queue.py::test_eig[opencl:gpu:0] tests/test_sycl_queue.py::test_eig[level_zero:gpu:0] tests/test_sycl_queue.py::test_eigh[opencl:gpu:0] @@ -778,14 +778,9 @@ tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_3 tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_31_{axis=1, func='nanprod', keepdims=False, shape=(20, 30, 40), transpose_axes=False}::test_nansum_axis_transposed tests/third_party/cupy/math_tests/test_sumprod.py::TestDiff::test_diff_2dim_with_scalar_append -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arccos -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arcsin -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arctan -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_arctan2 tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_deg2rad tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_hypot tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_rad2deg -tests/third_party/cupy/math_tests/test_trigonometric.py::TestTrigonometric::test_tan tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_6_{a_shape=(3, 2), b_shape=(3, 2), shape=(4, 3, 2)}::test_beta tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_7_{a_shape=(3, 2), b_shape=(3, 2), shape=(3, 2)}::test_beta diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index aaa7b3bae8e9..f64cd813c429 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -225,8 +225,14 @@ def test_meshgrid(device_x, device_y): "func,data", [ pytest.param("abs", [-1.2, 1.2]), + pytest.param("arccos", [-0.5, 0.0, 0.5]), + pytest.param("arcsin", [-0.5, 0.0, 0.5]), + pytest.param("arctan", [-1.0, 0.0, 1.0]), pytest.param("ceil", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("conjugate", [[1.0 + 1.0j, 0.0], [0.0, 1.0 + 1.0j]]), + pytest.param( + "cos", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] + ), pytest.param("copy", [1.0, 2.0, 3.0]), pytest.param("cumprod", [[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]), pytest.param("cumsum", [[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]), @@ -243,8 +249,14 @@ def test_meshgrid(device_x, device_y): pytest.param("prod", [1.0, 2.0]), pytest.param("sign", [-5.0, 0.0, 4.5]), pytest.param("signbit", [-5.0, 0.0, 4.5]), + pytest.param( + "sin", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] + ), pytest.param("sqrt", [1.0, 3.0, 9.0]), pytest.param("sum", [1.0, 2.0]), + pytest.param( + "tan", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] + ), pytest.param("trapz", [[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]]), pytest.param("trunc", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), ], @@ -255,13 +267,13 @@ def test_meshgrid(device_x, device_y): ids=[device.filter_string for device in valid_devices], ) def test_1in_1out(func, data, device): - x_orig = numpy.array(data) - expected = getattr(numpy, func)(x_orig) - x = dpnp.array(data, device=device) result = getattr(dpnp, func)(x) - assert_allclose(result, expected) + x_orig = dpnp.asnumpy(x) + expected = getattr(numpy, func)(x_orig) + + assert_allclose(result, expected, rtol=2e-07) expected_queue = x.get_array().sycl_queue result_queue = result.get_array().sycl_queue @@ -311,6 +323,11 @@ def test_proj(device): [1.0, dpnp.inf, -dpnp.inf], [1.0, dpnp.inf, -dpnp.inf], ), + pytest.param( + "arctan2", + [[-1, +1, +1, -1]], + [[-1, -1, +1, +1]], + ), pytest.param("copysign", [0.0, 1.0, 2.0], [-1.0, 0.0, 1.0]), pytest.param("cross", [1.0, 2.0, 3.0], [4.0, 5.0, 6.0]), pytest.param( diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 8eb8b1b779d8..ea52ba78b14e 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -275,14 +275,26 @@ def test_meshgrid(usm_type_x, usm_type_y): @pytest.mark.parametrize( "func,data", [ + pytest.param("arccos", [-0.5, 0.0, 0.5]), + pytest.param("arcsin", [-0.5, 0.0, 0.5]), + pytest.param("arctan", [-1.0, 0.0, 1.0]), pytest.param("ceil", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("conjugate", [[1.0 + 1.0j, 0.0], [0.0, 1.0 + 1.0j]]), + pytest.param( + "cos", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] + ), pytest.param("floor", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("negative", [1.0, 0.0, -1.0]), pytest.param("proj", [complex(1.0, 2.0), complex(dp.inf, -1.0)]), pytest.param("sign", [-5.0, 0.0, 4.5]), pytest.param("signbit", [-5.0, 0.0, 4.5]), + pytest.param( + "sin", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] + ), pytest.param("sqrt", [1.0, 3.0, 9.0]), + pytest.param( + "tan", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] + ), pytest.param("trunc", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), ], ) @@ -302,6 +314,11 @@ def test_1in_1out(func, data, usm_type): [[1.2, -0.0], [-7, 2.34567]], [[1.2, 0.0], [-7, 2.34567]], ), + pytest.param( + "arctan2", + [[-1, +1, +1, -1]], + [[-1, -1, +1, +1]], + ), pytest.param( "dot", [[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]], diff --git a/tests/third_party/cupy/math_tests/test_trigonometric.py b/tests/third_party/cupy/math_tests/test_trigonometric.py index 60533205d930..506dd6aae44d 100644 --- a/tests/third_party/cupy/math_tests/test_trigonometric.py +++ b/tests/third_party/cupy/math_tests/test_trigonometric.py @@ -1,17 +1,18 @@ import unittest +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing class TestTrigonometric(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose(atol=1e-5, type_check=False) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def check_binary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) b = testing.shaped_reverse_arange((2, 3), xp, dtype) From 1cc6dd519e04f23aea68f125fb6c0146df8e80af Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 29 Aug 2023 19:10:07 -0500 Subject: [PATCH 02/10] update tests and skipped test file --- tests/skipped_tests_gpu_no_fp64.tbl | 37 --------- tests/test_sycl_queue.py | 3 +- tests/test_umath.py | 112 +++++++++++++++------------- 3 files changed, 62 insertions(+), 90 deletions(-) diff --git a/tests/skipped_tests_gpu_no_fp64.tbl b/tests/skipped_tests_gpu_no_fp64.tbl index c92ef28d471d..38d088971f39 100644 --- a/tests/skipped_tests_gpu_no_fp64.tbl +++ b/tests/skipped_tests_gpu_no_fp64.tbl @@ -15,11 +15,8 @@ tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array0] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array1] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array2] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arccos] tests/test_strides.py::test_strides_1arg[(10,)-int32-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arcsin] tests/test_strides.py::test_strides_1arg[(10,)-int32-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arctan] tests/test_strides.py::test_strides_1arg[(10,)-int32-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-int32-cbrt] tests/test_strides.py::test_strides_1arg[(10,)-int32-cosh] @@ -34,11 +31,8 @@ tests/test_strides.py::test_strides_1arg[(10,)-int32-log2] tests/test_strides.py::test_strides_1arg[(10,)-int32-radians] tests/test_strides.py::test_strides_1arg[(10,)-int32-sinh] tests/test_strides.py::test_strides_1arg[(10,)-int32-tanh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arccos] tests/test_strides.py::test_strides_1arg[(10,)-int64-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arcsin] tests/test_strides.py::test_strides_1arg[(10,)-int64-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arctan] tests/test_strides.py::test_strides_1arg[(10,)-int64-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-int64-cbrt] tests/test_strides.py::test_strides_1arg[(10,)-int64-cosh] @@ -53,11 +47,8 @@ tests/test_strides.py::test_strides_1arg[(10,)-int64-log2] tests/test_strides.py::test_strides_1arg[(10,)-int64-radians] tests/test_strides.py::test_strides_1arg[(10,)-int64-sinh] tests/test_strides.py::test_strides_1arg[(10,)-int64-tanh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arccos] tests/test_strides.py::test_strides_1arg[(10,)-None-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arcsin] tests/test_strides.py::test_strides_1arg[(10,)-None-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arctan] tests/test_strides.py::test_strides_1arg[(10,)-None-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-None-cbrt] tests/test_strides.py::test_strides_1arg[(10,)-None-cosh] @@ -72,9 +63,6 @@ tests/test_strides.py::test_strides_1arg[(10,)-None-log2] tests/test_strides.py::test_strides_1arg[(10,)-None-radians] tests/test_strides.py::test_strides_1arg[(10,)-None-sinh] tests/test_strides.py::test_strides_1arg[(10,)-None-tanh] -tests/test_strides.py::test_strides_tan[(10,)-int32] -tests/test_strides.py::test_strides_tan[(10,)-int64] -tests/test_strides.py::test_strides_tan[(10,)-None] tests/test_sycl_queue.py::test_array_creation[opencl:gpu:0-arange-arg0-kwargs0] tests/test_sycl_queue.py::test_array_creation[level_zero:gpu:0-arange-arg0-kwargs0] @@ -100,31 +88,6 @@ tests/test_umath.py::TestExp::test_invalid_dtype[numpy.int32] tests/test_umath.py::TestExp::test_invalid_shape[(0,)] tests/test_umath.py::TestExp::test_invalid_shape[(15, )] tests/test_umath.py::TestExp::test_invalid_shape[(2,2)] -tests/test_umath.py::TestArcsin::test_arcsin -tests/test_umath.py::TestArcsin::test_invalid_dtype[numpy.float32] -tests/test_umath.py::TestArcsin::test_invalid_dtype[numpy.int64] -tests/test_umath.py::TestArcsin::test_invalid_dtype[numpy.int32] -tests/test_umath.py::TestArcsin::test_invalid_shape[(0,)] -tests/test_umath.py::TestArcsin::test_invalid_shape[(15, )] -tests/test_umath.py::TestArcsin::test_invalid_shape[(2,2)] -tests/test_umath.py::TestArctan::test_arctan -tests/test_umath.py::TestArctan::test_invalid_dtype[numpy.float32] -tests/test_umath.py::TestArctan::test_invalid_dtype[numpy.int64] -tests/test_umath.py::TestArctan::test_invalid_dtype[numpy.int32] -tests/test_umath.py::TestArctan::test_invalid_shape[(0,)] -tests/test_umath.py::TestArctan::test_invalid_shape[(15, )] -tests/test_umath.py::TestArctan::test_invalid_shape[(2,2)] -tests/test_umath.py::TestTan::test_tan -tests/test_umath.py::TestTan::test_invalid_dtype[numpy.float32] -tests/test_umath.py::TestTan::test_invalid_dtype[numpy.int64] -tests/test_umath.py::TestTan::test_invalid_dtype[numpy.int32] -tests/test_umath.py::TestTan::test_invalid_shape[(0,)] -tests/test_umath.py::TestTan::test_invalid_shape[(15, )] -tests/test_umath.py::TestTan::test_invalid_shape[(2,2)] -tests/test_umath.py::TestArctan2::test_arctan2 -tests/test_umath.py::TestArctan2::test_invalid_shape[(0,)] -tests/test_umath.py::TestArctan2::test_invalid_shape[(15, )] -tests/test_umath.py::TestArctan2::test_invalid_shape[(2,2)] tests/test_umath.py::TestSqrt::test_sqrt_complex[complex64] diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index f64cd813c429..37b678554345 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -273,7 +273,8 @@ def test_1in_1out(func, data, device): x_orig = dpnp.asnumpy(x) expected = getattr(numpy, func)(x_orig) - assert_allclose(result, expected, rtol=2e-07) + tol = numpy.finfo(x.dtype).resolution + assert_allclose(result, expected, rtol=tol) expected_queue = x.get_array().sycl_queue result_queue = result.get_array().sycl_queue diff --git a/tests/test_umath.py b/tests/test_umath.py index 8f1ac15d6ad8..bb055097eacc 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -11,6 +11,7 @@ from .helper import ( get_all_dtypes, get_complex_dtypes, + get_float_dtypes, has_support_aspect16, has_support_aspect64, ) @@ -381,137 +382,143 @@ def test_invalid_shape(self, shape): class TestArcsin: - def test_arcsin(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arcsin(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.arcsin(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arcsin(np_array, out=out) assert_array_equal(expected, result) @pytest.mark.parametrize( - "dtype", - [numpy.float32, numpy.int64, numpy.int32], - ids=["numpy.float32", "numpy.int64", "numpy.int32"], + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.arcsin(dp_array, out=dp_out) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.arcsin(dp_array, out=dp_out) class TestArctan: - def test_arctan(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_arctan(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.arctan(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arctan(np_array, out=out) - assert_array_equal(expected, result) + assert_allclose(expected, result) @pytest.mark.parametrize( - "dtype", - [numpy.float32, numpy.int64, numpy.int32], - ids=["numpy.float32", "numpy.int64", "numpy.int32"], + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.arctan(dp_array, out=dp_out) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.arctan(dp_array, out=dp_out) class TestTan: - def test_tan(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_tan(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.tan(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.tan(np_array, out=out) - assert_allclose(expected, result) + tol = numpy.finfo(dtype).resolution + assert_allclose(expected, result, rtol=tol) @pytest.mark.parametrize( - "dtype", - [numpy.float32, numpy.int64, numpy.int32], - ids=["numpy.float32", "numpy.int64", "numpy.int32"], + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.tan(dp_array, out=dp_out) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.tan(dp_array, out=dp_out) class TestArctan2: - def test_arctan2(self): + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_arctan2(self, dtype): array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + out = numpy.empty(10, dtype=dtype) # DPNP - dp_array = dpnp.array(array_data, dtype=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) result = dpnp.arctan2(dp_array, dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) + np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arctan2(np_array, np_array, out=out) - assert_array_equal(expected, result) + assert_allclose(expected, result) @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) @@ -520,21 +527,22 @@ def test_out_dtypes(self, dtype): size = 2 if dtype == dpnp.bool else 10 np_array = numpy.arange(size, dtype=dtype) - np_out = numpy.empty(size, dtype=numpy.complex64) + np_out = numpy.empty(size, dtype=numpy.float32) expected = numpy.arctan2(np_array, np_array, out=np_out) dp_array = dpnp.arange(size, dtype=dtype) - dp_out = dpnp.empty(size, dtype=dpnp.complex64) + dp_out = dpnp.empty(size, dtype=dpnp.float32) result = dpnp.arctan2(dp_array, dp_array, out=dp_out) assert_allclose(expected, result) + @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) - def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) with pytest.raises(ValueError): dpnp.arctan2(dp_array, dp_array, out=dp_out) From 1a951ad0c247ca1e477957a709c2bf24445938ba Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Tue, 29 Aug 2023 21:22:57 -0500 Subject: [PATCH 03/10] fix test_out_dtype arctan2 --- tests/test_umath.py | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/tests/test_umath.py b/tests/test_umath.py index bb055097eacc..805f1f4d2638 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -524,14 +524,18 @@ def test_arctan2(self, dtype): "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) ) def test_out_dtypes(self, dtype): + if has_support_aspect64() and dtype != numpy.float32: + dtype_out = numpy.float64 + else: + dtype_out = numpy.float32 size = 2 if dtype == dpnp.bool else 10 np_array = numpy.arange(size, dtype=dtype) - np_out = numpy.empty(size, dtype=numpy.float32) + np_out = numpy.empty(size, dtype=dtype_out) expected = numpy.arctan2(np_array, np_array, out=np_out) dp_array = dpnp.arange(size, dtype=dtype) - dp_out = dpnp.empty(size, dtype=dpnp.float32) + dp_out = dpnp.empty(size, dtype=dtype_out) result = dpnp.arctan2(dp_array, dp_array, out=dp_out) assert_allclose(expected, result) From ed4527080057263a5d8eac7e3f3cfc25e9cafe51 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 30 Aug 2023 13:23:54 -0500 Subject: [PATCH 04/10] address comments --- dpnp/dpnp_algo/dpnp_algo.pxd | 10 ----- dpnp/dpnp_iface_trigonometric.py | 2 +- tests/test_mathematical.py | 16 ++++--- tests/test_strides.py | 19 ++------ tests/test_umath.py | 44 ++++++++++++++++++- .../cupy/math_tests/test_trigonometric.py | 2 +- 6 files changed, 59 insertions(+), 34 deletions(-) diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 2197333af5c9..966df2993dde 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -38,18 +38,10 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_ALLCLOSE DPNP_FN_ALLCLOSE_EXT DPNP_FN_ARANGE - DPNP_FN_ARCCOS - DPNP_FN_ARCCOS_EXT DPNP_FN_ARCCOSH DPNP_FN_ARCCOSH_EXT - DPNP_FN_ARCSIN - DPNP_FN_ARCSIN_EXT DPNP_FN_ARCSINH DPNP_FN_ARCSINH_EXT - DPNP_FN_ARCTAN - DPNP_FN_ARCTAN_EXT - DPNP_FN_ARCTAN2 - DPNP_FN_ARCTAN2_EXT DPNP_FN_ARCTANH DPNP_FN_ARCTANH_EXT DPNP_FN_ARGMAX @@ -259,8 +251,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_SUM_EXT DPNP_FN_SVD DPNP_FN_SVD_EXT - DPNP_FN_TAN - DPNP_FN_TAN_EXT DPNP_FN_TANH DPNP_FN_TANH_EXT DPNP_FN_TRACE diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index 39dba2d118ae..c7b250249b38 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -305,7 +305,7 @@ def arctan( Parameters `where`, `dtype` and `subok` are supported with their default values. Keyword argument `kwargs` is currently unsupported. Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Input array data types are limited by supported real-valued floating-point data type. See Also -------- diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index cac1e3385c21..f47bca873474 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -137,16 +137,22 @@ def _test_mathematical(self, name, dtype, lhs, rhs): else: result = getattr(dpnp, name)(a_dpnp, b_dpnp) expected = getattr(numpy, name)(a_np, b_np) - assert_allclose(result, expected, rtol=1e-6) + if numpy.issubdtype(expected.dtype, numpy.floating): + tol = numpy.max( + [ + numpy.finfo(result.dtype).resolution, + numpy.finfo(expected.dtype).resolution, + ] + ) + else: + tol = 1e-06 + assert_allclose(result, expected, rtol=tol) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): self._test_mathematical("add", dtype, lhs, rhs) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") - @pytest.mark.parametrize( - "dtype", get_all_dtypes(no_bool=True, no_complex=True) - ) + @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_arctan2(self, dtype, lhs, rhs): self._test_mathematical("arctan2", dtype, lhs, rhs) diff --git a/tests/test_strides.py b/tests/test_strides.py index 4215e77e8cac..a6cfbfbde820 100644 --- a/tests/test_strides.py +++ b/tests/test_strides.py @@ -74,6 +74,7 @@ def test_strides(func_name, dtype): "sinh", "sqrt", "square", + "tan", "tanh", "trunc", ], @@ -93,7 +94,7 @@ def test_strides_1arg(func_name, dtype, shape): numpy_func = _getattr(numpy, func_name) expected = numpy_func(b) - assert_allclose(result, expected) + assert_allclose(result, expected, rtol=1e-06) @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) @@ -130,21 +131,6 @@ def test_strides_reciprocal(dtype, shape): assert_allclose(result, expected, rtol=1e-06) -@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) -@pytest.mark.parametrize("shape", [(10,)], ids=["(10,)"]) -def test_strides_tan(dtype, shape): - a = numpy.arange(numpy.prod(shape), dtype=dtype).reshape(shape) - b = a[::2] - - dpa = dpnp.reshape(dpnp.arange(numpy.prod(shape), dtype=dtype), shape) - dpb = dpa[::2] - - result = dpnp.tan(dpb) - expected = numpy.tan(b) - - assert_allclose(result, expected, rtol=1e-06) - - @pytest.mark.parametrize( "func_name", [ @@ -161,6 +147,7 @@ def test_strides_tan(dtype, shape): ) @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) @pytest.mark.parametrize("shape", [(3, 3)], ids=["(3, 3)"]) +@pytest.mark.usefixtures("suppress_invalid_numpy_warnings") def test_strides_2args(func_name, dtype, shape): a = numpy.arange(numpy.prod(shape), dtype=dtype).reshape(shape) b = a.T diff --git a/tests/test_umath.py b/tests/test_umath.py index 805f1f4d2638..ce918d849d53 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -381,6 +381,47 @@ def test_invalid_shape(self, shape): dpnp.exp(dp_array, out=dp_out) +class TestArccos: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arccos(self, dtype): + array_data = numpy.arange(10) + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arccos(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arccos(np_array, out=out) + + assert_array_equal(expected, result) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arccos(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arccos(dp_array, out=dp_out) + + class TestArcsin: @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") @@ -437,7 +478,8 @@ def test_arctan(self, dtype): np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arctan(np_array, out=out) - assert_allclose(expected, result) + tol = numpy.finfo(dtype).resolution + assert_allclose(expected, result, tol) @pytest.mark.parametrize( "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] diff --git a/tests/third_party/cupy/math_tests/test_trigonometric.py b/tests/third_party/cupy/math_tests/test_trigonometric.py index 506dd6aae44d..778dc461454d 100644 --- a/tests/third_party/cupy/math_tests/test_trigonometric.py +++ b/tests/third_party/cupy/math_tests/test_trigonometric.py @@ -18,7 +18,7 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) - @testing.for_dtypes(["f", "d"]) + @testing.for_dtypes(["e", "f", "d"]) @testing.numpy_cupy_allclose(atol=1e-5) def check_unary_unit(self, name, xp, dtype): a = xp.array([0.2, 0.4, 0.6, 0.8], dtype=dtype) From 4874bea23543a8f0ac2814f2251f13e51b2ea2f6 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 30 Aug 2023 15:25:44 -0500 Subject: [PATCH 05/10] use assert_dtype_allclose --- tests/test_mathematical.py | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index f47bca873474..4cc58f21b1cc 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -13,6 +13,7 @@ import dpnp from .helper import ( + assert_dtype_allclose, get_all_dtypes, get_complex_dtypes, get_float_dtypes, @@ -137,16 +138,7 @@ def _test_mathematical(self, name, dtype, lhs, rhs): else: result = getattr(dpnp, name)(a_dpnp, b_dpnp) expected = getattr(numpy, name)(a_np, b_np) - if numpy.issubdtype(expected.dtype, numpy.floating): - tol = numpy.max( - [ - numpy.finfo(result.dtype).resolution, - numpy.finfo(expected.dtype).resolution, - ] - ) - else: - tol = 1e-06 - assert_allclose(result, expected, rtol=tol) + assert_dtype_allclose(result, expected) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): From ec85cfe3a36937b29d35f481b927a0e2796f8acc Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 30 Aug 2023 19:53:38 -0500 Subject: [PATCH 06/10] add check_type --- tests/helper.py | 5 +++-- tests/test_mathematical.py | 2 +- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/tests/helper.py b/tests/helper.py index 42ada2949833..bb8ccd378a2b 100644 --- a/tests/helper.py +++ b/tests/helper.py @@ -7,7 +7,7 @@ import dpnp -def assert_dtype_allclose(dpnp_arr, numpy_arr): +def assert_dtype_allclose(dpnp_arr, numpy_arr, check_type=True): """ Assert DPNP and NumPy array based on maximum dtype resolution of input arrays for floating and complex types. @@ -24,7 +24,8 @@ def assert_dtype_allclose(dpnp_arr, numpy_arr): assert_allclose(dpnp_arr.asnumpy(), numpy_arr, atol=tol, rtol=tol) else: assert_array_equal(dpnp_arr.asnumpy(), numpy_arr) - assert dpnp_arr.dtype == numpy_arr.dtype + if check_type: + assert dpnp_arr.dtype == numpy_arr.dtype def get_complex_dtypes(device=None): diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 4cc58f21b1cc..705e85a6a62c 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -138,7 +138,7 @@ def _test_mathematical(self, name, dtype, lhs, rhs): else: result = getattr(dpnp, name)(a_dpnp, b_dpnp) expected = getattr(numpy, name)(a_np, b_np) - assert_dtype_allclose(result, expected) + assert_dtype_allclose(result, expected, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): From 7213d2541ed8d58e44c3419a4b530747827bd33c Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 31 Aug 2023 09:13:09 -0500 Subject: [PATCH 07/10] address new comments --- .github/workflows/generate_coverage.yaml | 3 +-- tests/skipped_tests.tbl | 19 +++++++++++++++++++ tests/test_mathematical.py | 10 +++++----- tests/test_strides.py | 2 -- 4 files changed, 25 insertions(+), 9 deletions(-) diff --git a/.github/workflows/generate_coverage.yaml b/.github/workflows/generate_coverage.yaml index 413835d336cf..776e521f66e0 100644 --- a/.github/workflows/generate_coverage.yaml +++ b/.github/workflows/generate_coverage.yaml @@ -49,8 +49,7 @@ jobs: conda list - name: Build dpnp with coverage run: | - python scripts/gen_coverage.py --pytest-opts="--ignore tests/test_random.py \ - --ignore tests/test_strides.py" + python scripts/gen_coverage.py --pytest-opts="--ignore tests/test_random.py" - name: Install coverall dependencies run: | sudo gem install coveralls-lcov diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index 9c5b855f3631..a8a41d619d0c 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -84,6 +84,25 @@ tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: (dpnp tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray([(i, i) for i in x], [("a", object), ("b", dpnp.int32)])]] tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray(x).astype(dpnp.int8)] +tests/test_strides.py::test_strides_1arg[(10,)-None-arccosh] +tests/test_strides.py::test_strides_1arg[(10,)-None-arcsinh] +tests/test_strides.py::test_strides_1arg[(10,)-None-arctanh] +tests/test_strides.py::test_strides_1arg[(10,)-None-cbrt] +tests/test_strides.py::test_strides_1arg[(10,)-None-cosh] +tests/test_strides.py::test_strides_1arg[(10,)-None-degrees] +tests/test_strides.py::test_strides_1arg[(10,)-None-exp] +tests/test_strides.py::test_strides_1arg[(10,)-None-exp2] +tests/test_strides.py::test_strides_1arg[(10,)-None-expm1] +tests/test_strides.py::test_strides_1arg[(10,)-None-fabs] +tests/test_strides.py::test_strides_1arg[(10,)-None-log1p] +tests/test_strides.py::test_strides_1arg[(10,)-None-radians] +tests/test_strides.py::test_strides_1arg[(10,)-None-log2] +tests/test_strides.py::test_strides_1arg[(10,)-None-log10] +tests/test_strides.py::test_strides_1arg[(10,)-None-sinh] +tests/test_strides.py::test_strides_1arg[(10,)-None-tanh] +tests/test_strides.py::test_strides_erf[(10,)-None] +tests/test_strides.py::test_strides_reciprocal[(10,)-None] + tests/test_umath.py::test_umaths[('divmod', 'ii')] tests/test_umath.py::test_umaths[('divmod', 'll')] tests/test_umath.py::test_umaths[('divmod', 'ff')] diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 705e85a6a62c..50db0d5ccacb 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -120,7 +120,7 @@ def array_or_scalar(xp, data, dtype=None): return xp.array(data, dtype=dtype) - def _test_mathematical(self, name, dtype, lhs, rhs): + def _test_mathematical(self, name, dtype, lhs, rhs, check_type=True): a_dpnp = self.array_or_scalar(dpnp, lhs, dtype=dtype) b_dpnp = self.array_or_scalar(dpnp, rhs, dtype=dtype) @@ -138,7 +138,7 @@ def _test_mathematical(self, name, dtype, lhs, rhs): else: result = getattr(dpnp, name)(a_dpnp, b_dpnp) expected = getattr(numpy, name)(a_np, b_np) - assert_dtype_allclose(result, expected, check_type=False) + assert_dtype_allclose(result, expected, check_type) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): @@ -175,7 +175,7 @@ def test_fmod(self, dtype, lhs, rhs): On a gpu without support for `float64`, dpnp produces results similar to the second one. """ pytest.skip("Due to accuracy reason, the results are different.") - self._test_mathematical("fmod", dtype, lhs, rhs) + self._test_mathematical("fmod", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_floor_divide(self, dtype, lhs, rhs): @@ -197,14 +197,14 @@ def test_hypot(self, dtype, lhs, rhs): "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_maximum(self, dtype, lhs, rhs): - self._test_mathematical("maximum", dtype, lhs, rhs) + self._test_mathematical("maximum", dtype, lhs, rhs, check_type=False) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_minimum(self, dtype, lhs, rhs): - self._test_mathematical("minimum", dtype, lhs, rhs) + self._test_mathematical("minimum", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_multiply(self, dtype, lhs, rhs): diff --git a/tests/test_strides.py b/tests/test_strides.py index a6cfbfbde820..1ddfcd14eb5a 100644 --- a/tests/test_strides.py +++ b/tests/test_strides.py @@ -315,7 +315,6 @@ def test_strided_in_out_2args_diff_out_dtype(func_name, dtype): @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) ) -@pytest.mark.skip("dpctl doesn't support overlap of arrays") def test_strided_in_2args_overlap(func_name, dtype): size = 5 @@ -337,7 +336,6 @@ def test_strided_in_2args_overlap(func_name, dtype): @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True, no_none=True) ) -@pytest.mark.skip("dpctl doesn't support overlap of arrays") def test_strided_in_out_2args_overlap(func_name, dtype): sh = (4, 3, 2) prod = numpy.prod(sh) From 8ad8114a4293dc79411211ebb0b21b740fed6dca Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Fri, 1 Sep 2023 18:02:07 -0500 Subject: [PATCH 08/10] bypass failed test on windows --- tests/test_mathematical.py | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 50db0d5ccacb..adb556fb79c7 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -142,7 +142,7 @@ def _test_mathematical(self, name, dtype, lhs, rhs, check_type=True): @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): - self._test_mathematical("add", dtype, lhs, rhs) + self._test_mathematical("add", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_arctan2(self, dtype, lhs, rhs): @@ -183,7 +183,9 @@ def test_floor_divide(self, dtype, lhs, rhs): pytest.skip( "In this case, a different result, but similar to xp.floor(xp.divide(lhs, rhs)." ) - self._test_mathematical("floor_divide", dtype, lhs, rhs) + self._test_mathematical( + "floor_divide", dtype, lhs, rhs, check_type=False + ) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize( @@ -208,7 +210,7 @@ def test_minimum(self, dtype, lhs, rhs): @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_multiply(self, dtype, lhs, rhs): - self._test_mathematical("multiply", dtype, lhs, rhs) + self._test_mathematical("multiply", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_remainder(self, dtype, lhs, rhs): @@ -227,15 +229,15 @@ def test_remainder(self, dtype, lhs, rhs): On a gpu without support for `float64`, dpnp produces results similar to the second one. """ pytest.skip("Due to accuracy reason, the results are different.") - self._test_mathematical("remainder", dtype, lhs, rhs) + self._test_mathematical("remainder", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_power(self, dtype, lhs, rhs): - self._test_mathematical("power", dtype, lhs, rhs) + self._test_mathematical("power", dtype, lhs, rhs, check_type=False) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_subtract(self, dtype, lhs, rhs): - self._test_mathematical("subtract", dtype, lhs, rhs) + self._test_mathematical("subtract", dtype, lhs, rhs, check_type=False) @pytest.mark.usefixtures("suppress_divide_invalid_numpy_warnings") From 605fac98896cf1823c2b48577f79786884b482c4 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Mon, 28 Aug 2023 16:25:13 -0500 Subject: [PATCH 09/10] use dpctl hyperbolic functions for dpnp --- dpnp/backend/extensions/vm/acosh.hpp | 78 +++++ dpnp/backend/extensions/vm/asinh.hpp | 78 +++++ dpnp/backend/extensions/vm/atanh.hpp | 78 +++++ dpnp/backend/extensions/vm/cosh.hpp | 78 +++++ dpnp/backend/extensions/vm/sinh.hpp | 78 +++++ dpnp/backend/extensions/vm/tanh.hpp | 78 +++++ dpnp/backend/extensions/vm/types_matrix.hpp | 102 ++++++ dpnp/backend/extensions/vm/vm_py.cpp | 180 ++++++++++ dpnp/backend/include/dpnp_iface_fptr.hpp | 16 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 54 --- dpnp/dpnp_algo/dpnp_algo.pxd | 18 - dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi | 30 -- dpnp/dpnp_algo/dpnp_elementwise_common.py | 327 ++++++++++++++++++ dpnp/dpnp_iface_trigonometric.py | 313 +++++++++++++---- tests/skipped_tests.tbl | 6 - tests/skipped_tests_gpu_no_fp64.tbl | 29 +- tests/test_sycl_queue.py | 6 + tests/test_usm_type.py | 6 + .../cupy/math_tests/test_hyperbolic.py | 7 +- 19 files changed, 1337 insertions(+), 225 deletions(-) create mode 100644 dpnp/backend/extensions/vm/acosh.hpp create mode 100644 dpnp/backend/extensions/vm/asinh.hpp create mode 100644 dpnp/backend/extensions/vm/atanh.hpp create mode 100644 dpnp/backend/extensions/vm/cosh.hpp create mode 100644 dpnp/backend/extensions/vm/sinh.hpp create mode 100644 dpnp/backend/extensions/vm/tanh.hpp diff --git a/dpnp/backend/extensions/vm/acosh.hpp b/dpnp/backend/extensions/vm/acosh.hpp new file mode 100644 index 000000000000..63d17cbd464b --- /dev/null +++ b/dpnp/backend/extensions/vm/acosh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event acosh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::acosh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AcoshContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AcoshOutputType::value_type, void>) + { + return nullptr; + } + else { + return acosh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/asinh.hpp b/dpnp/backend/extensions/vm/asinh.hpp new file mode 100644 index 000000000000..4e8e242257ac --- /dev/null +++ b/dpnp/backend/extensions/vm/asinh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event asinh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::asinh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AsinhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AsinhOutputType::value_type, void>) + { + return nullptr; + } + else { + return asinh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/atanh.hpp b/dpnp/backend/extensions/vm/atanh.hpp new file mode 100644 index 000000000000..6f7ee1064e8c --- /dev/null +++ b/dpnp/backend/extensions/vm/atanh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event atanh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::atanh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct AtanhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AtanhOutputType::value_type, void>) + { + return nullptr; + } + else { + return atanh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/cosh.hpp b/dpnp/backend/extensions/vm/cosh.hpp new file mode 100644 index 000000000000..4a02bf361a50 --- /dev/null +++ b/dpnp/backend/extensions/vm/cosh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event cosh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::cosh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct CoshContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::CoshOutputType::value_type, void>) + { + return nullptr; + } + else { + return cosh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/sinh.hpp b/dpnp/backend/extensions/vm/sinh.hpp new file mode 100644 index 000000000000..8b8114ad54ed --- /dev/null +++ b/dpnp/backend/extensions/vm/sinh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event sinh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::sinh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct SinhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::SinhOutputType::value_type, void>) + { + return nullptr; + } + else { + return sinh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/tanh.hpp b/dpnp/backend/extensions/vm/tanh.hpp new file mode 100644 index 000000000000..ed1be8f8d527 --- /dev/null +++ b/dpnp/backend/extensions/vm/tanh.hpp @@ -0,0 +1,78 @@ +//***************************************************************************** +// Copyright (c) 2023, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#pragma once + +#include + +#include "common.hpp" +#include "types_matrix.hpp" + +namespace dpnp +{ +namespace backend +{ +namespace ext +{ +namespace vm +{ +template +sycl::event tanh_contig_impl(sycl::queue exec_q, + const std::int64_t n, + const char *in_a, + char *out_y, + const std::vector &depends) +{ + type_utils::validate_type_for_device(exec_q); + + const T *a = reinterpret_cast(in_a); + T *y = reinterpret_cast(out_y); + + return mkl_vm::tanh(exec_q, + n, // number of elements to be calculated + a, // pointer `a` containing input vector of size n + y, // pointer `y` to the output vector of size n + depends); +} + +template +struct TanhContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::TanhOutputType::value_type, void>) + { + return nullptr; + } + else { + return tanh_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/types_matrix.hpp b/dpnp/backend/extensions/vm/types_matrix.hpp index 155dfc54d776..bbf74ed8c863 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -60,6 +60,23 @@ struct AcosOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::acosh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AcoshOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::add function. @@ -102,6 +119,23 @@ struct AsinOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::asinh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AsinhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::atan function. @@ -134,6 +168,23 @@ struct Atan2OutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::atanh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AtanhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::ceil function. @@ -183,6 +234,23 @@ struct CosOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::cosh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct CoshOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::div function. @@ -324,6 +392,23 @@ struct SinOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::sinh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct SinhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::sqr function. @@ -400,6 +485,23 @@ struct TanOutputType dpctl_td_ns::DefaultResultEntry>::result_type; }; +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::tanh function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct TanhOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry>, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::trunc function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index e769a40c2ba9..e4d89ff6e84f 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -31,14 +31,18 @@ #include #include "acos.hpp" +#include "acosh.hpp" #include "add.hpp" #include "asin.hpp" +#include "asinh.hpp" #include "atan.hpp" #include "atan2.hpp" +#include "atanh.hpp" #include "ceil.hpp" #include "common.hpp" #include "conj.hpp" #include "cos.hpp" +#include "cosh.hpp" #include "div.hpp" #include "floor.hpp" #include "ln.hpp" @@ -46,10 +50,12 @@ #include "pow.hpp" #include "round.hpp" #include "sin.hpp" +#include "sinh.hpp" #include "sqr.hpp" #include "sqrt.hpp" #include "sub.hpp" #include "tan.hpp" +#include "tanh.hpp" #include "trunc.hpp" #include "types_matrix.hpp" @@ -60,12 +66,16 @@ using vm_ext::binary_impl_fn_ptr_t; using vm_ext::unary_impl_fn_ptr_t; static unary_impl_fn_ptr_t acos_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t acosh_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t add_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t asin_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t asinh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t atan_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t atan2_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t atanh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ceil_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t cos_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t cosh_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t floor_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t conj_dispatch_vector[dpctl_td_ns::num_types]; @@ -74,10 +84,12 @@ static binary_impl_fn_ptr_t mul_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t pow_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t round_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t sinh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqr_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t sub_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t tan_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t tanh_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t trunc_dispatch_vector[dpctl_td_ns::num_types]; PYBIND11_MODULE(_vm_impl, m) @@ -113,6 +125,34 @@ PYBIND11_MODULE(_vm_impl, m) py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); } + // UnaryUfunc: ==== Acosh(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + acosh_dispatch_vector); + + auto acosh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + acosh_dispatch_vector); + }; + m.def("_acosh", acosh_pyapi, + "Call `acosh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto acosh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + acosh_dispatch_vector); + }; + m.def("_mkl_acosh_to_call", acosh_need_to_call_pyapi, + "Check input arguments to answer if `acosh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // BinaryUfunc: ==== Add(x1, x2) ==== { vm_ext::init_ufunc_dispatch_vector( + asinh_dispatch_vector); + + auto asinh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + asinh_dispatch_vector); + }; + m.def("_asinh", asinh_pyapi, + "Call `asinh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto asinh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + asinh_dispatch_vector); + }; + m.def("_mkl_asinh_to_call", asinh_need_to_call_pyapi, + "Check input arguments to answer if `asinh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Atan(x) ==== { vm_ext::init_ufunc_dispatch_vector( + atanh_dispatch_vector); + + auto atanh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + atanh_dispatch_vector); + }; + m.def("_atanh", atanh_pyapi, + "Call `atanh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto atanh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + atanh_dispatch_vector); + }; + m.def("_mkl_atanh_to_call", atanh_need_to_call_pyapi, + "Check input arguments to answer if `atanh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Ceil(x) ==== { vm_ext::init_ufunc_dispatch_vector( + cosh_dispatch_vector); + + auto cosh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + cosh_dispatch_vector); + }; + m.def("_cosh", cosh_pyapi, + "Call `cosh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto cosh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + cosh_dispatch_vector); + }; + m.def("_mkl_cosh_to_call", cosh_need_to_call_pyapi, + "Check input arguments to answer if `cosh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // BinaryUfunc: ==== Div(x1, x2) ==== { vm_ext::init_ufunc_dispatch_vector( + sinh_dispatch_vector); + + auto sinh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + sinh_dispatch_vector); + }; + m.def("_sinh", sinh_pyapi, + "Call `sinh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto sinh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + sinh_dispatch_vector); + }; + m.def("_mkl_sinh_to_call", sinh_need_to_call_pyapi, + "Check input arguments to answer if `sinh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Sqr(x) ==== { vm_ext::init_ufunc_dispatch_vector( + tanh_dispatch_vector); + + auto tanh_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + tanh_dispatch_vector); + }; + m.def("_tanh", tanh_pyapi, + "Call `tanh` function from OneMKL VM library to compute " + "inverse cosine of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto tanh_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + tanh_dispatch_vector); + }; + m.def("_mkl_tanh_to_call", tanh_need_to_call_pyapi, + "Check input arguments to answer if `tanh` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Trunc(x) ==== { vm_ext::init_ufunc_dispatch_vector}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCCOSH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_acosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_asin_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCSIN][eft_LNG][eft_LNG] = { @@ -269,15 +260,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCSINH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_asinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCSINH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_asinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_atan_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCTAN][eft_LNG][eft_LNG] = { @@ -296,15 +278,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCTANH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_atanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTANH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_atanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_CBRT][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_cbrt_c_default}; fmap[DPNPFuncName::DPNP_FN_CBRT][eft_LNG][eft_LNG] = { @@ -426,15 +399,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_COSH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_cosh_c_default}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cosh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_degrees_c_default}; fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_LNG][eft_LNG] = { @@ -633,15 +597,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_SINH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_sinh_c_default}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sinh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SQRT][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_sqrt_c_default}; fmap[DPNPFuncName::DPNP_FN_SQRT][eft_LNG][eft_LNG] = { @@ -675,15 +630,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_TANH][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_tanh_c_default}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TANH_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_tanh_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_trunc_c_default}; fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_LNG][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 966df2993dde..c3ba9338cee9 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -38,12 +38,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_ALLCLOSE DPNP_FN_ALLCLOSE_EXT DPNP_FN_ARANGE - DPNP_FN_ARCCOSH - DPNP_FN_ARCCOSH_EXT - DPNP_FN_ARCSINH - DPNP_FN_ARCSINH_EXT - DPNP_FN_ARCTANH - DPNP_FN_ARCTANH_EXT DPNP_FN_ARGMAX DPNP_FN_ARGMAX_EXT DPNP_FN_ARGMIN @@ -64,8 +58,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_COPYSIGN_EXT DPNP_FN_CORRELATE DPNP_FN_CORRELATE_EXT - DPNP_FN_COSH - DPNP_FN_COSH_EXT DPNP_FN_COUNT_NONZERO DPNP_FN_COUNT_NONZERO_EXT DPNP_FN_CROSS @@ -241,8 +233,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_RNG_ZIPF_EXT DPNP_FN_SEARCHSORTED DPNP_FN_SEARCHSORTED_EXT - DPNP_FN_SINH - DPNP_FN_SINH_EXT DPNP_FN_SORT DPNP_FN_SORT_EXT DPNP_FN_STD @@ -251,8 +241,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_SUM_EXT DPNP_FN_SVD DPNP_FN_SVD_EXT - DPNP_FN_TANH - DPNP_FN_TANH_EXT DPNP_FN_TRACE DPNP_FN_TRACE_EXT DPNP_FN_TRANSPOSE @@ -432,11 +420,7 @@ cpdef dpnp_descriptor dpnp_argmin(dpnp_descriptor array1) """ Trigonometric functions """ -cpdef dpnp_descriptor dpnp_arccosh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arcsinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_arctanh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_cbrt(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_cosh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_degrees(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_exp(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_exp2(dpnp_descriptor array1) @@ -446,5 +430,3 @@ cpdef dpnp_descriptor dpnp_log1p(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_log2(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_radians(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_recip(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_sinh(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_tanh(dpnp_descriptor array1) diff --git a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi index 12995c346fde..55e7e9c8dc53 100644 --- a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi @@ -36,11 +36,7 @@ and the rest of the library # NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file __all__ += [ - 'dpnp_arccosh', - 'dpnp_arcsinh', - 'dpnp_arctanh', 'dpnp_cbrt', - 'dpnp_cosh', 'dpnp_degrees', 'dpnp_exp', 'dpnp_exp2', @@ -50,32 +46,14 @@ __all__ += [ 'dpnp_log2', 'dpnp_radians', 'dpnp_recip', - 'dpnp_sinh', - 'dpnp_tanh', 'dpnp_unwrap' ] -cpdef utils.dpnp_descriptor dpnp_arccosh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCCOSH_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_arcsinh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCSINH_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_arctanh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_ARCTANH_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_cbrt(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_CBRT_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_cosh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_COSH_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_degrees(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_DEGREES_EXT, x1) @@ -112,14 +90,6 @@ cpdef utils.dpnp_descriptor dpnp_radians(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_RADIANS_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_sinh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_SINH_EXT, x1) - - -cpdef utils.dpnp_descriptor dpnp_tanh(utils.dpnp_descriptor x1): - return call_fptr_1in_1out_strides(DPNP_FN_TANH_EXT, x1) - - cpdef utils.dpnp_descriptor dpnp_unwrap(utils.dpnp_descriptor array1): result_type = dpnp.float64 diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 3da994049deb..ffc1f20065c4 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -43,16 +43,20 @@ __all__ = [ "check_nd_call_func", "dpnp_acos", + "dpnp_acosh", "dpnp_add", "dpnp_asin", + "dpnp_asinh", "dpnp_atan", "dpnp_atan2", + "dpnp_atanh", "dpnp_bitwise_and", "dpnp_bitwise_or", "dpnp_bitwise_xor", "dpnp_ceil", "dpnp_conj", "dpnp_cos", + "dpnp_cosh", "dpnp_divide", "dpnp_equal", "dpnp_floor", @@ -82,10 +86,12 @@ "dpnp_sign", "dpnp_signbit", "dpnp_sin", + "dpnp_sinh", "dpnp_sqrt", "dpnp_square", "dpnp_subtract", "dpnp_tan", + "dpnp_tanh", "dpnp_trunc", ] @@ -210,6 +216,60 @@ def dpnp_acos(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_acosh_docstring = """ +acosh(x, out=None, order='K') + +Computes hyperbolic inverse cosine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse hyperbolic cosine. + The data type of the returned array is determined by + the Type Promotion Rules. +""" + + +def _call_acosh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_acosh_to_call(sycl_queue, src, dst): + # call pybind11 extension for acosh() function from OneMKL VM + return vmi._acosh(sycl_queue, src, dst, depends) + return ti._acosh(src, dst, sycl_queue, depends) + + +acosh_func = UnaryElementwiseFunc( + "acosh", ti._acosh_result_type, _call_acosh, _acosh_docstring +) + + +def dpnp_acosh(x, out=None, order="K"): + """ + Invokes acosh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for acosh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = acosh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _add_docstring_ = """ add(x1, x2, out=None, order="K") @@ -323,6 +383,60 @@ def dpnp_asin(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_asinh_docstring = """ +asinh(x, out=None, order='K') + +Computes inverse hyperbolic sine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise inverse hyperbolic sine. + The data type of the returned array is determined by + the Type Promotion Rules. +""" + + +def _call_asinh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_asinh_to_call(sycl_queue, src, dst): + # call pybind11 extension for asinh() function from OneMKL VM + return vmi._asinh(sycl_queue, src, dst, depends) + return ti._asinh(src, dst, sycl_queue, depends) + + +asinh_func = UnaryElementwiseFunc( + "asinh", ti._asinh_result_type, _call_asinh, _asinh_docstring +) + + +def dpnp_asinh(x, out=None, order="K"): + """ + Invokes asinh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for asinh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = asinh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _atan_docstring = """ atan(x, out=None, order='K') @@ -443,6 +557,60 @@ def dpnp_atan2(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_atanh_docstring = """ +atanh(x, out=None, order='K') + +Computes hyperbolic inverse tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic inverse tangent. + The data type of the returned array is determined by + the Type Promotion Rules. +""" + + +def _call_atanh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_atanh_to_call(sycl_queue, src, dst): + # call pybind11 extension for atanh() function from OneMKL VM + return vmi._atanh(sycl_queue, src, dst, depends) + return ti._atanh(src, dst, sycl_queue, depends) + + +atanh_func = UnaryElementwiseFunc( + "atanh", ti._atanh_result_type, _call_atanh, _atanh_docstring +) + + +def dpnp_atanh(x, out=None, order="K"): + """ + Invokes atanh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for atanh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = atanh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _bitwise_and_docstring_ = """ bitwise_and(x1, x2, out=None, order='K') @@ -693,6 +861,59 @@ def dpnp_cos(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_cosh_docstring = """ +cosh(x, out=None, order='K') + +Computes hyperbolic cosine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic cosine. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_cosh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_cosh_to_call(sycl_queue, src, dst): + # call pybind11 extension for cosh() function from OneMKL VM + return vmi._cosh(sycl_queue, src, dst, depends) + return ti._cosh(src, dst, sycl_queue, depends) + + +cosh_func = UnaryElementwiseFunc( + "cosh", ti._cosh_result_type, _call_cosh, _cosh_docstring +) + + +def dpnp_cosh(x, out=None, order="K"): + """ + Invokes cosh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for cosh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = cosh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _conj_docstring = """ conj(x, out=None, order='K') @@ -2077,6 +2298,59 @@ def dpnp_sin(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_sinh_docstring = """ +sinh(x, out=None, order='K') + +Computes hyperbolic sine for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic sine. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_sinh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_sinh_to_call(sycl_queue, src, dst): + # call pybind11 extension for sinh() function from OneMKL VM + return vmi._sinh(sycl_queue, src, dst, depends) + return ti._sinh(src, dst, sycl_queue, depends) + + +sinh_func = UnaryElementwiseFunc( + "sinh", ti._sinh_result_type, _call_sinh, _sinh_docstring +) + + +def dpnp_sinh(x, out=None, order="K"): + """ + Invokes sinh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for sinh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = sinh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _sqrt_docstring_ = """ sqrt(x, out=None, order='K') Computes the non-negative square-root for each element `x_i` for input array `x`. @@ -2310,6 +2584,59 @@ def dpnp_tan(x, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_tanh_docstring = """ +tanh(x, out=None, order='K') + +Computes hyperbolic tangent for each element `x_i` for input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise hyperbolic tangent. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_tanh(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_tanh_to_call(sycl_queue, src, dst): + # call pybind11 extension for tanh() function from OneMKL VM + return vmi._tanh(sycl_queue, src, dst, depends) + return ti._tanh(src, dst, sycl_queue, depends) + + +tanh_func = UnaryElementwiseFunc( + "tanh", ti._tanh_result_type, _call_tanh, _tanh_docstring +) + + +def dpnp_tanh(x, out=None, order="K"): + """ + Invokes tanh() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for tanh() function. + + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = tanh_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _trunc_docstring = """ trunc(x, out=None, order='K') diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index c7b250249b38..3ed6de42e998 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -49,15 +49,21 @@ from .dpnp_algo.dpnp_elementwise_common import ( check_nd_call_func, dpnp_acos, + dpnp_acosh, dpnp_asin, + dpnp_asinh, dpnp_atan, dpnp_atan2, + dpnp_atanh, dpnp_cos, + dpnp_cosh, dpnp_log, dpnp_sin, + dpnp_sinh, dpnp_sqrt, dpnp_square, dpnp_tan, + dpnp_tanh, ) __all__ = [ @@ -152,42 +158,62 @@ def arccos( ) -def arccosh(x1): +def arccosh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Inverse hyperbolic cosine, element-wise. For full documentation refer to :obj:`numpy.arccosh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic inverse cosine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also -------- :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. - :obj:`dpnp.arcsinh` : Inverse hyperbolic sine element-wise. - :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. - :obj:`dpnp.arctanh` : Inverse hyperbolic tangent element-wise. - :obj:`dpnp.tanh` : Compute hyperbolic tangent element-wise. + :obj:`dpnp.arctanh` : Hyperbolic inverse tangent, element-wise. + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. + :obj:`dpnp.arccos` : Trigonometric inverse cosine, element-wise. Examples -------- >>> import dpnp as np - >>> x = np.array([np.e, 10.0]) - >>> out = np.arccosh(x) - >>> [i for i in out] - [1.65745445, 2.99322285] + >>> x = np.array([1.0, np.e, 10.0]) + >>> np.arccosh(x) + array([0.0, 1.65745445, 2.99322285]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arccosh, + dpnp_acosh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arccosh(x1_desc).get_pyobj() - - return call_origin(numpy.arccosh, x1, **kwargs) def arcsin( @@ -248,34 +274,62 @@ def arcsin( ) -def arcsinh(x1): +def arcsinh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Inverse hyperbolic sine, element-wise. For full documentation refer to :obj:`numpy.arcsinh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic inverse sine of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. + :obj:`dpnp.arctanh` : Hyperbolic inverse tangent, element-wise. + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. + :obj:`dpnp.arcsin` : Trigonometric inverse sine, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([np.e, 10.0]) - >>> out = np.arcsinh(x) - >>> [i for i in out] - [1.72538256, 2.99822295] + >>> np.arcsinh(x) + array([1.72538256, 2.99822295]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arcsinh, + dpnp_asinh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arcsinh(x1_desc).get_pyobj() - - return call_origin(numpy.arcsinh, x1, **kwargs) def arctan( @@ -412,34 +466,62 @@ def arctan2( ) -def arctanh(x1): +def arctanh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Hyperbolic inverse tangent, element-wise. For full documentation refer to :obj:`numpy.arctanh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic inverse tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. + :obj:`dpnp.arctan` : Trigonometric inverse tangent, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([0, -0.5]) - >>> out = np.arctanh(x) - >>> [i for i in out] - [0.0, -0.54930614] + >>> np.arctanh(x) + array([0.0, -0.54930614]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.arctanh, + dpnp_atanh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_arctanh(x1_desc).get_pyobj() - - return call_origin(numpy.arctanh, x1, **kwargs) def cbrt(x1): @@ -530,34 +612,63 @@ def cos( ) -def cosh(x1): +def cosh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Hyperbolic cosine, element-wise. For full documentation refer to :obj:`numpy.cosh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic cosine of each element of `x`. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arccosh` : Hyperbolic inverse cosine, element-wise. + :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + :obj:`dpnp.cos` : Trigonometric cosine, element-wise. + + Examples -------- >>> import dpnp as np - >>> x = np.array([0]) - >>> out = np.cosh(x) - >>> [i for i in out] - [1.0] + >>> x = np.array([0, np.pi/2, np.pi]) + >>> np.cosh(x) + array([1.0, 2.5091786, 11.591953]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.cosh, + dpnp_cosh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_cosh(x1_desc).get_pyobj() - - return call_origin(numpy.cosh, x1, **kwargs) def deg2rad(x1): @@ -1120,34 +1231,62 @@ def sin( ) -def sinh(x1): +def sinh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Hyperbolic sine, element-wise. For full documentation refer to :obj:`numpy.sinh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic sine of each element of `x`. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arcsinh` : Hyperbolic inverse sine, element-wise. + :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. + :obj:`dpnp.tanh` : Hyperbolic tangent, element-wise. + :obj:`dpnp.sin` : Trigonometric sine, element-wise. + Examples -------- >>> import dpnp as np >>> x = np.array([0, np.pi/2, np.pi]) - >>> out = np.sinh(x) - >>> [i for i in out] - [0.0, 2.3012989, 11.548739] + >>> np.sinh(x) + array([0.0, 2.3012989, 11.548739]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.sinh, + dpnp_sinh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_sinh(x1_desc).get_pyobj() - - return call_origin(numpy.sinh, x1, **kwargs) def sqrt( @@ -1329,34 +1468,62 @@ def tan( ) -def tanh(x1): +def tanh( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Compute hyperbolic tangent element-wise. For full documentation refer to :obj:`numpy.tanh`. + Returns + ------- + out : dpnp.ndarray + The hyperbolic tangent of each element of `x`. + Limitations ----------- - Input array is supported as :class:`dpnp.ndarray`. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Keyword argument `kwargs` is currently unsupported. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. + See Also + -------- + :obj:`dpnp.arctanh` : Hyperbolic inverse tangent, element-wise. + :obj:`dpnp.sinh` : Hyperbolic sine, element-wise. + :obj:`dpnp.cosh` : Hyperbolic cosine, element-wise. + :obj:`dpnp.tan` : Trigonometric tangent, element-wise. + Examples -------- >>> import dpnp as np - >>> x = np.array([-np.pi, np.pi/2, np.pi]) - >>> out = np.tanh(x) - >>> [i for i in out] - [-0.996272, 0.917152, 0.996272] + >>> x = np.array([0, -np.pi, np.pi/2, np.pi]) + >>> np.tanh(x) + array([0.0, -0.996272, 0.917152, 0.996272]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.tanh, + dpnp_tanh, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - return dpnp_tanh(x1_desc).get_pyobj() - - return call_origin(numpy.tanh, x1, **kwargs) def unwrap(x1): diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index 6d4e4139dfb1..128b7f7f5da2 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -83,11 +83,7 @@ tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: (dpnp tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray([(i, i) for i in x], [("a", object), ("b", dpnp.int32)])]] tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray(x).astype(dpnp.int8)] -tests/test_strides.py::test_strides_1arg[(10,)-None-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-None-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-None-cosh] tests/test_strides.py::test_strides_1arg[(10,)-None-degrees] tests/test_strides.py::test_strides_1arg[(10,)-None-exp] tests/test_strides.py::test_strides_1arg[(10,)-None-exp2] @@ -97,8 +93,6 @@ tests/test_strides.py::test_strides_1arg[(10,)-None-log1p] tests/test_strides.py::test_strides_1arg[(10,)-None-radians] tests/test_strides.py::test_strides_1arg[(10,)-None-log2] tests/test_strides.py::test_strides_1arg[(10,)-None-log10] -tests/test_strides.py::test_strides_1arg[(10,)-None-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-tanh] tests/test_strides.py::test_strides_erf[(10,)-None] tests/test_strides.py::test_strides_reciprocal[(10,)-None] diff --git a/tests/skipped_tests_gpu_no_fp64.tbl b/tests/skipped_tests_gpu_no_fp64.tbl index 38d088971f39..5319e959fade 100644 --- a/tests/skipped_tests_gpu_no_fp64.tbl +++ b/tests/skipped_tests_gpu_no_fp64.tbl @@ -15,11 +15,7 @@ tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array0] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array1] tests/test_mathematical.py::TestGradient::test_gradient_y1_dx[3.5-array2] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-int32-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-int32-cosh] tests/test_strides.py::test_strides_1arg[(10,)-int32-degrees] tests/test_strides.py::test_strides_1arg[(10,)-int32-exp] tests/test_strides.py::test_strides_1arg[(10,)-int32-exp2] @@ -29,13 +25,7 @@ tests/test_strides.py::test_strides_1arg[(10,)-int32-log10] tests/test_strides.py::test_strides_1arg[(10,)-int32-log1p] tests/test_strides.py::test_strides_1arg[(10,)-int32-log2] tests/test_strides.py::test_strides_1arg[(10,)-int32-radians] -tests/test_strides.py::test_strides_1arg[(10,)-int32-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-int32-tanh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-int64-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-int64-cosh] tests/test_strides.py::test_strides_1arg[(10,)-int64-degrees] tests/test_strides.py::test_strides_1arg[(10,)-int64-exp] tests/test_strides.py::test_strides_1arg[(10,)-int64-exp2] @@ -45,13 +35,7 @@ tests/test_strides.py::test_strides_1arg[(10,)-int64-log10] tests/test_strides.py::test_strides_1arg[(10,)-int64-log1p] tests/test_strides.py::test_strides_1arg[(10,)-int64-log2] tests/test_strides.py::test_strides_1arg[(10,)-int64-radians] -tests/test_strides.py::test_strides_1arg[(10,)-int64-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-int64-tanh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arccosh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arcsinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-arctanh] tests/test_strides.py::test_strides_1arg[(10,)-None-cbrt] -tests/test_strides.py::test_strides_1arg[(10,)-None-cosh] tests/test_strides.py::test_strides_1arg[(10,)-None-degrees] tests/test_strides.py::test_strides_1arg[(10,)-None-exp] tests/test_strides.py::test_strides_1arg[(10,)-None-exp2] @@ -61,13 +45,11 @@ tests/test_strides.py::test_strides_1arg[(10,)-None-log10] tests/test_strides.py::test_strides_1arg[(10,)-None-log1p] tests/test_strides.py::test_strides_1arg[(10,)-None-log2] tests/test_strides.py::test_strides_1arg[(10,)-None-radians] -tests/test_strides.py::test_strides_1arg[(10,)-None-sinh] -tests/test_strides.py::test_strides_1arg[(10,)-None-tanh] tests/test_sycl_queue.py::test_array_creation[opencl:gpu:0-arange-arg0-kwargs0] tests/test_sycl_queue.py::test_array_creation[level_zero:gpu:0-arange-arg0-kwargs0] -tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-gradient-data14] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data14] +tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-gradient-data18] +tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data18] tests/test_sycl_queue.py::test_eig[opencl:gpu:0] tests/test_sycl_queue.py::test_eig[level_zero:gpu:0] tests/test_sycl_queue.py::test_eigh[opencl:gpu:0] @@ -676,13 +658,6 @@ tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_log2 tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_copysign_combination -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_arccosh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_arcsinh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_arctanh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_cosh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_sinh -tests/third_party/cupy/math_tests/test_hyperbolic.py::TestHyperbolic::test_tanh - tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_0_{axis=0, func='nansum', keepdims=True, shape=(2, 3, 4), transpose_axes=True}::test_nansum_all tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_0_{axis=0, func='nansum', keepdims=True, shape=(2, 3, 4), transpose_axes=True}::test_nansum_axis_transposed tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_1_{axis=0, func='nansum', keepdims=True, shape=(2, 3, 4), transpose_axes=False}::test_nansum_all diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index f3b1bc80270a..4fec38e50c45 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -226,13 +226,17 @@ def test_meshgrid(device_x, device_y): [ pytest.param("abs", [-1.2, 1.2]), pytest.param("arccos", [-0.5, 0.0, 0.5]), + pytest.param("arccosh", [1.5, 3.5, 5.0]), pytest.param("arcsin", [-0.5, 0.0, 0.5]), + pytest.param("arcsinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("arctan", [-1.0, 0.0, 1.0]), + pytest.param("arctanh", [-0.5, 0.0, 0.5]), pytest.param("ceil", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("conjugate", [[1.0 + 1.0j, 0.0], [0.0, 1.0 + 1.0j]]), pytest.param( "cos", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] ), + pytest.param("cosh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("copy", [1.0, 2.0, 3.0]), pytest.param("cumprod", [[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]), pytest.param("cumsum", [[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]), @@ -252,11 +256,13 @@ def test_meshgrid(device_x, device_y): pytest.param( "sin", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] ), + pytest.param("sinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("sqrt", [1.0, 3.0, 9.0]), pytest.param("sum", [1.0, 2.0]), pytest.param( "tan", [-dpnp.pi / 2, -dpnp.pi / 4, 0.0, dpnp.pi / 4, dpnp.pi / 2] ), + pytest.param("tanh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("trapz", [[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]]), pytest.param("trunc", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), ], diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 2937eaf9e0f5..17ccb55b6fd2 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -283,13 +283,17 @@ def test_meshgrid(usm_type_x, usm_type_y): "func,data", [ pytest.param("arccos", [-0.5, 0.0, 0.5]), + pytest.param("arccosh", [1.5, 3.5, 5.0]), pytest.param("arcsin", [-0.5, 0.0, 0.5]), + pytest.param("arcsinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("arctan", [-1.0, 0.0, 1.0]), + pytest.param("arctanh", [-0.5, 0.0, 0.5]), pytest.param("ceil", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("conjugate", [[1.0 + 1.0j, 0.0], [0.0, 1.0 + 1.0j]]), pytest.param( "cos", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] ), + pytest.param("cosh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("floor", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), pytest.param("negative", [1.0, 0.0, -1.0]), pytest.param("proj", [complex(1.0, 2.0), complex(dp.inf, -1.0)]), @@ -298,10 +302,12 @@ def test_meshgrid(usm_type_x, usm_type_y): pytest.param( "sin", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] ), + pytest.param("sinh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("sqrt", [1.0, 3.0, 9.0]), pytest.param( "tan", [-dp.pi / 2, -dp.pi / 4, 0.0, dp.pi / 4, dp.pi / 2] ), + pytest.param("tanh", [-5.0, -3.5, 0.0, 3.5, 5.0]), pytest.param("trunc", [-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]), ], ) diff --git a/tests/third_party/cupy/math_tests/test_hyperbolic.py b/tests/third_party/cupy/math_tests/test_hyperbolic.py index 6fa732c0200b..6f35337f2619 100644 --- a/tests/third_party/cupy/math_tests/test_hyperbolic.py +++ b/tests/third_party/cupy/math_tests/test_hyperbolic.py @@ -1,17 +1,18 @@ import unittest +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @testing.gpu class TestHyperbolic(unittest.TestCase): @testing.for_all_dtypes() - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) - @testing.for_dtypes(["f", "d"]) + @testing.for_dtypes(["e", "f", "d"]) @testing.numpy_cupy_allclose(atol=1e-5) def check_unary_unit(self, name, xp, dtype): a = xp.array([0.2, 0.4, 0.6, 0.8], dtype=dtype) @@ -29,7 +30,7 @@ def test_tanh(self): def test_arcsinh(self): self.check_unary("arcsinh") - @testing.for_dtypes(["f", "d"]) + @testing.for_dtypes(["e", "f", "d"]) @testing.numpy_cupy_allclose(atol=1e-5) def test_arccosh(self, xp, dtype): a = xp.array([1, 2, 3], dtype=dtype) From f391367241c7d792b5a8bb5476ef37349e59bde4 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 7 Sep 2023 10:38:33 -0600 Subject: [PATCH 10/10] address comments --- tests/test_umath.py | 301 +++++++++++++++++- .../cupy/math_tests/test_hyperbolic.py | 1 - 2 files changed, 297 insertions(+), 5 deletions(-) diff --git a/tests/test_umath.py b/tests/test_umath.py index ce918d849d53..791e9eb30223 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -176,6 +176,89 @@ def test_invalid_shape(self, shape): dpnp.sin(dp_array, out=dp_out) +class TestSinh: + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_sinh(self, dtype): + np_array = numpy.arange(10, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.float64) + + # DPNP + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.sinh(dp_array, out=dp_out) + + # original + expected = numpy.sinh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_sinh_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) + + # DPNP + dp_out_dtype = dpnp.complex64 + if has_support_aspect64() and dtype != dpnp.complex64: + dp_out_dtype = dpnp.complex128 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.sinh(dp_array, out=dp_out) + + # original + expected = numpy.sinh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_sinh_bool(self): + np_array = numpy.arange(2, dtype=numpy.bool_) + np_out = numpy.empty(2, dtype=numpy.float16) + + # DPNP + dp_array = dpnp.array(np_array, dtype=np_array.dtype) + dp_out = dpnp.array(np_out, dtype=np_out.dtype) + result = dpnp.sinh(dp_array, out=dp_out) + + # original + expected = numpy.sinh(np_array, out=np_out) + assert_allclose(expected, result) + + @pytest.mark.parametrize( + "dtype", + [numpy.float32, numpy.int64, numpy.int32], + ids=["numpy.float32", "numpy.int64", "numpy.int32"], + ) + def test_invalid_dtype(self, dtype): + dp_array = dpnp.arange(10, dtype=dpnp.complex64) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.sinh(dp_array, out=dp_out) + + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape): + dp_array = dpnp.arange(10) + dp_out = dpnp.empty(shape, dtype=dp_array.dtype) + + with pytest.raises(ValueError): + dpnp.sinh(dp_array, out=dp_out) + + class TestCos: @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) @@ -259,6 +342,89 @@ def test_invalid_shape(self, shape): dpnp.cos(dp_array, out=dp_out) +class TestCosh: + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_cosh(self, dtype): + np_array = numpy.arange(10, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.float64) + + # DPNP + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.cosh(dp_array, out=dp_out) + + # original + expected = numpy.cosh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_cosh_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) + + # DPNP + dp_out_dtype = dpnp.complex64 + if has_support_aspect64() and dtype != dpnp.complex64: + dp_out_dtype = dpnp.complex128 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) + result = dpnp.cosh(dp_array, out=dp_out) + + # original + expected = numpy.cosh(np_array, out=np_out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_cosh_bool(self): + np_array = numpy.arange(2, dtype=numpy.bool_) + np_out = numpy.empty(2, dtype=numpy.float16) + + # DPNP + dp_array = dpnp.array(np_array, dtype=np_array.dtype) + dp_out = dpnp.array(np_out, dtype=np_out.dtype) + result = dpnp.cosh(dp_array, out=dp_out) + + # original + expected = numpy.cosh(np_array, out=np_out) + assert_allclose(expected, result) + + @pytest.mark.parametrize( + "dtype", + [numpy.float32, numpy.int64, numpy.int32], + ids=["numpy.float32", "numpy.int64", "numpy.int32"], + ) + def test_invalid_dtype(self, dtype): + dp_array = dpnp.arange(10, dtype=dpnp.complex64) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.cosh(dp_array, out=dp_out) + + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape): + dp_array = dpnp.arange(10) + dp_out = dpnp.empty(shape, dtype=dp_array.dtype) + + with pytest.raises(ValueError): + dpnp.cosh(dp_array, out=dp_out) + + class TestsLog: @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) @@ -385,7 +551,7 @@ class TestArccos: @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") def test_arccos(self, dtype): - array_data = numpy.arange(10) + array_data = numpy.arange(-9, 10, 2) / 10 out = numpy.empty(10, dtype=dtype) # DPNP @@ -397,7 +563,8 @@ def test_arccos(self, dtype): np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arccos(np_array, out=out) - assert_array_equal(expected, result) + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) @pytest.mark.parametrize( "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] @@ -422,11 +589,53 @@ def test_invalid_shape(self, shape, dtype): dpnp.arccos(dp_array, out=dp_out) +class TestArccosh: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arccosh(self, dtype): + array_data = numpy.arange(2, 12) + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arccosh(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arccosh(np_array, out=out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arccosh(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arccosh(dp_array, out=dp_out) + + class TestArcsin: @pytest.mark.parametrize("dtype", get_float_dtypes()) @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") def test_arcsin(self, dtype): - array_data = numpy.arange(10) + array_data = numpy.arange(-9, 10, 2) / 10 out = numpy.empty(10, dtype=dtype) # DPNP @@ -438,7 +647,8 @@ def test_arcsin(self, dtype): np_array = numpy.array(array_data, dtype=dtype) expected = numpy.arcsin(np_array, out=out) - assert_array_equal(expected, result) + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) @pytest.mark.parametrize( "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] @@ -463,6 +673,48 @@ def test_invalid_shape(self, shape, dtype): dpnp.arcsin(dp_array, out=dp_out) +class TestArcsinh: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.usefixtures("suppress_invalid_numpy_warnings") + def test_arcsinh(self, dtype): + array_data = numpy.arange(10) + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arcsinh(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arcsinh(np_array, out=out) + + tol = numpy.finfo(dtype=result.dtype).resolution + assert_allclose(expected, result.asnumpy(), rtol=tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arcsinh(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arcsinh(dp_array, out=dp_out) + + class TestArctan: @pytest.mark.parametrize("dtype", get_float_dtypes()) def test_arctan(self, dtype): @@ -504,6 +756,47 @@ def test_invalid_shape(self, shape, dtype): dpnp.arctan(dp_array, out=dp_out) +class TestArctanh: + @pytest.mark.parametrize("dtype", get_float_dtypes()) + def test_arctanh(self, dtype): + array_data = numpy.arange(-9, 10, 2) / 10 + out = numpy.empty(10, dtype=dtype) + + # DPNP + dp_array = dpnp.array(array_data, dtype=dtype) + dp_out = dpnp.array(out, dtype=dtype) + result = dpnp.arctanh(dp_array, out=dp_out) + + # original + np_array = numpy.array(array_data, dtype=dtype) + expected = numpy.arctanh(np_array, out=out) + + tol = numpy.finfo(dtype).resolution + assert_allclose(expected, result, tol) + + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_complex=True, no_none=True)[:-1] + ) + def test_invalid_dtype(self, dtype): + dpnp_dtype = get_all_dtypes(no_complex=True, no_none=True)[-1] + dp_array = dpnp.arange(10, dtype=dpnp_dtype) + dp_out = dpnp.empty(10, dtype=dtype) + + with pytest.raises(TypeError): + dpnp.arctanh(dp_array, out=dp_out) + + @pytest.mark.parametrize("dtype", get_float_dtypes()) + @pytest.mark.parametrize( + "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] + ) + def test_invalid_shape(self, shape, dtype): + dp_array = dpnp.arange(10, dtype=dtype) + dp_out = dpnp.empty(shape, dtype=dtype) + + with pytest.raises(ValueError): + dpnp.arctanh(dp_array, out=dp_out) + + class TestTan: @pytest.mark.parametrize("dtype", get_float_dtypes()) def test_tan(self, dtype): diff --git a/tests/third_party/cupy/math_tests/test_hyperbolic.py b/tests/third_party/cupy/math_tests/test_hyperbolic.py index 6f35337f2619..7d7fa6a6b13d 100644 --- a/tests/third_party/cupy/math_tests/test_hyperbolic.py +++ b/tests/third_party/cupy/math_tests/test_hyperbolic.py @@ -4,7 +4,6 @@ from tests.third_party.cupy import testing -@testing.gpu class TestHyperbolic(unittest.TestCase): @testing.for_all_dtypes() @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64())