implement dpnp.cbrt, dpnp.exp2, dpnp.copysign, dpnp.rsqrt

dpnp/backend/extensions/vm/cbrt.hpp

@@ -0,0 +1,79 @@
+//*****************************************************************************
+// 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 cbrt_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);
+    using resTy = typename types::CbrtOutputType<T>::value_type;
+    resTy *y = reinterpret_cast<resTy *>(out_y);
+
+    return mkl_vm::cbrt(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 CbrtContigFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_same_v<
+                          typename types::CbrtOutputType<T>::value_type, void>)
+        {
+            return nullptr;
+        }
+        else {
+            return cbrt_contig_impl<T>;
+        }
+    }
+};
+} // namespace vm
+} // namespace ext
+} // namespace backend
+} // namespace dpnp

dpnp/backend/extensions/vm/exp2.hpp

@@ -0,0 +1,79 @@
+//*****************************************************************************
+// 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 exp2_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);
+    using resTy = typename types::Exp2OutputType<T>::value_type;
+    resTy *y = reinterpret_cast<resTy *>(out_y);
+
+    return mkl_vm::exp2(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 Exp2ContigFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_same_v<
+                          typename types::Exp2OutputType<T>::value_type, void>)
+        {
+            return nullptr;
+        }
+        else {
+            return exp2_contig_impl<T>;
+        }
+    }
+};
+} // namespace vm
+} // namespace ext
+} // namespace backend
+} // namespace dpnp

dpnp/backend/extensions/vm/types_matrix.hpp

@@ -202,6 +202,21 @@ struct AtanhOutputType
         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::cbrt<T> function.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct CbrtOutputType
+{
+    using value_type = typename std::disjunction<
+        dpctl_td_ns::TypeMapResultEntry<T, double>,
+        dpctl_td_ns::TypeMapResultEntry<T, 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::ceil<T> function.
@@ -308,6 +323,21 @@ struct ExpOutputType
         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::exp2<T> function.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct Exp2OutputType
+{
+    using value_type = typename std::disjunction<
+        dpctl_td_ns::TypeMapResultEntry<T, double>,
+        dpctl_td_ns::TypeMapResultEntry<T, 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::expm1<T> function.

dpnp/backend/extensions/vm/vm_py.cpp

@@ -39,13 +39,15 @@
 #include "atan.hpp"
 #include "atan2.hpp"
 #include "atanh.hpp"
+#include "cbrt.hpp"
 #include "ceil.hpp"
 #include "common.hpp"
 #include "conj.hpp"
 #include "cos.hpp"
 #include "cosh.hpp"
 #include "div.hpp"
 #include "exp.hpp"
+#include "exp2.hpp"
 #include "expm1.hpp"
 #include "floor.hpp"
 #include "hypot.hpp"
@@ -81,12 +83,14 @@ static unary_impl_fn_ptr_t asinh_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t atan_dispatch_vector[dpctl_td_ns::num_types];
 static binary_impl_fn_ptr_t atan2_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t atanh_dispatch_vector[dpctl_td_ns::num_types];
+static unary_impl_fn_ptr_t cbrt_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t ceil_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t conj_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 cosh_dispatch_vector[dpctl_td_ns::num_types];
 static binary_impl_fn_ptr_t div_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t exp_dispatch_vector[dpctl_td_ns::num_types];
+static unary_impl_fn_ptr_t exp2_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t expm1_dispatch_vector[dpctl_td_ns::num_types];
 static unary_impl_fn_ptr_t floor_dispatch_vector[dpctl_td_ns::num_types];
 static binary_impl_fn_ptr_t hypot_dispatch_vector[dpctl_td_ns::num_types];
@@ -366,6 +370,34 @@ PYBIND11_MODULE(_vm_impl, m)
               py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
     }
 
+    // UnaryUfunc: ==== Cbrt(x) ====
+    {
+        vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,
+                                           vm_ext::CbrtContigFactory>(
+            cbrt_dispatch_vector);
+
+        auto cbrt_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst,
+                              const event_vecT &depends = {}) {
+            return vm_ext::unary_ufunc(exec_q, src, dst, depends,
+                                       cbrt_dispatch_vector);
+        };
+        m.def("_cbrt", cbrt_pyapi,
+              "Call `cbrt` function from OneMKL VM library to compute "
+              "the element-wise cube root of vector elements",
+              py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),
+              py::arg("depends") = py::list());
+
+        auto cbrt_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src,
+                                           arrayT dst) {
+            return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst,
+                                                    cbrt_dispatch_vector);
+        };
+        m.def("_mkl_cbrt_to_call", cbrt_need_to_call_pyapi,
+              "Check input arguments to answer if `cbrt` function from "
+              "OneMKL VM library can be used",
+              py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
+    }
+
     // UnaryUfunc: ==== Ceil(x) ====
     {
         vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,
@@ -536,6 +568,34 @@ PYBIND11_MODULE(_vm_impl, m)
               py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
     }
 
+    // UnaryUfunc: ==== exp2(x) ====
+    {
+        vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,
+                                           vm_ext::Exp2ContigFactory>(
+            exp2_dispatch_vector);
+
+        auto exp2_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst,
+                              const event_vecT &depends = {}) {
+            return vm_ext::unary_ufunc(exec_q, src, dst, depends,
+                                       exp2_dispatch_vector);
+        };
+        m.def("_exp2", exp2_pyapi,
+              "Call `exp2` function from OneMKL VM library to compute "
+              "the element-wise base-2 exponential of vector elements",
+              py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),
+              py::arg("depends") = py::list());
+
+        auto exp2_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src,
+                                           arrayT dst) {
+            return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst,
+                                                    exp2_dispatch_vector);
+        };
+        m.def("_mkl_exp2_to_call", exp2_need_to_call_pyapi,
+              "Check input arguments to answer if `exp2` function from "
+              "OneMKL VM library can be used",
+              py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
+    }
+
     // UnaryUfunc: ==== expm1(x) ====
     {
         vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -86,10 +86,8 @@ enum class DPNPFuncName : size_t
     DPNP_FN_BITWISE_OR,   /**< Used in numpy.bitwise_or() impl  */
     DPNP_FN_BITWISE_XOR,  /**< Used in numpy.bitwise_xor() impl  */
     DPNP_FN_CBRT,         /**< Used in numpy.cbrt() impl  */
-    DPNP_FN_CBRT_EXT, /**< Used in numpy.cbrt() impl, requires extra parameters
-                       */
-    DPNP_FN_CEIL,     /**< Used in numpy.ceil() impl  */
-    DPNP_FN_CHOLESKY, /**< Used in numpy.linalg.cholesky() impl  */
+    DPNP_FN_CEIL,         /**< Used in numpy.ceil() impl  */
+    DPNP_FN_CHOLESKY,     /**< Used in numpy.linalg.cholesky() impl  */
     DPNP_FN_CHOLESKY_EXT, /**< Used in numpy.linalg.cholesky() impl, requires
                              extra parameters */
     DPNP_FN_CONJUGATE,    /**< Used in numpy.conjugate() impl  */
@@ -100,9 +98,7 @@ enum class DPNPFuncName : size_t
     DPNP_FN_COPY_EXT, /**< Used in numpy.copy() impl, requires extra parameters
                        */
     DPNP_FN_COPYSIGN, /**< Used in numpy.copysign() impl  */
-    DPNP_FN_COPYSIGN_EXT,  /**< Used in numpy.copysign() impl, requires extra
-                              parameters */
-    DPNP_FN_COPYTO,        /**< Used in numpy.copyto() impl  */
+    DPNP_FN_COPYTO,   /**< Used in numpy.copyto() impl  */
     DPNP_FN_COPYTO_EXT,    /**< Used in numpy.copyto() impl, requires extra
                               parameters */
     DPNP_FN_CORRELATE,     /**< Used in numpy.correlate() impl  */
@@ -154,10 +150,8 @@ enum class DPNPFuncName : size_t
     DPNP_FN_EYE,         /**< Used in numpy.eye() impl  */
     DPNP_FN_EXP,         /**< Used in numpy.exp() impl  */
     DPNP_FN_EXP2,        /**< Used in numpy.exp2() impl  */
-    DPNP_FN_EXP2_EXT, /**< Used in numpy.exp2() impl, requires extra parameters
-                       */
-    DPNP_FN_EXPM1,    /**< Used in numpy.expm1() impl  */
-    DPNP_FN_FABS,     /**< Used in numpy.fabs() impl  */
+    DPNP_FN_EXPM1,       /**< Used in numpy.expm1() impl  */
+    DPNP_FN_FABS,        /**< Used in numpy.fabs() impl  */
     DPNP_FN_FABS_EXT, /**< Used in numpy.fabs() impl, requires extra parameters
                        */
     DPNP_FN_FFT_FFT,  /**< Used in numpy.fft.fft() impl  */