Skip to content

Commit

Permalink
Update dpnp.ediff1d() function (#1970)
Browse files Browse the repository at this point in the history 
* 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
  • Loading branch information
@vlad-perevezentsev
vlad-perevezentsev authored Aug 12, 2024
1 parent 5fda819 commit 05196fe
Show file tree
Hide file tree
Showing 10 changed files with 276 additions and 293 deletions.
47 changes: 0 additions & 47 deletions dpnp/backend/include/dpnp_iface.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -192,53 +192,6 @@ INP_DLLEXPORT void dpnp_dot_c(void *result_out,
const shape_elem_type *input2_shape,
const shape_elem_type *input2_strides);

/**
* @ingroup BACKEND_API
* @brief The differences between consecutive elements of an array.
*
* @param [in] q_ref Reference to SYCL queue.
* @param [out] result_out Output array.
* @param [in] result_size Size of output array.
* @param [in] result_ndim Number of output array dimensions.
* @param [in] result_shape Shape of output array.
* @param [in] result_strides Strides of output array.
* @param [in] input1_in First input array.
* @param [in] input1_size Size of first input array.
* @param [in] input1_ndim Number of first input array dimensions.
* @param [in] input1_shape Shape of first input array.
* @param [in] input1_strides Strides of first input array.
* @param [in] where Mask array.
* @param [in] dep_event_vec_ref Reference to vector of SYCL events.
*/
template <typename _DataType_input, typename _DataType_output>
INP_DLLEXPORT DPCTLSyclEventRef
dpnp_ediff1d_c(DPCTLSyclQueueRef q_ref,
void *result_out,
const size_t result_size,
const size_t result_ndim,
const shape_elem_type *result_shape,
const shape_elem_type *result_strides,
const void *input1_in,
const size_t input1_size,
const size_t input1_ndim,
const shape_elem_type *input1_shape,
const shape_elem_type *input1_strides,
const size_t *where,
const DPCTLEventVectorRef dep_event_vec_ref);

template <typename _DataType_input, typename _DataType_output>
INP_DLLEXPORT void dpnp_ediff1d_c(void *result_out,
const size_t result_size,
const size_t result_ndim,
const shape_elem_type *result_shape,
const shape_elem_type *result_strides,
const void *input1_in,
const size_t input1_size,
const size_t input1_ndim,
const shape_elem_type *input1_shape,
const shape_elem_type *input1_strides,
const size_t *where);

/**
* @ingroup BACKEND_API
* @brief Compute summary of input array elements.
Expand Down
9 changes: 3 additions & 6 deletions dpnp/backend/include/dpnp_iface_fptr.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -72,12 +72,9 @@ enum class DPNPFuncName : size_t
DPNP_FN_COV, /**< Used in numpy.cov() impl */
DPNP_FN_DOT, /**< Used in numpy.dot() impl */
DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */
DPNP_FN_EDIFF1D, /**< Used in numpy.ediff1d() impl */
DPNP_FN_EDIFF1D_EXT, /**< Used in numpy.ediff1d() impl, requires extra
parameters */
DPNP_FN_ERF, /**< Used in scipy.special.erf impl */
DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra
parameters */
DPNP_FN_ERF, /**< Used in scipy.special.erf impl */
DPNP_FN_ERF_EXT, /**< Used in scipy.special.erf impl, requires extra
parameters */
DPNP_FN_INITVAL, /**< Used in numpy ones, ones_like, zeros, zeros_like impls
*/
DPNP_FN_INITVAL_EXT, /**< Used in numpy ones, ones_like, zeros, zeros_like
Expand Down
142 changes: 0 additions & 142 deletions dpnp/backend/kernels/dpnp_krnl_mathematical.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -44,130 +44,6 @@ using dpctl::tensor::kernels::alignment_utils::required_alignment;
static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VECTOR_ABS_CHANGED,
"SYCL DPC++ compiler does not meet minimum version requirement");

template <typename _KernelNameSpecialization1,
typename _KernelNameSpecialization2>
class dpnp_ediff1d_c_kernel;

template <typename _DataType_input, typename _DataType_output>
DPCTLSyclEventRef dpnp_ediff1d_c(DPCTLSyclQueueRef q_ref,
void *result_out,
const size_t result_size,
const size_t result_ndim,
const shape_elem_type *result_shape,
const shape_elem_type *result_strides,
const void *input1_in,
const size_t input1_size,
const size_t input1_ndim,
const shape_elem_type *input1_shape,
const shape_elem_type *input1_strides,
const size_t *where,
const DPCTLEventVectorRef dep_event_vec_ref)
{
/* avoid warning unused variable*/
(void)result_ndim;
(void)result_shape;
(void)result_strides;
(void)input1_ndim;
(void)input1_shape;
(void)input1_strides;
(void)where;
(void)dep_event_vec_ref;

DPCTLSyclEventRef event_ref = nullptr;

if (!input1_size) {
return event_ref;
}

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

DPNPC_ptr_adapter<_DataType_input> input1_ptr(q_ref, input1_in,
input1_size);
DPNPC_ptr_adapter<_DataType_output> result_ptr(q_ref, result_out,
result_size, false, true);

_DataType_input *input1_data = input1_ptr.get_ptr();
_DataType_output *result = result_ptr.get_ptr();

sycl::event event;
sycl::range<1> gws(result_size);

auto kernel_parallel_for_func = [=](sycl::id<1> global_id) {
size_t output_id =
global_id[0]; /*for (size_t i = 0; i < result_size; ++i)*/
{
const _DataType_output curr_elem = input1_data[output_id];
const _DataType_output next_elem = input1_data[output_id + 1];
result[output_id] = next_elem - curr_elem;
}
};
auto kernel_func = [&](sycl::handler &cgh) {
cgh.parallel_for<
class dpnp_ediff1d_c_kernel<_DataType_input, _DataType_output>>(
gws, kernel_parallel_for_func);
};
event = q.submit(kernel_func);

input1_ptr.depends_on(event);
result_ptr.depends_on(event);
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);

return DPCTLEvent_Copy(event_ref);
}

template <typename _DataType_input, typename _DataType_output>
void dpnp_ediff1d_c(void *result_out,
const size_t result_size,
const size_t result_ndim,
const shape_elem_type *result_shape,
const shape_elem_type *result_strides,
const void *input1_in,
const size_t input1_size,
const size_t input1_ndim,
const shape_elem_type *input1_shape,
const shape_elem_type *input1_strides,
const size_t *where)
{
DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
DPCTLEventVectorRef dep_event_vec_ref = nullptr;
DPCTLSyclEventRef event_ref =
dpnp_ediff1d_c<_DataType_input, _DataType_output>(
q_ref, result_out, result_size, result_ndim, result_shape,
result_strides, input1_in, input1_size, input1_ndim, input1_shape,
input1_strides, where, dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
}

template <typename _DataType_input, typename _DataType_output>
void (*dpnp_ediff1d_default_c)(void *,
const size_t,
const size_t,
const shape_elem_type *,
const shape_elem_type *,
const void *,
const size_t,
const size_t,
const shape_elem_type *,
const shape_elem_type *,
const size_t *) =
dpnp_ediff1d_c<_DataType_input, _DataType_output>;

template <typename _DataType_input, typename _DataType_output>
DPCTLSyclEventRef (*dpnp_ediff1d_ext_c)(DPCTLSyclQueueRef,
void *,
const size_t,
const size_t,
const shape_elem_type *,
const shape_elem_type *,
const void *,
const size_t,
const size_t,
const shape_elem_type *,
const shape_elem_type *,
const size_t *,
const DPCTLEventVectorRef) =
dpnp_ediff1d_c<_DataType_input, _DataType_output>;

template <typename _KernelNameSpecialization1,
typename _KernelNameSpecialization2>
class dpnp_modf_c_kernel;
Expand Down Expand Up @@ -260,24 +136,6 @@ DPCTLSyclEventRef (*dpnp_modf_ext_c)(DPCTLSyclQueueRef,
void func_map_init_mathematical(func_map_t &fmap)
{

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

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

fmap[DPNPFuncName::DPNP_FN_MODF][eft_INT][eft_INT] = {
eft_DBL, (void *)dpnp_modf_default_c<int32_t, double>};
fmap[DPNPFuncName::DPNP_FN_MODF][eft_LNG][eft_LNG] = {
Expand Down
1 change: 0 additions & 1 deletion dpnp/dpnp_algo/dpnp_algo.pxd
View file
Original file line number Diff line number Diff line change
Expand Up @@ -35,7 +35,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
cdef enum DPNPFuncName "DPNPFuncName":
DPNP_FN_CHOOSE_EXT
DPNP_FN_CORRELATE_EXT
DPNP_FN_EDIFF1D_EXT
DPNP_FN_ERF_EXT
DPNP_FN_MEDIAN_EXT
DPNP_FN_MODF_EXT
Expand Down
57 changes: 0 additions & 57 deletions dpnp/dpnp_algo/dpnp_algo_mathematical.pxi
View file
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,6 @@ and the rest of the library
# NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file

__all__ += [
"dpnp_ediff1d",
"dpnp_modf",
]

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


cpdef utils.dpnp_descriptor dpnp_ediff1d(utils.dpnp_descriptor x1):

if x1.size <= 1:
return utils.dpnp_descriptor(dpnp.empty(0, dtype=x1.dtype)) # TODO need to call dpnp_empty instead

# Convert type (x1.dtype) to C enum DPNPFuncType
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)

# get the FPTR data structure
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_EDIFF1D_EXT, param1_type, param1_type)

result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)

# Currently shape and strides of the input array are not took into account for the function ediff1d
cdef shape_type_c x1_shape = (x1.size,)
cdef shape_type_c x1_strides = utils.strides_to_vector(None, x1_shape)

x1_obj = x1.get_array()

cdef shape_type_c result_shape = (x1.size - 1,)
cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape,
kernel_data.return_type,
None,
device=x1_obj.sycl_device,
usm_type=x1_obj.usm_type,
sycl_queue=x1_obj.sycl_queue)

cdef shape_type_c result_strides = utils.strides_to_vector(result.strides, result_shape)

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()

# Call FPTR function
cdef fptr_1in_1out_strides_t func = <fptr_1in_1out_strides_t > kernel_data.ptr
cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref,
result.get_data(),
result.size,
result.ndim,
result_shape.data(),
result_strides.data(),
x1.get_data(),
x1.size,
x1.ndim,
x1_shape.data(),
x1_strides.data(),
NULL,
NULL) # dep_events_ref

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

return result


cpdef tuple dpnp_modf(utils.dpnp_descriptor x1):
""" Convert string type names (array.dtype) to C enum DPNPFuncType """
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
Expand Down
Loading

0 comments on commit 05196fe

Please sign in to comment.