Skip to content
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

Update dpnp.ediff1d() function #1970

Merged
merged 15 commits into from
Aug 12, 2024
Merged

Update dpnp.ediff1d() function #1970

Show file tree
Hide file tree
Changes from 11 commits
Commits
Show all changes
15 commits
Select commit Hold shift + click to select a range
5d79f76
Update dpnp.ediff1d() function
vlad-perevezentsev Aug 5, 2024
33149e5
Remove old implementation of dpnp_ediff1d
vlad-perevezentsev Aug 5, 2024
67fd9a5
Update cupy tests
vlad-perevezentsev Aug 5, 2024
5048a9e
Add new tests for dpnp.ediff1d()
vlad-perevezentsev Aug 5, 2024
5a11403
Merge master into impl_ediff1d
vlad-perevezentsev Aug 6, 2024
d1f89c5
Remove old implementation of dpnp_ediff1d_c
vlad-perevezentsev Aug 6, 2024
58bffa2
Update docstrings for ediff1d
vlad-perevezentsev Aug 6, 2024
cb02fe4
Update TestEdiff1d to increase coverage
vlad-perevezentsev Aug 6, 2024
68e3589
Applied review comments
vlad-perevezentsev Aug 6, 2024
df2c705
Add support usm_ndarray
vlad-perevezentsev Aug 6, 2024
fce3ded
Dont use asarray for supported array type
vlad-perevezentsev Aug 6, 2024
7e7703d
Applied review comments
vlad-perevezentsev Aug 7, 2024
5e52bbf
Merge master into impl_ediff1d
vlad-perevezentsev Aug 12, 2024
22d9cf9
Merge master into impl_ediff1d
vlad-perevezentsev Aug 12, 2024
dcc6d3c
Merge master into impl_ediff1d
vlad-perevezentsev Aug 12, 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
47 changes: 0 additions & 47 deletions dpnp/backend/include/dpnp_iface.hpp
View file
Open in desktop
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
Open in desktop
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
Open in desktop
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
Open in desktop
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
Open in desktop
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
Loading