Kernel abs

dpnp/backend.pxd

@@ -31,6 +31,7 @@ from dpnp.dparray cimport dparray, dparray_shape_type
 
 cdef extern from "backend/backend_iface_fptr.hpp" namespace "DPNPFuncName":  # need this namespace for Enum import
     cdef enum DPNPFuncName "DPNPFuncName":
+        DPNP_FN_ABSOLUTE
         DPNP_FN_ADD
         DPNP_FN_ARCCOS
         DPNP_FN_ARCCOSH

dpnp/backend/backend_iface.hpp

@@ -97,6 +97,26 @@ template <typename _DataType>
 INP_DLLEXPORT void
     custom_blas_gemm_c(void* array1, void* array2, void* result1, size_t size_m, size_t size_n, size_t size_k);
 
+
+/**
+ * @ingroup BACKEND_API
+ * @brief absolute function.
+ *
+ * @param [in]  array1_in    Input array.
+ *
+ * @param [in]  input_shape  Input shape.
+ *
+ * @param [out] result1      Output array.
+ *
+ * @param [in]  size         Number of elements in input arrays.
+ *
+ */
+template <typename _DataType>
+INP_DLLEXPORT void custom_elemwise_absolute_c(void* array1_in,
+                                               const std::vector<long>& input_shape,
+                                               void* result1,
+                                               size_t size);
+
 /**
  * @ingroup BACKEND_API
  * @brief Custom implementation of dot function

dpnp/backend/backend_iface_fptr.cpp

@@ -143,6 +143,11 @@ static func_map_t func_map_init()
     const DPNPFuncType eft_DBL = DPNPFuncType::DPNP_FT_DOUBLE;
     func_map_t fmap;
 
+    fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_INT][eft_INT] = {eft_INT, (void*)custom_elemwise_absolute_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_LNG][eft_LNG] = {eft_LNG, (void*)custom_elemwise_absolute_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_FLT][eft_FLT] = {eft_FLT, (void*)custom_elemwise_absolute_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_DBL][eft_DBL] = {eft_DBL, (void*)custom_elemwise_absolute_c<double>};
+
     fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_INT] = {eft_INT, (void*)custom_elemwise_add_c<int, int, int>};
     fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_LNG] = {eft_LNG, (void*)custom_elemwise_add_c<int, long, long>};
     fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_FLT] = {eft_DBL, (void*)custom_elemwise_add_c<int, float, double>};

dpnp/backend/backend_iface_fptr.hpp

@@ -59,7 +59,8 @@
  */
 enum class DPNPFuncName : size_t
 {
-    DPNP_FN_NONE,      /**< Very first element of the enumeration */
+    DPNP_FN_NONE,     /**< Very first element of the enumeration */
+    DPNP_FN_ABSOLUTE,      /**< Used in numpy.absolute() implementation  */
     DPNP_FN_ADD,       /**< Used in numpy.add() implementation  */
     DPNP_FN_ARCCOS,    /**< Used in numpy.arccos() implementation  */
     DPNP_FN_ARCCOSH,   /**< Used in numpy.arccosh() implementation  */

dpnp/backend/custom_kernels_mathematical.cpp

@@ -0,0 +1,97 @@
+//*****************************************************************************
+// 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 <cmath>
+#include <iostream>
+#include <mkl_blas_sycl.hpp>
+#include <vector>
+
+#include <backend_iface.hpp>
+#include "backend_utils.hpp"
+#include "queue_sycl.hpp"
+
+template <typename _KernelNameSpecialization>
+class custom_elemwise_absolute_c_kernel;
+
+template <typename _DataType>
+void custom_elemwise_absolute_c(void* array1_in,
+                                 const std::vector<long>& input_shape,
+                                 void* result1,
+                                 size_t size)
+{
+    if (!size)
+    {
+        return;
+    }
+
+    cl::sycl::event event;
+    _DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
+    _DataType* result = reinterpret_cast<_DataType*>(result1);
+
+    const size_t input_shape_size = input_shape.size();
+    size_t* input_offset_shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(input_shape_size * sizeof(long)));
+    size_t* result_offset_shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(input_shape_size * sizeof(long)));
+
+    cl::sycl::range<1> gws(size);
+    event = DPNP_QUEUE.submit([&](cl::sycl::handler& cgh) {
+            cgh.parallel_for<class custom_elemwise_absolute_c_kernel<_DataType> >(
+                gws,
+                [=](cl::sycl::id<1> global_id)
+            {
+                const size_t idx = global_id[0];
+
+                if (array1[idx] >= 0) {
+                    result[idx] = array1[idx];
+                }
+                else{
+                    result[idx] = -1 * array1[idx];
+                }
+
+            }); // parallel_for
+    });         // queue.submit
+
+    event.wait();
+
+    free(input_offset_shape, DPNP_QUEUE);
+    free(result_offset_shape, DPNP_QUEUE);
+
+}
+
+template void custom_elemwise_absolute_c<double>(void* array1_in,
+                                                  const std::vector<long>& input_shape,
+                                                  void* result1,
+                                                  size_t size);
+template void custom_elemwise_absolute_c<float>(void* array1_in,
+                                                 const std::vector<long>& input_shape,
+                                                 void* result1,
+                                                 size_t size);
+template void custom_elemwise_absolute_c<long>(void* array1_in,
+                                                const std::vector<long>& input_shape,
+                                                void* result1,
+                                                size_t size);
+template void custom_elemwise_absolute_c<int>(void* array1_in,
+                                               const std::vector<long>& input_shape,
+                                               void* result1,
+                                               size_t size);

dpnp/backend_mathematical.pyx

@@ -63,54 +63,29 @@ __all__ += [
 ]
 
 
+ctypedef void(*fptr_custom_elemwise_absolute_1in_1out_t)(void *, dparray_shape_type &,
+                                                          void *, size_t)
+
+
 cpdef dparray dpnp_absolute(dparray input):
-    cdef dparray_shape_type shape_input = input.shape
-    cdef long size_input = input.size
-    output_shape = dparray(len(shape_input), dtype=numpy.int64)
-    for id, shape_ in enumerate(shape_input):
-        output_shape[id] = shape_
-    cdef long prod = 1
-    for i in range(len(output_shape)):
-        if output_shape[i] != 0:
-            prod *= output_shape[i]
-    result_array = [None] * prod
-    input_shape_offsets = [None] * len(shape_input)
-    acc = 1
-    for i in range(len(shape_input)):
-        ind = len(shape_input) - 1 - i
-        input_shape_offsets[ind] = acc
-        acc *= shape_input[ind]
-    output_shape_offsets = [None] * len(shape_input)
-    acc = 1
-    for i in range(len(output_shape)):
-        ind = len(output_shape) - 1 - i
-        output_shape_offsets[ind] = acc
-        acc *= output_shape[ind]
-        result_offsets = input_shape_offsets[:]  # need copy. not a reference
-
-    for source_idx in range(size_input):
-
-        # reconstruct x,y,z from linear source_idx
-        xyz = []
-        remainder = source_idx
-        for i in input_shape_offsets:
-            quotient, remainder = divmod(remainder, i)
-            xyz.append(quotient)
-
-        result_indexes = []
-        for idx, offset in enumerate(xyz):
-            result_indexes.append(offset)
-
-        result_offset = 0
-        for i, result_indexes_val in enumerate(result_indexes):
-            result_offset += (output_shape_offsets[i] * result_indexes_val)
-
-        input_elem = input.item(source_idx)
-        result_array[result_offset] = input_elem if input_elem >= 0 else -1 * input_elem
-
-    dpnp_array = dpnp.array(result_array, dtype=input.dtype)
-    dpnp_result_array = dpnp_array.reshape(output_shape)
-    return dpnp_result_array
+    cdef dparray_shape_type input_shape = input.shape
+    cdef size_t input_shape_size = input.ndim
+
+    # convert string type names (dparray.dtype) to C enum DPNPFuncType
+    cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(input.dtype)
+
+    # get the FPTR data structure
+    cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_ABSOLUTE, param1_type, param1_type)
+
+    result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
+    # ceate result array with type given by FPTR data
+    cdef dparray result = dparray(input_shape, dtype=result_type)
+
+    cdef fptr_custom_elemwise_absolute_1in_1out_t func = <fptr_custom_elemwise_absolute_1in_1out_t > kernel_data.ptr
+    # call FPTR function
+    func(input.get_data(), input_shape, result.get_data(), input.size)
+
+    return result
 
 
 cpdef dparray dpnp_add(dparray x1, dparray x2):

setup.py

@@ -251,6 +251,7 @@
                 "dpnp/backend/custom_kernels.cpp",
                 "dpnp/backend/custom_kernels_elemwise.cpp",
                 "dpnp/backend/custom_kernels_manipulation.cpp",
+                "dpnp/backend/custom_kernels_mathematical.cpp",
                 "dpnp/backend/custom_kernels_reduction.cpp",
                 "dpnp/backend/custom_kernels_searching.cpp",
                 "dpnp/backend/custom_kernels_sorting.cpp",