FULL add kernel (#558)

* FULL add kernel

Author: densmirn

dpnp/backend/CMakeLists.txt

@@ -164,6 +164,7 @@ project(dpnp_project
 # Building logic...
 # -----------------------------------------------------------------------------------------------
 set(DPNP_SRC
+    kernels/dpnp_krnl_arraycreation.cpp
     kernels/dpnp_krnl_bitwise.cpp
     kernels/dpnp_krnl_common.cpp
     kernels/dpnp_krnl_elemwise.cpp

dpnp/backend/include/dpnp_iface.hpp

@@ -118,6 +118,17 @@ void dpnp_memory_memcpy_c(void* dst, const void* src, size_t size_in_bytes);
 template <typename _DataType>
 INP_DLLEXPORT void dpnp_arange_c(size_t start, size_t step, void* result1, size_t size);
 
+/**
+ * @ingroup BACKEND_API
+ * @brief Implementation of full function
+ *
+ * @param [in]  array_in  Input one-element array.
+ * @param [out] result    Output array.
+ * @param [in]  size      Number of elements in the output array.
+ */
+template <typename _DataType>
+INP_DLLEXPORT void dpnp_full_c(void* array_in, void* result, const size_t size);
+
 /**
  * @ingroup BACKEND_API
  * @brief Matrix multiplication.

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -102,6 +102,7 @@ enum class DPNPFuncName : size_t
     DPNP_FN_FLOOR,                    /**< Used in numpy.floor() implementation  */
     DPNP_FN_FLOOR_DIVIDE,             /**< Used in numpy.floor_divide() implementation  */
     DPNP_FN_FMOD,                     /**< Used in numpy.fmod() implementation  */
+    DPNP_FN_FULL,                     /**< Used in numpy.full() implementation  */
     DPNP_FN_HYPOT,                    /**< Used in numpy.hypot() implementation  */
     DPNP_FN_INITVAL,                  /**< Used in numpy ones, ones_like, zeros, zeros_like, full, full_like impl  */
     DPNP_FN_INV,                      /**< Used in numpy.linalg.inv() implementation  */

dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp

@@ -0,0 +1,88 @@
+//*****************************************************************************
+// Copyright (c) 2016-2020, 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.
+//*****************************************************************************
+
+#include <iostream>
+
+#include "dpnp_fptr.hpp"
+#include "dpnp_iface.hpp"
+#include "queue_sycl.hpp"
+
+template <typename _KernelNameSpecialization>
+class dpnp_arange_c_kernel;
+
+template <typename _DataType>
+void dpnp_arange_c(size_t start, size_t step, void* result1, size_t size)
+{
+    // parameter `size` used instead `stop` to avoid dependency on array length calculation algorithm
+    // TODO: floating point (and negatives) types from `start` and `step`
+
+    if (!size)
+    {
+        return;
+    }
+
+    cl::sycl::event event;
+
+    _DataType* result = reinterpret_cast<_DataType*>(result1);
+
+    cl::sycl::range<1> gws(size);
+    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
+        size_t i = global_id[0];
+
+        result[i] = start + i * step;
+    };
+
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.parallel_for<class dpnp_arange_c_kernel<_DataType>>(gws, kernel_parallel_for_func);
+    };
+
+    event = DPNP_QUEUE.submit(kernel_func);
+
+    event.wait();
+}
+
+template <typename _KernelNameSpecialization>
+class dpnp_full_c_kernel;
+
+template <typename _DataType>
+void dpnp_full_c(void* array_in, void* result, const size_t size)
+{    
+    dpnp_initval_c<_DataType>(result, array_in, size);
+}
+
+void func_map_init_arraycreation(func_map_t& fmap)
+{
+    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_arange_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_arange_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_arange_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_arange_c<double>};
+
+    fmap[DPNPFuncName::DPNP_FN_FULL][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_full_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_FULL][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_full_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_FULL][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_full_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_FULL][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_full_c<double>};
+
+    return;
+}

dpnp/backend/kernels/dpnp_krnl_elemwise.cpp

@@ -254,47 +254,8 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
 
 #include <dpnp_gen_1arg_1type_tbl.hpp>
 
-template <typename _KernelNameSpecialization>
-class dpnp_arange_c_kernel;
-
-template <typename _DataType>
-void dpnp_arange_c(size_t start, size_t step, void* result1, size_t size)
-{
-    // parameter `size` used instead `stop` to avoid dependency on array length calculation algorithm
-    // TODO: floating point (and negatives) types from `start` and `step`
-
-    if (!size)
-    {
-        return;
-    }
-
-    cl::sycl::event event;
-
-    _DataType* result = reinterpret_cast<_DataType*>(result1);
-
-    cl::sycl::range<1> gws(size);
-    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
-        size_t i = global_id[0];
-
-        result[i] = start + i * step;
-    };
-
-    auto kernel_func = [&](cl::sycl::handler& cgh) {
-        cgh.parallel_for<class dpnp_arange_c_kernel<_DataType>>(gws, kernel_parallel_for_func);
-    };
-
-    event = DPNP_QUEUE.submit(kernel_func);
-
-    event.wait();
-}
-
 static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
 {
-    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_arange_c<int>};
-    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_arange_c<long>};
-    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_arange_c<float>};
-    fmap[DPNPFuncName::DPNP_FN_ARANGE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_arange_c<double>};
-
     fmap[DPNPFuncName::DPNP_FN_CONJIGUATE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_copy_c<int>};
     fmap[DPNPFuncName::DPNP_FN_CONJIGUATE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_copy_c<long>};
     fmap[DPNPFuncName::DPNP_FN_CONJIGUATE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_copy_c<float>};

dpnp/backend/src/dpnp_fptr.hpp

@@ -66,6 +66,7 @@ const DPNPFuncType eft_BOOL = DPNPFuncType::DPNP_FT_BOOL;
 /**
  * FPTR interface initialization functions
  */
+void func_map_init_arraycreation(func_map_t& fmap);
 void func_map_init_bitwise(func_map_t& fmap);
 void func_map_init_elemwise(func_map_t& fmap);
 void func_map_init_fft_func(func_map_t& fmap);

dpnp/backend/src/dpnp_iface_fptr.cpp

@@ -134,6 +134,7 @@ static func_map_t func_map_init()
 {
     func_map_t fmap;
 
+    func_map_init_arraycreation(fmap);
     func_map_init_bitwise(fmap);
     func_map_init_elemwise(fmap);
     func_map_init_fft_func(fmap);

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -75,6 +75,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this na
         DPNP_FN_FLOOR
         DPNP_FN_FLOOR_DIVIDE
         DPNP_FN_FMOD
+        DPNP_FN_FULL
         DPNP_FN_HYPOT
         DPNP_FN_INITVAL
         DPNP_FN_INV

dpnp/dpnp_algo/dpnp_algo_arraycreation.pyx

@@ -43,6 +43,7 @@ from dpnp.dpnp_algo cimport *
 __all__ += [
     "dpnp_copy",
     "dpnp_diag",
+    "dpnp_full",
     "dpnp_geomspace",
     "dpnp_linspace",
     "dpnp_logspace",
@@ -89,6 +90,28 @@ cpdef dparray dpnp_diag(v, k):
     return result
 
 
+cpdef dparray dpnp_full(result_shape, value_in, result_dtype):
+    # Convert string type names (dparray.dtype) to C enum DPNPFuncType
+    cdef DPNPFuncType dtype_in = dpnp_dtype_to_DPNPFuncType(result_dtype)
+
+    # get the FPTR data structure
+    cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_FULL, dtype_in, DPNP_FT_NONE)
+
+    result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
+    # Create single-element input array with type given by FPTR data
+    cdef dparray_shape_type shape_in = (1,)
+    cdef dparray array_in = dparray(shape_in, dtype=result_type)
+    array_in[0] = value_in
+    # Create result array with type given by FPTR data
+    cdef dparray result = dparray(result_shape, dtype=result_type)
+
+    cdef fptr_1in_1out_t func = <fptr_1in_1out_t > kernel_data.ptr
+    # Call FPTR function
+    func(array_in.get_data(), result.get_data(), result.size)
+
+    return result
+
+
 cpdef dparray dpnp_geomspace(start, stop, num, endpoint, dtype, axis):
     cdef dparray result = dparray(num, dtype=dtype)
 

dpnp/dpnp_iface_arraycreation.py

@@ -611,16 +611,16 @@ def full(shape, fill_value, dtype=None, order='C'):
     [10, 10, 10, 10]
 
     """
-
-    if (not use_origin_backend()):
+    if not use_origin_backend():
         if order not in ('C', 'c', None):
-            checker_throw_value_error("full", "order", order, 'C')
-
-        _dtype = dtype if dtype is not None else dpnp.dtype(type(fill_value))
+            pass
+        else:
+            if dtype is None:
+                dtype = numpy.array(fill_value).dtype.type
 
-        return dpnp_init_val(shape, _dtype, fill_value)
+            return dpnp_full(shape, fill_value, dtype)
 
-    return numpy.full(shape, fill_value, dtype, order)
+    return call_origin(numpy.full, shape, fill_value, dtype, order)
 
 
 # numpy.full_like(a, fill_value, dtype=None, order='K', subok=True, shape=None)