From 6ab618d0aad61ece28dae925300582547dd0a512 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Mon, 7 Aug 2023 19:00:53 -0500 Subject: [PATCH 1/5] use_dpctl_remainder_func --- dpnp/backend/extensions/vm/remainder.hpp | 83 ++++++ dpnp/backend/extensions/vm/types_matrix.hpp | 15 ++ dpnp/backend/extensions/vm/vm_py.cpp | 33 +++ dpnp/backend/include/dpnp_iface_fptr.hpp | 36 ++- .../kernels/dpnp_krnl_mathematical.cpp | 50 ---- dpnp/dpnp_algo/dpnp_algo.pxd | 5 - dpnp/dpnp_algo/dpnp_algo_mathematical.pxi | 9 - dpnp/dpnp_algo/dpnp_elementwise_common.py | 46 +++- dpnp/dpnp_iface_mathematical.py | 244 +++++++++--------- tests/skipped_tests.tbl | 7 - tests/skipped_tests_gpu.tbl | 5 - tests/test_mathematical.py | 4 - 12 files changed, 308 insertions(+), 229 deletions(-) create mode 100644 dpnp/backend/extensions/vm/remainder.hpp diff --git a/dpnp/backend/extensions/vm/remainder.hpp b/dpnp/backend/extensions/vm/remainder.hpp new file mode 100644 index 000000000000..ac77dbe7f628 --- /dev/null +++ b/dpnp/backend/extensions/vm/remainder.hpp @@ -0,0 +1,83 @@ +//***************************************************************************** +// 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 remainder_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::remainder( + 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 RemainderContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::RemainderOutputType::value_type, + void>) + { + return nullptr; + } + else { + return remainder_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 667e27f0c9ad..1ca03e3b3a7b 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::remainder function. + * + * @tparam T Type of input vectors `a` and `b` and of result vector `y`. + */ +template +struct RemainderOutputType +{ + 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::cos function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 924db66025b3..2c2578105b6d 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -34,6 +34,7 @@ #include "cos.hpp" #include "div.hpp" #include "ln.hpp" +#include "remainder.hpp" #include "sin.hpp" #include "sqr.hpp" #include "sqrt.hpp" @@ -46,6 +47,7 @@ using vm_ext::binary_impl_fn_ptr_t; using vm_ext::unary_impl_fn_ptr_t; static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types]; +static binary_impl_fn_ptr_t remainder_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 ln_dispatch_vector[dpctl_td_ns::num_types]; @@ -88,6 +90,37 @@ PYBIND11_MODULE(_vm_impl, m) py::arg("dst")); } + // BinaryUfunc: ==== REMAINDER(x1, x2) ==== + { + vm_ext::init_ufunc_dispatch_vector( + remainder_dispatch_vector); + + auto remainder_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, + remainder_dispatch_vector); + }; + m.def("_remainder", remainder_pyapi, + "Call `remainder` function from OneMKL VM library to performs " + "element " + "by element remainder of vector `src1` by vector `src2` " + "to resulting vector `dst`", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst"), py::arg("depends") = py::list()); + + auto remainder_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, + remainder_dispatch_vector); + }; + m.def("_mkl_remainder_to_call", remainder_need_to_call_pyapi, + "Check input arguments to answer if `remainder` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), + py::arg("dst")); + } + // UnaryUfunc: ==== Cos(x) ==== { vm_ext::init_ufunc_dispatch_vector; -template -DPCTLSyclEventRef (*dpnp_remainder_ext_c)(DPCTLSyclQueueRef, - void *, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const void *, - const size_t, - const shape_elem_type *, - const size_t, - const size_t *, - const DPCTLEventVectorRef) = - dpnp_remainder_c<_DataType_output, _DataType_input1, _DataType_input2>; - template @@ -1385,39 +1368,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_REMAINDER][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_remainder_default_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_REMAINDER_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_remainder_ext_c}; - fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_trapz_default_c}; fmap[DPNPFuncName::DPNP_FN_TRAPZ][eft_INT][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 5bbb0dc4a010..7427c6c7ce44 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -188,8 +188,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_QR_EXT DPNP_FN_RADIANS DPNP_FN_RADIANS_EXT - DPNP_FN_REMAINDER - DPNP_FN_REMAINDER_EXT DPNP_FN_RECIP DPNP_FN_RECIP_EXT DPNP_FN_REPEAT @@ -442,9 +440,6 @@ cpdef dpnp_descriptor dpnp_minimum(dpnp_descriptor x1_obj, dpnp_descriptor x2_ob cpdef dpnp_descriptor dpnp_negative(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_power(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, dpnp_descriptor out=*, object where=*) -cpdef dpnp_descriptor dpnp_remainder(dpnp_descriptor x1_obj, dpnp_descriptor x2_obj, object dtype=*, - dpnp_descriptor out=*, object where=*) - """ Array manipulation routines diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index ae03fa54051c..1685c770d4ee 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -62,7 +62,6 @@ __all__ += [ "dpnp_negative", "dpnp_power", "dpnp_prod", - "dpnp_remainder", "dpnp_sign", "dpnp_sum", "dpnp_trapz", @@ -546,14 +545,6 @@ cpdef utils.dpnp_descriptor dpnp_prod(utils.dpnp_descriptor x1, return result -cpdef utils.dpnp_descriptor dpnp_remainder(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(DPNP_FN_REMAINDER_EXT, x1_obj, x2_obj, dtype, out, where) - - cpdef utils.dpnp_descriptor dpnp_sign(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_SIGN_EXT, x1) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 410f5542f7ba..c9a0796c8652 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -66,6 +66,7 @@ "dpnp_logical_xor", "dpnp_multiply", "dpnp_not_equal", + "dpnp_remainder", "dpnp_right_shift", "dpnp_sin", "dpnp_sqrt", @@ -86,7 +87,7 @@ def check_nd_call_func( **kwargs, ): """ - Checks arguments and calls function with a single input array. + Checks arguments and calls a function. Chooses a common internal elementwise function to call in DPNP based on input arguments or to fallback on NumPy call if any passed argument is not currently supported. @@ -127,7 +128,6 @@ def check_nd_call_func( order ) ) - return dpnp_func(*x_args, out=out, order=order) return call_origin( origin_func, @@ -1174,6 +1174,48 @@ def dpnp_not_equal(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_remainder_docstring_ = """ +remainder(x1, x2, out=None, order='K') +Calculates the remainder of division for each element `x1_i` of the input array +`x1` with the respective element `x2_i` of the input array `x2`. +This function is equivalent to the Python modulus operator. +Args: + x1 (dpnp.ndarray): + First input array, expected to have a real-valued data type. + x2 (dpnp.ndarray): + Second input array, also expected to have a real-valued data type. + out ({None, usm_ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + dpnp.ndarray: + an array containing the element-wise remainders. The data type of + the returned array is determined by the Type Promotion Rules. +""" + + +remainder_func = BinaryElementwiseFunc( + "remainder", + ti._remainder_result_type, + ti._remainder, + _remainder_docstring_, +) + + +def dpnp_remainder(x1, x2, out=None, order="K"): + # 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 = remainder_func( + x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order + ) + return dpnp_array._create_from_usm_ndarray(res_usm) + _right_shift_docstring_ = """ right_shift(x1, x2, out=None, order='K') diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 36e9804618f8..114b2097aee9 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -49,10 +49,12 @@ from .dpnp_algo import * from .dpnp_algo.dpnp_elementwise_common import ( + check_nd_call_func, dpnp_add, dpnp_divide, dpnp_floor_divide, dpnp_multiply, + dpnp_remainder, dpnp_subtract, ) from .dpnp_utils import * @@ -103,63 +105,6 @@ ] -def _check_nd_call( - origin_func, - dpnp_func, - x1, - x2, - out=None, - where=True, - order="K", - dtype=None, - subok=True, - **kwargs, -): - """ - Checks arguments and calls a function. - - Chooses a common internal elementwise function to call in DPNP based on input arguments - or to fallback on NumPy call if any passed argument is not currently supported. - - """ - - if kwargs: - pass - elif where is not True: - pass - elif dtype is not None: - pass - elif subok is not True: - pass - elif dpnp.isscalar(x1) and dpnp.isscalar(x2): - # at least either x1 or x2 has to be an array - pass - else: - if order in "afkcAFKC": - order = order.upper() - elif order is None: - order = "K" - else: - raise ValueError( - "order must be one of 'C', 'F', 'A', or 'K' (got '{}')".format( - order - ) - ) - - return dpnp_func(x1, x2, out=out, order=order) - return call_origin( - origin_func, - x1, - x2, - out=out, - where=where, - order=order, - dtype=dtype, - subok=subok, - **kwargs, - ) - - def abs(*args, **kwargs): """ Calculate the absolute value element-wise. @@ -281,16 +226,29 @@ def add( Examples -------- - >>> import dpnp as dp - >>> a = dp.array([1, 2, 3]) - >>> b = dp.array([1, 2, 3]) - >>> result = dp.add(a, b) - >>> print(result) - [2, 4, 6] + >>> import dpnp as np + >>> a = np.array([1, 2, 3]) + >>> b = np.array([1, 2, 3]) + >>> np.add(a, b) + array([2, 4, 6]) + >>> x1 = np.arange(9.0).reshape((3, 3)) + >>> x2 = np.arange(3.0) + >>> np.add(x1, x2) + array([[ 0., 2., 4.], + [ 3., 5., 7.], + [ 6., 8., 10.]]) + + The ``+`` operator can be used as a shorthand for ``add`` on + :class:`dpnp.ndarray`. + + >>> x1 + x2 + array([[ 0., 2., 4.], + [ 3., 5., 7.], + [ 6., 8., 10.]]) """ - return _check_nd_call( + return check_nd_call_func( numpy.add, dpnp_add, x1, @@ -681,14 +639,29 @@ def divide( Examples -------- - >>> import dpnp as dp - >>> result = dp.divide(dp.array([1, -2, 6, -9]), dp.array([-2, -2, -2, -2])) - >>> print(result) - [-0.5, 1.0, -3.0, 4.5] + >>> import dpnp as np + >>> np.divide(dp.array([1, -2, 6, -9]), np.array([-2, -2, -2, -2])) + array([-0.5, 1. , -3. , 4.5]) + >>> x1 = np.arange(9.0).reshape((3, 3)) + >>> x2 = np.arange(3.0) + >>> np.divide(x1, x2) + array([[nan, 1. , 1. ], + [inf, 4. , 2.5], + [inf, 7. , 4. ]]) + + The ``/`` operator can be used as a shorthand for ``divide`` on + :class:`dpnp.ndarray`. + + >>> x1 = np.arange(9.0).reshape((3, 3)) + >>> x2 = 2 * np.ones(3) + >>> x1/x2 + array([[0. , 0.5, 1. ], + [1.5, 2. , 2.5], + [3. , 3.5, 4. ]]) """ - return _check_nd_call( + return check_nd_call_func( numpy.divide, dpnp_divide, x1, @@ -858,12 +831,14 @@ def floor_divide( Examples -------- >>> import dpnp as np - >>> np.floor_divide(np.array([1, -1, -2, -9]), np.array([-2, -2, -2, -2])) + >>> np.floor_divide(np.array([1, -1, -2, -9]), -2) array([-1, 0, 1, 4]) + >>> np.floor_divide(np.array([1., 2., 3., 4.]), 2.5) + array([ 0., 0., 1., 1.]) """ - return _check_nd_call( + return check_nd_call_func( numpy.floor_divide, dpnp_floor_divide, x1, @@ -1247,15 +1222,28 @@ def multiply( Examples -------- - >>> import dpnp as dp - >>> a = dp.array([1, 2, 3, 4, 5]) - >>> result = dp.multiply(a, a) - >>> print(result) - [1, 4, 9, 16, 25] + >>> import dpnp as np + >>> a = np.array([1, 2, 3, 4, 5]) + >>> np.multiply(a, a) + array([ 1, 4, 9, 16, 25])] + + >>> x1 = np.arange(9.0).reshape((3, 3)) + >>> x2 = np.arange(3.0) + >>> np.multiply(x1, x2) + array([[ 0., 1., 4.], + [ 0., 4., 10.], + [ 0., 7., 16.]]) + The ``*`` operator can be used as a shorthand for ``multiply`` on + :class:`dpnp.ndarray`. + + >>> x1 * x2 + array([[ 0., 1., 4.], + [ 0., 4., 10.], + [ 0., 7., 16.]]) """ - return _check_nd_call( + return check_nd_call_func( numpy.multiply, dpnp_multiply, x1, @@ -1595,7 +1583,18 @@ def prod( ) -def remainder(x1, x2, out=None, where=True, dtype=None, **kwargs): +def remainder( + x1, + x2, + /, + out=None, + *, + where=True, + order="K", + dtype=None, + subok=True, + **kwargs, +): """ Return element-wise remainder of division. @@ -1619,55 +1618,31 @@ def remainder(x1, x2, out=None, where=True, dtype=None, **kwargs): Example ------- >>> import dpnp as np - >>> result = np.remainder(np.array([4, 7]), np.array([2, 3])) - >>> [x for x in result] - [0, 1] + >>> np.remainder(np.array([4, 7]), np.array([2, 3])) + array([0, 1]) - """ + >>> np.remainder(np.arange(7), 5) + array([0, 1, 2, 3, 4, 0, 1]) - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_nondefault_queue=False) + The ``%`` operator can be used as a shorthand for ``remainder`` on + :class:`dpnp.ndarray`. - 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 x2_is_scalar and not x2_desc: - pass - elif x1_desc and x2_desc and x1_desc.size != x2_desc.size: - # TODO: enable broadcasting - pass - elif x1_desc and x2_desc and x1_desc.shape != x2_desc.shape: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - elif x1_is_scalar and x2_desc.ndim > 1: - pass - else: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_remainder( - x1_desc, x2_desc, dtype, out_desc, where - ).get_pyobj() + >>> x1 = np.arange(7) + >>> x1 % 5 + array([0, 1, 2, 3, 4, 0, 1]) + """ - return call_origin( - numpy.remainder, x1, x2, out=out, where=where, dtype=dtype, **kwargs + return check_nd_call_func( + numpy.remainder, + dpnp_remainder, + x1, + x2, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) @@ -1750,14 +1725,27 @@ def subtract( Example ------- - >>> import dpnp as dp - >>> result = dp.subtract(dp.array([4, 3]), dp.array([2, 7])) - >>> print(result) - [2, -4] + >>> import dpnp as np + >>> np.subtract(dp.array([4, 3]), np.array([2, 7])) + array([ 2, -4]) + + >>> x1 = np.arange(9.0).reshape((3, 3)) + >>> x2 = np.arange(3.0) + >>> np.subtract(x1, x2) + array([[ 0., 0., 0.], + [ 3., 3., 3.], + [ 6., 6., 6.]]) + + The ``-`` operator can be used as a shorthand for ``subtract`` on + :class:`dpnp.ndarray`. + >>> x1 - x2 + array([[ 0., 0., 0.], + [ 3., 3., 3.], + [ 6., 6., 6.]]) """ - return _check_nd_call( + return check_nd_call_func( numpy.subtract, dpnp_subtract, x1, diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index cd00837d0379..07df45e2915c 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -7,8 +7,6 @@ tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp. tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-trapz-data19] tests/test_sycl_queue.py::test_1in_1out[opencl:cpu:0-trapz-data19] -tests/test_sycl_queue.py::test_broadcasting[opencl:gpu:0-remainder-data15-data25] -tests/test_sycl_queue.py::test_broadcasting[opencl:cpu:0-remainder-data15-data25] tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_1_{axes=None, norm=None, s=(1, None), shape=(3, 4)}::test_fft2 tests/third_party/cupy/fft_tests/test_fft.py::TestFft2_param_7_{axes=(), norm=None, s=None, shape=(3, 4)}::test_fft2 @@ -577,20 +575,15 @@ tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_3_{reps tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_4_{reps=(2, 3)}::test_array_tile tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_5_{reps=(2, 3, 4, 5)}::test_array_tile tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_457_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_459_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='remainder', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_461_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_465_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_467_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='remainder', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_469_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_537_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_539_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='remainder', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_541_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_545_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_547_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='remainder', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_549_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticModf::test_modf -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_10_{name='remainder', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_11_{name='mod', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_1_{name='angle', nargs=1}::test_raises_with_numpy_input diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index d0a477666efa..22b409df03d4 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -40,10 +40,6 @@ 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] -tests/test_sycl_queue.py::test_broadcasting[opencl:cpu:0-remainder-data15-data25] - tests/test_sycl_queue.py::test_modf[level_zero:gpu:0] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-trapz-data19] @@ -770,7 +766,6 @@ tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_3_{reps tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_4_{reps=(2, 3)}::test_array_tile tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_5_{reps=(2, 3, 4, 5)}::test_array_tile -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_10_{name='remainder', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_11_{name='mod', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_1_{name='angle', nargs=1}::test_raises_with_numpy_input diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index e0a168695671..49fb6101b622 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -210,10 +210,6 @@ def test_minimum(self, dtype, lhs, rhs): def test_multiply(self, dtype, lhs, rhs): self._test_mathematical("multiply", dtype, lhs, rhs) - @pytest.mark.skipif( - not has_support_aspect64(), reason="Aborted on Iris Xe: SAT-6039" - ) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) From 2550ab6a64921449a58477ec2798fc4c5b01bfa0 Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Thu, 10 Aug 2023 11:24:02 -0500 Subject: [PATCH 2/5] remove mkl::remainder implementation --- dpnp/backend/extensions/vm/remainder.hpp | 83 --------------------- dpnp/backend/extensions/vm/types_matrix.hpp | 15 ---- dpnp/backend/extensions/vm/vm_py.cpp | 33 -------- dpnp/dpnp_iface_mathematical.py | 8 +- 4 files changed, 4 insertions(+), 135 deletions(-) delete mode 100644 dpnp/backend/extensions/vm/remainder.hpp diff --git a/dpnp/backend/extensions/vm/remainder.hpp b/dpnp/backend/extensions/vm/remainder.hpp deleted file mode 100644 index ac77dbe7f628..000000000000 --- a/dpnp/backend/extensions/vm/remainder.hpp +++ /dev/null @@ -1,83 +0,0 @@ -//***************************************************************************** -// 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 remainder_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::remainder( - 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 RemainderContigFactory -{ - fnT get() - { - if constexpr (std::is_same_v< - typename types::RemainderOutputType::value_type, - void>) - { - return nullptr; - } - else { - return remainder_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 1ca03e3b3a7b..667e27f0c9ad 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -68,21 +68,6 @@ 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::remainder function. - * - * @tparam T Type of input vectors `a` and `b` and of result vector `y`. - */ -template -struct RemainderOutputType -{ - 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::cos function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 2c2578105b6d..924db66025b3 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -34,7 +34,6 @@ #include "cos.hpp" #include "div.hpp" #include "ln.hpp" -#include "remainder.hpp" #include "sin.hpp" #include "sqr.hpp" #include "sqrt.hpp" @@ -47,7 +46,6 @@ using vm_ext::binary_impl_fn_ptr_t; using vm_ext::unary_impl_fn_ptr_t; static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types]; -static binary_impl_fn_ptr_t remainder_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 ln_dispatch_vector[dpctl_td_ns::num_types]; @@ -90,37 +88,6 @@ PYBIND11_MODULE(_vm_impl, m) py::arg("dst")); } - // BinaryUfunc: ==== REMAINDER(x1, x2) ==== - { - vm_ext::init_ufunc_dispatch_vector( - remainder_dispatch_vector); - - auto remainder_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, - remainder_dispatch_vector); - }; - m.def("_remainder", remainder_pyapi, - "Call `remainder` function from OneMKL VM library to performs " - "element " - "by element remainder of vector `src1` by vector `src2` " - "to resulting vector `dst`", - py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), - py::arg("dst"), py::arg("depends") = py::list()); - - auto remainder_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, - remainder_dispatch_vector); - }; - m.def("_mkl_remainder_to_call", remainder_need_to_call_pyapi, - "Check input arguments to answer if `remainder` function from " - "OneMKL VM library can be used", - py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), - py::arg("dst")); - } - // UnaryUfunc: ==== Cos(x) ==== { vm_ext::init_ufunc_dispatch_vector Date: Fri, 11 Aug 2023 10:53:46 -0500 Subject: [PATCH 3/5] address reviewer's comments --- dpnp/dpnp_algo/dpnp_elementwise_common.py | 4 +- dpnp/dpnp_iface_mathematical.py | 71 +++++++++++++------ tests/test_mathematical.py | 10 ++- tests/test_strides.py | 21 +++++- tests/test_usm_type.py | 14 ++++ .../cupy/math_tests/test_arithmetic.py | 10 --- 6 files changed, 93 insertions(+), 37 deletions(-) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index c9a0796c8652..d1ffc6694c92 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -1213,6 +1213,7 @@ def dpnp_remainder(x1, x2, out=None, order="K"): res_usm = remainder_func( x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order +<<<<<<< HEAD ) return dpnp_array._create_from_usm_ndarray(res_usm) @@ -1254,8 +1255,9 @@ def dpnp_right_shift(x1, x2, out=None, order="K"): ti._bitwise_right_shift_result_type, ti._bitwise_right_shift, _right_shift_docstring_, +======= +>>>>>>> address reviewer's comments ) - 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) diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 800ffb251dbb..f5c9530046f2 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -220,7 +220,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`. @@ -626,14 +626,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. + 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`. @@ -823,7 +823,7 @@ def floor_divide( See Also -------- - :obj:`dpnp.reminder` : Remainder complementary to floor_divide. + :obj:`dpnp.remainder` : Remainder complementary to floor_divide. :obj:`dpnp.divide` : Standard division. :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. @@ -910,7 +910,7 @@ def fmod(x1, x2, dtype=None, out=None, where=True, **kwargs): See Also -------- - :obj:`dpnp.reminder` : Remainder complementary to floor_divide. + :obj:`dpnp.remainder` : Remainder complementary to floor_divide. :obj:`dpnp.divide` : Standard division. Examples @@ -1137,16 +1137,36 @@ def minimum(x1, x2, dtype=None, out=None, where=True, **kwargs): ) -def mod(*args, **kwargs): +def mod( + x1, + x2, + /, + out=None, + *, + where=True, + order="K", + dtype=None, + subok=True, + **kwargs, +): """ Compute element-wise remainder of division. For full documentation refer to :obj:`numpy.mod`. + 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.fmod` : Calculate the element-wise remainder of division - :obj:`dpnp.reminder` : Remainder complementary to floor_divide. + :obj:`dpnp.remainder` : Remainder complementary to floor_divide. :obj:`dpnp.divide` : Standard division. Notes @@ -1155,7 +1175,16 @@ def mod(*args, **kwargs): """ - return dpnp.remainder(*args, **kwargs) + return dpnp.remainder( + x1, + x2, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, + ) def modf(x1, **kwargs): @@ -1602,18 +1631,20 @@ def remainder( Limitations ----------- - Parameters ``x1`` and ``x2`` are supported as either :obj:`dpnp.ndarray`, - :class:`dpctl.tensor.usm_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 -------- - :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. + :obj:`dpnp.floor_divide` : Compute the largest integer smaller or equal to the division of the inputs. + :obj:`dpnp.mod` : Calculate the element-wise remainder of division. Example ------- @@ -1718,8 +1749,8 @@ 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. + 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/test_mathematical.py b/tests/test_mathematical.py index 49fb6101b622..c5b5ea336ba6 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -136,6 +136,12 @@ 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 ( + name == "remainder" + and result.dtype != expected.dtype + and not has_support_aspect64() + ): + pytest.skip("skipping since output is promoted differently") assert_allclose(result, expected, rtol=1e-6) @pytest.mark.parametrize("dtype", get_all_dtypes()) @@ -210,9 +216,7 @@ def test_minimum(self, dtype, lhs, rhs): def test_multiply(self, dtype, lhs, rhs): self._test_mathematical("multiply", dtype, lhs, rhs) - @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_remainder(self, dtype, lhs, rhs): self._test_mathematical("remainder", dtype, lhs, rhs) diff --git a/tests/test_strides.py b/tests/test_strides.py index 68e9dc4ea37f..e39414a7bd7f 100644 --- a/tests/test_strides.py +++ b/tests/test_strides.py @@ -233,15 +233,30 @@ def test_strides_fmod(dtype, shape): @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) @pytest.mark.parametrize("shape", [(3, 3)], ids=["(3, 3)"]) -def test_strides_true_devide(dtype, shape): +def test_strides_true_divide(dtype, shape): a = numpy.arange(numpy.prod(shape), dtype=dtype).reshape(shape) b = a.T + 1 dpa = dpnp.reshape(dpnp.arange(numpy.prod(shape), dtype=dtype), shape) dpb = dpa.T + 1 - result = dpnp.fmod(dpa, dpb) - expected = numpy.fmod(a, b) + result = dpnp.true_divide(dpa, dpb) + expected = numpy.true_divide(a, b) + + assert_allclose(result, expected, atol=1e-08) + + +@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) +@pytest.mark.parametrize("shape", [(3, 3)], ids=["(3, 3)"]) +def test_strides_remainder(dtype, shape): + a = numpy.arange(numpy.prod(shape), dtype=dtype).reshape(shape) + b = a.T + 1 + + dpa = dpnp.reshape(dpnp.arange(numpy.prod(shape), dtype=dtype), shape) + dpb = dpa.T + 1 + + result = dpnp.remainder(dpa, dpb) + expected = numpy.remainder(a, b) assert_allclose(result, expected) diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index f6e0f38a4e7d..bda39505a1c3 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -74,6 +74,20 @@ def test_coerced_usm_types_divide(usm_type_x, usm_type_y): assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) +@pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) +@pytest.mark.parametrize("usm_type_y", list_of_usm_types, ids=list_of_usm_types) +def test_coerced_usm_types_remainder(usm_type_x, usm_type_y): + x = dp.arange(100, usm_type=usm_type_x).reshape(10, 10) + y = dp.arange(100, usm_type=usm_type_y).reshape(10, 10) + y = y.T + 1 + + z = x % y + + assert x.usm_type == usm_type_x + assert y.usm_type == usm_type_y + assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) + + @pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) @pytest.mark.parametrize("usm_type_y", list_of_usm_types, ids=list_of_usm_types) def test_coerced_usm_types_power(usm_type_x, usm_type_y): diff --git a/tests/third_party/cupy/math_tests/test_arithmetic.py b/tests/third_party/cupy/math_tests/test_arithmetic.py index 955bd9e2247f..b61c3df9a3a2 100644 --- a/tests/third_party/cupy/math_tests/test_arithmetic.py +++ b/tests/third_party/cupy/math_tests/test_arithmetic.py @@ -199,16 +199,6 @@ def check_binary(self, xp): elif is_array_arg2 and not is_array_arg1: y = y.astype(dtype2) - # NumPy returns different values (nan/inf) on division by zero - # depending on the architecture. - # As it is not possible for CuPy to replicate this behavior, we ignore - # the difference here. - # if self.name in ('floor_divide', 'remainder', 'mod'): - # if y.dtype in (float_types + complex_types) and (np2 == 0).any(): - # y = xp.asarray(y) - # y[y == numpy.inf] = numpy.nan - # y[y == -numpy.inf] = numpy.nan - return y From 307db4d014c8d52e9df761b082fefd2230e9e00b Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Mon, 14 Aug 2023 15:18:06 -0500 Subject: [PATCH 4/5] remove skipping tests for mod function --- tests/skipped_tests.tbl | 5 ----- 1 file changed, 5 deletions(-) diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index 07df45e2915c..fbaa9eae002c 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -575,16 +575,11 @@ tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_3_{reps tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_4_{reps=(2, 3)}::test_array_tile tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_5_{reps=(2, 3, 4, 5)}::test_array_tile tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_457_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_461_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_465_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_469_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_537_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_541_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_545_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='fmod', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_549_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticModf::test_modf -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_11_{name='mod', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_1_{name='angle', nargs=1}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp From 988ec2091de4c9f4d0e3e7a1653d4117d5d860fd Mon Sep 17 00:00:00 2001 From: Vahid Tavanashad Date: Wed, 16 Aug 2023 13:23:22 -0500 Subject: [PATCH 5/5] rebase and resolve conflicts --- dpnp/backend/include/dpnp_iface_fptr.hpp | 28 +++++++++++------------ dpnp/dpnp_algo/dpnp_elementwise_common.py | 5 ++-- 2 files changed, 16 insertions(+), 17 deletions(-) diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index d56f3440c6a9..e2c29a0f4a88 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -313,20 +313,20 @@ enum class DPNPFuncName : size_t DPNP_FN_PUT_ALONG_AXIS_EXT, /**< Used in numpy.put_along_axis() impl, requires extra parameters */ DPNP_FN_QR, /**< Used in numpy.linalg.qr() impl */ - DPNP_FN_QR_EXT, /**< Used in numpy.linalg.qr() impl, requires extra - parameters */ - DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */ - DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra - parameters */ - DPNP_FN_REMAINDER, /**< Used in numpy.remainder() impl */ - DPNP_FN_RECIP, /**< Used in numpy.recip() impl */ - DPNP_FN_RECIP_EXT, /**< Used in numpy.recip() impl, requires extra - parameters */ - DPNP_FN_REPEAT, /**< Used in numpy.repeat() impl */ - DPNP_FN_REPEAT_EXT, /**< Used in numpy.repeat() impl, requires extra - parameters */ - DPNP_FN_RIGHT_SHIFT, /**< Used in numpy.right_shift() impl */ - DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ + DPNP_FN_QR_EXT, /**< Used in numpy.linalg.qr() impl, requires extra + parameters */ + DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */ + DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra + parameters */ + DPNP_FN_REMAINDER, /**< Used in numpy.remainder() impl */ + DPNP_FN_RECIP, /**< Used in numpy.recip() impl */ + DPNP_FN_RECIP_EXT, /**< Used in numpy.recip() impl, requires extra + parameters */ + DPNP_FN_REPEAT, /**< Used in numpy.repeat() impl */ + DPNP_FN_REPEAT_EXT, /**< Used in numpy.repeat() impl, requires extra + parameters */ + DPNP_FN_RIGHT_SHIFT, /**< Used in numpy.right_shift() impl */ + DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */ DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra parameters */ DPNP_FN_RNG_BINOMIAL, /**< Used in numpy.random.binomial() impl */ diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index d1ffc6694c92..33776e26a012 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -1213,10 +1213,10 @@ def dpnp_remainder(x1, x2, out=None, order="K"): res_usm = remainder_func( x1_usm_or_scalar, x2_usm_or_scalar, out=out_usm, order=order -<<<<<<< HEAD ) return dpnp_array._create_from_usm_ndarray(res_usm) + _right_shift_docstring_ = """ right_shift(x1, x2, out=None, order='K') @@ -1255,9 +1255,8 @@ def dpnp_right_shift(x1, x2, out=None, order="K"): ti._bitwise_right_shift_result_type, ti._bitwise_right_shift, _right_shift_docstring_, -======= ->>>>>>> address reviewer's comments ) + 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)