use_dpctl_round_func_in_dpnp

Author: vtavana

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 <CL/sycl.hpp>
+
+#include "common.hpp"
+#include "types_matrix.hpp"
+
+namespace dpnp
+{
+namespace backend
+{
+namespace ext
+{
+namespace vm
+{
+template <typename T>
+sycl::event round_contig_impl(sycl::queue exec_q,
+                              const std::int64_t n,
+                              const char *in_a,
+                              char *out_y,
+                              const std::vector<sycl::event> &depends)
+{
+    type_utils::validate_type_for_device<T>(exec_q);
+
+    const T *a = reinterpret_cast<const T *>(in_a);
+    T *y = reinterpret_cast<T *>(out_y);
+
+    return mkl_vm::round(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 <typename fnT, typename T>
+struct RoundContigFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_same_v<
+                          typename types::RoundOutputType<T>::value_type, void>)
+        {
+            return nullptr;
+        }
+        else {
+            return round_contig_impl<T>;
+        }
+    }
+};
+} // namespace vm
+} // namespace ext
+} // namespace backend
+} // namespace dpnp

dpnp/backend/extensions/vm/types_matrix.hpp

@@ -136,6 +136,21 @@ struct LnOutputType
         dpctl_td_ns::DefaultResultEntry<void>>::result_type;
 };
 
+/**
+ * @brief A factory to define pairs of supported types for which
+ * MKL VM library provides support in oneapi::mkl::vm::round<T> function.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct RoundOutputType
+{
+    using value_type = typename std::disjunction<
+        dpctl_td_ns::TypeMapResultEntry<T, double, double>,
+        dpctl_td_ns::TypeMapResultEntry<T, float, float>,
+        dpctl_td_ns::DefaultResultEntry<void>>::result_type;
+};
+
 /**
  * @brief A factory to define pairs of supported types for which
  * MKL VM library provides support in oneapi::mkl::vm::sin<T> function.

dpnp/backend/extensions/vm/vm_py.cpp

@@ -36,6 +36,7 @@
 #include "div.hpp"
 #include "floor.hpp"
 #include "ln.hpp"
+#include "round.hpp"
 #include "sin.hpp"
 #include "sqr.hpp"
 #include "sqrt.hpp"
@@ -54,6 +55,7 @@ static unary_impl_fn_ptr_t ceil_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t cos_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t floor_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types];
+static unary_impl_fn_ptr_t 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];
@@ -206,6 +208,34 @@ PYBIND11_MODULE(_vm_impl, m)
               py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
     }
 
+    // UnaryUfunc: ==== Round(x) ====
+    {
+        vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,
+                                           vm_ext::RoundContigFactory>(
+            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 `round` 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 `round` 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<unary_impl_fn_ptr_t,

dpnp/dpnp_algo/dpnp_elementwise_common.py

@@ -63,6 +63,7 @@
     "dpnp_logical_xor",
     "dpnp_multiply",
     "dpnp_not_equal",
+    "dpnp_round",
     "dpnp_sin",
     "dpnp_sqrt",
     "dpnp_square",
@@ -1062,6 +1063,58 @@ def dpnp_not_equal(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 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.
+
+    """
+
+    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)
+
+    # dpctl.tensor only works with usm_ndarray
+    x1_usm = dpnp.get_usm_ndarray(x)
+    out_usm = None if out is None else dpnp.get_usm_ndarray(out)
+
+    func = UnaryElementwiseFunc(
+        "round", ti._round_result_type, _call_round, _round_docstring
+    )
+    res_usm = func(x1_usm, out=out_usm, order=order)
+    return dpnp_array._create_from_usm_ndarray(res_usm)
+
+
 _sin_docstring = """
 sin(x, out=None, order='K')
 Computes sine for each element `x_i` of input array `x`.

dpnp/dpnp_iface_mathematical.py

@@ -372,11 +372,19 @@ def ceil(
 
     Limitations
     -----------
+<<<<<<< HEAD
     Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`.
     Parameters `where`, `dtype`, and `subok` are supported with their default values.
     Keyword arguments `kwargs` are currently unsupported.
     Otherwise the function will be executed sequentially on CPU.
     Input array data types are limited by real-value data types.
+=======
+        Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`.
+        Parameters `where`, `dtype`, and `subok` are supported with their default values.
+        Keyword arguments `kwargs` are currently unsupported.
+        Otherwise the functions will be executed sequentially on CPU.
+        Input array data types are limited by real-value data types.
+>>>>>>> use_dpctl_rounding_funcs_in_dpnp
 
     See Also
     --------
@@ -821,6 +829,7 @@ def floor(
 
     Limitations
     -----------
+<<<<<<< HEAD
     Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`.
     Parameters `where`, `dtype`, and `subok` are supported with their default values.
     Keyword arguments `kwargs` are currently unsupported.
@@ -836,6 +845,23 @@ def floor(
     -----
     Some spreadsheet programs calculate the "floor-towards-zero", in other words floor(-2.5) == -2.
     DPNP instead uses the definition of floor where floor(-2.5) == -3.
+=======
+        Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`.
+        Parameters `where`, `dtype`, and `subok` are supported with their default values.
+        Keyword arguments `kwargs` are currently unsupported.
+        Otherwise the functions will be executed sequentially on CPU.
+        Input array data types are limited by real-value data types.
+
+    See Also
+    --------
+        :obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
+        :obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
+
+    Notes
+    -----
+        Some spreadsheet programs calculate the "floor-towards-zero", in other words floor(-2.5) == -2.
+        DPNP instead uses the definition of floor where floor(-2.5) == -3.
+>>>>>>> use_dpctl_rounding_funcs_in_dpnp
 
     Examples
     --------
@@ -2028,11 +2054,19 @@ def trunc(
 
     Limitations
     -----------
+<<<<<<< HEAD
     Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`.
     Parameters `where`, `dtype`, and `subok` are supported with their default values.
     Keyword arguments `kwargs` are currently unsupported.
     Otherwise the function will be executed sequentially on CPU.
     Input array data types are limited by real-value data types.
+=======
+        Parameter `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`.
+        Parameters `where`, `dtype`, and `subok` are supported with their default values.
+        Keyword arguments `kwargs` are currently unsupported.
+        Otherwise the functions will be executed sequentially on CPU.
+        Input array data types are limited by real-value data types.
+>>>>>>> use_dpctl_rounding_funcs_in_dpnp
 
     See Also
     --------