diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index e99921baa2d..6ceddf88962 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -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 -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 -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. diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index bc5d0042a35..8652a88b9b6 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -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 diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 7e358a8a710..7bb70bc63a6 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -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 -class dpnp_ediff1d_c_kernel; - -template -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(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(&event); - - return DPCTLEvent_Copy(event_ref); -} - -template -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(&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 -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 -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 class dpnp_modf_c_kernel; @@ -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}; - fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_ediff1d_default_c}; - fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ediff1d_default_c}; - fmap[DPNPFuncName::DPNP_FN_EDIFF1D][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ediff1d_default_c}; - - fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_INT][eft_INT] = { - eft_LNG, (void *)dpnp_ediff1d_ext_c}; - fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_ediff1d_ext_c}; - fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_ediff1d_ext_c}; - fmap[DPNPFuncName::DPNP_FN_EDIFF1D_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_ediff1d_ext_c}; - fmap[DPNPFuncName::DPNP_FN_MODF][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_modf_default_c}; fmap[DPNPFuncName::DPNP_FN_MODF][eft_LNG][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 8fb06cc58e4..489a20f064c 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -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 diff --git a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi index 84b004856bd..75a1141f3b0 100644 --- a/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_mathematical.pxi @@ -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", ] @@ -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 = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - # Call FPTR function - cdef fptr_1in_1out_strides_t func = 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) diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index e3c8efefdb2..98f4c6c8fa8 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -57,10 +57,7 @@ import dpnp.backend.extensions.ufunc._ufunc_impl as ufi from .backend.extensions.sycl_ext import _sycl_ext_impl -from .dpnp_algo import ( - dpnp_ediff1d, - dpnp_modf, -) +from .dpnp_algo import dpnp_modf from .dpnp_algo.dpnp_elementwise_common import ( DPNPAngle, DPNPBinaryFunc, @@ -1279,49 +1276,116 @@ def diff(a, n=1, axis=-1, prepend=None, append=None): ) -def ediff1d(x1, to_end=None, to_begin=None): +def ediff1d(ary, to_end=None, to_begin=None): """ The differences between consecutive elements of an array. For full documentation refer to :obj:`numpy.ediff1d`. - Limitations - ----------- - Parameter `x1`is supported as :class:`dpnp.ndarray`. - Keyword arguments `to_end` and `to_begin` are currently supported only - with default values `None`. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters + ---------- + ary : {dpnp.ndarray, usm_ndarray} + If necessary, will be flattened before the differences are taken. + to_end : array_like, optional + Number(s) to append at the end of the returned differences. + Default: ``None``. + to_begin : array_like, optional + Number(s) to prepend at the beginning of the returned differences. + Default: ``None``. + + Returns + ------- + out : dpnp.ndarray + New array consisting differences among succeeding elements. + Loosely, this is ``ary.flat[1:] - ary.flat[:-1]``. See Also -------- :obj:`dpnp.diff` : Calculate the n-th discrete difference along the given axis. + :obj:`dpnp.gradient` : Return the gradient of an N-dimensional array. Examples -------- >>> import dpnp as np - >>> a = np.array([1, 2, 4, 7, 0]) - >>> result = np.ediff1d(a) - >>> [x for x in result] - [1, 2, 3, -7] - >>> b = np.array([[1, 2, 4], [1, 6, 24]]) - >>> result = np.ediff1d(b) - >>> [x for x in result] - [1, 2, -3, 5, 18] + >>> x = np.array([1, 2, 4, 7, 0]) + >>> np.ediff1d(x) + array([ 1, 2, 3, -7]) + + >>> np.ediff1d(x, to_begin=-99, to_end=np.array([88, 99])) + array([-99, 1, 2, 3, -7, 88, 99]) + + The returned array is always 1D. + + >>> y = np.array([[1, 2, 4], [1, 6, 24]]) + >>> np.ediff1d(y) + array([ 1, 2, -3, 5, 18]) """ - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - if x1_desc: - if to_begin is not None: - pass - elif to_end is not None: - pass - else: - return dpnp_ediff1d(x1_desc).get_pyobj() + dpnp.check_supported_arrays_type(ary) + if ary.ndim > 1: + ary = dpnp.ravel(ary) + + # fast track default case + if to_begin is None and to_end is None: + return ary[1:] - ary[:-1] + + ary_dtype = ary.dtype + ary_usm_type = ary.usm_type + ary_sycl_queue = ary.sycl_queue + + if to_begin is None: + l_begin = 0 + else: + if not dpnp.is_supported_array_type(to_begin): + to_begin = dpnp.asarray( + to_begin, usm_type=ary_usm_type, sycl_queue=ary_sycl_queue + ) + if not dpnp.can_cast(to_begin, ary_dtype, casting="same_kind"): + raise TypeError( + "dtype of `to_begin` must be compatible " + "with input `ary` under the `same_kind` rule." + ) + + to_begin_ndim = to_begin.ndim + + if to_begin_ndim > 1: + to_begin = dpnp.ravel(to_begin) + + l_begin = to_begin.size + + if to_end is None: + l_end = 0 + else: + if not dpnp.is_supported_array_type(to_end): + to_end = dpnp.asarray( + to_end, usm_type=ary_usm_type, sycl_queue=ary_sycl_queue + ) + if not dpnp.can_cast(to_end, ary_dtype, casting="same_kind"): + raise TypeError( + "dtype of `to_end` must be compatible " + "with input `ary` under the `same_kind` rule." + ) + + to_end_ndim = to_end.ndim + + if to_end_ndim > 1: + to_end = dpnp.ravel(to_end) + + l_end = to_end.size + + # calculating using in place operation + l_diff = max(len(ary) - 1, 0) + result = dpnp.empty_like(ary, shape=l_diff + l_begin + l_end) + + if l_begin > 0: + result[:l_begin] = to_begin + if l_end > 0: + result[l_begin + l_diff :] = to_end + dpnp.subtract(ary[1:], ary[:-1], out=result[l_begin : l_begin + l_diff]) - return call_origin(numpy.ediff1d, x1, to_end=to_end, to_begin=to_begin) + return result _FABS_DOCSTRING = """ diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index b160601e939..b78371033ad 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -2011,9 +2011,7 @@ def test_power(array, val, data_type, val_type): class TestEdiff1d: - @pytest.mark.parametrize( - "data_type", get_all_dtypes(no_bool=True, no_complex=True) - ) + @pytest.mark.parametrize("data_type", get_all_dtypes(no_bool=True)) @pytest.mark.parametrize( "array", [ @@ -2023,7 +2021,7 @@ class TestEdiff1d: [[1, 2, 3], [5, 2, 8], [7, 3, 4]], ], ) - def test_ediff1d_int(self, array, data_type): + def test_ediff1d(self, array, data_type): np_a = numpy.array(array, dtype=data_type) dpnp_a = dpnp.array(array, dtype=data_type) @@ -2031,17 +2029,128 @@ def test_ediff1d_int(self, array, data_type): expected = numpy.ediff1d(np_a) assert_array_equal(expected, result) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") - def test_ediff1d_args(self): + @pytest.mark.parametrize( + "to_begin", + [ + -20, + numpy.array([-20, -30]), + dpnp.array([-20, -30]), + dpnp.array([[-20], [-30]]), + [1, 2], + (1, 2), + ], + ) + def test_ediff1d_to_begin(self, to_begin): + np_a = numpy.array([1, 2, 4, 7, 0]) + dpnp_a = dpnp.array([1, 2, 4, 7, 0]) + + if isinstance(to_begin, dpnp.ndarray): + np_to_begin = dpnp.asnumpy(to_begin) + else: + np_to_begin = to_begin + + result = dpnp.ediff1d(dpnp_a, to_begin=to_begin) + expected = numpy.ediff1d(np_a, to_begin=np_to_begin) + assert_array_equal(expected, result) + + @pytest.mark.parametrize( + "to_end", + [ + 20, + numpy.array([20, 15]), + dpnp.array([20, 15]), + dpnp.array([[-20], [-30]]), + [3, 4], + (3, 4), + ], + ) + def test_ediff1d_to_end(self, to_end): + np_a = numpy.array([1, 2, 4, 7, 0]) + dpnp_a = dpnp.array([1, 2, 4, 7, 0]) + + if isinstance(to_end, dpnp.ndarray): + np_to_end = dpnp.asnumpy(to_end) + else: + np_to_end = to_end + + result = dpnp.ediff1d(dpnp_a, to_end=to_end) + expected = numpy.ediff1d(np_a, to_end=np_to_end) + assert_array_equal(expected, result) + + @pytest.mark.parametrize( + "to_begin, to_end", + [ + (-20, 20), + (numpy.array([-20, -30]), numpy.array([20, 15])), + (dpnp.array([-20, -30]), dpnp.array([20, 15])), + (dpnp.array([[-20], [-30]]), dpnp.array([[20], [15]])), + ([1, 2], [3, 4]), + ((1, 2), (3, 4)), + ], + ) + def test_ediff1d_to_begin_to_end(self, to_begin, to_end): np_a = numpy.array([1, 2, 4, 7, 0]) + dpnp_a = dpnp.array([1, 2, 4, 7, 0]) + + if isinstance(to_begin, dpnp.ndarray): + np_to_begin = dpnp.asnumpy(to_begin) + else: + np_to_begin = to_begin - to_begin = numpy.array([-20, -30]) - to_end = numpy.array([20, 15]) + if isinstance(to_end, dpnp.ndarray): + np_to_end = dpnp.asnumpy(to_end) + else: + np_to_end = to_end - result = dpnp.ediff1d(np_a, to_end=to_end, to_begin=to_begin) - expected = numpy.ediff1d(np_a, to_end=to_end, to_begin=to_begin) + result = dpnp.ediff1d(dpnp_a, to_end=to_end, to_begin=to_begin) + expected = numpy.ediff1d(np_a, to_end=np_to_end, to_begin=np_to_begin) assert_array_equal(expected, result) + @pytest.mark.parametrize( + "to_begin, to_end", + [ + (-20, 20), + (dpt.asarray([-20, -30]), dpt.asarray([20, 15])), + (dpt.asarray([[-20, -30]]), dpt.asarray([[20, 15]])), + ([1, 2], [3, 4]), + ((1, 2), (3, 4)), + ], + ) + def test_ediff1d_usm_ndarray(self, to_begin, to_end): + np_a = numpy.array([[1, 2, 0]]) + dpt_a = dpt.asarray(np_a) + + if isinstance(to_begin, dpt.usm_ndarray): + np_to_begin = dpt.asnumpy(to_begin) + else: + np_to_begin = to_begin + + if isinstance(to_end, dpt.usm_ndarray): + np_to_end = dpt.asnumpy(to_end) + else: + np_to_end = to_end + + result = dpnp.ediff1d(dpt_a, to_end=to_end, to_begin=to_begin) + expected = numpy.ediff1d(np_a, to_end=np_to_end, to_begin=np_to_begin) + + assert_array_equal(expected, result) + assert isinstance(result, dpnp.ndarray) + + def test_ediff1d_errors(self): + a_dp = dpnp.array([[1, 2], [2, 5]]) + + # unsupported type + a_np = dpnp.asnumpy(a_dp) + assert_raises(TypeError, dpnp.ediff1d, a_np) + + # unsupported `to_begin` type according to the `same_kind` rules + to_begin = dpnp.array([-5], dtype="f4") + assert_raises(TypeError, dpnp.ediff1d, a_dp, to_begin=to_begin) + + # unsupported `to_end` type according to the `same_kind` rules + to_end = dpnp.array([5], dtype="f4") + assert_raises(TypeError, dpnp.ediff1d, a_dp, to_end=to_end) + @pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestTrapz: diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index 9a7899659d2..9c97c226c31 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -2385,3 +2385,36 @@ def test_nan_to_num(copy, device): assert_sycl_queue_equal(result.sycl_queue, a.sycl_queue) assert copy == (result is not a) + + +@pytest.mark.parametrize( + "device_x", + valid_devices, + ids=[device.filter_string for device in valid_devices], +) +@pytest.mark.parametrize( + "device_args", + valid_devices, + ids=[device.filter_string for device in valid_devices], +) +@pytest.mark.parametrize( + ["to_end", "to_begin"], + [ + (10, None), + (None, -10), + (10, -10), + ], +) +def test_ediff1d(device_x, device_args, to_end, to_begin): + data = [1, 3, 5, 7] + + x = dpnp.array(data, device=device_x) + if to_end: + to_end = dpnp.array(to_end, device=device_args) + + if to_begin: + to_begin = dpnp.array(to_begin, device=device_args) + + res = dpnp.ediff1d(x, to_end=to_end, to_begin=to_begin) + + assert_sycl_queue_equal(res.sycl_queue, x.sycl_queue) diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index a682ee9dce4..d440c4d5573 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -571,6 +571,7 @@ def test_norm(usm_type, ord, axis): pytest.param("degrees", [numpy.pi, numpy.pi / 2, 0]), pytest.param("diagonal", [[[1, 2], [3, 4]]]), pytest.param("diff", [1.0, 2.0, 4.0, 7.0, 0.0]), + pytest.param("ediff1d", [1.0, 2.0, 4.0, 7.0, 0.0]), pytest.param("exp", [1.0, 2.0, 4.0, 7.0]), pytest.param("exp2", [0.0, 1.0, 2.0]), pytest.param("expm1", [1.0e-10, 1.0, 2.0, 4.0, 7.0]), @@ -1392,3 +1393,30 @@ def test_nan_to_num(copy, usm_type_a): assert result.usm_type == usm_type_a assert copy == (result is not a) + + +@pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) +@pytest.mark.parametrize( + "usm_type_args", list_of_usm_types, ids=list_of_usm_types +) +@pytest.mark.parametrize( + ["to_end", "to_begin"], + [ + (10, None), + (None, -10), + (10, -10), + ], +) +def test_ediff1d(usm_type_x, usm_type_args, to_end, to_begin): + data = [1, 3, 5, 7] + + x = dp.array(data, usm_type=usm_type_x) + if to_end: + to_end = dp.array(to_end, usm_type=usm_type_args) + + if to_begin: + to_begin = dp.array(to_begin, usm_type=usm_type_args) + + res = dp.ediff1d(x, to_end=to_end, to_begin=to_begin) + + assert res.usm_type == x.usm_type diff --git a/tests/third_party/cupy/math_tests/test_sumprod.py b/tests/third_party/cupy/math_tests/test_sumprod.py index 5d8c4053cd5..ba57994ca6c 100644 --- a/tests/third_party/cupy/math_tests/test_sumprod.py +++ b/tests/third_party/cupy/math_tests/test_sumprod.py @@ -860,7 +860,6 @@ def test_gradient_bool_input(self): xp.gradient(x) -@pytest.mark.skip("ediff1d() is not implemented yet") class TestEdiff1d: @testing.for_all_dtypes(no_bool=True) @testing.numpy_cupy_allclose()