From cf7362d93b906e0057747335c47ba5f748f2bfab Mon Sep 17 00:00:00 2001 From: vtavana <120411540+vtavana@users.noreply.github.com> Date: Tue, 10 Oct 2023 08:26:51 -0500 Subject: [PATCH] re-write dpnp.abs (#1575) * re-write dpnp.abs * address comments --- dpnp/backend/extensions/vm/abs.hpp | 78 ++++++++++++++++ dpnp/backend/extensions/vm/types_matrix.hpp | 17 ++++ dpnp/backend/extensions/vm/vm_py.cpp | 30 ++++++ dpnp/backend/include/dpnp_iface_fptr.hpp | 8 +- .../kernels/dpnp_krnl_mathematical.cpp | 23 ----- dpnp/dpnp_algo/dpnp_algo.pxd | 2 - dpnp/dpnp_algo/dpnp_algo_mathematical.pxi | 39 -------- dpnp/dpnp_algo/dpnp_elementwise_common.py | 58 ++++++++++++ dpnp/dpnp_iface_mathematical.py | 93 +++++++------------ tests/skipped_tests.tbl | 1 - tests/skipped_tests_gpu.tbl | 1 - tests/test_usm_type.py | 1 + 12 files changed, 223 insertions(+), 128 deletions(-) create mode 100644 dpnp/backend/extensions/vm/abs.hpp diff --git a/dpnp/backend/extensions/vm/abs.hpp b/dpnp/backend/extensions/vm/abs.hpp new file mode 100644 index 00000000000..1cae10f70ad --- /dev/null +++ b/dpnp/backend/extensions/vm/abs.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 abs_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::abs(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 AbsContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::AbsOutputType::value_type, void>) + { + return nullptr; + } + else { + return abs_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 3d91c95bfbb..172b74b4e28 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -43,6 +43,23 @@ namespace vm { namespace types { +/** + * @brief A factory to define pairs of supported types for which + * MKL VM library provides support in oneapi::mkl::vm::abs function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct AbsOutputType +{ + using value_type = typename std::disjunction< + // TODO: Add complex type here after updating the dispatching to allow + // output type to be different than input + 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::acos function. diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp index 1cf4fd7d854..0b31901569a 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -30,6 +30,7 @@ #include #include +#include "abs.hpp" #include "acos.hpp" #include "acosh.hpp" #include "add.hpp" @@ -66,6 +67,7 @@ namespace vm_ext = dpnp::backend::ext::vm; using vm_ext::binary_impl_fn_ptr_t; using vm_ext::unary_impl_fn_ptr_t; +static unary_impl_fn_ptr_t abs_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t acos_dispatch_vector[dpctl_td_ns::num_types]; static unary_impl_fn_ptr_t acosh_dispatch_vector[dpctl_td_ns::num_types]; static binary_impl_fn_ptr_t add_dispatch_vector[dpctl_td_ns::num_types]; @@ -99,6 +101,34 @@ PYBIND11_MODULE(_vm_impl, m) using arrayT = dpctl::tensor::usm_ndarray; using event_vecT = std::vector; + // UnaryUfunc: ==== Abs(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + abs_dispatch_vector); + + auto abs_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + abs_dispatch_vector); + }; + m.def("_abs", abs_pyapi, + "Call `abs` function from OneMKL VM library to compute " + "the absolute value of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto abs_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + abs_dispatch_vector); + }; + m.def("_mkl_abs_to_call", abs_need_to_call_pyapi, + "Check input arguments to answer if `abs` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Acos(x) ==== { vm_ext::init_ufunc_dispatch_vector void (*dpnp_elemwise_absolute_default_c)(const void *, void *, size_t) = dpnp_elemwise_absolute_c<_DataType>; -template -DPCTLSyclEventRef (*dpnp_elemwise_absolute_ext_c)(DPCTLSyclQueueRef, - const void *, - void *, - size_t, - const DPCTLEventVectorRef) = - dpnp_elemwise_absolute_c<_DataType_input, _DataType_output>; - template @@ -1151,21 +1143,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_elemwise_absolute_default_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_elemwise_absolute_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_elemwise_absolute_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_elemwise_absolute_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_elemwise_absolute_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE_EXT][eft_C64][eft_C64] = { - eft_FLT, - (void *)dpnp_elemwise_absolute_ext_c, float>}; - fmap[DPNPFuncName::DPNP_FN_ABSOLUTE_EXT][eft_C128][eft_C128] = { - eft_DBL, - (void *)dpnp_elemwise_absolute_ext_c, double>}; - fmap[DPNPFuncName::DPNP_FN_AROUND][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_around_default_c}; fmap[DPNPFuncName::DPNP_FN_AROUND][eft_LNG][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 9770d3abe6e..442209621bd 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -33,8 +33,6 @@ from dpnp.dpnp_utils.dpnp_algo_utils cimport dpnp_descriptor cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this namespace for Enum import cdef enum DPNPFuncName "DPNPFuncName": - DPNP_FN_ABSOLUTE - DPNP_FN_ABSOLUTE_EXT DPNP_FN_ALLCLOSE DPNP_FN_ALLCLOSE_EXT DPNP_FN_ARANGE diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index da5f2ec1040..425b277d330 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -36,7 +36,6 @@ and the rest of the library # NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file __all__ += [ - "dpnp_absolute", "dpnp_copysign", "dpnp_cross", "dpnp_cumprod", @@ -59,9 +58,6 @@ __all__ += [ ] -ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_custom_elemwise_absolute_1in_1out_t)(c_dpctl.DPCTLSyclQueueRef, - void * , void * , size_t, - const c_dpctl.DPCTLEventVectorRef) ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef, void * , void * , void * , size_t, const c_dpctl.DPCTLEventVectorRef) @@ -70,41 +66,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_trapz_2in_1out_with_2size_t)(c_d const c_dpctl.DPCTLEventVectorRef) -cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1): - cdef shape_type_c x1_shape = x1.shape - cdef size_t x1_shape_size = x1.ndim - - # convert string type names (array.dtype) to C enum DPNPFuncType - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype) - - # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_ABSOLUTE_EXT, param1_type, param1_type) - - x1_obj = x1.get_array() - - # ceate result array with type given by FPTR data - cdef utils.dpnp_descriptor result = utils.create_output_descriptor(x1_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 fptr_custom_elemwise_absolute_1in_1out_t func = kernel_data.ptr - # call FPTR function - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, x1.get_data(), result.get_data(), x1.size, 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 69f043b9a8d..629c1e0a4c9 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -42,6 +42,7 @@ __all__ = [ "check_nd_call_func", + "dpnp_abs", "dpnp_acos", "dpnp_acosh", "dpnp_add", @@ -169,6 +170,63 @@ def check_nd_call_func( ) +_abs_docstring = """ +abs(x, out=None, order='K') + +Calculates the absolute value for each element `x_i` of input array `x`. + +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + out ({None, dpnp.ndarray}, optional): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + dpnp.ndarray: + An array containing the element-wise absolute values. + For complex input, the absolute value is its magnitude. + If `x` has a real-valued data type, the returned array has the + same data type as `x`. If `x` has a complex floating-point data type, + the returned array has a real-valued floating-point data type whose + precision matches the precision of `x`. +""" + + +def _call_abs(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_abs_to_call(sycl_queue, src, dst): + # call pybind11 extension for abs() function from OneMKL VM + return vmi._abs(sycl_queue, src, dst, depends) + return ti._abs(src, dst, sycl_queue, depends) + + +abs_func = UnaryElementwiseFunc( + "abs", ti._abs_result_type, _call_abs, _abs_docstring +) + + +def dpnp_abs(x, out=None, order="K"): + """ + Invokes abs() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for abs() 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 = abs_func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _acos_docstring = """ acos(x, out=None, order='K') diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 40a4d1ef73b..650c0478e17 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -50,6 +50,7 @@ from .dpnp_algo import * from .dpnp_algo.dpnp_elementwise_common import ( check_nd_call_func, + dpnp_abs, dpnp_add, dpnp_ceil, dpnp_conj, @@ -126,37 +127,22 @@ ] -def abs(*args, **kwargs): - """ - Calculate the absolute value element-wise. - - For full documentation refer to :obj:`numpy.absolute`. - - Notes - ----- - :obj:`dpnp.abs` is a shorthand for :obj:`dpnp.absolute`. - - Examples - -------- - >>> import dpnp as np - >>> a = np.array([-1.2, 1.2]) - >>> result = np.abs(a) - >>> [x for x in result] - [1.2, 1.2] - - """ - - return dpnp.absolute(*args, **kwargs) - - -def absolute(x, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): +def absolute( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Calculate the absolute value element-wise. For full documentation refer to :obj:`numpy.absolute`. - .. seealso:: :obj:`dpnp.abs` : Calculate the absolute value element-wise. - Returns ------- out : dpnp.ndarray @@ -170,50 +156,43 @@ def absolute(x, /, out=None, *, where=True, dtype=None, subok=True, **kwargs): 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.fabs` : Calculate the absolute value element-wise excluding complex types. + + Notes + ----- + ``dpnp.abs`` is a shorthand for this function. + Examples -------- - >>> import dpnp as dp - >>> a = dp.array([-1.2, 1.2]) - >>> result = dp.absolute(a) - >>> [x for x in result] - [1.2, 1.2] + >>> import dpnp as np + >>> a = np.array([-1.2, 1.2]) + >>> np.absolute(a) + array([1.2, 1.2]) - """ + >>> a = np.array(1.2 + 1j) + >>> np.absolute(a) + array(1.5620499351813308) - if out is not None: - pass - elif where is not True: - pass - elif dtype is not None: - pass - elif subok is not True: - pass - elif dpnp.isscalar(x): - pass - else: - x_desc = dpnp.get_dpnp_descriptor(x, copy_when_nondefault_queue=False) - if x_desc: - if x_desc.dtype == dpnp.bool: - # return a copy of input array "x" - return dpnp.array( - x, - dtype=x.dtype, - sycl_queue=x.sycl_queue, - usm_type=x.usm_type, - ) - return dpnp_absolute(x_desc).get_pyobj() + """ - return call_origin( + return check_nd_call_func( numpy.absolute, + dpnp_abs, x, out=out, where=where, + order=order, dtype=dtype, subok=subok, **kwargs, ) +abs = absolute + + def add( x1, x2, @@ -826,7 +805,7 @@ def fabs(x1, **kwargs): See Also -------- - :obj:`dpnp.abs` : Calculate the absolute value element-wise. + :obj:`dpnp.absolute` : Calculate the absolute value element-wise. Examples -------- @@ -2143,7 +2122,7 @@ def proj( See Also -------- - :obj:`dpnp.abs` : Returns the magnitude of a complex number, element-wise. + :obj:`dpnp.absolute` : Returns the magnitude of a complex number, element-wise. :obj:`dpnp.conj` : Return the complex conjugate, element-wise. Examples diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index e3a7d60bb24..ccf926dc770 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -498,7 +498,6 @@ tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_clip_min_max_none tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_external_clip4 -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_absolute_negative tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_fmax_nan tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_fmin_nan tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index b9da380e113..e36491861a4 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -633,7 +633,6 @@ tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_clip_min_max_none tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_external_clip4 -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_absolute_negative tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_fmax_nan tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_fmin_nan tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index dcd29b7af9e..83c0be25473 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -325,6 +325,7 @@ def test_meshgrid(usm_type_x, usm_type_y): @pytest.mark.parametrize( "func,data", [ + pytest.param("abs", [-1.2, 1.2]), pytest.param("arccos", [-0.5, 0.0, 0.5]), pytest.param("arccosh", [1.5, 3.5, 5.0]), pytest.param("arcsin", [-0.5, 0.0, 0.5]),