Native repeat (#644)

* add kernel dpnp.repeat

Author: Rubtsowa

dpnp/backend/include/dpnp_iface.hpp

@@ -821,6 +821,18 @@ INP_DLLEXPORT void dpnp_ones_like_c(void* result, size_t size);
 template <typename _DataType_input1, typename _DataType_input2, typename _DataType_output>
 INP_DLLEXPORT void dpnp_remainder_c(void* array1_in, void* array2_in, void* result1, size_t size);
 
+/**
+ * @ingroup BACKEND_API
+ * @brief repeat elements of an array.
+ *
+ * @param [in]  array_in    Input array.
+ * @param [out] result      Output array.
+ * @param [in]  repeats      The number of repetitions for each element.
+ * @param [in]  size         Number of elements in input arrays.
+ */
+template <typename _DataType>
+INP_DLLEXPORT void dpnp_repeat_c(const void* array_in, void* result, const size_t repeats, const size_t size);
+
 /**
  * @ingroup BACKEND_API
  * @brief copyto function.

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -144,6 +144,7 @@ enum class DPNPFuncName : size_t
     DPNP_FN_RADIANS,                  /**< Used in numpy.radians() implementation  */
     DPNP_FN_REMAINDER,                /**< Used in numpy.remainder() implementation  */
     DPNP_FN_RECIP,                    /**< Used in numpy.recip() implementation  */
+    DPNP_FN_REPEAT,                   /**< Used in numpy.repeat() implementation  */
     DPNP_FN_RIGHT_SHIFT,              /**< Used in numpy.right_shift() implementation  */
     DPNP_FN_RNG_BETA,                 /**< Used in numpy.random.beta() implementation  */
     DPNP_FN_RNG_BINOMIAL,             /**< Used in numpy.random.binomial() implementation  */

dpnp/backend/kernels/dpnp_krnl_manipulation.cpp

@@ -39,6 +39,43 @@ void dpnp_copyto_c(void* destination, void* source, const size_t size)
     __dpnp_copyto_c<_DataType_src, _DataType_dst>(source, destination, size);
 }
 
+template <typename _DataType>
+class dpnp_repeat_c_kernel;
+
+template <typename _DataType>
+void dpnp_repeat_c(const void* array1_in, void* result1, const size_t repeats, const size_t size)
+{
+    cl::sycl::event event;
+
+    const _DataType* array_in = reinterpret_cast<const _DataType*>(array1_in);
+    _DataType* result = reinterpret_cast<_DataType*>(result1);
+
+    if (!array_in || !result)
+    {
+        return;
+    }
+
+    if (!size || !repeats)
+    {
+        return;
+    }
+
+    cl::sycl::range<2> gws(size, repeats);
+    auto kernel_parallel_for_func = [=](cl::sycl::id<2> global_id) {
+        size_t idx1 = global_id[0];
+        size_t idx2 = global_id[1];
+        result[(idx1 * repeats) + idx2] = array_in[idx1];
+    };
+
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.parallel_for<class dpnp_repeat_c_kernel<_DataType>>(gws, kernel_parallel_for_func);
+    };
+
+    event = DPNP_QUEUE.submit(kernel_func);
+
+    event.wait();
+}
+
 template <typename _KernelNameSpecialization>
 class dpnp_elemwise_transpose_c_kernel;
 
@@ -134,6 +171,11 @@ void func_map_init_manipulation(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_COPYTO][eft_C128][eft_C128] = {
         eft_C128, (void*)dpnp_copyto_c<std::complex<double>, std::complex<double>>};
 
+    fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_repeat_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_repeat_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_repeat_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_REPEAT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_repeat_c<double>};
+
     fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_elemwise_transpose_c<int>};
     fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_elemwise_transpose_c<long>};
     fmap[DPNPFuncName::DPNP_FN_TRANSPOSE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_elemwise_transpose_c<float>};

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -116,6 +116,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this na
         DPNP_FN_RADIANS
         DPNP_FN_REMAINDER
         DPNP_FN_RECIP
+        DPNP_FN_REPEAT
         DPNP_FN_RIGHT_SHIFT
         DPNP_FN_RNG_BETA
         DPNP_FN_RNG_BINOMIAL

dpnp/dpnp_algo/dpnp_algo_manipulation.pyx

@@ -50,6 +50,7 @@ __all__ += [
 # C function pointer to the C library template functions
 ctypedef void(*fptr_custom_elemwise_transpose_1in_1out_t)(void * , size_t * , size_t * ,
                                                           size_t * , size_t, void * , size_t)
+ctypedef void(*fptr_dpnp_repeat_t)(const void *, void * , const size_t , const size_t)
 
 
 cpdef dparray dpnp_atleast_2d(dparray arr):
@@ -129,12 +130,16 @@ cpdef dparray dpnp_expand_dims(dparray in_array, axis):
 
 
 cpdef dparray dpnp_repeat(dparray array1, repeats, axes=None):
+    cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype)
+
+    cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_REPEAT, param1_type, param1_type)
+
+    result_type = dpnp_DPNPFuncType_to_dtype(< size_t > kernel_data.return_type)
     cdef long new_size = array1.size * repeats
     cdef dparray result = dparray((new_size, ), dtype=array1.dtype)
 
-    for idx2 in range(array1.size):
-        for idx1 in range(repeats):
-            result[(idx2 * repeats) + idx1] = array1[idx2]
+    cdef fptr_dpnp_repeat_t func = <fptr_dpnp_repeat_t > kernel_data.ptr
+    func(array1.get_data(), result.get_data(), repeats, array1.size)
 
     return result
 

dpnp/dpnp_iface_manipulation.py

@@ -403,30 +403,20 @@ def repeat(x1, repeats, axis=None):
 
     """
 
-    is_x1_dparray = isinstance(x1, dparray)
-
-    if (not use_origin_backend(x1) and is_x1_dparray and (axis is None or axis == 0) and (x1.ndim < 2)):
-
-        repeat_val = repeats
-        if isinstance(repeats, (tuple, list)):
-            if (len(repeats) > 1):
-                checker_throw_value_error("repeat", "len(repeats)", len(repeats), 1)
-
-            repeat_val = repeats[0]
-
-        return dpnp_repeat(x1, repeat_val, axis)
-
-    input1 = dpnp.asnumpy(x1) if is_x1_dparray else x1
-
-    # TODO need to put dparray memory into NumPy call
-    result_numpy = numpy.repeat(input1, repeats, axis=axis)
-    result = result_numpy
-    if isinstance(result, numpy.ndarray):
-        result = dparray(result_numpy.shape, dtype=result_numpy.dtype)
-        for i in range(result.size):
-            result._setitem_scalar(i, result_numpy.item(i))
+    if not use_origin_backend(x1):
+        if not isinstance(x1, dparray):
+            pass
+        elif axis is not None and axis != 0:
+            pass
+        elif x1.ndim >= 2:
+            pass
+        elif not dpnp.isscalar(repeats) and len(repeats) > 1:
+            pass
+        else:
+            repeat_val = repeats if dpnp.isscalar(repeats) else repeats[0]
+            return dpnp_repeat(x1, repeat_val, axis)
 
-    return result
+    return call_origin(numpy.repeat, x1, repeats, axis)
 
 
 def rollaxis(a, axis, start=0):

tests/skipped_tests.tbl

@@ -974,8 +974,6 @@ tests/third_party/cupy/manipulation_tests/test_shape.py::TestShape_param_1_{shap
 tests/third_party/cupy/manipulation_tests/test_shape.py::TestShape_param_1_{shape=()}::test_shape_list
 tests/third_party/cupy/manipulation_tests/test_shape.py::TestShape_param_2_{shape=(4,)}::test_shape
 tests/third_party/cupy/manipulation_tests/test_shape.py::TestShape_param_2_{shape=(4,)}::test_shape_list
-tests/third_party/cupy/manipulation_tests/test_tiling.py::TestRepeat1D_param_3_{axis=None, repeats=[1, 2, 3, 4]}::test_array_repeat
-tests/third_party/cupy/manipulation_tests/test_tiling.py::TestRepeat1D_param_4_{axis=0, repeats=[1, 2, 3, 4]}::test_array_repeat
 tests/third_party/cupy/manipulation_tests/test_tiling.py::TestRepeatRepeatsNdarray::test_func
 tests/third_party/cupy/manipulation_tests/test_tiling.py::TestRepeatRepeatsNdarray::test_method
 tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTileFailure_param_0_{reps=-1}::test_tile_failure

tests/test_manipulation.py

@@ -21,3 +21,14 @@ def test_copyto_dtype(in_obj, out_dtype):
     dpnp.copyto(result, dparr)
 
     numpy.testing.assert_array_equal(result, expected)
+
+
+@pytest.mark.parametrize("arr",
+                         [[], [1, 2, 3, 4], [[1, 2], [3, 4]], [[[1], [2]], [[3], [4]]]],
+                         ids=['[]', '[1, 2, 3, 4]', '[[1, 2], [3, 4]]', '[[[1], [2]], [[3], [4]]]'])
+def test_repeat(arr):
+    a = numpy.array(arr)
+    dpnp_a = dpnp.array(arr)
+    expected = numpy.repeat(a, 2)
+    result = dpnp.repeat(dpnp_a, 2)
+    numpy.testing.assert_array_equal(expected, result)