Skip to content

Commit 05196fe

Browse files
vlad-perevezentsevvlad-perevezentsev
authored
Update dpnp.ediff1d() function (#1970)
* Update dpnp.ediff1d() function * Remove old implementation of dpnp_ediff1d * Update cupy tests * Add new tests for dpnp.ediff1d() * Remove old implementation of dpnp_ediff1d_c * Update docstrings for ediff1d * Update TestEdiff1d to increase coverage * Add support usm_ndarray * Dont use asarray for supported array type
1 parent 5fda819 commit 05196fe

File tree

10 files changed

+276
-293
lines changed

10 files changed

+276
-293
lines changed

‎dpnp/backend/include/dpnp_iface.hpp‎

Lines changed: 0 additions & 47 deletions
Original file line numberDiff line numberDiff line change
@@ -192,53 +192,6 @@ INP_DLLEXPORT void dpnp_dot_c(void *result_out,
192192
const shape_elem_type *input2_shape,
193193
const shape_elem_type *input2_strides);
194194

195-
/**
196-
* @ingroup BACKEND_API
197-
* @brief The differences between consecutive elements of an array.
198-
*
199-
* @param [in] q_ref Reference to SYCL queue.
200-
* @param [out] result_out Output array.
201-
* @param [in] result_size Size of output array.
202-
* @param [in] result_ndim Number of output array dimensions.
203-
* @param [in] result_shape Shape of output array.
204-
* @param [in] result_strides Strides of output array.
205-
* @param [in] input1_in First input array.
206-
* @param [in] input1_size Size of first input array.
207-
* @param [in] input1_ndim Number of first input array dimensions.
208-
* @param [in] input1_shape Shape of first input array.
209-
* @param [in] input1_strides Strides of first input array.
210-
* @param [in] where Mask array.
211-
* @param [in] dep_event_vec_ref Reference to vector of SYCL events.
212-
*/
213-
template <typename _DataType_input, typename _DataType_output>
214-
INP_DLLEXPORT DPCTLSyclEventRef
215-
dpnp_ediff1d_c(DPCTLSyclQueueRef q_ref,
216-
void *result_out,
217-
const size_t result_size,
218-
const size_t result_ndim,
219-
const shape_elem_type *result_shape,
220-
const shape_elem_type *result_strides,
221-
const void *input1_in,
222-
const size_t input1_size,
223-
const size_t input1_ndim,
224-
const shape_elem_type *input1_shape,
225-
const shape_elem_type *input1_strides,
226-
const size_t *where,
227-
const DPCTLEventVectorRef dep_event_vec_ref);
228-
229-
template <typename _DataType_input, typename _DataType_output>
230-
INP_DLLEXPORT void dpnp_ediff1d_c(void *result_out,
231-
const size_t result_size,
232-
const size_t result_ndim,
233-
const shape_elem_type *result_shape,
234-
const shape_elem_type *result_strides,
235-
const void *input1_in,
236-
const size_t input1_size,
237-
const size_t input1_ndim,
238-
const shape_elem_type *input1_shape,
239-
const shape_elem_type *input1_strides,
240-
const size_t *where);
241-
242195
/**
243196
* @ingroup BACKEND_API
244197
* @brief Compute summary of input array elements.

‎dpnp/backend/include/dpnp_iface_fptr.hpp‎

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -72,12 +72,9 @@ enum class DPNPFuncName : size_t
7272
DPNP_FN_COV, /**< Used in numpy.cov() impl */
7373
DPNP_FN_DOT, /**< Used in numpy.dot() impl */
7474
DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */
75-
DPNP_FN_EDIFF1D, /**< Used in numpy.ediff1d() impl */
76-
DPNP_FN_EDIFF1D_EXT, /**< Used in numpy.ediff1d() impl, requires extra
77-
parameters */
78-
DPNP_FN_ERF, /**< Used in scipy.special.erf impl */
79-
DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra
80-
parameters */
75+
DPNP_FN_ERF, /**< Used in scipy.special.erf impl */
76+
DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra
77+
parameters */
8178
DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls
8279
*/
8380
DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like

‎dpnp/backend/kernels/dpnp_krnl_mathematical.cpp‎

Lines changed: 0 additions & 142 deletions
Original file line numberDiff line numberDiff line change
@@ -44,130 +44,6 @@ using dpctl::tensor::kernels::alignment_utils::required_alignment;
4444
static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VECTOR_ABS_CHANGED,
4545
"SYCL DPC++ compiler does not meet minimum version requirement");
4646

47-
template <typename _KernelNameSpecialization1,
48-
typename _KernelNameSpecialization2>
49-
class dpnp_ediff1d_c_kernel;
50-
51-
template <typename _DataType_input, typename _DataType_output>
52-
DPCTLSyclEventRef dpnp_ediff1d_c(DPCTLSyclQueueRef q_ref,
53-
void *result_out,
54-
const size_t result_size,
55-
const size_t result_ndim,
56-
const shape_elem_type *result_shape,
57-
const shape_elem_type *result_strides,
58-
const void *input1_in,
59-
const size_t input1_size,
60-
const size_t input1_ndim,
61-
const shape_elem_type *input1_shape,
62-
const shape_elem_type *input1_strides,
63-
const size_t *where,
64-
const DPCTLEventVectorRef dep_event_vec_ref)
65-
{
66-
/* avoid warning unused variable*/
67-
(void)result_ndim;
68-
(void)result_shape;
69-
(void)result_strides;
70-
(void)input1_ndim;
71-
(void)input1_shape;
72-
(void)input1_strides;
73-
(void)where;
74-
(void)dep_event_vec_ref;
75-
76-
DPCTLSyclEventRef event_ref = nullptr;
77-
78-
if (!input1_size) {
79-
return event_ref;
80-
}
81-
82-
sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));
83-
84-
DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, input1_in,
85-
input1_size);
86-
DPNPC_ptr_adapter<_DataType_output> result_ptr(q_ref, result_out,
87-
result_size, false, true);
88-
89-
_DataType_input *input1_data = input1_ptr.get_ptr();
90-
_DataType_output *result = result_ptr.get_ptr();
91-
92-
sycl::event event;
93-
sycl::range<1> gws(result_size);
94-
95-
auto kernel_parallel_for_func = [=](sycl::id<1> global_id) {
96-
size_t output_id =
97-
global_id[0]; /*for (size_t i = 0; i < result_size; ++i)*/
98-
{
99-
const _DataType_output curr_elem = input1_data[output_id];
100-
const _DataType_output next_elem = input1_data[output_id + 1];
101-
result[output_id] = next_elem - curr_elem;
102-
}
103-
};
104-
auto kernel_func = [&](sycl::handler &cgh) {
105-
cgh.parallel_for<
106-
class dpnp_ediff1d_c_kernel<_DataType_input, _DataType_output>>(
107-
gws, kernel_parallel_for_func);
108-
};
109-
event = q.submit(kernel_func);
110-
111-
input1_ptr.depends_on(event);
112-
result_ptr.depends_on(event);
113-
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
114-
115-
return DPCTLEvent_Copy(event_ref);
116-
}
117-
118-
template <typename _DataType_input, typename _DataType_output>
119-
void dpnp_ediff1d_c(void *result_out,
120-
const size_t result_size,
121-
const size_t result_ndim,
122-
const shape_elem_type *result_shape,
123-
const shape_elem_type *result_strides,
124-
const void *input1_in,
125-
const size_t input1_size,
126-
const size_t input1_ndim,
127-
const shape_elem_type *input1_shape,
128-
const shape_elem_type *input1_strides,
129-
const size_t *where)
130-
{
131-
DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
132-
DPCTLEventVectorRef dep_event_vec_ref = nullptr;
133-
DPCTLSyclEventRef event_ref =
134-
dpnp_ediff1d_c<_DataType_input, _DataType_output>(
135-
q_ref, result_out, result_size, result_ndim, result_shape,
136-
result_strides, input1_in, input1_size, input1_ndim, input1_shape,
137-
input1_strides, where, dep_event_vec_ref);
138-
DPCTLEvent_WaitAndThrow(event_ref);
139-
}
140-
141-
template <typename _DataType_input, typename _DataType_output>
142-
void (*dpnp_ediff1d_default_c)(void *,
143-
const size_t,
144-
const size_t,
145-
const shape_elem_type *,
146-
const shape_elem_type *,
147-
const void *,
148-
const size_t,
149-
const size_t,
150-
const shape_elem_type *,
151-
const shape_elem_type *,
152-
const size_t *) =
153-
dpnp_ediff1d_c<_DataType_input, _DataType_output>;
154-
155-
template <typename _DataType_input, typename _DataType_output>
156-
DPCTLSyclEventRef (*dpnp_ediff1d_ext_c)(DPCTLSyclQueueRef,
157-
void *,
158-
const size_t,
159-
const size_t,
160-
const shape_elem_type *,
161-
const shape_elem_type *,
162-
const void *,
163-
const size_t,
164-
const size_t,
165-
const shape_elem_type *,
166-
const shape_elem_type *,
167-
const size_t *,
168-
const DPCTLEventVectorRef) =
169-
dpnp_ediff1d_c<_DataType_input, _DataType_output>;
170-
17147
template <typename _KernelNameSpecialization1,
17248
typename _KernelNameSpecialization2>
17349
class dpnp_modf_c_kernel;
@@ -260,24 +136,6 @@ DPCTLSyclEventRef (*dpnp_modf_ext_c)(DPCTLSyclQueueRef,
260136
void func_map_init_mathematical(func_map_t &fmap)
261137
{
262138

263-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_INT][eft_INT] = {
264-
eft_LNG, (void *)dpnp_ediff1d_default_c<int32_t, int64_t>};
265-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_LNG][eft_LNG] = {
266-
eft_LNG, (void *)dpnp_ediff1d_default_c<int64_t, int64_t>};
267-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_FLT][eft_FLT] = {
268-
eft_FLT, (void *)dpnp_ediff1d_default_c<float, float>};
269-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_DBL][eft_DBL] = {
270-
eft_DBL, (void *)dpnp_ediff1d_default_c<double, double>};
271-
272-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_INT][eft_INT] = {
273-
eft_LNG, (void *)dpnp_ediff1d_ext_c<int32_t, int64_t>};
274-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_LNG][eft_LNG] = {
275-
eft_LNG, (void *)dpnp_ediff1d_ext_c<int64_t, int64_t>};
276-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_FLT][eft_FLT] = {
277-
eft_FLT, (void *)dpnp_ediff1d_ext_c<float, float>};
278-
fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_DBL][eft_DBL] = {
279-
eft_DBL, (void *)dpnp_ediff1d_ext_c<double, double>};
280-
281139
fmap[DPNPFuncName::DPNP_FN_MODF][eft_INT][eft_INT] = {
282140
eft_DBL, (void *)dpnp_modf_default_c<int32_t, double>};
283141
fmap[DPNPFuncName::DPNP_FN_MODF][eft_LNG][eft_LNG] = {

‎dpnp/dpnp_algo/dpnp_algo.pxd‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
3535
cdef enum DPNPFuncName "DPNPFuncName":
3636
DPNP_FN_CHOOSE_EXT
3737
DPNP_FN_CORRELATE_EXT
38-
DPNP_FN_EDIFF1D_EXT
3938
DPNP_FN_ERF_EXT
4039
DPNP_FN_MEDIAN_EXT
4140
DPNP_FN_MODF_EXT

‎dpnp/dpnp_algo/dpnp_algo_mathematical.pxi‎

Lines changed: 0 additions & 57 deletions
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,6 @@ and the rest of the library
3636
# NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file
3737

3838
__all__ += [
39-
"dpnp_ediff1d",
4039
"dpnp_modf",
4140
]
4241

@@ -46,62 +45,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef,
4645
const c_dpctl.DPCTLEventVectorRef)
4746

4847

49-
cpdef utils.dpnp_descriptor dpnp_ediff1d(utils.dpnp_descriptor x1):
50-
51-
if x1.size <= 1:
52-
return utils.dpnp_descriptor(dpnp.empty(0, dtype=x1.dtype)) # TODO need to call dpnp_empty instead
53-
54-
# Convert type (x1.dtype) to C enum DPNPFuncType
55-
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
56-
57-
# get the FPTR data structure
58-
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_EDIFF1D_EXT, param1_type, param1_type)
59-
60-
result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
61-
62-
# Currently shape and strides of the input array are not took into account for the function ediff1d
63-
cdef shape_type_c x1_shape = (x1.size,)
64-
cdef shape_type_c x1_strides = utils.strides_to_vector(None, x1_shape)
65-
66-
x1_obj = x1.get_array()
67-
68-
cdef shape_type_c result_shape = (x1.size - 1,)
69-
cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape,
70-
kernel_data.return_type,
71-
None,
72-
device=x1_obj.sycl_device,
73-
usm_type=x1_obj.usm_type,
74-
sycl_queue=x1_obj.sycl_queue)
75-
76-
cdef shape_type_c result_strides = utils.strides_to_vector(result.strides, result_shape)
77-
78-
result_sycl_queue = result.get_array().sycl_queue
79-
80-
cdef c_dpctl.SyclQueue q = <c_dpctl.SyclQueue> result_sycl_queue
81-
cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref()
82-
83-
# Call FPTR function
84-
cdef fptr_1in_1out_strides_t func = <fptr_1in_1out_strides_t > kernel_data.ptr
85-
cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref,
86-
result.get_data(),
87-
result.size,
88-
result.ndim,
89-
result_shape.data(),
90-
result_strides.data(),
91-
x1.get_data(),
92-
x1.size,
93-
x1.ndim,
94-
x1_shape.data(),
95-
x1_strides.data(),
96-
NULL,
97-
NULL) # dep_events_ref
98-
99-
with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref)
100-
c_dpctl.DPCTLEvent_Delete(event_ref)
101-
102-
return result
103-
104-
10548
cpdef tuple dpnp_modf(utils.dpnp_descriptor x1):
10649
""" Convert string type names (array.dtype) to C enum DPNPFuncType """
10750
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)

0 commit comments

Comments
 (0)