Skip to content

Commit 28de9c3

Browse files
densmirndensmirn
authored
Implement strides for 29 element-wise functions part 4 (#994)
* Implement strides for 29 element-wise functions part 4
1 parent 57b6052 commit 28de9c3

11 files changed

+370
-185
lines changed

‎dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp

Lines changed: 83 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -47,58 +47,110 @@
4747
/** */ \
4848
/** Function "__name__" executes operator "__operation1__" over each element of the array */ \
4949
/** */ \
50-
/** @param[in] array1 Input array. */ \
51-
/** @param[out] result1 Output array. */ \
52-
/** @param[in] size Number of elements in the input array. */ \
50+
/** @param[out] result_out Output array. */ \
51+
/** @param[in] result_size Output array size. */ \
52+
/** @param[in] result_ndim Number of output array dimensions. */ \
53+
/** @param[in] result_shape Output array shape. */ \
54+
/** @param[in] result_strides Output array strides. */ \
55+
/** @param[in] input1_in Input array 1. */ \
56+
/** @param[in] input1_size Input array 1 size. */ \
57+
/** @param[in] input1_ndim Number of input array 1 dimensions. */ \
58+
/** @param[in] input1_shape Input array 1 shape. */ \
59+
/** @param[in] input1_strides Input array 1 strides. */ \
60+
/** @param[in] where Where condition. */ \
5361
template <typename _DataType_input, typename _DataType_output> \
54-
void __name__(void* array1, void* result1, size_t size);
62+
void __name__(void* result_out, \
63+
const size_t result_size, \
64+
const size_t result_ndim, \
65+
const size_t* result_shape, \
66+
const size_t* result_strides, \
67+
const void* input1_in, \
68+
const size_t input1_size, \
69+
const size_t input1_ndim, \
70+
const size_t* input1_shape, \
71+
const size_t* input1_strides, \
72+
const size_t* where);
5573

5674
#endif
5775

58-
MACRO_1ARG_2TYPES_OP(dpnp_acos_c, cl::sycl::acos(input_elem), oneapi::mkl::vm::acos(DPNP_QUEUE, size, array1, result))
76+
MACRO_1ARG_2TYPES_OP(dpnp_acos_c,
77+
cl::sycl::acos(input_elem),
78+
oneapi::mkl::vm::acos(DPNP_QUEUE, input1_size, input1_data, result))
5979
MACRO_1ARG_2TYPES_OP(dpnp_acosh_c,
6080
cl::sycl::acosh(input_elem),
61-
oneapi::mkl::vm::acosh(DPNP_QUEUE, size, array1, result))
62-
MACRO_1ARG_2TYPES_OP(dpnp_asin_c, cl::sycl::asin(input_elem), oneapi::mkl::vm::asin(DPNP_QUEUE, size, array1, result))
81+
oneapi::mkl::vm::acosh(DPNP_QUEUE, input1_size, input1_data, result))
82+
MACRO_1ARG_2TYPES_OP(dpnp_asin_c,
83+
cl::sycl::asin(input_elem),
84+
oneapi::mkl::vm::asin(DPNP_QUEUE, input1_size, input1_data, result))
6385
MACRO_1ARG_2TYPES_OP(dpnp_asinh_c,
6486
cl::sycl::asinh(input_elem),
65-
oneapi::mkl::vm::asinh(DPNP_QUEUE, size, array1, result))
66-
MACRO_1ARG_2TYPES_OP(dpnp_atan_c, cl::sycl::atan(input_elem), oneapi::mkl::vm::atan(DPNP_QUEUE, size, array1, result))
87+
oneapi::mkl::vm::asinh(DPNP_QUEUE, input1_size, input1_data, result))
88+
MACRO_1ARG_2TYPES_OP(dpnp_atan_c,
89+
cl::sycl::atan(input_elem),
90+
oneapi::mkl::vm::atan(DPNP_QUEUE, input1_size, input1_data, result))
6791
MACRO_1ARG_2TYPES_OP(dpnp_atanh_c,
6892
cl::sycl::atanh(input_elem),
69-
oneapi::mkl::vm::atanh(DPNP_QUEUE, size, array1, result))
70-
MACRO_1ARG_2TYPES_OP(dpnp_cbrt_c, cl::sycl::cbrt(input_elem), oneapi::mkl::vm::cbrt(DPNP_QUEUE, size, array1, result))
71-
MACRO_1ARG_2TYPES_OP(dpnp_ceil_c, cl::sycl::ceil(input_elem), oneapi::mkl::vm::ceil(DPNP_QUEUE, size, array1, result))
72-
MACRO_1ARG_2TYPES_OP(__dpnp_copyto_c, input_elem, DPNP_QUEUE.submit(kernel_func))
73-
MACRO_1ARG_2TYPES_OP(dpnp_cos_c, cl::sycl::cos(input_elem), oneapi::mkl::vm::cos(DPNP_QUEUE, size, array1, result))
74-
MACRO_1ARG_2TYPES_OP(dpnp_cosh_c, cl::sycl::cosh(input_elem), oneapi::mkl::vm::cosh(DPNP_QUEUE, size, array1, result))
93+
oneapi::mkl::vm::atanh(DPNP_QUEUE, input1_size, input1_data, result))
94+
MACRO_1ARG_2TYPES_OP(dpnp_cbrt_c,
95+
cl::sycl::cbrt(input_elem),
96+
oneapi::mkl::vm::cbrt(DPNP_QUEUE, input1_size, input1_data, result))
97+
MACRO_1ARG_2TYPES_OP(dpnp_ceil_c,
98+
cl::sycl::ceil(input_elem),
99+
oneapi::mkl::vm::ceil(DPNP_QUEUE, input1_size, input1_data, result))
100+
MACRO_1ARG_2TYPES_OP(dpnp_copyto_c, input_elem, DPNP_QUEUE.submit(kernel_func))
101+
MACRO_1ARG_2TYPES_OP(dpnp_cos_c,
102+
cl::sycl::cos(input_elem),
103+
oneapi::mkl::vm::cos(DPNP_QUEUE, input1_size, input1_data, result))
104+
MACRO_1ARG_2TYPES_OP(dpnp_cosh_c,
105+
cl::sycl::cosh(input_elem),
106+
oneapi::mkl::vm::cosh(DPNP_QUEUE, input1_size, input1_data, result))
75107
MACRO_1ARG_2TYPES_OP(dpnp_degrees_c, cl::sycl::degrees(input_elem), DPNP_QUEUE.submit(kernel_func))
76-
MACRO_1ARG_2TYPES_OP(dpnp_ediff1d_c, array1[i + 1] - input_elem, DPNP_QUEUE.submit(kernel_func))
77-
MACRO_1ARG_2TYPES_OP(dpnp_exp2_c, cl::sycl::exp2(input_elem), oneapi::mkl::vm::exp2(DPNP_QUEUE, size, array1, result))
78-
MACRO_1ARG_2TYPES_OP(dpnp_exp_c, cl::sycl::exp(input_elem), oneapi::mkl::vm::exp(DPNP_QUEUE, size, array1, result))
108+
MACRO_1ARG_2TYPES_OP(dpnp_ediff1d_c, input1_data[output_id + 1] - input_elem, DPNP_QUEUE.submit(kernel_func))
109+
MACRO_1ARG_2TYPES_OP(dpnp_exp2_c,
110+
cl::sycl::exp2(input_elem),
111+
oneapi::mkl::vm::exp2(DPNP_QUEUE, input1_size, input1_data, result))
112+
MACRO_1ARG_2TYPES_OP(dpnp_exp_c,
113+
cl::sycl::exp(input_elem),
114+
oneapi::mkl::vm::exp(DPNP_QUEUE, input1_size, input1_data, result))
79115
MACRO_1ARG_2TYPES_OP(dpnp_expm1_c,
80116
cl::sycl::expm1(input_elem),
81-
oneapi::mkl::vm::expm1(DPNP_QUEUE, size, array1, result))
82-
MACRO_1ARG_2TYPES_OP(dpnp_fabs_c, cl::sycl::fabs(input_elem), oneapi::mkl::vm::abs(DPNP_QUEUE, size, array1, result))
117+
oneapi::mkl::vm::expm1(DPNP_QUEUE, input1_size, input1_data, result))
118+
MACRO_1ARG_2TYPES_OP(dpnp_fabs_c,
119+
cl::sycl::fabs(input_elem),
120+
oneapi::mkl::vm::abs(DPNP_QUEUE, input1_size, input1_data, result))
83121
MACRO_1ARG_2TYPES_OP(dpnp_floor_c,
84122
cl::sycl::floor(input_elem),
85-
oneapi::mkl::vm::floor(DPNP_QUEUE, size, array1, result))
123+
oneapi::mkl::vm::floor(DPNP_QUEUE, input1_size, input1_data, result))
86124
MACRO_1ARG_2TYPES_OP(dpnp_log10_c,
87125
cl::sycl::log10(input_elem),
88-
oneapi::mkl::vm::log10(DPNP_QUEUE, size, array1, result))
126+
oneapi::mkl::vm::log10(DPNP_QUEUE, input1_size, input1_data, result))
89127
MACRO_1ARG_2TYPES_OP(dpnp_log1p_c,
90128
cl::sycl::log1p(input_elem),
91-
oneapi::mkl::vm::log1p(DPNP_QUEUE, size, array1, result))
92-
MACRO_1ARG_2TYPES_OP(dpnp_log2_c, cl::sycl::log2(input_elem), oneapi::mkl::vm::log2(DPNP_QUEUE, size, array1, result))
93-
MACRO_1ARG_2TYPES_OP(dpnp_log_c, cl::sycl::log(input_elem), oneapi::mkl::vm::ln(DPNP_QUEUE, size, array1, result))
129+
oneapi::mkl::vm::log1p(DPNP_QUEUE, input1_size, input1_data, result))
130+
MACRO_1ARG_2TYPES_OP(dpnp_log2_c,
131+
cl::sycl::log2(input_elem),
132+
oneapi::mkl::vm::log2(DPNP_QUEUE, input1_size, input1_data, result))
133+
MACRO_1ARG_2TYPES_OP(dpnp_log_c,
134+
cl::sycl::log(input_elem),
135+
oneapi::mkl::vm::ln(DPNP_QUEUE, input1_size, input1_data, result))
94136
MACRO_1ARG_2TYPES_OP(dpnp_radians_c, cl::sycl::radians(input_elem), DPNP_QUEUE.submit(kernel_func))
95-
MACRO_1ARG_2TYPES_OP(dpnp_sin_c, cl::sycl::sin(input_elem), oneapi::mkl::vm::sin(DPNP_QUEUE, size, array1, result))
96-
MACRO_1ARG_2TYPES_OP(dpnp_sinh_c, cl::sycl::sinh(input_elem), oneapi::mkl::vm::sinh(DPNP_QUEUE, size, array1, result))
97-
MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c, cl::sycl::sqrt(input_elem), oneapi::mkl::vm::sqrt(DPNP_QUEUE, size, array1, result))
98-
MACRO_1ARG_2TYPES_OP(dpnp_tan_c, cl::sycl::tan(input_elem), oneapi::mkl::vm::tan(DPNP_QUEUE, size, array1, result))
99-
MACRO_1ARG_2TYPES_OP(dpnp_tanh_c, cl::sycl::tanh(input_elem), oneapi::mkl::vm::tanh(DPNP_QUEUE, size, array1, result))
137+
MACRO_1ARG_2TYPES_OP(dpnp_sin_c,
138+
cl::sycl::sin(input_elem),
139+
oneapi::mkl::vm::sin(DPNP_QUEUE, input1_size, input1_data, result))
140+
MACRO_1ARG_2TYPES_OP(dpnp_sinh_c,
141+
cl::sycl::sinh(input_elem),
142+
oneapi::mkl::vm::sinh(DPNP_QUEUE, input1_size, input1_data, result))
143+
MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c,
144+
cl::sycl::sqrt(input_elem),
145+
oneapi::mkl::vm::sqrt(DPNP_QUEUE, input1_size, input1_data, result))
146+
MACRO_1ARG_2TYPES_OP(dpnp_tan_c,
147+
cl::sycl::tan(input_elem),
148+
oneapi::mkl::vm::tan(DPNP_QUEUE, input1_size, input1_data, result))
149+
MACRO_1ARG_2TYPES_OP(dpnp_tanh_c,
150+
cl::sycl::tanh(input_elem),
151+
oneapi::mkl::vm::tanh(DPNP_QUEUE, input1_size, input1_data, result))
100152
MACRO_1ARG_2TYPES_OP(dpnp_trunc_c,
101153
cl::sycl::trunc(input_elem),
102-
oneapi::mkl::vm::trunc(DPNP_QUEUE, size, array1, result))
154+
oneapi::mkl::vm::trunc(DPNP_QUEUE, input1_size, input1_data, result))
103155

104156
#undef MACRO_1ARG_2TYPES_OP

‎dpnp/backend/include/dpnp_iface.hpp

Lines changed: 11 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -975,7 +975,17 @@ INP_DLLEXPORT void dpnp_invert_c(void* array1_in, void* result, size_t size);
975975

976976
#define MACRO_1ARG_2TYPES_OP(__name__, __operation1__, __operation2__) \
977977
template <typename _DataType_input, typename _DataType_output> \
978-
INP_DLLEXPORT void __name__(void* array1, void* result1, size_t size);
978+
INP_DLLEXPORT void __name__(void* result_out, \
979+
const size_t result_size, \
980+
const size_t result_ndim, \
981+
const size_t* result_shape, \
982+
const size_t* result_strides, \
983+
const void* input1_in, \
984+
const size_t input1_size, \
985+
const size_t input1_ndim, \
986+
const size_t* input1_shape, \
987+
const size_t* input1_strides, \
988+
const size_t* where);
979989

980990
#include <dpnp_gen_1arg_2type_tbl.hpp>
981991

@@ -1114,17 +1124,6 @@ INP_DLLEXPORT void dpnp_remainder_c(void* result_out,
11141124
template <typename _DataType>
11151125
INP_DLLEXPORT void dpnp_repeat_c(const void* array_in, void* result, const size_t repeats, const size_t size);
11161126

1117-
/**
1118-
* @ingroup BACKEND_API
1119-
* @brief copyto function.
1120-
*
1121-
* @param [out] destination Destination array.
1122-
* @param [in] source Source array.
1123-
* @param [in] size Number of elements in destination array.
1124-
*/
1125-
template <typename _DataType_dst, typename _DataType_src>
1126-
INP_DLLEXPORT void dpnp_copyto_c(void* destination, void* source, const size_t size);
1127-
11281127
/**
11291128
* @ingroup BACKEND_API
11301129
* @brief transpose function. Permute axes of the input to the output with elements permutation.

0 commit comments

Comments
 (0)