From 342b3a4240772a96a0c9a0979c5affe263a723ff Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 10 Aug 2023 11:19:52 -0500 Subject: [PATCH] use_dpctl_rounding_funcs_in_dpnp --- dpnp/backend/extensions/vm/ceil.hpp | 78 ++++ dpnp/backend/extensions/vm/floor.hpp | 78 ++++ dpnp/backend/extensions/vm/trunc.hpp | 78 ++++ dpnp/backend/extensions/vm/types_matrix.hpp | 46 ++ dpnp/backend/extensions/vm/vm_py.cpp | 90 ++++ dpnp/backend/include/dpnp_iface_fptr.hpp | 8 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 27 -- dpnp/dpnp_algo/dpnp_algo.pxd | 6 - dpnp/dpnp_algo/dpnp_algo_mathematical.pxi | 15 - dpnp/dpnp_algo/dpnp_elementwise_common.py | 161 +++++++ dpnp/dpnp_iface_mathematical.py | 411 ++++++++++-------- dpnp/dpnp_iface_trigonometric.py | 6 +- tests/skipped_tests.tbl | 3 - tests/skipped_tests_gpu.tbl | 8 - .../cupy/math_tests/test_rounding.py | 2 +- 15 files changed, 770 insertions(+), 247 deletions(-) create mode 100644 dpnp/backend/extensions/vm/ceil.hpp create mode 100644 dpnp/backend/extensions/vm/floor.hpp create mode 100644 dpnp/backend/extensions/vm/trunc.hpp diff --git a/dpnp/backend/extensions/vm/ceil.hpp b/dpnp/backend/extensions/vm/ceil.hpp new file mode 100644 index 00000000000..b1354b884c3 --- /dev/null +++ b/dpnp/backend/extensions/vm/ceil.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 ceil_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::ceil(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 CeilContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::CeilOutputType::value_type, void>) + { + return nullptr; + } + else { + return ceil_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/floor.hpp b/dpnp/backend/extensions/vm/floor.hpp new file mode 100644 index 00000000000..66707ba2979 --- /dev/null +++ b/dpnp/backend/extensions/vm/floor.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 floor_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::floor(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 FloorContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::FloorOutputType::value_type, void>) + { + return nullptr; + } + else { + return floor_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/trunc.hpp b/dpnp/backend/extensions/vm/trunc.hpp new file mode 100644 index 00000000000..d8ecadff363 --- /dev/null +++ b/dpnp/backend/extensions/vm/trunc.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 trunc_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::trunc(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 TruncContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::TruncOutputType::value_type, void>) + { + return nullptr; + } + else { + return trunc_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 667e27f0c9a..cd4fd76d4be 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -68,6 +68,21 @@ struct DivOutputType 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. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct CeilOutputType +{ + using value_type = typename std::disjunction< + 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::cos function. @@ -87,6 +102,21 @@ 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::floor function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct FloorOutputType +{ + using value_type = typename std::disjunction< + 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::ln function. @@ -158,6 +188,22 @@ struct SqrtOutputType 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. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct TruncOutputType +{ + using value_type = typename std::disjunction< + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::TypeMapResultEntry, + dpctl_td_ns::DefaultResultEntry>::result_type; +}; + } // namespace types } // namespace vm } // namespace ext diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 924db66025b..c7435ae9e2e 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -30,13 +30,16 @@ #include #include +#include "ceil.hpp" #include "common.hpp" #include "cos.hpp" #include "div.hpp" +#include "floor.hpp" #include "ln.hpp" #include "sin.hpp" #include "sqr.hpp" #include "sqrt.hpp" +#include "trunc.hpp" #include "types_matrix.hpp" namespace py = pybind11; @@ -47,11 +50,14 @@ using vm_ext::unary_impl_fn_ptr_t; static binary_impl_fn_ptr_t div_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 floor_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ln_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 sqr_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqrt_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) { @@ -88,6 +94,34 @@ PYBIND11_MODULE(_vm_impl, m) py::arg("dst")); } + // UnaryUfunc: ==== Ceil(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + ceil_dispatch_vector); + + auto ceil_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + ceil_dispatch_vector); + }; + m.def("_ceil", ceil_pyapi, + "Call `ceil` function from OneMKL VM library to compute " + "ceiling of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto ceil_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + ceil_dispatch_vector); + }; + m.def("_mkl_ceil_to_call", ceil_need_to_call_pyapi, + "Check input arguments to answer if `ceil` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Cos(x) ==== { vm_ext::init_ufunc_dispatch_vector( + floor_dispatch_vector); + + auto floor_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + floor_dispatch_vector); + }; + m.def("_floor", floor_pyapi, + "Call `floor` function from OneMKL VM library to compute " + "floor of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto floor_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + floor_dispatch_vector); + }; + m.def("_mkl_floor_to_call", floor_need_to_call_pyapi, + "Check input arguments to answer if `floor` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Ln(x) ==== { vm_ext::init_ufunc_dispatch_vector( + trunc_dispatch_vector); + + auto trunc_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + trunc_dispatch_vector); + }; + m.def("_trunc", trunc_pyapi, + "Call `trunc` function from OneMKL VM library to compute " + "the truncated value of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto trunc_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + trunc_dispatch_vector); + }; + m.def("_mkl_trunc_to_call", trunc_need_to_call_pyapi, + "Check input arguments to answer if `trunc` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } } diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 96b85885473..6831e61aef0 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -121,8 +121,6 @@ enum class DPNPFuncName : size_t DPNP_FN_CBRT_EXT, /**< Used in numpy.cbrt() impl, requires extra parameters */ DPNP_FN_CEIL, /**< Used in numpy.ceil() impl */ - DPNP_FN_CEIL_EXT, /**< Used in numpy.ceil() impl, requires extra parameters - */ DPNP_FN_CHOLESKY, /**< Used in numpy.linalg.cholesky() impl */ DPNP_FN_CHOLESKY_EXT, /**< Used in numpy.linalg.cholesky() impl, requires extra parameters */ @@ -220,8 +218,6 @@ enum class DPNPFuncName : size_t DPNP_FN_FLATTEN_EXT, /**< Used in numpy.flatten() impl, requires extra parameters */ DPNP_FN_FLOOR, /**< Used in numpy.floor() impl */ - DPNP_FN_FLOOR_EXT, /**< Used in numpy.floor() impl, requires extra - parameters */ DPNP_FN_FLOOR_DIVIDE, /**< Used in numpy.floor_divide() impl */ DPNP_FN_FLOOR_DIVIDE_EXT, /**< Used in numpy.floor_divide() impl, requires extra parameters */ @@ -504,9 +500,7 @@ enum class DPNPFuncName : size_t DPNP_FN_TRIL, /**< Used in numpy.tril() impl */ DPNP_FN_TRIU, /**< Used in numpy.triu() impl */ DPNP_FN_TRUNC, /**< Used in numpy.trunc() impl */ - DPNP_FN_TRUNC_EXT, /**< Used in numpy.trunc() impl, requires extra - parameters */ - DPNP_FN_VANDER, /**< Used in numpy.vander() impl */ + DPNP_FN_VANDER, /**< Used in numpy.vander() impl */ DPNP_FN_VANDER_EXT, /**< Used in numpy.vander() impl, requires extra parameters */ DPNP_FN_VAR, /**< Used in numpy.var() impl */ diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 1bd0d1922e0..a382d85b573 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -359,15 +359,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_CEIL][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_ceil_c_default}; - fmap[DPNPFuncName::DPNP_FN_CEIL_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_ceil_c_ext}; - fmap[DPNPFuncName::DPNP_FN_CEIL_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_ceil_c_ext}; - fmap[DPNPFuncName::DPNP_FN_CEIL_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ceil_c_ext}; - fmap[DPNPFuncName::DPNP_FN_CEIL_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ceil_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYTO][eft_BLN][eft_BLN] = { eft_BLN, (void *)dpnp_copyto_c_default}; fmap[DPNPFuncName::DPNP_FN_COPYTO][eft_BLN][eft_INT] = { @@ -603,15 +594,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_FLOOR][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_floor_c_default}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_floor_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_floor_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_floor_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_c_ext}; - fmap[DPNPFuncName::DPNP_FN_LOG10][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_log10_c_default}; fmap[DPNPFuncName::DPNP_FN_LOG10][eft_LNG][eft_LNG] = { @@ -780,15 +762,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_trunc_c_default}; - fmap[DPNPFuncName::DPNP_FN_TRUNC_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_trunc_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TRUNC_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_trunc_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TRUNC_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_trunc_c_ext}; - fmap[DPNPFuncName::DPNP_FN_TRUNC_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_trunc_c_ext}; - return; } diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 37d4a7d3694..438ea8f1c65 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -74,8 +74,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_BITWISE_XOR_EXT DPNP_FN_CBRT DPNP_FN_CBRT_EXT - DPNP_FN_CEIL - DPNP_FN_CEIL_EXT DPNP_FN_CHOLESKY DPNP_FN_CHOLESKY_EXT DPNP_FN_CHOOSE @@ -138,8 +136,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_FILL_DIAGONAL_EXT DPNP_FN_FLATTEN DPNP_FN_FLATTEN_EXT - DPNP_FN_FLOOR - DPNP_FN_FLOOR_EXT DPNP_FN_FMOD DPNP_FN_FMOD_EXT DPNP_FN_FULL @@ -310,8 +306,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_TRIL_EXT DPNP_FN_TRIU DPNP_FN_TRIU_EXT - DPNP_FN_TRUNC - DPNP_FN_TRUNC_EXT DPNP_FN_VANDER DPNP_FN_VANDER_EXT DPNP_FN_VAR diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index ae03fa54051..040ae129af9 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -39,7 +39,6 @@ __all__ += [ "dpnp_absolute", "dpnp_arctan2", "dpnp_around", - "dpnp_ceil", "dpnp_conjugate", "dpnp_copysign", "dpnp_cross", @@ -48,7 +47,6 @@ __all__ += [ "dpnp_diff", "dpnp_ediff1d", "dpnp_fabs", - "dpnp_floor", "dpnp_fmod", "dpnp_gradient", 'dpnp_hypot', @@ -66,7 +64,6 @@ __all__ += [ "dpnp_sign", "dpnp_sum", "dpnp_trapz", - "dpnp_trunc" ] @@ -159,10 +156,6 @@ cpdef utils.dpnp_descriptor dpnp_around(utils.dpnp_descriptor x1, int decimals): return result -cpdef utils.dpnp_descriptor dpnp_ceil(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_CEIL_EXT, x1, dtype=None, out=out, where=True, func_name='ceil') - - cpdef utils.dpnp_descriptor dpnp_conjugate(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_CONJIGUATE_EXT, x1) @@ -296,10 +289,6 @@ cpdef utils.dpnp_descriptor dpnp_fabs(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_FABS_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_floor(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_FLOOR_EXT, x1, dtype=None, out=out, where=True, func_name='floor') - - cpdef utils.dpnp_descriptor dpnp_fmod(utils.dpnp_descriptor x1_obj, utils.dpnp_descriptor x2_obj, object dtype=None, @@ -650,7 +639,3 @@ cpdef utils.dpnp_descriptor dpnp_trapz(utils.dpnp_descriptor y1, utils.dpnp_desc c_dpctl.DPCTLEvent_Delete(event_ref) return result - - -cpdef utils.dpnp_descriptor dpnp_trunc(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_TRUNC_EXT, x1, dtype=None, out=out, where=True, func_name='trunc') diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index c97c79ec0ad..b2f7844b1d5 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -43,9 +43,11 @@ __all__ = [ "check_nd_call_func", "dpnp_add", + "dpnp_ceil", "dpnp_cos", "dpnp_divide", "dpnp_equal", + "dpnp_floor", "dpnp_floor_divide", "dpnp_greater", "dpnp_greater_equal", @@ -65,6 +67,7 @@ "dpnp_sqrt", "dpnp_square", "dpnp_subtract", + "dpnp_trunc", ] @@ -180,6 +183,58 @@ def dpnp_add(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_ceil_docstring = """ +ceil(x, out=None, order='K') + +Returns the ceiling for each element `x_i` for input array `x`. +The ceil of the scalar `x` is the smallest integer `i`, such that `i >= x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have a real-valued 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 ceiling of input array. + The returned array has the same data type as `x`. +""" + + +def dpnp_ceil(x, out=None, order="K"): + """ + Invokes ceil() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for ceil() function. + + """ + + def _call_ceil(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_ceil_to_call(sycl_queue, src, dst): + # call pybind11 extension for ceil() function from OneMKL VM + return vmi._ceil(sycl_queue, src, dst, depends) + return ti._ceil(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( + "ceil", ti._ceil_result_type, _call_ceil, _ceil_docstring + ) + res_usm = func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _cos_docstring = """ cos(x, out=None, order='K') Computes cosine for each element `x_i` for input array `x`. @@ -346,6 +401,58 @@ def dpnp_equal(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_floor_docstring = """ +floor(x, out=None, order='K') + +Returns the floor for each element `x_i` for input array `x`. +The floor of the scalar `x` is the largest integer `i`, such that `i <= x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have a real-valued 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 floor of input array. + The returned array has the same data type as `x`. +""" + + +def dpnp_floor(x, out=None, order="K"): + """ + Invokes floor() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for floor() function. + + """ + + def _call_floor(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_floor_to_call(sycl_queue, src, dst): + # call pybind11 extension for floor() function from OneMKL VM + return vmi._floor(sycl_queue, src, dst, depends) + return ti._floor(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( + "floor", ti._floor_result_type, _call_floor, _floor_docstring + ) + res_usm = func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _floor_divide_docstring_ = """ floor_divide(x1, x2, out=None, order="K") Calculates the ratio for each element `x1_i` of the input array `x1` with @@ -1162,3 +1269,57 @@ def dpnp_subtract(x1, x2, out=None, order="K"): ) res_usm = func(x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order) return dpnp_array._create_from_usm_ndarray(res_usm) + + +_trunc_docstring = """ +trunc(x, out=None, order='K') + +Returns the truncated value for each element `x_i` for input array `x`. +The truncated value of the scalar `x` is the nearest integer `i` which is +closer to zero than `x` is. In short, the fractional part of the +signed number `x` is discarded. + +Args: + x (dpnp.ndarray): + Input array, expected to have a real-valued 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 truncated value of each element in `x`. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def dpnp_trunc(x, out=None, order="K"): + """ + Invokes trunc() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for trunc() function. + + """ + + def _call_trunc(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_trunc_to_call(sycl_queue, src, dst): + # call pybind11 extension for trunc() function from OneMKL VM + return vmi._trunc(sycl_queue, src, dst, depends) + return ti._trunc(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( + "trunc", ti._trunc_result_type, _call_trunc, _trunc_docstring + ) + res_usm = func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 36e9804618f..0054f78fb49 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -49,11 +49,15 @@ from .dpnp_algo import * from .dpnp_algo.dpnp_elementwise_common import ( + check_nd_call_func, dpnp_add, + dpnp_ceil, dpnp_divide, + dpnp_floor, dpnp_floor_divide, dpnp_multiply, dpnp_subtract, + dpnp_trunc, ) from .dpnp_utils import * @@ -200,7 +204,7 @@ def absolute(x, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): ----------- Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. Parameters `out`, `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments ``kwargs`` are currently unsupported. + 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`. @@ -275,7 +279,7 @@ def add( 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. + 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`. @@ -312,9 +316,9 @@ def around(x1, decimals=0, out=None): Limitations ----------- - Parameters ``x1`` is supported as :obj:`dpnp.ndarray`. - Parameters ``decimals`` and ``out`` are supported with their default values. - Otherwise the functions will be executed sequentially on CPU. + Parameters `x1` is supported as :class:`dpnp.ndarray`. + Parameters `decimals` and `out` are supported with their default values. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. Examples @@ -345,18 +349,34 @@ def around(x1, decimals=0, out=None): return call_origin(numpy.around, x1, decimals=decimals, out=out) -def ceil(x1, out=None, **kwargs): +def ceil( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ - Compute the ceiling of the input, element-wise. + Compute the ceiling of the input, element-wise. For full documentation refer to :obj:`numpy.ceil`. + Returns + ------- + out : dpnp.ndarray + The ceiling of each element of `x`. + Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + 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 arguments `kwargs` are currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by real-value data types. See Also -------- @@ -367,24 +387,22 @@ def ceil(x1, out=None, **kwargs): -------- >>> import dpnp as np >>> a = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) - >>> result = np.ceil(a) - >>> [x for x in result] - [-1.0, -1.0, -0.0, 1.0, 2.0, 2.0, 2.0] + >>> np.ceil(a) + array([-1.0, -1.0, -0.0, 1.0, 2.0, 2.0, 2.0]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.ceil, + dpnp_ceil, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc and not kwargs: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_ceil(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.ceil, x1, out=out, **kwargs) def conjugate(x1, **kwargs): @@ -446,10 +464,10 @@ def copysign(x1, x2, dtype=None, out=None, where=True, **kwargs): 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 arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. + Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. + Parameters `dtype`, `out` and `where` 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`. Examples @@ -505,12 +523,12 @@ def cross(x1, x2, axisa=-1, axisb=-1, axisc=-1, axis=None): Limitations ----------- - Parameters ``x1`` and ``x2`` are supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Sizes of input arrays are limited by ``x1.size == 3 and x2.size == 3``. - Shapes of input arrays are limited by ``x1.shape == (3,) and x2.shape == (3,)``. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters `x1` and `x2` are supported as :class:`dpnp.ndarray`. + Keyword arguments `kwargs` are currently unsupported. + Sizes of input arrays are limited by `x1.size == 3 and x2.size == 3`. + Shapes of input arrays are limited by `x1.shape == (3,) and x2.shape == (3,)`. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. Examples -------- @@ -553,10 +571,10 @@ def cumprod(x1, **kwargs): Limitations ----------- - Parameter ``x`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x` is supported as :class:`dpnp.ndarray`. + 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`. Examples -------- @@ -587,10 +605,10 @@ def cumsum(x1, **kwargs): Limitations ----------- - Parameter ``x`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x` is supported as :obj:`dpnp.ndarray`. + 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`. Examples -------- @@ -622,7 +640,7 @@ def diff(x1, n=1, axis=-1, prepend=numpy._NoValue, append=numpy._NoValue): Limitations ----------- Input array is supported as :obj:`dpnp.ndarray`. - Parameters ``axis``, ``prepend`` and ``append`` are supported only with default values. + Parameters `axis`, `prepend` and `append` are supported only with default values. Otherwise the function will be executed sequentially on CPU. """ @@ -668,14 +686,14 @@ def divide( Returns ------- y : dpnp.ndarray - The quotient ``x1/x2``, element-wise. + The quotient `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 `out`, `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments ``kwargs`` are currently unsupported. + 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`. @@ -710,12 +728,14 @@ def ediff1d(x1, to_end=None, to_begin=None): Limitations ----------- - Parameter ``x1``is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``to_end`` and ``to_begin`` are currently supported only with default values `None`. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x1`is supported as :class:`dpnp.ndarray`. + Keyword arguments `to_end` and `to_begin` are currently supported only with default values `None`. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. - ..seealso:: :obj:`dpnp.diff` : Calculate the n-th discrete difference along the given axis. + See Also + -------- + :obj:`dpnp.diff` : Calculate the n-th discrete difference along the given axis. Examples -------- @@ -751,12 +771,14 @@ def fabs(x1, **kwargs): Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x1` is supported as :class:`dpnp.ndarray`. + 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`. - .. seealso:: :obj:`dpnp.abs` : Calculate the absolute value element-wise. + See Also + -------- + :obj:`dpnp.abs` : Calculate the absolute value element-wise. Examples -------- @@ -776,51 +798,65 @@ def fabs(x1, **kwargs): return call_origin(numpy.fabs, x1, **kwargs) -def floor(x1, out=None, **kwargs): +def floor( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Round a number to the nearest integer toward minus infinity. For full documentation refer to :obj:`numpy.floor`. + Returns + ------- + out : dpnp.ndarray + The floor of each element of `x`. + Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + 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 arguments `kwargs` are currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by real-value data types. See Also -------- - :obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise. - :obj:`dpnp.trunc` : Return the truncated value of the input, element-wise. + :obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise. + :obj:`dpnp.trunc` : Return the truncated value of the input, element-wise. Notes ----- - Some spreadsheet programs calculate the "floor-towards-zero", in other words floor(-2.5) == -2. - dpNP instead uses the definition of floor where floor(-2.5) == -3. + Some spreadsheet programs calculate the "floor-towards-zero", in other words floor(-2.5) == -2. + DPNP instead uses the definition of floor where floor(-2.5) == -3. Examples -------- >>> import dpnp as np >>> a = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) - >>> result = np.floor(a) - >>> [x for x in result] - [-2.0, -2.0, -1.0, 0.0, 1.0, 1.0, 2.0] + >>> np.floor(a) + array([-2.0, -2.0, -1.0, 0.0, 1.0, 1.0, 2.0]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.floor, + dpnp_floor, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc and not kwargs: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_floor(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.floor, x1, out=out, **kwargs) def floor_divide( @@ -842,11 +878,12 @@ def floor_divide( Limitations ----------- - Parameters ``x1`` and ``x2`` are supported as either :obj:`dpnp.ndarray` or scalar. - Parameters ``where``, ``dtype`` and ``subok`` are supported with their default values. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + 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 -------- @@ -927,10 +964,10 @@ def fmod(x1, x2, dtype=None, out=None, where=True, **kwargs): 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 arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. + Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. + Parameters `dtype`, `out` and `where` 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 @@ -998,11 +1035,11 @@ def gradient(x1, *varargs, **kwargs): Limitations ----------- - Parameter ``y1`` is supported as :obj:`dpnp.ndarray`. - Argument ``varargs[0]`` is supported as `int`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `y1` is supported as :class:`dpnp.ndarray`. + Argument `varargs[0]` is supported as `int`. + 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`. Example ------- @@ -1040,10 +1077,10 @@ def maximum(x1, x2, dtype=None, out=None, where=True, **kwargs): 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 arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. + Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. + Parameters `dtype`, `out` and `where` 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 @@ -1105,10 +1142,10 @@ def minimum(x1, x2, dtype=None, out=None, where=True, **kwargs): 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 arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. + Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. + Parameters `dtype`, `out` and `where` 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 @@ -1191,10 +1228,10 @@ def modf(x1, **kwargs): Limitations ----------- - Parameter ``x`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x` is supported as :obj:`dpnp.ndarray`. + 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`. Examples -------- @@ -1241,8 +1278,8 @@ def multiply( 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 functions will be executed sequentially on CPU. + 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`. Examples @@ -1277,10 +1314,10 @@ def nancumprod(x1, **kwargs): Limitations ----------- - Parameter ``x`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x` is supported as :class:`dpnp.ndarray`. + 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`. .. seealso:: :obj:`dpnp.cumprod` : Return the cumulative product of elements along a given axis. @@ -1314,12 +1351,14 @@ def nancumsum(x1, **kwargs): Limitations ----------- - Parameter ``x`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x` is supported as :class:`dpnp.ndarray`. + 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`. - .. seealso:: :obj:`dpnp.cumsum` : Return the cumulative sum of the elements along a given axis. + See Also + -------- + :obj:`dpnp.cumsum` : Return the cumulative sum of the elements along a given axis. Examples -------- @@ -1350,10 +1389,10 @@ def nanprod(x1, **kwargs): Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x1` is supported as :obj:`dpnp.ndarray`. + 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`. Examples -------- @@ -1380,10 +1419,10 @@ def nansum(x1, **kwargs): Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x1` is supported as :class:`dpnp.ndarray`. + 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`. Examples -------- @@ -1412,12 +1451,14 @@ def negative(x1, **kwargs): Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x1` is supported as :class:`dpnp.ndarray`. + 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.copysign` : Change the sign of x1 to that of x2, element-wise. + See Also + -------- + :obj:`dpnp.copysign` : Change the sign of x1 to that of x2, element-wise. Examples -------- @@ -1456,7 +1497,7 @@ def power(x1, x2, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): 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. + 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`. @@ -1551,8 +1592,8 @@ def prod( Limitations ----------- - Parameter ``where`` is unsupported. - Input array data types are limited by DPNP :ref:`Data types`. + Parameter `where` is unsupported. + Input array data types are limited by DPNP :ref:`Data types`. Examples -------- @@ -1603,18 +1644,18 @@ def remainder(x1, x2, out=None, where=True, dtype=None, **kwargs): 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 arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. - Parameters ``x1`` and ``x2`` are supported with equal sizes and shapes. + 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 arguments `kwargs` are currently unsupported. + Otherwise the functions will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters `x1` and `x2` are supported with equal sizes and shapes. See Also -------- - :obj:`dpnp.fmod` : Calculate the element-wise remainder of division. - :obj:`dpnp.divide` : Standard division. - :obj:`dpnp.floor` : Round a number to the nearest integer toward minus infinity. + :obj:`dpnp.fmod` : Calculate the element-wise remainder of division. + :obj:`dpnp.divide` : Standard division. + :obj:`dpnp.floor` : Round a number to the nearest integer toward minus infinity. Example ------- @@ -1679,7 +1720,7 @@ def round_(a, decimals=0, out=None): See Also -------- - :obj:`dpnp.around` : equivalent function; see for details. + :obj:`dpnp.around` : equivalent function; see for details. """ @@ -1694,10 +1735,10 @@ def sign(x1, **kwargs): Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameter `x1` is supported as :class:`dpnp.ndarray`. + 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`. Examples -------- @@ -1744,7 +1785,7 @@ def subtract( 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 `out`, `where`, `dtype` and `subok` are supported with their default values. - Keyword arguments ``kwargs`` are currently unsupported. + 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`. @@ -1797,11 +1838,11 @@ def sum( Limitations ----------- - Parameters `x` is supported as either :class:`dpnp.ndarray` - or :class:`dpctl.tensor.usm_ndarray`. - Parameters `out`, `initial` and `where` are supported with their default values. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters `x` is supported as either :class:`dpnp.ndarray` + or :class:`dpctl.tensor.usm_ndarray`. + Parameters `out`, `initial` and `where` are supported with their default values. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. Examples -------- @@ -1892,10 +1933,10 @@ def trapz(y1, x1=None, dx=1.0, axis=-1): Limitations ----------- - Parameters ``y`` and ``x`` are supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters `y` and `x` are supported as :class:`dpnp.ndarray`. + 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`. Examples -------- @@ -1952,7 +1993,7 @@ def true_divide(*args, **kwargs): See Also -------- - .. seealso:: :obj:`dpnp.divide` : Standard division. + :obj:`dpnp.divide` : Standard division. Notes ----- @@ -1964,43 +2005,57 @@ def true_divide(*args, **kwargs): return dpnp.divide(*args, **kwargs) -def trunc(x1, out=None, **kwargs): +def trunc( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Compute the truncated value of the input, element-wise. For full documentation refer to :obj:`numpy.trunc`. + Returns + ------- + out : dpnp.ndarray + The truncated value of each element of `x`. + Limitations ----------- - Parameter ``x1`` is supported as :obj:`dpnp.ndarray`. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + 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 arguments `kwargs` are currently unsupported. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by real-value data types. See Also -------- - :obj:`dpnp.floor` : Round a number to the nearest integer toward minus infinity. - :obj:`dpnp.ceil` : Round a number to the nearest integer toward infinity. + :obj:`dpnp.floor` : Round a number to the nearest integer toward minus infinity. + :obj:`dpnp.ceil` : Round a number to the nearest integer toward infinity. Examples -------- >>> import dpnp as np >>> a = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) - >>> result = np.trunc(a) - >>> [x for x in result] - [-1.0, -1.0, -0.0, 0.0, 1.0, 1.0, 2.0] + >>> np.trunc(a) + array([-1.0, -1.0, -0.0, 0.0, 1.0, 1.0, 2.0]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.trunc, + dpnp_trunc, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc and not kwargs: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_trunc(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.trunc, x1, out=out, **kwargs) diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index 904e1e3a1a1..d64359e1db8 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -434,8 +434,9 @@ def cos( Limitations ----------- - Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_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 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`. @@ -986,8 +987,9 @@ def sin( Limitations ----------- - Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_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 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`. diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index cd00837d037..a02f2078b51 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -784,13 +784,10 @@ tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{de tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_large tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_small tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_around -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_ceil tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_fix -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_floor tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint_negative tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_round_ -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_trunc tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_sum_out tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_sum_out_wrong_shape tests/third_party/cupy/math_tests/test_sumprod.py::TestCumprod::test_ndarray_cumprod_2dim_with_axis diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index d0a477666ef..32b7a3ed04b 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -12,23 +12,19 @@ tests/test_random.py::TestPermutationsTestShuffle::test_no_miss_numbers[int64] tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.array([])] tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.astype(dpnp.asarray(x), dpnp.float32)] -tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-ceil-data1] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-conjugate-data2] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-copy-data3] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-cumprod-data4] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-cumsum-data5] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-ediff1d-data7] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-fabs-data8] -tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-floor-data9] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-ceil-data1] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-conjugate-data2] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-copy-data3] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-cumprod-data4] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-cumsum-data5] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-ediff1d-data7] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-fabs-data8] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-floor-data9] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-nancumprod-data11] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-nancumsum-data12] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-nanprod-data13] @@ -38,7 +34,6 @@ tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-prod-data16] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-sign-data17] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-sum-data18] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-trapz-data19] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-trunc-data20] tests/test_sycl_queue.py::test_broadcasting[level_zero:gpu:0-remainder-data15-data25] tests/test_sycl_queue.py::test_broadcasting[opencl:gpu:0-remainder-data15-data25] @@ -964,13 +959,10 @@ tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{de tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_large tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_6_{decimals=100}::test_round_small tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_around -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_ceil tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_fix -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_floor tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_rint_negative tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_round_ -tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_trunc tests/third_party/cupy/math_tests/test_sumprod.py::TestCumprod::test_ndarray_cumprod_2dim_with_axis tests/third_party/cupy/math_tests/test_sumprod.py::TestDiff::test_diff_1dim tests/third_party/cupy/math_tests/test_sumprod.py::TestDiff::test_diff_1dim_with_n diff --git a/tests/third_party/cupy/math_tests/test_rounding.py b/tests/third_party/cupy/math_tests/test_rounding.py index e4927495cc5..e17a48aa9e3 100644 --- a/tests/third_party/cupy/math_tests/test_rounding.py +++ b/tests/third_party/cupy/math_tests/test_rounding.py @@ -10,7 +10,7 @@ @testing.gpu class TestRounding(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(type_check=False, atol=1e-5) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a)