diff --git a/dpnp/backend/extensions/vm/round.hpp b/dpnp/backend/extensions/vm/round.hpp new file mode 100644 index 00000000000..0760db4d3fc --- /dev/null +++ b/dpnp/backend/extensions/vm/round.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 round_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::rint(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 RoundContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::RoundOutputType::value_type, void>) + { + return nullptr; + } + else { + return round_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 e710a5b4673..df02f409b73 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -203,6 +203,21 @@ struct MulOutputType 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::rint function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct RoundOutputType +{ + 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::sin function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index da088582af1..d2949d14bb2 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -39,6 +39,7 @@ #include "floor.hpp" #include "ln.hpp" #include "mul.hpp" +#include "round.hpp" #include "sin.hpp" #include "sqr.hpp" #include "sqrt.hpp" @@ -60,6 +61,7 @@ static unary_impl_fn_ptr_t floor_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t conj_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t mul_dispatch_vector[dpctl_td_ns::num_types]; +static unary_impl_fn_ptr_t round_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqr_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types]; @@ -301,6 +303,34 @@ PYBIND11_MODULE(_vm_impl, m) py::arg("dst")); } + // UnaryUfunc: ==== Round(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + round_dispatch_vector); + + auto round_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + round_dispatch_vector); + }; + m.def("_round", round_pyapi, + "Call `rint` function from OneMKL VM library to compute " + "the rounded value of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto round_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + round_dispatch_vector); + }; + m.def("_mkl_round_to_call", round_need_to_call_pyapi, + "Check input arguments to answer if `rint` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Sin(x) ==== { vm_ext::init_ufunc_dispatch_vector void (*dpnp_around_default_c)(const void *, void *, const size_t, const int) = dpnp_around_c<_DataType>; -template -DPCTLSyclEventRef (*dpnp_around_ext_c)(DPCTLSyclQueueRef, - const void *, - void *, - const size_t, - const int, - const DPCTLEventVectorRef) = - dpnp_around_c<_DataType>; - template class dpnp_elemwise_absolute_c_kernel; @@ -1184,15 +1175,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_AROUND][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_around_default_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_around_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_cross_default_c}; fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 9a4a556435a..568fc76abc7 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -58,8 +58,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_ARGMIN_EXT DPNP_FN_ARGSORT DPNP_FN_ARGSORT_EXT - DPNP_FN_AROUND - DPNP_FN_AROUND_EXT DPNP_FN_ASTYPE DPNP_FN_ASTYPE_EXT DPNP_FN_CBRT diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index c697e76c682..9fd17b6c91e 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -170,7 +170,7 @@ cpdef utils.dpnp_descriptor dpnp_flatten(utils.dpnp_descriptor x1): cpdef utils.dpnp_descriptor dpnp_init_val(shape, dtype, value): """ - same as dpnp_full(). TODO remove code dumplication + same as dpnp_full(). TODO remove code duplication """ cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index 86579933b01..d2ec12b4c3d 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -38,7 +38,6 @@ and the rest of the library __all__ += [ "dpnp_absolute", "dpnp_arctan2", - "dpnp_around", "dpnp_copysign", "dpnp_cross", "dpnp_cumprod", @@ -72,9 +71,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef, ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_trapz_2in_1out_with_2size_t)(c_dpctl.DPCTLSyclQueueRef, void *, void * , void * , double, size_t, size_t, const c_dpctl.DPCTLEventVectorRef) -ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_around_1in_1out_t)(c_dpctl.DPCTLSyclQueueRef, - const void * , void * , const size_t, const int, - const c_dpctl.DPCTLEventVectorRef) cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1): @@ -120,38 +116,6 @@ cpdef utils.dpnp_descriptor dpnp_arctan2(utils.dpnp_descriptor x1_obj, return call_fptr_2in_1out_strides(DPNP_FN_ARCTAN2_EXT, x1_obj, x2_obj, dtype, out, where, func_name="arctan2") -cpdef utils.dpnp_descriptor dpnp_around(utils.dpnp_descriptor x1, int decimals): - - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype) - - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_AROUND_EXT, param1_type, param1_type) - - x1_obj = x1.get_array() - - # ceate result array with type given by FPTR data - cdef shape_type_c result_shape = x1.shape - cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, - kernel_data.return_type, - None, - device=x1_obj.sycl_device, - usm_type=x1_obj.usm_type, - sycl_queue=x1_obj.sycl_queue) - - result_sycl_queue = result.get_array().sycl_queue - - cdef c_dpctl.SyclQueue q = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef ftpr_custom_around_1in_1out_t func = kernel_data.ptr - - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, x1.get_data(), result.get_data(), x1.size, decimals, NULL) - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return result - - cpdef utils.dpnp_descriptor dpnp_copysign(utils.dpnp_descriptor x1_obj, utils.dpnp_descriptor x2_obj, object dtype=None, diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index dd3a34369e3..733b9086022 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -71,6 +71,7 @@ "dpnp_not_equal", "dpnp_remainder", "dpnp_right_shift", + "dpnp_round", "dpnp_sin", "dpnp_sqrt", "dpnp_square", @@ -1568,6 +1569,57 @@ def dpnp_right_shift(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_round_docstring = """ +round(x, out=None, order='K') +Rounds each element `x_i` of the input array `x` to +the nearest integer-valued number. +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise rounded value. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def _call_round(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_round_to_call(sycl_queue, src, dst): + # call pybind11 extension for round() function from OneMKL VM + return vmi._round(sycl_queue, src, dst, depends) + return ti._round(src, dst, sycl_queue, depends) + + +round_func = UnaryElementwiseFunc( + "round", ti._round_result_type, _call_round, _round_docstring +) + + +def dpnp_round(x, out=None, order="K"): + """ + Invokes round() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for round() function. + """ + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + res_usm = round_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _sign_docstring = """ sign(x, out=None, order="K") diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 7552bcd2130..8dacfcf515f 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -59,6 +59,7 @@ dpnp_multiply, dpnp_negative, dpnp_remainder, + dpnp_round, dpnp_sign, dpnp_subtract, dpnp_trunc, @@ -101,7 +102,7 @@ "power", "prod", "remainder", - "round_", + "round", "sign", "subtract", "sum", @@ -268,45 +269,29 @@ def add( ) -def around(x1, decimals=0, out=None): +def around(x, /, decimals=0, out=None): """ - Evenly round to the given number of decimals. + Round an array to the given number of decimals. For full documentation refer to :obj:`numpy.around`. Limitations ----------- - Parameters `x1` is supported as :class:`dpnp.ndarray`. - Parameters `decimals` and `out` are supported with their default values. + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `decimals` is supported with its default value. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. - Examples + See Also -------- - >>> import dpnp as np - >>> np.around([0.37, 1.64]) - array([0., 2.]) - >>> np.around([0.37, 1.64], decimals=1) - array([0.4, 1.6]) - >>> np.around([.5, 1.5, 2.5, 3.5, 4.5]) # rounds to nearest even value - array([0., 2., 2., 4., 4.]) - >>> np.around([1,2,3,11], decimals=1) # ndarray of ints is returned - array([ 1, 2, 3, 11]) - >>> np.around([1,2,3,11], decimals=-1) - array([ 0, 0, 0, 10]) + :obj:`dpnp.round` : equivalent function; see for details. + Notes + ----- + This function works the same as :obj:`dpnp.round`. """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if out is not None: - pass - elif decimals != 0: - pass - else: - return dpnp_around(x1_desc, decimals).get_pyobj() - - return call_origin(numpy.around, x1, decimals=decimals, out=out) + return round(x, decimals, out) def ceil( @@ -1771,19 +1756,43 @@ def remainder( ) -def round_(a, decimals=0, out=None): +def round(x, decimals=0, out=None): """ - Round an array to the given number of decimals. + Evenly round to the given number of decimals. - For full documentation refer to :obj:`numpy.round_`. + For full documentation refer to :obj:`numpy.round`. - See Also + Limitations + ----------- + Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `decimals` is supported with its default value. + Otherwise the function will be executed sequentially on CPU. + Input array data types are limited by supported DPNP :ref:`Data types`. + + Examples -------- - :obj:`dpnp.around` : equivalent function; see for details. + >>> import dpnp as np + >>> np.round(np.array([0.37, 1.64])) + array([0., 2.]) + >>> np.round(np.array([0.37, 1.64]), decimals=1) + array([0.4, 1.6]) + >>> np.round(np.array([.5, 1.5, 2.5, 3.5, 4.5])) # rounds to nearest even value + array([0., 2., 2., 4., 4.]) + >>> np.round(np.array([1,2,3,11]), decimals=1) # ndarray of ints is returned + array([ 1, 2, 3, 11]) + >>> np.round(np.array([1,2,3,11]), decimals=-1) + array([ 0, 0, 0, 10]) """ - return around(a, decimals, out) + if decimals != 0: + pass + elif dpnp.isscalar(x): + # input has to be an array + pass + else: + return dpnp_round(x, out=out) + return call_origin(numpy.round, x, decimals=decimals, out=out) def sign( diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index a78de781e3b..90c4159c629 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -161,14 +161,6 @@ tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAn tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAndView::test_view_non_contiguous_raise tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_0_{src_order='C'}::test_isinstance_numpy_view_copy_f tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_1_{src_order='F'}::test_isinstance_numpy_view_copy_f -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_uint tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestArrayReduction::test_min_nan @@ -720,43 +712,9 @@ tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[same] tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[full] -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_small -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_fix 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_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 0b95dde9bb5..d8099f933ff 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -156,10 +156,6 @@ tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_2_{shape=(2, 3)}::test_item tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_3_{order='C', shape=(2, 3)}::test_item tests/third_party/cupy/core_tests/test_ndarray_conversion.py::TestNdarrayToBytes_param_4_{order='F', shape=(2, 3)}::test_item -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_float -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_float tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_construction tests/third_party/cupy/creation_tests/test_matrix.py::TestMatrix::test_diag_construction_from_list @@ -295,10 +291,6 @@ tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAn tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAndView::test_view_non_contiguous_raise tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_0_{src_order='C'}::test_isinstance_numpy_view_copy_f tests/third_party/cupy/core_tests/test_ndarray_copy_and_view.py::TestNumPyArrayCopyView_param_1_{src_order='F'}::test_isinstance_numpy_view_copy_f -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_0_{decimals=-3}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_1_{decimals=-2}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_2_{decimals=-1}::test_round_halfway_uint -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRoundHalfway_param_3_{decimals=0}::test_round_halfway_uint tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestArrayReduction::test_ptp_all tests/third_party/cupy/core_tests/test_ndarray_reduction.py::TestArrayReduction::test_ptp_all_keepdims @@ -864,43 +856,9 @@ tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[same] tests/third_party/cupy/math_tests/test_misc.py::TestConvolve::test_convolve_diff_types[full] -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_0_{value=(14, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_1_{value=(15, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_2_{value=(16, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_3_{value=(14.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_4_{value=(15.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_5_{value=(16.0, -1)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_negative2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_positive2 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_small -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_large -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_small -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_fix 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_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/skipped_tests_gpu_no_fp64.tbl b/tests/skipped_tests_gpu_no_fp64.tbl index 31c4b499cd0..15b28cd06d5 100644 --- a/tests/skipped_tests_gpu_no_fp64.tbl +++ b/tests/skipped_tests_gpu_no_fp64.tbl @@ -497,7 +497,6 @@ tests/test_umath.py::TestArctan2::test_invalid_shape[(2,2)] tests/test_umath.py::TestSqrt::test_sqrt_complex[complex64] -tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_2_{decimals=0}::test_round tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_0_{decimals=-2}::test_round_out tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_1_{decimals=-1}::test_round_out tests/third_party/cupy/core_tests/test_ndarray_math.py::TestRound_param_2_{decimals=0}::test_round_out @@ -1192,16 +1191,23 @@ tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_para tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_0_{decimals=-2}::test_round_out tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_1_{decimals=-1}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_2_{decimals=0}::test_round tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_2_{decimals=0}::test_round_out tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_3_{decimals=1}::test_round_out tests/third_party/cupy/math_tests/test_rounding.py::TestRound_param_4_{decimals=2}::test_round_out -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_negative1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_6_{value=(1.4, 0)}::test_around_positive1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_negative1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_7_{value=(1.5, 0)}::test_around_positive1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_negative1 -tests/third_party/cupy/math_tests/test_rounding.py::TestRoundBorder_param_8_{value=(1.6, 0)}::test_around_positive1 +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_large +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_0_{decimals=-100}::test_round_small +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_large +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_1_{decimals=-99}::test_round_small +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_large +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_2_{decimals=-90}::test_round_small +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_large +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_3_{decimals=0}::test_round_small +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_large +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_4_{decimals=90}::test_round_small +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_large +tests/third_party/cupy/math_tests/test_rounding.py::TestRoundExtreme_param_5_{decimals=99}::test_round_small +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_explog.py::TestExplog::test_exp tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_exp2 diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 2b5c98083e2..a45aeff5941 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -137,12 +137,7 @@ def _test_mathematical(self, name, dtype, lhs, rhs): else: result = getattr(dpnp, name)(a_dpnp, b_dpnp) expected = getattr(numpy, name)(a_np, b_np) - if ( - 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()) @@ -174,14 +169,17 @@ def test_divide(self, dtype, lhs, rhs): def test_fmod(self, dtype, lhs, rhs): if dtype == None and rhs == 0.3 and not has_support_aspect64(): """ - Due to accuracy reason NumPy behaves differently, when: + Due to accuracy reason, the results are different for `float32` and `float64` >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float32), 0.3) array([0.29999995], dtype=float32) - while numpy with float64 returns something around zero which is aligned with dpnp: + >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float64), 0.3) array([9.53674318e-08]) + On a gpu without support for `float64`, dpnp produces results similar to the second one. """ - pytest.skip("missaligned between numpy results") + pytest.skip( + "Due to accuracy reason, the results are different for numpy and dpnp" + ) self._test_mathematical("fmod", dtype, lhs, rhs) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) @@ -219,6 +217,23 @@ def test_multiply(self, dtype, lhs, rhs): @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) def test_remainder(self, dtype, lhs, rhs): + if ( + dtype in [dpnp.int32, dpnp.int64, None] + and rhs == 0.3 + and not has_support_aspect64() + ): + """ + Due to accuracy reason, the results are different for `float32` and `float64` + >>> numpy.remainder(numpy.array([6, 3], dtype='i4'), 0.3, dtype='f8') + array([2.22044605e-16, 1.11022302e-16]) + + >>> numpy.remainder(numpy.array([6, 3], dtype='i4'), 0.3, dtype='f4') + usm_ndarray([0.29999977, 0.2999999 ], dtype=float32) + On a gpu without support for `float64`, dpnp produces results similar to the second one. + """ + pytest.skip( + "Due to accuracy reason, the results are different for numpy and dpnp" + ) self._test_mathematical("remainder", dtype, lhs, rhs) @pytest.mark.parametrize("dtype", get_all_dtypes()) diff --git a/tests/third_party/cupy/core_tests/test_ndarray_math.py b/tests/third_party/cupy/core_tests/test_ndarray_math.py index 9b16b76e4c8..cff07e5c104 100644 --- a/tests/third_party/cupy/core_tests/test_ndarray_math.py +++ b/tests/third_party/cupy/core_tests/test_ndarray_math.py @@ -50,6 +50,7 @@ def test_round_out(self, xp): } ) ) +@pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestRoundHalfway(unittest.TestCase): shape = (20,) diff --git a/tests/third_party/cupy/math_tests/test_rounding.py b/tests/third_party/cupy/math_tests/test_rounding.py index e17a48aa9e3..73dead998b9 100644 --- a/tests/third_party/cupy/math_tests/test_rounding.py +++ b/tests/third_party/cupy/math_tests/test_rounding.py @@ -4,6 +4,7 @@ import pytest import dpnp as cupy +from tests.helper import has_support_aspect64 from tests.third_party.cupy import testing @@ -70,8 +71,8 @@ def test_around(self): self.check_unary("around") self.check_unary_complex("around") - def test_round_(self): - self.check_unary("round_") + def test_round(self): + self.check_unary("round") self.check_unary_complex("around") @@ -115,6 +116,7 @@ def test_round_out(self, xp): } ) ) +@pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestRoundExtreme(unittest.TestCase): shape = (20,) @@ -150,23 +152,23 @@ def test_round_small(self, xp, dtype): ) @pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestRoundBorder(unittest.TestCase): - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_positive1(self, xp): a, decimals = self.value return xp.around(a, decimals) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_positive2(self, xp): a, decimals = self.value a = xp.asarray(a) return xp.around(a, decimals) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_negative1(self, xp): a, decimals = self.value return xp.around(-a, decimals) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=has_support_aspect64()) def test_around_negative2(self, xp): a, decimals = self.value a = xp.asarray(a)