Skip to content

implement dpnp.median #2117

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 8 commits into from
Oct 22, 2024
Merged

implement dpnp.median #2117

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
33fda9d
implement dpnp.median
vtavana Oct 14, 2024
ba2e4c8
use assert_dtype_allclose
vtavana Oct 18, 2024
356553e
add type_check=has_support_aspect64()
vtavana Oct 18, 2024
d5dfee4
address comments
vtavana Oct 21, 2024
744a988
Merge branch 'master' into impl_median
vtavana Oct 21, 2024
9d3ede1
Update dpnp/dpnp_iface_statistics.py
vtavana Oct 21, 2024
a4b6818
remove a test
vtavana Oct 21, 2024
74d1184
add new tests
vtavana Oct 22, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
53 changes: 0 additions & 53 deletions dpnp/backend/include/dpnp_iface.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -195,27 +195,6 @@ INP_DLLEXPORT void dpnp_partition_c(void *array,
const shape_elem_type *shape,
const size_t ndim);

/**
* @ingroup BACKEND_API
* @brief math library implementation of sort function
*
* @param [in] q_ref Reference to SYCL queue.
* @param [in] array Input array with data.
* @param [out] result Output array with indices.
* @param [in] size Number of elements in input arrays.
* @param [in] dep_event_vec_ref Reference to vector of SYCL events.
*/
template <typename _DataType>
INP_DLLEXPORT DPCTLSyclEventRef
dpnp_sort_c(DPCTLSyclQueueRef q_ref,
void *array,
void *result,
size_t size,
const DPCTLEventVectorRef dep_event_vec_ref);

template <typename _DataType>
INP_DLLEXPORT void dpnp_sort_c(void *array, void *result, size_t size);

/**
* @ingroup BACKEND_API
* @brief correlate function
Expand Down Expand Up @@ -318,38 +297,6 @@ INP_DLLEXPORT DPCTLSyclEventRef
template <typename _DataType>
INP_DLLEXPORT void dpnp_initval_c(void *result1, void *value, size_t size);

/**
* @ingroup BACKEND_API
* @brief math library implementation of median function
*
* @param [in] q_ref Reference to SYCL queue.
* @param [in] array Input array with data.
* @param [out] result Output array.
* @param [in] shape Shape of input array.
* @param [in] ndim Number of elements in shape.
* @param [in] axis Axis.
* @param [in] naxis Number of elements in axis.
* @param [in] dep_event_vec_ref Reference to vector of SYCL events.
*/
template <typename _DataType, typename _ResultType>
INP_DLLEXPORT DPCTLSyclEventRef
dpnp_median_c(DPCTLSyclQueueRef q_ref,
void *array,
void *result,
const shape_elem_type *shape,
size_t ndim,
const shape_elem_type *axis,
size_t naxis,
const DPCTLEventVectorRef dep_event_vec_ref);

template <typename _DataType, typename _ResultType>
INP_DLLEXPORT void dpnp_median_c(void *array,
void *result,
const shape_elem_type *shape,
size_t ndim,
const shape_elem_type *axis,
size_t naxis);

#define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \
template <typename _DataType> \
INP_DLLEXPORT DPCTLSyclEventRef __name__( \
Expand Down
4 changes: 0 additions & 4 deletions dpnp/backend/include/dpnp_iface_fptr.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,9 +74,6 @@ enum class DPNPFuncName : size_t
*/
DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like
impls */
DPNP_FN_MEDIAN, /**< Used in numpy.median() impl */
DPNP_FN_MEDIAN_EXT, /**< Used in numpy.median() impl, requires extra
parameters */
DPNP_FN_MODF, /**< Used in numpy.modf() impl */
DPNP_FN_MODF_EXT, /**< Used in numpy.modf() impl, requires extra parameters
*/
Expand Down Expand Up @@ -203,7 +200,6 @@ enum class DPNPFuncName : size_t
DPNP_FN_RNG_ZIPF, /**< Used in numpy.random.zipf() impl */
DPNP_FN_RNG_ZIPF_EXT, /**< Used in numpy.random.zipf() impl, requires extra
parameters */
DPNP_FN_SORT, /**< Used in numpy.sort() impl */
DPNP_FN_ZEROS, /**< Used in numpy.zeros() impl */
DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */
DPNP_FN_LAST, /**< The latest element of the enumeration */
Expand Down
67 changes: 0 additions & 67 deletions dpnp/backend/kernels/dpnp_krnl_sorting.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,15 +30,6 @@
#include "queue_sycl.hpp"
#include <dpnp_iface.hpp>

template <typename _DataType>
struct _sort_less
{
inline bool operator()(const _DataType &val1, const _DataType &val2)
{
return (val1 < val2);
}
};

template <typename _DataType>
class dpnp_partition_c_kernel;

Expand Down Expand Up @@ -199,55 +190,6 @@ DPCTLSyclEventRef (*dpnp_partition_ext_c)(DPCTLSyclQueueRef,
const DPCTLEventVectorRef) =
dpnp_partition_c<_DataType>;

template <typename _DataType>
class dpnp_sort_c_kernel;

template <typename _DataType>
DPCTLSyclEventRef dpnp_sort_c(DPCTLSyclQueueRef q_ref,
void *array1_in,
void *result1,
size_t size,
const DPCTLEventVectorRef dep_event_vec_ref)
{
// avoid warning unused variable
(void)dep_event_vec_ref;

DPCTLSyclEventRef event_ref = nullptr;
sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));

DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size, true);
DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result1, size, true, true);
_DataType *array_1 = input1_ptr.get_ptr();
_DataType *result = result1_ptr.get_ptr();

std::copy(array_1, array_1 + size, result);

auto policy = oneapi::dpl::execution::make_device_policy<
class dpnp_sort_c_kernel<_DataType>>(q);

// fails without explicitly specifying of comparator or with std::less
// during kernels compilation affects other kernels
std::sort(policy, result, result + size, _sort_less<_DataType>());

policy.queue().wait();

return event_ref;
}

template <typename _DataType>
void dpnp_sort_c(void *array1_in, void *result1, size_t size)
{
DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
DPCTLEventVectorRef dep_event_vec_ref = nullptr;
DPCTLSyclEventRef event_ref = dpnp_sort_c<_DataType>(
q_ref, array1_in, result1, size, dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType>
void (*dpnp_sort_default_c)(void *, void *, size_t) = dpnp_sort_c<_DataType>;

void func_map_init_sorting(func_map_t &fmap)
{
fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_INT][eft_INT] = {
Expand All @@ -274,14 +216,5 @@ void func_map_init_sorting(func_map_t &fmap)
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_C128][eft_C128] = {
eft_C128, (void *)dpnp_partition_ext_c<std::complex<double>>};

fmap[DPNPFuncName::DPNP_FN_SORT][eft_INT][eft_INT] = {
eft_INT, (void *)dpnp_sort_default_c<int32_t>};
fmap[DPNPFuncName::DPNP_FN_SORT][eft_LNG][eft_LNG] = {
eft_LNG, (void *)dpnp_sort_default_c<int64_t>};
fmap[DPNPFuncName::DPNP_FN_SORT][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_sort_default_c<float>};
fmap[DPNPFuncName::DPNP_FN_SORT][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_sort_default_c<double>};

return;
}
112 changes: 0 additions & 112 deletions dpnp/backend/kernels/dpnp_krnl_statistics.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -128,88 +128,6 @@ DPCTLSyclEventRef (*dpnp_correlate_ext_c)(DPCTLSyclQueueRef,
const DPCTLEventVectorRef) =
dpnp_correlate_c<_DataType_output, _DataType_input1, _DataType_input2>;

template <typename _DataType, typename _ResultType>
DPCTLSyclEventRef dpnp_median_c(DPCTLSyclQueueRef q_ref,
void *array1_in,
void *result1,
const shape_elem_type *shape,
size_t ndim,
const shape_elem_type *axis,
size_t naxis,
const DPCTLEventVectorRef dep_event_vec_ref)
{
// avoid warning unused variable
(void)dep_event_vec_ref;

__attribute__((unused)) void *tmp = (void *)(axis + naxis);

DPCTLSyclEventRef event_ref = nullptr;

const size_t size = std::accumulate(shape, shape + ndim, 1,
std::multiplies<shape_elem_type>());
if (!size) {
return event_ref;
}

sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));

DPNPC_ptr_adapter<_ResultType> result_ptr(q_ref, result1, 1, true, true);
_ResultType *result = result_ptr.get_ptr();

_DataType *sorted = reinterpret_cast<_DataType *>(
sycl::malloc_shared(size * sizeof(_DataType), q));

dpnp_sort_c<_DataType>(array1_in, sorted, size);

if (size % 2 == 0) {
result[0] =
static_cast<_ResultType>(sorted[size / 2] + sorted[size / 2 - 1]) /
2;
}
else {
result[0] = sorted[(size - 1) / 2];
}

sycl::free(sorted, q);

return event_ref;
}

template <typename _DataType, typename _ResultType>
void dpnp_median_c(void *array1_in,
void *result1,
const shape_elem_type *shape,
size_t ndim,
const shape_elem_type *axis,
size_t naxis)
{
DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
DPCTLEventVectorRef dep_event_vec_ref = nullptr;
DPCTLSyclEventRef event_ref = dpnp_median_c<_DataType, _ResultType>(
q_ref, array1_in, result1, shape, ndim, axis, naxis, dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType, typename _ResultType>
void (*dpnp_median_default_c)(void *,
void *,
const shape_elem_type *,
size_t,
const shape_elem_type *,
size_t) = dpnp_median_c<_DataType, _ResultType>;

template <typename _DataType, typename _ResultType>
DPCTLSyclEventRef (*dpnp_median_ext_c)(DPCTLSyclQueueRef,
void *,
void *,
const shape_elem_type *,
size_t,
const shape_elem_type *,
size_t,
const DPCTLEventVectorRef) =
dpnp_median_c<_DataType, _ResultType>;

void func_map_init_statistics(func_map_t &fmap)
{
fmap[DPNPFuncName::DPNP_FN_CORRELATE][eft_INT][eft_INT] = {
Expand Down Expand Up @@ -278,35 +196,5 @@ void func_map_init_statistics(func_map_t &fmap)
fmap[DPNPFuncName::DPNP_FN_CORRELATE_EXT][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_correlate_ext_c<double, double, double>};

fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_INT][eft_INT] = {
eft_DBL, (void *)dpnp_median_default_c<int32_t, double>};
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_LNG][eft_LNG] = {
eft_DBL, (void *)dpnp_median_default_c<int64_t, double>};
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_median_default_c<float, float>};
fmap[DPNPFuncName::DPNP_FN_MEDIAN][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_median_default_c<double, double>};

fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_INT][eft_INT] = {
get_default_floating_type(),
(void *)dpnp_median_ext_c<
int32_t, func_type_map_t::find_type<get_default_floating_type()>>,
get_default_floating_type<std::false_type>(),
(void *)dpnp_median_ext_c<
int32_t, func_type_map_t::find_type<
get_default_floating_type<std::false_type>()>>};
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_LNG][eft_LNG] = {
get_default_floating_type(),
(void *)dpnp_median_ext_c<
int64_t, func_type_map_t::find_type<get_default_floating_type()>>,
get_default_floating_type<std::false_type>(),
(void *)dpnp_median_ext_c<
int64_t, func_type_map_t::find_type<
get_default_floating_type<std::false_type>()>>};
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_median_ext_c<float, float>};
fmap[DPNPFuncName::DPNP_FN_MEDIAN_EXT][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_median_ext_c<double, double>};

return;
}
11 changes: 0 additions & 11 deletions dpnp/backend/src/dpnp_fptr.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -260,17 +260,6 @@ class dpnp_less_comp
}
};

/**
* A template function that determines the default floating-point type
* based on the value of the template parameter has_fp64.
*/
template <typename has_fp64 = std::true_type>
static constexpr DPNPFuncType get_default_floating_type()
{
return has_fp64::value ? DPNPFuncType::DPNP_FT_DOUBLE
: DPNPFuncType::DPNP_FT_FLOAT;
}

/**
* FPTR interface initialization functions
*/
Expand Down
1 change: 0 additions & 1 deletion dpnp/dpnp_algo/dpnp_algo.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
DPNP_FN_CHOOSE_EXT
DPNP_FN_CORRELATE_EXT
DPNP_FN_ERF_EXT
DPNP_FN_MEDIAN_EXT
DPNP_FN_MODF_EXT
DPNP_FN_PARTITION_EXT
DPNP_FN_RNG_BETA_EXT
Expand Down
52 changes: 0 additions & 52 deletions dpnp/dpnp_algo/dpnp_algo_statistics.pxi
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,17 +37,9 @@ and the rest of the library

__all__ += [
"dpnp_correlate",
"dpnp_median",
]


# C function pointer to the C library template functions
ctypedef c_dpctl.DPCTLSyclEventRef(*custom_statistic_1in_1out_func_ptr_t)(c_dpctl.DPCTLSyclQueueRef,
void *, void * , shape_elem_type * , size_t,
shape_elem_type * , size_t,
const c_dpctl.DPCTLEventVectorRef)


cpdef utils.dpnp_descriptor dpnp_correlate(utils.dpnp_descriptor x1, utils.dpnp_descriptor x2):
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(x2.dtype)
Expand Down Expand Up @@ -90,47 +82,3 @@ cpdef utils.dpnp_descriptor dpnp_correlate(utils.dpnp_descriptor x1, utils.dpnp_
c_dpctl.DPCTLEvent_Delete(event_ref)

return result


cpdef utils.dpnp_descriptor dpnp_median(utils.dpnp_descriptor array1):
cdef shape_type_c x1_shape = array1.shape
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype)

cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_MEDIAN_EXT, param1_type, param1_type)

array1_obj = array1.get_array()

cdef (DPNPFuncType, void *) ret_type_and_func = utils.get_ret_type_and_func(kernel_data,
array1_obj.sycl_device.has_aspect_fp64)
cdef DPNPFuncType return_type = ret_type_and_func[0]
cdef custom_statistic_1in_1out_func_ptr_t func = < custom_statistic_1in_1out_func_ptr_t > ret_type_and_func[1]

cdef utils.dpnp_descriptor result = utils.create_output_descriptor((1,),
return_type,
None,
device=array1_obj.sycl_device,
usm_type=array1_obj.usm_type,
sycl_queue=array1_obj.sycl_queue)

result_sycl_queue = result.get_array().sycl_queue

cdef c_dpctl.SyclQueue q = <c_dpctl.SyclQueue> result_sycl_queue
cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref()

# stub for interface support
cdef shape_type_c axis
cdef Py_ssize_t axis_size = 0

cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref,
array1.get_data(),
result.get_data(),
x1_shape.data(),
array1.ndim,
axis.data(),
axis_size,
NULL) # dep_events_ref

with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref)
c_dpctl.DPCTLEvent_Delete(event_ref)

return result
Loading
Loading