Skip to content

Commit 7c45c10

Browse files
vtavanaantonwolfy
vtavana
and
antonwolfy
authored
implement dpnp.median (#2117)
* implement dpnp.median * use assert_dtype_allclose * add type_check=has_support_aspect64() * address comments * Update dpnp/dpnp_iface_statistics.py Co-authored-by: Anton <[email protected]> * remove a test * add new tests --------- Co-authored-by: Anton <[email protected]>
1 parent a42a87f commit 7c45c10

17 files changed

+501
-595
lines changed

‎dpnp/backend/include/dpnp_iface.hpp

Lines changed: 0 additions & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -195,27 +195,6 @@ INP_DLLEXPORT void dpnp_partition_c(void *array,
195195
const shape_elem_type *shape,
196196
const size_t ndim);
197197

198-
/**
199-
* @ingroup BACKEND_API
200-
* @brief math library implementation of sort function
201-
*
202-
* @param [in] q_ref Reference to SYCL queue.
203-
* @param [in] array Input array with data.
204-
* @param [out] result Output array with indices.
205-
* @param [in] size Number of elements in input arrays.
206-
* @param [in] dep_event_vec_ref Reference to vector of SYCL events.
207-
*/
208-
template <typename _DataType>
209-
INP_DLLEXPORT DPCTLSyclEventRef
210-
dpnp_sort_c(DPCTLSyclQueueRef q_ref,
211-
void *array,
212-
void *result,
213-
size_t size,
214-
const DPCTLEventVectorRef dep_event_vec_ref);
215-
216-
template <typename _DataType>
217-
INP_DLLEXPORT void dpnp_sort_c(void *array, void *result, size_t size);
218-
219198
/**
220199
* @ingroup BACKEND_API
221200
* @brief correlate function
@@ -318,38 +297,6 @@ INP_DLLEXPORT DPCTLSyclEventRef
318297
template <typename _DataType>
319298
INP_DLLEXPORT void dpnp_initval_c(void *result1, void *value, size_t size);
320299

321-
/**
322-
* @ingroup BACKEND_API
323-
* @brief math library implementation of median function
324-
*
325-
* @param [in] q_ref Reference to SYCL queue.
326-
* @param [in] array Input array with data.
327-
* @param [out] result Output array.
328-
* @param [in] shape Shape of input array.
329-
* @param [in] ndim Number of elements in shape.
330-
* @param [in] axis Axis.
331-
* @param [in] naxis Number of elements in axis.
332-
* @param [in] dep_event_vec_ref Reference to vector of SYCL events.
333-
*/
334-
template <typename _DataType, typename _ResultType>
335-
INP_DLLEXPORT DPCTLSyclEventRef
336-
dpnp_median_c(DPCTLSyclQueueRef q_ref,
337-
void *array,
338-
void *result,
339-
const shape_elem_type *shape,
340-
size_t ndim,
341-
const shape_elem_type *axis,
342-
size_t naxis,
343-
const DPCTLEventVectorRef dep_event_vec_ref);
344-
345-
template <typename _DataType, typename _ResultType>
346-
INP_DLLEXPORT void dpnp_median_c(void *array,
347-
void *result,
348-
const shape_elem_type *shape,
349-
size_t ndim,
350-
const shape_elem_type *axis,
351-
size_t naxis);
352-
353300
#define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \
354301
template <typename _DataType> \
355302
INP_DLLEXPORT DPCTLSyclEventRef __name__( \

‎dpnp/backend/include/dpnp_iface_fptr.hpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -74,9 +74,6 @@ enum class DPNPFuncName : size_t
7474
*/
7575
DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like
7676
impls */
77-
DPNP_FN_MEDIAN, /**< Used in numpy.median() impl */
78-
DPNP_FN_MEDIAN_EXT, /**< Used in numpy.median() impl, requires extra
79-
parameters */
8077
DPNP_FN_MODF, /**< Used in numpy.modf() impl */
8178
DPNP_FN_MODF_EXT, /**< Used in numpy.modf() impl, requires extra parameters
8279
*/
@@ -203,7 +200,6 @@ enum class DPNPFuncName : size_t
203200
DPNP_FN_RNG_ZIPF, /**< Used in numpy.random.zipf() impl */
204201
DPNP_FN_RNG_ZIPF_EXT, /**< Used in numpy.random.zipf() impl, requires extra
205202
parameters */
206-
DPNP_FN_SORT, /**< Used in numpy.sort() impl */
207203
DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */
208204
DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */
209205
DPNP_FN_LAST, /**< The latest element of the enumeration */

‎dpnp/backend/kernels/dpnp_krnl_sorting.cpp

Lines changed: 0 additions & 67 deletions
Original file line numberDiff line numberDiff line change
@@ -30,15 +30,6 @@
3030
#include "queue_sycl.hpp"
3131
#include <dpnp_iface.hpp>
3232

33-
template <typename _DataType>
34-
struct _sort_less
35-
{
36-
inline bool operator()(const _DataType &val1, const _DataType &val2)
37-
{
38-
return (val1 < val2);
39-
}
40-
};
41-
4233
template <typename _DataType>
4334
class dpnp_partition_c_kernel;
4435

@@ -199,55 +190,6 @@ DPCTLSyclEventRef (*dpnp_partition_ext_c)(DPCTLSyclQueueRef,
199190
const DPCTLEventVectorRef) =
200191
dpnp_partition_c<_DataType>;
201192

202-
template <typename _DataType>
203-
class dpnp_sort_c_kernel;
204-
205-
template <typename _DataType>
206-
DPCTLSyclEventRef dpnp_sort_c(DPCTLSyclQueueRef q_ref,
207-
void *array1_in,
208-
void *result1,
209-
size_t size,
210-
const DPCTLEventVectorRef dep_event_vec_ref)
211-
{
212-
// avoid warning unused variable
213-
(void)dep_event_vec_ref;
214-
215-
DPCTLSyclEventRef event_ref = nullptr;
216-
sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));
217-
218-
DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size, true);
219-
DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result1, size, true, true);
220-
_DataType *array_1 = input1_ptr.get_ptr();
221-
_DataType *result = result1_ptr.get_ptr();
222-
223-
std::copy(array_1, array_1 + size, result);
224-
225-
auto policy = oneapi::dpl::execution::make_device_policy<
226-
class dpnp_sort_c_kernel<_DataType>>(q);
227-
228-
// fails without explicitly specifying of comparator or with std::less
229-
// during kernels compilation affects other kernels
230-
std::sort(policy, result, result + size, _sort_less<_DataType>());
231-
232-
policy.queue().wait();
233-
234-
return event_ref;
235-
}
236-
237-
template <typename _DataType>
238-
void dpnp_sort_c(void *array1_in, void *result1, size_t size)
239-
{
240-
DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
241-
DPCTLEventVectorRef dep_event_vec_ref = nullptr;
242-
DPCTLSyclEventRef event_ref = dpnp_sort_c<_DataType>(
243-
q_ref, array1_in, result1, size, dep_event_vec_ref);
244-
DPCTLEvent_WaitAndThrow(event_ref);
245-
DPCTLEvent_Delete(event_ref);
246-
}
247-
248-
template <typename _DataType>
249-
void (*dpnp_sort_default_c)(void *, void *, size_t) = dpnp_sort_c<_DataType>;
250-
251193
void func_map_init_sorting(func_map_t &fmap)
252194
{
253195
fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_INT][eft_INT] = {
@@ -274,14 +216,5 @@ void func_map_init_sorting(func_map_t &fmap)
274216
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_C128][eft_C128] = {
275217
eft_C128, (void *)dpnp_partition_ext_c<std::complex<double>>};
276218

277-
fmap[DPNPFuncName::DPNP_FN_SORT][eft_INT][eft_INT] = {
278-
eft_INT, (void *)dpnp_sort_default_c<int32_t>};
279-
fmap[DPNPFuncName::DPNP_FN_SORT][eft_LNG][eft_LNG] = {
280-
eft_LNG, (void *)dpnp_sort_default_c<int64_t>};
281-
fmap[DPNPFuncName::DPNP_FN_SORT][eft_FLT][eft_FLT] = {
282-
eft_FLT, (void *)dpnp_sort_default_c<float>};
283-
fmap[DPNPFuncName::DPNP_FN_SORT][eft_DBL][eft_DBL] = {
284-
eft_DBL, (void *)dpnp_sort_default_c<double>};
285-
286219
return;
287220
}

‎dpnp/backend/kernels/dpnp_krnl_statistics.cpp

Lines changed: 0 additions & 112 deletions
Original file line numberDiff line numberDiff line change
@@ -128,88 +128,6 @@ DPCTLSyclEventRef (*dpnp_correlate_ext_c)(DPCTLSyclQueueRef,
128128
const DPCTLEventVectorRef) =
129129
dpnp_correlate_c<_DataType_output, _DataType_input1, _DataType_input2>;
130130

131-
template <typename _DataType, typename _ResultType>
132-
DPCTLSyclEventRef dpnp_median_c(DPCTLSyclQueueRef q_ref,
133-
void *array1_in,
134-
void *result1,
135-
const shape_elem_type *shape,
136-
size_t ndim,
137-
const shape_elem_type *axis,
138-
size_t naxis,
139-
const DPCTLEventVectorRef dep_event_vec_ref)
140-
{
141-
// avoid warning unused variable
142-
(void)dep_event_vec_ref;
143-
144-
__attribute__((unused)) void *tmp = (void *)(axis + naxis);
145-
146-
DPCTLSyclEventRef event_ref = nullptr;
147-
148-
const size_t size = std::accumulate(shape, shape + ndim, 1,
149-
std::multiplies<shape_elem_type>());
150-
if (!size) {
151-
return event_ref;
152-
}
153-
154-
sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));
155-
156-
DPNPC_ptr_adapter<_ResultType> result_ptr(q_ref, result1, 1, true, true);
157-
_ResultType *result = result_ptr.get_ptr();
158-
159-
_DataType *sorted = reinterpret_cast<_DataType *>(
160-
sycl::malloc_shared(size * sizeof(_DataType), q));
161-
162-
dpnp_sort_c<_DataType>(array1_in, sorted, size);
163-
164-
if (size % 2 == 0) {
165-
result[0] =
166-
static_cast<_ResultType>(sorted[size / 2] + sorted[size / 2 - 1]) /
167-
2;
168-
}
169-
else {
170-
result[0] = sorted[(size - 1) / 2];
171-
}
172-
173-
sycl::free(sorted, q);
174-
175-
return event_ref;
176-
}
177-
178-
template <typename _DataType, typename _ResultType>
179-
void dpnp_median_c(void *array1_in,
180-
void *result1,
181-
const shape_elem_type *shape,
182-
size_t ndim,
183-
const shape_elem_type *axis,
184-
size_t naxis)
185-
{
186-
DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
187-
DPCTLEventVectorRef dep_event_vec_ref = nullptr;
188-
DPCTLSyclEventRef event_ref = dpnp_median_c<_DataType, _ResultType>(
189-
q_ref, array1_in, result1, shape, ndim, axis, naxis, dep_event_vec_ref);
190-
DPCTLEvent_WaitAndThrow(event_ref);
191-
DPCTLEvent_Delete(event_ref);
192-
}
193-
194-
template <typename _DataType, typename _ResultType>
195-
void (*dpnp_median_default_c)(void *,
196-
void *,
197-
const shape_elem_type *,
198-
size_t,
199-
const shape_elem_type *,
200-
size_t) = dpnp_median_c<_DataType, _ResultType>;
201-
202-
template <typename _DataType, typename _ResultType>
203-
DPCTLSyclEventRef (*dpnp_median_ext_c)(DPCTLSyclQueueRef,
204-
void *,
205-
void *,
206-
const shape_elem_type *,
207-
size_t,
208-
const shape_elem_type *,
209-
size_t,
210-
const DPCTLEventVectorRef) =
211-
dpnp_median_c<_DataType, _ResultType>;
212-
213131
void func_map_init_statistics(func_map_t &fmap)
214132
{
215133
fmap[DPNPFuncName::DPNP_FN_CORRELATE][eft_INT][eft_INT] = {
@@ -278,35 +196,5 @@ void func_map_init_statistics(func_map_t &fmap)
278196
fmap[DPNPFuncName::DPNP_FN_CORRELATE_EXT][eft_DBL][eft_DBL] = {
279197
eft_DBL, (void *)dpnp_correlate_ext_c<double, double, double>};
280198

281-
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_INT][eft_INT] = {
282-
eft_DBL, (void *)dpnp_median_default_c<int32_t, double>};
283-
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_LNG][eft_LNG] = {
284-
eft_DBL, (void *)dpnp_median_default_c<int64_t, double>};
285-
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_FLT][eft_FLT] = {
286-
eft_FLT, (void *)dpnp_median_default_c<float, float>};
287-
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_DBL][eft_DBL] = {
288-
eft_DBL, (void *)dpnp_median_default_c<double, double>};
289-
290-
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_INT][eft_INT] = {
291-
get_default_floating_type(),
292-
(void *)dpnp_median_ext_c<
293-
int32_t, func_type_map_t::find_type<get_default_floating_type()>>,
294-
get_default_floating_type<std::false_type>(),
295-
(void *)dpnp_median_ext_c<
296-
int32_t, func_type_map_t::find_type<
297-
get_default_floating_type<std::false_type>()>>};
298-
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_LNG][eft_LNG] = {
299-
get_default_floating_type(),
300-
(void *)dpnp_median_ext_c<
301-
int64_t, func_type_map_t::find_type<get_default_floating_type()>>,
302-
get_default_floating_type<std::false_type>(),
303-
(void *)dpnp_median_ext_c<
304-
int64_t, func_type_map_t::find_type<
305-
get_default_floating_type<std::false_type>()>>};
306-
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_FLT][eft_FLT] = {
307-
eft_FLT, (void *)dpnp_median_ext_c<float, float>};
308-
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_DBL][eft_DBL] = {
309-
eft_DBL, (void *)dpnp_median_ext_c<double, double>};
310-
311199
return;
312200
}

‎dpnp/backend/src/dpnp_fptr.hpp

Lines changed: 0 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -260,17 +260,6 @@ class dpnp_less_comp
260260
}
261261
};
262262

263-
/**
264-
* A template function that determines the default floating-point type
265-
* based on the value of the template parameter has_fp64.
266-
*/
267-
template <typename has_fp64 = std::true_type>
268-
static constexpr DPNPFuncType get_default_floating_type()
269-
{
270-
return has_fp64::value ? DPNPFuncType::DPNP_FT_DOUBLE
271-
: DPNPFuncType::DPNP_FT_FLOAT;
272-
}
273-
274263
/**
275264
* FPTR interface initialization functions
276265
*/

‎dpnp/dpnp_algo/dpnp_algo.pxd

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
3636
DPNP_FN_CHOOSE_EXT
3737
DPNP_FN_CORRELATE_EXT
3838
DPNP_FN_ERF_EXT
39-
DPNP_FN_MEDIAN_EXT
4039
DPNP_FN_MODF_EXT
4140
DPNP_FN_PARTITION_EXT
4241
DPNP_FN_RNG_BETA_EXT

‎dpnp/dpnp_algo/dpnp_algo_statistics.pxi

Lines changed: 0 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -37,17 +37,9 @@ and the rest of the library
3737

3838
__all__ += [
3939
"dpnp_correlate",
40-
"dpnp_median",
4140
]
4241

4342

44-
# C function pointer to the C library template functions
45-
ctypedef c_dpctl.DPCTLSyclEventRef(*custom_statistic_1in_1out_func_ptr_t)(c_dpctl.DPCTLSyclQueueRef,
46-
void *, void * , shape_elem_type * , size_t,
47-
shape_elem_type * , size_t,
48-
const c_dpctl.DPCTLEventVectorRef)
49-
50-
5143
cpdef utils.dpnp_descriptor dpnp_correlate(utils.dpnp_descriptor x1, utils.dpnp_descriptor x2):
5244
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
5345
cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(x2.dtype)
@@ -90,47 +82,3 @@ cpdef utils.dpnp_descriptor dpnp_correlate(utils.dpnp_descriptor x1, utils.dpnp_
9082
c_dpctl.DPCTLEvent_Delete(event_ref)
9183

9284
return result
93-
94-
95-
cpdef utils.dpnp_descriptor dpnp_median(utils.dpnp_descriptor array1):
96-
cdef shape_type_c x1_shape = array1.shape
97-
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype)
98-
99-
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_MEDIAN_EXT, param1_type, param1_type)
100-
101-
array1_obj = array1.get_array()
102-
103-
cdef (DPNPFuncType, void *) ret_type_and_func = utils.get_ret_type_and_func(kernel_data,
104-
array1_obj.sycl_device.has_aspect_fp64)
105-
cdef DPNPFuncType return_type = ret_type_and_func[0]
106-
cdef custom_statistic_1in_1out_func_ptr_t func = < custom_statistic_1in_1out_func_ptr_t > ret_type_and_func[1]
107-
108-
cdef utils.dpnp_descriptor result = utils.create_output_descriptor((1,),
109-
return_type,
110-
None,
111-
device=array1_obj.sycl_device,
112-
usm_type=array1_obj.usm_type,
113-
sycl_queue=array1_obj.sycl_queue)
114-
115-
result_sycl_queue = result.get_array().sycl_queue
116-
117-
cdef c_dpctl.SyclQueue q = <c_dpctl.SyclQueue> result_sycl_queue
118-
cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref()
119-
120-
# stub for interface support
121-
cdef shape_type_c axis
122-
cdef Py_ssize_t axis_size = 0
123-
124-
cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref,
125-
array1.get_data(),
126-
result.get_data(),
127-
x1_shape.data(),
128-
array1.ndim,
129-
axis.data(),
130-
axis_size,
131-
NULL) # dep_events_ref
132-
133-
with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref)
134-
c_dpctl.DPCTLEvent_Delete(event_ref)
135-
136-
return result

0 commit comments

Comments
 (0)