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

Reuse OneDPL implementation of std::nth_element() for partition of 1D array #1406

Merged
merged 4 commits into from
May 30, 2023
Merged

Reuse OneDPL implementation of std::nth_element() for partition of 1D array #1406

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
5e1d133
Reuse OneDPL implementation of std::nth_element() for partition of 1D…
antonwolfy May 22, 2023
bd8be5b
Add complex type in dpnp.copy() and tune tests
antonwolfy May 22, 2023
3b29372
Added dpnp.ndarray.partition() & unmuted cupy tests
antonwolfy May 23, 2023
15ae06c
Merge branch 'master' into fix_crash_in_partition
antonwolfy May 30, 2023
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
2 changes: 2 additions & 0 deletions .github/workflows/conda-package.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,8 +18,10 @@ env:
test_linalg.py
test_mathematical.py
test_random_state.py
test_sort.py
test_special.py
test_usm_type.py
third_party/cupy/sorting_tests/test_sort.py
VER_JSON_NAME: 'version.json'
VER_SCRIPT1: "import json; f = open('version.json', 'r'); j = json.load(f); f.close(); "
VER_SCRIPT2: "d = j['dpnp'][0]; print('='.join((d[s] for s in ('version', 'build'))))"
Expand Down
4 changes: 2 additions & 2 deletions dpnp/backend/kernels/dpnp_krnl_elemwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -811,8 +811,8 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_copy_c_ext<int64_t>};
fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_copy_c_ext<float>};
fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_copy_c_ext<double>};
fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_C128][eft_C128] = {eft_C128,
(void*)dpnp_copy_c_ext<std::complex<double>>};
fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_C64][eft_C64] = {eft_C64, (void*)dpnp_copy_c_ext<std::complex<float>>};
fmap[DPNPFuncName::DPNP_FN_COPY_EXT][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_copy_c_ext<std::complex<double>>};

fmap[DPNPFuncName::DPNP_FN_ERF][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_erf_c_default<int32_t>};
fmap[DPNPFuncName::DPNP_FN_ERF][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_erf_c_default<int64_t>};
Expand Down
25 changes: 23 additions & 2 deletions dpnp/backend/kernels/dpnp_krnl_sorting.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
//*****************************************************************************
// Copyright (c) 2016-2020, Intel Corporation
// Copyright (c) 2016-2023, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
Expand Down Expand Up @@ -160,6 +160,24 @@ DPCTLSyclEventRef dpnp_partition_c(DPCTLSyclQueueRef q_ref,

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

if (ndim == 1) // 1d array with C-contiguous data
{
_DataType* arr = static_cast<_DataType*>(array1_in);
_DataType* result = static_cast<_DataType*>(result1);

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

// fill the result array with data from input one
q.memcpy(result, arr, size * sizeof(_DataType)).wait();

// make a partial sorting such that:
// 1. result[0 <= i < kth] <= result[kth]
// 2. result[kth <= i < size] >= result[kth]
// event-blocking call, no need for wait()
std::nth_element(policy, result, result + kth, result + size, dpnp_less_comp());
return event_ref;
}

DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size, true);
DPNPC_ptr_adapter<_DataType> input2_ptr(q_ref, array2_in, size, true);
DPNPC_ptr_adapter<_DataType> result1_ptr(q_ref, result1, size, true, true);
Expand All @@ -181,7 +199,7 @@ DPCTLSyclEventRef dpnp_partition_c(DPCTLSyclQueueRef q_ref,
size_t ind = j - ind_begin;
matrix[ind] = arr2[j];
}
std::partial_sort(matrix, matrix + shape_[ndim - 1], matrix + shape_[ndim - 1]);
std::partial_sort(matrix, matrix + shape_[ndim - 1], matrix + shape_[ndim - 1], dpnp_less_comp());
for (size_t j = ind_begin; j < ind_end + 1; ++j)
{
size_t ind = j - ind_begin;
Expand Down Expand Up @@ -492,10 +510,13 @@ void func_map_init_sorting(func_map_t& fmap)
fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_partition_default_c<float>};
fmap[DPNPFuncName::DPNP_FN_PARTITION][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_partition_default_c<double>};

fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_BLN][eft_BLN] = {eft_BLN, (void*)dpnp_partition_ext_c<bool>};
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_partition_ext_c<int32_t>};
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_partition_ext_c<int64_t>};
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_partition_ext_c<float>};
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_partition_ext_c<double>};
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_C64][eft_C64] = {eft_C64, (void*)dpnp_partition_ext_c<std::complex<float>>};
fmap[DPNPFuncName::DPNP_FN_PARTITION_EXT][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_partition_ext_c<std::complex<double>>};

fmap[DPNPFuncName::DPNP_FN_SEARCHSORTED][eft_INT][eft_INT] = {
eft_INT, (void*)dpnp_searchsorted_default_c<int32_t, int64_t>};
Expand Down
53 changes: 53 additions & 0 deletions dpnp/backend/src/dpnp_fptr.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -187,6 +187,59 @@ constexpr auto both_types_are_any_of = std::conjunction_v<is_any<T1, Ts...>, is_
template <typename T1, typename T2, typename... Ts>
constexpr auto none_of_both_types = !std::disjunction_v<is_any<T1, Ts...>, is_any<T2, Ts...>>;


/**
* @brief If the type _Tp is a reference type, provides the member typedef type which is the type referred to by _Tp
* with its topmost cv-qualifiers removed. Otherwise type is _Tp with its topmost cv-qualifiers removed.
*
* @note std::remove_cvref is only available since c++20
*/
template<typename _Tp>
using dpnp_remove_cvref_t = typename std::remove_cv_t<typename std::remove_reference_t<_Tp>>;


/**
* @brief "<" comparison with complex types support.
*
* @note return a result of lexicographical "<" comparison for complex types.
*/
class dpnp_less_comp
{
public:
template <typename _Xp, typename _Yp>
bool operator()(_Xp&& __x, _Yp&& __y) const
{
if constexpr (both_types_are_same<dpnp_remove_cvref_t<_Xp>, dpnp_remove_cvref_t<_Yp>, std::complex<float>, std::complex<double>>)
{
bool ret = false;
_Xp a = std::forward<_Xp>(__x);
_Yp b = std::forward<_Yp>(__y);

if (a.real() < b.real())
{
ret = (a.imag() == a.imag() || b.imag() != b.imag());
}
else if (a.real() > b.real())
{
ret = (b.imag() != b.imag() && a.imag() == a.imag());
}
else if (a.real() == b.real() || (a.real() != a.real() && b.real() != b.real()))
{
ret = (a.imag() < b.imag() || (b.imag() != b.imag() && a.imag() == a.imag()));
}
else
{
ret = (b.real() != b.real());
}
return ret;
}
else
{
return std::forward<_Xp>(__x) < std::forward<_Yp>(__y);
}
}
};

/**
* FPTR interface initialization functions
*/
Expand Down
27 changes: 26 additions & 1 deletion dpnp/dpnp_array.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -832,7 +832,32 @@ def ndim(self):
def nonzero(self):
return dpnp.nonzero(self)

# 'partition',
def partition(self, kth, axis=-1, kind='introselect', order=None):
"""
Rearranges the elements in the array in such a way that the value of the
element in kth position is in the position it would be in a sorted array.

All elements smaller than the kth element are moved before this element and
all equal or greater are moved behind it. The ordering of the elements in
the two partitions is undefined.

Refer to `dpnp.partition` for full documentation.

See Also
--------
:obj:`dpnp.partition` : Return a partitioned copy of an array.

Examples
--------
>>> import dpnp as np
>>> a = np.array([3, 4, 2, 1])
>>> a.partition(3)
>>> a
array([1, 2, 3, 4])

"""

self._array_obj = dpnp.partition(self, kth, axis=axis, kind=kind, order=order).get_array()

def prod(self, axis=None, dtype=None, out=None, keepdims=False, initial=None, where=True):
"""
Expand Down
40 changes: 2 additions & 38 deletions tests/skipped_tests.tbl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -202,8 +202,6 @@ tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: (dpnp
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray([(i, i) for i in x], [("a", object), ("b", dpnp.int32)])]]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[lambda x: dpnp.asarray(x).astype(dpnp.int8)]

tests/test_sort.py::test_partition[[[1, 0], [3, 0]]-float32-1]

tests/third_party/cupy/core_tests/test_ndarray_complex_ops.py::TestAngle::test_angle
tests/third_party/cupy/core_tests/test_ndarray_complex_ops.py::TestRealImag::test_imag
tests/third_party/cupy/core_tests/test_ndarray_complex_ops.py::TestRealImag::test_imag_inplace
Expand Down Expand Up @@ -1181,18 +1179,7 @@ tests/third_party/cupy/sorting_tests/test_sort.py::TestArgsort_param_0_{external
tests/third_party/cupy/sorting_tests/test_sort.py::TestArgsort_param_0_{external=False}::test_argsort_original_array_not_modified_multi_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestArgsort_param_0_{external=False}::test_argsort_original_array_not_modified_one_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestArgsort_param_0_{external=False}::test_argsort_zero_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_axis1
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_axis2
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_negative_axis1
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_negative_axis2
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_invalid_negative_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_axis1
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_axis2
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_negative_axis1
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_negative_axis2
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_invalid_negative_kth

tests/third_party/cupy/sorting_tests/test_sort.py::TestLexsort::test_F_order
tests/third_party/cupy/sorting_tests/test_sort.py::TestLexsort::test_lexsort_dtype
tests/third_party/cupy/sorting_tests/test_sort.py::TestLexsort::test_lexsort_three_or_more_dim
Expand All @@ -1202,30 +1189,7 @@ tests/third_party/cupy/sorting_tests/test_sort.py::TestLexsort::test_nan3
tests/third_party/cupy/sorting_tests/test_sort.py::TestLexsort::test_view
tests/third_party/cupy/sorting_tests/test_sort.py::TestMsort::test_msort_multi_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestMsort::test_msort_one_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_axis
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_multi_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_negative_axis
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_negative_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_non_contiguous
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_one_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_sequence_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_0_{external=False, length=10}::test_partition_zero_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_axis
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_multi_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_negative_axis
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_negative_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_non_contiguous
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_one_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_sequence_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_1_{external=False, length=20000}::test_partition_zero_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_multi_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_negative_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_one_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_2_{external=True, length=10}::test_partition_zero_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_multi_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_negative_kth
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_one_dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestPartition_param_3_{external=True, length=20000}::test_partition_zero_dim

tests/third_party/cupy/sorting_tests/test_sort.py::TestSort_complex::test_sort_complex_1dim
tests/third_party/cupy/sorting_tests/test_sort.py::TestSort_complex::test_sort_complex_nan
tests/third_party/cupy/sorting_tests/test_sort.py::TestSort_complex::test_sort_complex_ndim
Expand Down
Loading