Skip to content

Commit 605416c

Browse files
RubtsowaRubtsowa
authored
add abs kernel (#104)
1 parent 80d9fca commit 605416c

File tree

7 files changed

+148
-48
lines changed

7 files changed

+148
-48
lines changed

‎dpnp/backend.pxd‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@ from dpnp.dparray cimport dparray, dparray_shape_type
3131

3232
cdef extern from "backend/backend_iface_fptr.hpp" namespace "DPNPFuncName": # need this namespace for Enum import
3333
cdef enum DPNPFuncName "DPNPFuncName":
34+
DPNP_FN_ABSOLUTE
3435
DPNP_FN_ADD
3536
DPNP_FN_ARCCOS
3637
DPNP_FN_ARCCOSH

‎dpnp/backend/backend_iface.hpp‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -97,6 +97,26 @@ template <typename _DataType>
9797
INP_DLLEXPORT void
9898
custom_blas_gemm_c(void* array1, void* array2, void* result1, size_t size_m, size_t size_n, size_t size_k);
9999

100+
101+
/**
102+
* @ingroup BACKEND_API
103+
* @brief absolute function.
104+
*
105+
* @param [in] array1_in Input array.
106+
*
107+
* @param [in] input_shape Input shape.
108+
*
109+
* @param [out] result1 Output array.
110+
*
111+
* @param [in] size Number of elements in input arrays.
112+
*
113+
*/
114+
template <typename _DataType>
115+
INP_DLLEXPORT void custom_elemwise_absolute_c(void* array1_in,
116+
const std::vector<long>& input_shape,
117+
void* result1,
118+
size_t size);
119+
100120
/**
101121
* @ingroup BACKEND_API
102122
* @brief Custom implementation of dot function

‎dpnp/backend/backend_iface_fptr.cpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -143,6 +143,11 @@ static func_map_t func_map_init()
143143
const DPNPFuncType eft_DBL = DPNPFuncType::DPNP_FT_DOUBLE;
144144
func_map_t fmap;
145145

146+
fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_INT][eft_INT] = {eft_INT, (void*)custom_elemwise_absolute_c<int>};
147+
fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_LNG][eft_LNG] = {eft_LNG, (void*)custom_elemwise_absolute_c<long>};
148+
fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_FLT][eft_FLT] = {eft_FLT, (void*)custom_elemwise_absolute_c<float>};
149+
fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_DBL][eft_DBL] = {eft_DBL, (void*)custom_elemwise_absolute_c<double>};
150+
146151
fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_INT] = {eft_INT, (void*)custom_elemwise_add_c<int, int, int>};
147152
fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_LNG] = {eft_LNG, (void*)custom_elemwise_add_c<int, long, long>};
148153
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‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -59,7 +59,8 @@
5959
*/
6060
enum class DPNPFuncName : size_t
6161
{
62-
DPNP_FN_NONE, /**< Very first element of the enumeration */
62+
DPNP_FN_NONE, /**< Very first element of the enumeration */
63+
DPNP_FN_ABSOLUTE, /**< Used in numpy.absolute() implementation */
6364
DPNP_FN_ADD, /**< Used in numpy.add() implementation */
6465
DPNP_FN_ARCCOS, /**< Used in numpy.arccos() implementation */
6566
DPNP_FN_ARCCOSH, /**< Used in numpy.arccosh() implementation */

‎dpnp/backend/custom_kernels_mathematical.cpp‎

Lines changed: 97 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,97 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2016-2020, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#include <cmath>
27+
#include <iostream>
28+
#include <mkl_blas_sycl.hpp>
29+
#include <vector>
30+
31+
#include <backend_iface.hpp>
32+
#include "backend_utils.hpp"
33+
#include "queue_sycl.hpp"
34+
35+
template <typename _KernelNameSpecialization>
36+
class custom_elemwise_absolute_c_kernel;
37+
38+
template <typename _DataType>
39+
void custom_elemwise_absolute_c(void* array1_in,
40+
const std::vector<long>& input_shape,
41+
void* result1,
42+
size_t size)
43+
{
44+
if (!size)
45+
{
46+
return;
47+
}
48+
49+
cl::sycl::event event;
50+
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
51+
_DataType* result = reinterpret_cast<_DataType*>(result1);
52+
53+
const size_t input_shape_size = input_shape.size();
54+
size_t* input_offset_shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(input_shape_size * sizeof(long)));
55+
size_t* result_offset_shape = reinterpret_cast<size_t*>(dpnp_memory_alloc_c(input_shape_size * sizeof(long)));
56+
57+
cl::sycl::range<1> gws(size);
58+
event = DPNP_QUEUE.submit([&](cl::sycl::handler& cgh) {
59+
cgh.parallel_for<class custom_elemwise_absolute_c_kernel<_DataType> >(
60+
gws,
61+
[=](cl::sycl::id<1> global_id)
62+
{
63+
const size_t idx = global_id[0];
64+
65+
if (array1[idx] >= 0) {
66+
result[idx] = array1[idx];
67+
}
68+
else{
69+
result[idx] = -1 * array1[idx];
70+
}
71+
72+
}); // parallel_for
73+
}); // queue.submit
74+
75+
event.wait();
76+
77+
free(input_offset_shape, DPNP_QUEUE);
78+
free(result_offset_shape, DPNP_QUEUE);
79+
80+
}
81+
82+
template void custom_elemwise_absolute_c<double>(void* array1_in,
83+
const std::vector<long>& input_shape,
84+
void* result1,
85+
size_t size);
86+
template void custom_elemwise_absolute_c<float>(void* array1_in,
87+
const std::vector<long>& input_shape,
88+
void* result1,
89+
size_t size);
90+
template void custom_elemwise_absolute_c<long>(void* array1_in,
91+
const std::vector<long>& input_shape,
92+
void* result1,
93+
size_t size);
94+
template void custom_elemwise_absolute_c<int>(void* array1_in,
95+
const std::vector<long>& input_shape,
96+
void* result1,
97+
size_t size);

‎dpnp/backend_mathematical.pyx‎

Lines changed: 22 additions & 47 deletions
Original file line numberDiff line numberDiff line change
@@ -63,54 +63,29 @@ __all__ += [
6363
]
6464

6565

66+
ctypedef void(*fptr_custom_elemwise_absolute_1in_1out_t)(void *, dparray_shape_type &,
67+
void *, size_t)
68+
69+
6670
cpdef dparray dpnp_absolute(dparray input):
67-
cdef dparray_shape_type shape_input = input.shape
68-
cdef long size_input = input.size
69-
output_shape = dparray(len(shape_input), dtype=numpy.int64)
70-
for id, shape_ in enumerate(shape_input):
71-
output_shape[id] = shape_
72-
cdef long prod = 1
73-
for i in range(len(output_shape)):
74-
if output_shape[i] != 0:
75-
prod *= output_shape[i]
76-
result_array = [None] * prod
77-
input_shape_offsets = [None] * len(shape_input)
78-
acc = 1
79-
for i in range(len(shape_input)):
80-
ind = len(shape_input) - 1 - i
81-
input_shape_offsets[ind] = acc
82-
acc *= shape_input[ind]
83-
output_shape_offsets = [None] * len(shape_input)
84-
acc = 1
85-
for i in range(len(output_shape)):
86-
ind = len(output_shape) - 1 - i
87-
output_shape_offsets[ind] = acc
88-
acc *= output_shape[ind]
89-
result_offsets = input_shape_offsets[:] # need copy. not a reference
90-
91-
for source_idx in range(size_input):
92-
93-
# reconstruct x,y,z from linear source_idx
94-
xyz = []
95-
remainder = source_idx
96-
for i in input_shape_offsets:
97-
quotient, remainder = divmod(remainder, i)
98-
xyz.append(quotient)
99-
100-
result_indexes = []
101-
for idx, offset in enumerate(xyz):
102-
result_indexes.append(offset)
103-
104-
result_offset = 0
105-
for i, result_indexes_val in enumerate(result_indexes):
106-
result_offset += (output_shape_offsets[i] * result_indexes_val)
107-
108-
input_elem = input.item(source_idx)
109-
result_array[result_offset] = input_elem if input_elem >= 0 else -1 * input_elem
110-
111-
dpnp_array = dpnp.array(result_array, dtype=input.dtype)
112-
dpnp_result_array = dpnp_array.reshape(output_shape)
113-
return dpnp_result_array
71+
cdef dparray_shape_type input_shape = input.shape
72+
cdef size_t input_shape_size = input.ndim
73+
74+
# convert string type names (dparray.dtype) to C enum DPNPFuncType
75+
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(input.dtype)
76+
77+
# get the FPTR data structure
78+
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_ABSOLUTE, param1_type, param1_type)
79+
80+
result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
81+
# ceate result array with type given by FPTR data
82+
cdef dparray result = dparray(input_shape, dtype=result_type)
83+
84+
cdef fptr_custom_elemwise_absolute_1in_1out_t func = <fptr_custom_elemwise_absolute_1in_1out_t > kernel_data.ptr
85+
# call FPTR function
86+
func(input.get_data(), input_shape, result.get_data(), input.size)
87+
88+
return result
11489

11590

11691
cpdef dparray dpnp_add(dparray x1, dparray x2):

‎setup.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -251,6 +251,7 @@
251251
"dpnp/backend/custom_kernels.cpp",
252252
"dpnp/backend/custom_kernels_elemwise.cpp",
253253
"dpnp/backend/custom_kernels_manipulation.cpp",
254+
"dpnp/backend/custom_kernels_mathematical.cpp",
254255
"dpnp/backend/custom_kernels_reduction.cpp",
255256
"dpnp/backend/custom_kernels_searching.cpp",
256257
"dpnp/backend/custom_kernels_sorting.cpp",

0 commit comments

Comments
 (0)