add backend at around function (#726)

* add backend for around func

* add limitations

Author: KsanaKozlova

dpnp/backend/include/dpnp_iface.hpp

@@ -677,6 +677,21 @@ INP_DLLEXPORT void dpnp_argmax_c(void* array, void* result, size_t size);
 template <typename _DataType, typename _idx_DataType>
 INP_DLLEXPORT void dpnp_argmin_c(void* array, void* result, size_t size);
 
+/**
+ * @ingroup BACKEND_API
+ * @brief math library implementation of around function
+ *
+ * @param [in]  input_in     Input array with data.
+ * @param [out] result_out   Output array with indeces.
+ * @param [in]  input_size   Number of elements in input arrays.
+ * @param [in]  decimals     Number of decimal places to round. Support only with default value 0.
+ */
+template <typename _DataType>
+INP_DLLEXPORT void dpnp_around_c(const void* input_in,
+                                  void* result_out,
+                                  const size_t input_size,
+                                  const int decimals);
+
 /**
  * @ingroup BACKEND_API
  * @brief math library implementation of std function

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -73,6 +73,7 @@ enum class DPNPFuncName : size_t
     DPNP_FN_ARGMAX,                   /**< Used in numpy.argmax() implementation  */
     DPNP_FN_ARGMIN,                   /**< Used in numpy.argmin() implementation  */
     DPNP_FN_ARGSORT,                  /**< Used in numpy.argsort() implementation  */
+    DPNP_FN_AROUND,                   /**< Used in numpy.around() implementation  */
     DPNP_FN_ASTYPE,                   /**< Used in numpy.astype() implementation  */
     DPNP_FN_BITWISE_AND,              /**< Used in numpy.bitwise_and() implementation  */
     DPNP_FN_BITWISE_OR,               /**< Used in numpy.bitwise_or() implementation  */

dpnp/backend/kernels/dpnp_krnl_mathematical.cpp

@@ -32,6 +32,51 @@
 #include "dpnp_utils.hpp"
 #include "queue_sycl.hpp"
 
+template <typename _KernelNameSpecialization>
+class dpnp_around_c_kernel;
+
+template <typename _DataType>
+void dpnp_around_c(const void* input_in,
+                   void* result_out,
+                   const size_t input_size,
+                   const int decimals)
+{
+    (void)decimals;
+
+    if (!input_size)
+    {
+        return;
+    }
+
+    cl::sycl::event event;
+    _DataType* input = reinterpret_cast<_DataType*>(const_cast<void*>(input_in));
+    _DataType* result = reinterpret_cast<_DataType*>(result_out);
+
+    if constexpr(std::is_same<_DataType, double>::value || std::is_same<_DataType, float>::value)
+    {
+        event = oneapi::mkl::vm::rint(DPNP_QUEUE, input_size, input, result);
+    }
+    else
+    {
+        cl::sycl::range<1> gws(input_size);
+        auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
+            size_t i = global_id[0];
+            {
+                result[i] = std::rint(input[i]);
+            }
+        };
+
+        auto kernel_func = [&](cl::sycl::handler& cgh) {
+            cgh.parallel_for<class dpnp_around_c_kernel<_DataType>>(
+                gws, kernel_parallel_for_func);
+        };
+
+        event = DPNP_QUEUE.submit(kernel_func);
+    }
+
+    event.wait();
+}
+
 template <typename _KernelNameSpecialization>
 class dpnp_elemwise_absolute_c_kernel;
 
@@ -394,6 +439,11 @@ void func_map_init_mathematical(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_elemwise_absolute_c<float>};
     fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_elemwise_absolute_c<double>};
 
+    fmap[DPNPFuncName::DPNP_FN_AROUND][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_around_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_AROUND][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_around_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_AROUND][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_around_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_AROUND][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_around_c<double>};
+
     fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_cross_c<int, int, int>};
     fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_LNG] = {eft_LNG, (void*)dpnp_cross_c<long, int, long>};
     fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_FLT] = {eft_DBL, (void*)dpnp_cross_c<double, int, float>};

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -46,6 +46,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this na
         DPNP_FN_ARGMAX
         DPNP_FN_ARGMIN
         DPNP_FN_ARGSORT
+        DPNP_FN_AROUND
         DPNP_FN_ASTYPE
         DPNP_FN_BITWISE_AND
         DPNP_FN_BITWISE_OR

dpnp/dpnp_algo/dpnp_algo_mathematical.pyx

@@ -80,6 +80,7 @@ __all__ += [
 ctypedef void(*fptr_custom_elemwise_absolute_1in_1out_t)(void * , void * , size_t)
 ctypedef void(*fptr_1in_2out_t)(void * , void * , void * , size_t)
 ctypedef void(*ftpr_custom_trapz_2in_1out_with_2size_t)(void *, void * , void * , double, size_t, size_t)
+ctypedef void(*ftpr_custom_around_1in_1out_t)(const void *, void *, const size_t, const int)
 
 
 cpdef dparray dpnp_absolute(dparray input):
@@ -111,16 +112,19 @@ cpdef dparray dpnp_arctan2(object x1_obj, object x2_obj, object dtype=None, dpar
     return call_fptr_2in_1out(DPNP_FN_ARCTAN2, x1_obj, x2_obj, dtype=dtype, out=out, where=where)
 
 
-cpdef dpnp_around(dparray a, decimals, out):
-    cdef dparray result
+cpdef dpnp_around(dparray x1, int decimals):
 
-    if out is None:
-        result = dparray(a.shape, dtype=a.dtype)
-    else:
-        result = out
+    cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
 
-    for i in range(result.size):
-        result[i] = round(a[i], decimals)
+    cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_AROUND, param1_type, param1_type)
+
+    result_type = dpnp_DPNPFuncType_to_dtype(< size_t > kernel_data.return_type)
+
+    cdef dparray result = dparray(x1.shape, dtype=result_type)
+
+    cdef ftpr_custom_around_1in_1out_t func = <ftpr_custom_around_1in_1out_t > kernel_data.ptr
+
+    func(x1.get_data(), result.get_data(), x1.size, decimals)
 
     return result
 

dpnp/dpnp_iface_mathematical.py

@@ -213,12 +213,19 @@ def add(x1, x2, dtype=None, out=None, where=True, **kwargs):
     return call_origin(numpy.add, x1, x2, dtype=dtype, out=out, where=where, **kwargs)
 
 
-def around(a, decimals=0, out=None):
+def around(x1, decimals=0, out=None):
     """
     Evenly round to the given number of decimals.
 
     For full documentation refer to :obj:`numpy.around`.
 
+    Limitations
+    -----------
+    Parameters ``x1`` is supported as :obj:`dpnp.ndarray`.
+    Parameters ``decimals`` and ``out`` are supported with their default values.
+    Otherwise the functions will be executed sequentially on CPU.
+    Input array data types are limited by supported DPNP :ref:`Data types`.
+
     Examples
     --------
     >>> import dpnp as np
@@ -235,13 +242,17 @@ def around(a, decimals=0, out=None):
 
     """
 
-    if not use_origin_backend(a):
-        if not isinstance(a, dparray):
+    if not use_origin_backend(x1):
+        if not isinstance(x1, dparray):
+            pass
+        elif out is not None:
+            pass
+        elif decimals != 0:
             pass
         else:
-            return dpnp_around(a, decimals, out)
+            return dpnp_around(x1, decimals)
 
-    return call_origin(numpy.around, a, decimals, out)
+    return call_origin(numpy.around, x1, decimals=decimals, out=out)
 
 
 def ceil(x1, **kwargs):