From 208d2acb1b0660480c7ffeebae8d3e1d9525e17a Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 19 Mar 2023 15:26:27 -0700 Subject: [PATCH 1/4] Changes to integer indexing modes - Re-implements ClipIndex to clip indices from 0 <= i < n - Adds FancyIndex class, which clips indices to -n <= i < n then wraps negative indices (old ClipIndex) - Integer indexing functions have "fancy" mode (old "clip") and "clip" mode (uses new behavior) - dpctl.tensor.take and dpctl.tensor.put now default to "fancy" mode rather than "clip" mode - Indexing mode tests now test for new clip behavior --- dpctl/tensor/_indexing_functions.py | 46 +++++++++-------- .../kernels/integer_advanced_indexing.hpp | 49 ++++++++++++++++++- .../source/integer_advanced_indexing.cpp | 24 ++++++--- dpctl/tests/test_usm_ndarray_indexing.py | 15 ++++-- 4 files changed, 102 insertions(+), 32 deletions(-) diff --git a/dpctl/tensor/_indexing_functions.py b/dpctl/tensor/_indexing_functions.py index f166e94f49..b54fb9b133 100644 --- a/dpctl/tensor/_indexing_functions.py +++ b/dpctl/tensor/_indexing_functions.py @@ -26,8 +26,19 @@ from ._copy_utils import _extract_impl, _nonzero_impl -def take(x, indices, /, *, axis=None, mode="clip"): - """take(x, indices, axis=None, mode="clip") +def _get_indexing_mode(name): + modes = {"fancy": 0, "clip": 1, "wrap": 2} + try: + return modes[name] + except KeyError: + raise ValueError( + "`mode` must be `fancy`, `clip`, or `wrap`." + "Got `{}`.".format(name) + ) + + +def take(x, indices, /, *, axis=None, mode="fancy"): + """take(x, indices, axis=None, mode="fancy") Takes elements from array along a given axis. @@ -42,15 +53,16 @@ def take(x, indices, /, *, axis=None, mode="clip"): Default: `None`. mode: How out-of-bounds indices will be handled. - "clip" - clamps indices to (-n <= i < n), then wraps + "fancy" - clamps indices to (-n <= i < n), then wraps negative indices. + "clip" - clips indices to (0 <= i < n) "wrap" - wraps both negative and positive indices. - Default: `"clip"`. + Default: `"fancy"`. Returns: out: usm_ndarray Array with shape x.shape[:axis] + indices.shape + x.shape[axis + 1:] - filled with elements . + filled with elements from x. """ if not isinstance(x, dpt.usm_ndarray): raise TypeError( @@ -80,11 +92,7 @@ def take(x, indices, /, *, axis=None, mode="clip"): [x.usm_type, indices.usm_type] ) - modes = {"clip": 0, "wrap": 1} - try: - mode = modes[mode] - except KeyError: - raise ValueError("`mode` must be `clip` or `wrap`.") + mode = _get_indexing_mode(mode) x_ndim = x.ndim if axis is None: @@ -114,8 +122,8 @@ def take(x, indices, /, *, axis=None, mode="clip"): return res -def put(x, indices, vals, /, *, axis=None, mode="clip"): - """put(x, indices, vals, axis=None, mode="clip") +def put(x, indices, vals, /, *, axis=None, mode="fancy"): + """put(x, indices, vals, axis=None, mode="fancy") Puts values of an array into another array along a given axis. @@ -134,10 +142,11 @@ def put(x, indices, vals, /, *, axis=None, mode="clip"): Default: `None`. mode: How out-of-bounds indices will be handled. - "clip" - clamps indices to (-axis_size <= i < axis_size), - then wraps negative indices. + "fancy" - clamps indices to (-n <= i < n), then wraps + negative indices. + "clip" - clips indices to (0 <= i < n) "wrap" - wraps both negative and positive indices. - Default: `"clip"`. + Default: `"fancy"`. """ if not isinstance(x, dpt.usm_ndarray): raise TypeError( @@ -175,11 +184,8 @@ def put(x, indices, vals, /, *, axis=None, mode="clip"): if exec_q is None: raise dpctl.utils.ExecutionPlacementError vals_usm_type = dpctl.utils.get_coerced_usm_type(usm_types_) - modes = {"clip": 0, "wrap": 1} - try: - mode = modes[mode] - except KeyError: - raise ValueError("`mode` must be `clip` or `wrap`.") + + mode = _get_indexing_mode(mode) x_ndim = x.ndim if axis is None: diff --git a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp index 07b0a4e9b8..1c9fd39af2 100644 --- a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp @@ -46,10 +46,10 @@ namespace py = pybind11; template class take_kernel; template class put_kernel; -class ClipIndex +class FancyIndex { public: - ClipIndex() = default; + FancyIndex() = default; void operator()(py::ssize_t max_item, py::ssize_t &ind) const { @@ -60,6 +60,19 @@ class ClipIndex } }; +class ClipIndex +{ +public: + ClipIndex() = default; + + void operator()(py::ssize_t max_item, py::ssize_t &ind) const + { + max_item = std::max(max_item, 1); + ind = std::clamp(ind, 0, max_item - 1); + return; + } +}; + class WrapIndex { public: @@ -348,6 +361,22 @@ sycl::event put_impl(sycl::queue q, return put_ev; } +template struct TakeFancyFactory +{ + fnT get() + { + if constexpr (std::is_integral::value && + !std::is_same::value) { + fnT fn = take_impl; + return fn; + } + else { + fnT fn = nullptr; + return fn; + } + } +}; + template struct TakeWrapFactory { fnT get() @@ -380,6 +409,22 @@ template struct TakeClipFactory } }; +template struct PutFancyFactory +{ + fnT get() + { + if constexpr (std::is_integral::value && + !std::is_same::value) { + fnT fn = put_impl; + return fn; + } + else { + fnT fn = nullptr; + return fn; + } + } +}; + template struct PutWrapFactory { fnT get() diff --git a/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp b/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp index cb148c7df3..936a13e2bc 100644 --- a/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp +++ b/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp @@ -39,9 +39,10 @@ #include "integer_advanced_indexing.hpp" -#define INDEXING_MODES 2 -#define CLIP_MODE 0 -#define WRAP_MODE 1 +#define INDEXING_MODES 3 +#define FANCY_MODE 0 +#define CLIP_MODE 1 +#define WRAP_MODE 2 namespace dpctl { @@ -252,8 +253,8 @@ usm_ndarray_take(dpctl::tensor::usm_ndarray src, throw py::value_error("Axis cannot be negative."); } - if (mode != 0 && mode != 1) { - throw py::value_error("Mode must be 0 or 1."); + if (mode != 0 && mode != 1 && mode != 2) { + throw py::value_error("Mode must be 0, 1, or 2."); } const dpctl::tensor::usm_ndarray ind_rep = ind[0]; @@ -575,8 +576,8 @@ usm_ndarray_put(dpctl::tensor::usm_ndarray dst, throw py::value_error("Axis cannot be negative."); } - if (mode != 0 && mode != 1) { - throw py::value_error("Mode must be 0 or 1."); + if (mode != 0 && mode != 1 && mode != 2) { + throw py::value_error("Mode must be 0, 1, or 2."); } if (!dst.is_writable()) { @@ -883,6 +884,11 @@ void init_advanced_indexing_dispatch_tables(void) { using namespace dpctl::tensor::detail; + using dpctl::tensor::kernels::indexing::TakeFancyFactory; + DispatchTableBuilder + dtb_takefancy; + dtb_takefancy.populate_dispatch_table(take_dispatch_table[FANCY_MODE]); + using dpctl::tensor::kernels::indexing::TakeClipFactory; DispatchTableBuilder dtb_takeclip; @@ -893,6 +899,10 @@ void init_advanced_indexing_dispatch_tables(void) dtb_takewrap; dtb_takewrap.populate_dispatch_table(take_dispatch_table[WRAP_MODE]); + using dpctl::tensor::kernels::indexing::PutFancyFactory; + DispatchTableBuilder dtb_putfancy; + dtb_putfancy.populate_dispatch_table(put_dispatch_table[FANCY_MODE]); + using dpctl::tensor::kernels::indexing::PutClipFactory; DispatchTableBuilder dtb_putclip; dtb_putclip.populate_dispatch_table(put_dispatch_table[CLIP_MODE]); diff --git a/dpctl/tests/test_usm_ndarray_indexing.py b/dpctl/tests/test_usm_ndarray_indexing.py index 5d160bf0f1..d6d427fcf8 100644 --- a/dpctl/tests/test_usm_ndarray_indexing.py +++ b/dpctl/tests/test_usm_ndarray_indexing.py @@ -895,11 +895,20 @@ def test_integer_indexing_modes(): q = get_queue_or_skip() x = dpt.arange(5, sycl_queue=q) + x_np = dpt.asnumpy(x) - # wrapping ind = dpt.asarray([-6, -3, 0, 2, 6], dtype=np.intp, sycl_queue=q) + ind_np = dpt.asnumpy(ind) + + # wrapping res = dpt.take(x, ind, mode="wrap") - expected_arr = np.take(dpt.asnumpy(x), dpt.asnumpy(ind), mode="wrap") + expected_arr = np.take(x_np, ind_np, mode="wrap") + + assert (dpt.asnumpy(res) == expected_arr).all() + + # clipping to 0 (disabling negative indices) + res = dpt.take(x, ind, mode="clip") + expected_arr = np.take(x_np, ind_np, mode="clip") assert (dpt.asnumpy(res) == expected_arr).all() @@ -907,7 +916,7 @@ def test_integer_indexing_modes(): # where n is the axis length ind = dpt.asarray([-4, -3, 0, 2, 4], dtype=np.intp, sycl_queue=q) - res = dpt.take(x, ind, mode="clip") + res = dpt.take(x, ind, mode="fancy") expected_arr = np.take(dpt.asnumpy(x), dpt.asnumpy(ind), mode="raise") assert (dpt.asnumpy(res) == expected_arr).all() From 762042213f22b7a7df5bb9b583fb69ec9785a7fa Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Mon, 20 Mar 2023 06:57:20 -0500 Subject: [PATCH 2/4] Improve _indexing_functions.py coverage to 100% --- dpctl/tests/test_usm_ndarray_indexing.py | 49 ++++++++++++++++++++++++ 1 file changed, 49 insertions(+) diff --git a/dpctl/tests/test_usm_ndarray_indexing.py b/dpctl/tests/test_usm_ndarray_indexing.py index d6d427fcf8..5637ccb580 100644 --- a/dpctl/tests/test_usm_ndarray_indexing.py +++ b/dpctl/tests/test_usm_ndarray_indexing.py @@ -20,6 +20,7 @@ from helper import get_queue_or_skip, skip_if_dtype_not_supported from numpy.testing import assert_array_equal +import dpctl import dpctl.tensor as dpt from dpctl.utils import ExecutionPlacementError @@ -948,6 +949,10 @@ def test_take_arg_validation(): dpt.take(dpt.reshape(x, (2, 2)), ind0, axis=None) with pytest.raises(ValueError): dpt.take(x, dpt.reshape(ind0, (2, 2))) + with pytest.raises(ValueError): + dpt.take(x[0], ind0, axis=2) + with pytest.raises(ValueError): + dpt.take(x[:, dpt.newaxis, dpt.newaxis], ind0, axis=None) def test_put_arg_validation(): @@ -977,6 +982,10 @@ def test_put_arg_validation(): dpt.put(x, ind0, val, mode=0) with pytest.raises(ValueError): dpt.put(x, dpt.reshape(ind0, (2, 2)), val) + with pytest.raises(ValueError): + dpt.put(x[0], ind0, val, axis=2) + with pytest.raises(ValueError): + dpt.put(x[:, dpt.newaxis, dpt.newaxis], ind0, val, axis=None) def test_advanced_indexing_compute_follows_data(): @@ -1278,3 +1287,43 @@ def test_nonzero_large(): m = dpt.full((30, 60, 80), True) assert m[m].size == m.size + + +def test_extract_arg_validation(): + get_queue_or_skip() + with pytest.raises(TypeError): + dpt.extract(None, None) + cond = dpt.ones(10, dtype="?") + with pytest.raises(TypeError): + dpt.extract(cond, None) + q1 = dpctl.SyclQueue() + with pytest.raises(ExecutionPlacementError): + dpt.extract(cond.to_device(q1), dpt.zeros_like(cond, dtype="u1")) + with pytest.raises(ValueError): + dpt.extract(dpt.ones((2, 3), dtype="?"), dpt.ones((3, 2), dtype="i1")) + + +def test_place_arg_validation(): + get_queue_or_skip() + with pytest.raises(TypeError): + dpt.place(None, None, None) + arr = dpt.zeros(8, dtype="i1") + with pytest.raises(TypeError): + dpt.place(arr, None, None) + cond = dpt.ones(8, dtype="?") + with pytest.raises(TypeError): + dpt.place(arr, cond, None) + vals = dpt.ones_like(arr) + q1 = dpctl.SyclQueue() + with pytest.raises(ExecutionPlacementError): + dpt.place(arr.to_device(q1), cond, vals) + with pytest.raises(ValueError): + dpt.place(dpt.reshape(arr, (2, 2, 2)), cond, vals) + + +def test_nonzero_arg_validation(): + get_queue_or_skip() + with pytest.raises(TypeError): + dpt.nonzero(list()) + with pytest.raises(ValueError): + dpt.nonzero(dpt.asarray(1)) From 4e06ba951bb278ebbe6e0803221251e8ac4346cd Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 20 Mar 2023 09:09:27 -0700 Subject: [PATCH 3/4] Renamed "fancy" indexing mode to "default" --- dpctl/tensor/_indexing_functions.py | 20 ++++++++++---------- dpctl/tests/test_usm_ndarray_indexing.py | 2 +- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/dpctl/tensor/_indexing_functions.py b/dpctl/tensor/_indexing_functions.py index b54fb9b133..86067be1f2 100644 --- a/dpctl/tensor/_indexing_functions.py +++ b/dpctl/tensor/_indexing_functions.py @@ -27,18 +27,18 @@ def _get_indexing_mode(name): - modes = {"fancy": 0, "clip": 1, "wrap": 2} + modes = {"default": 0, "clip": 1, "wrap": 2} try: return modes[name] except KeyError: raise ValueError( - "`mode` must be `fancy`, `clip`, or `wrap`." + "`mode` must be `default`, `clip`, or `wrap`." "Got `{}`.".format(name) ) -def take(x, indices, /, *, axis=None, mode="fancy"): - """take(x, indices, axis=None, mode="fancy") +def take(x, indices, /, *, axis=None, mode="default"): + """take(x, indices, axis=None, mode="default") Takes elements from array along a given axis. @@ -53,11 +53,11 @@ def take(x, indices, /, *, axis=None, mode="fancy"): Default: `None`. mode: How out-of-bounds indices will be handled. - "fancy" - clamps indices to (-n <= i < n), then wraps + "default" - clamps indices to (-n <= i < n), then wraps negative indices. "clip" - clips indices to (0 <= i < n) "wrap" - wraps both negative and positive indices. - Default: `"fancy"`. + Default: `"default"`. Returns: out: usm_ndarray @@ -122,8 +122,8 @@ def take(x, indices, /, *, axis=None, mode="fancy"): return res -def put(x, indices, vals, /, *, axis=None, mode="fancy"): - """put(x, indices, vals, axis=None, mode="fancy") +def put(x, indices, vals, /, *, axis=None, mode="default"): + """put(x, indices, vals, axis=None, mode="default") Puts values of an array into another array along a given axis. @@ -142,11 +142,11 @@ def put(x, indices, vals, /, *, axis=None, mode="fancy"): Default: `None`. mode: How out-of-bounds indices will be handled. - "fancy" - clamps indices to (-n <= i < n), then wraps + "default" - clamps indices to (-n <= i < n), then wraps negative indices. "clip" - clips indices to (0 <= i < n) "wrap" - wraps both negative and positive indices. - Default: `"fancy"`. + Default: `"default"`. """ if not isinstance(x, dpt.usm_ndarray): raise TypeError( diff --git a/dpctl/tests/test_usm_ndarray_indexing.py b/dpctl/tests/test_usm_ndarray_indexing.py index 5637ccb580..d2d26e6498 100644 --- a/dpctl/tests/test_usm_ndarray_indexing.py +++ b/dpctl/tests/test_usm_ndarray_indexing.py @@ -917,7 +917,7 @@ def test_integer_indexing_modes(): # where n is the axis length ind = dpt.asarray([-4, -3, 0, 2, 4], dtype=np.intp, sycl_queue=q) - res = dpt.take(x, ind, mode="fancy") + res = dpt.take(x, ind, mode="default") expected_arr = np.take(dpt.asnumpy(x), dpt.asnumpy(ind), mode="raise") assert (dpt.asnumpy(res) == expected_arr).all() From a1078c7449a6e774ed1a6ccb8719efb104e8d438 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 20 Mar 2023 14:00:28 -0700 Subject: [PATCH 4/4] Integer indexing "wrap" mode now default - For performance reasons, "wrap" now clips positive indices and wraps negative indices --- dpctl/tensor/_indexing_functions.py | 23 ++++----- .../kernels/integer_advanced_indexing.hpp | 50 +------------------ .../source/integer_advanced_indexing.cpp | 14 +----- dpctl/tests/test_usm_ndarray_indexing.py | 20 +++----- 4 files changed, 20 insertions(+), 87 deletions(-) diff --git a/dpctl/tensor/_indexing_functions.py b/dpctl/tensor/_indexing_functions.py index 86067be1f2..a2301b5fb2 100644 --- a/dpctl/tensor/_indexing_functions.py +++ b/dpctl/tensor/_indexing_functions.py @@ -27,18 +27,17 @@ def _get_indexing_mode(name): - modes = {"default": 0, "clip": 1, "wrap": 2} + modes = {"wrap": 0, "clip": 1} try: return modes[name] except KeyError: raise ValueError( - "`mode` must be `default`, `clip`, or `wrap`." - "Got `{}`.".format(name) + "`mode` must be `wrap` or `clip`." "Got `{}`.".format(name) ) -def take(x, indices, /, *, axis=None, mode="default"): - """take(x, indices, axis=None, mode="default") +def take(x, indices, /, *, axis=None, mode="wrap"): + """take(x, indices, axis=None, mode="wrap") Takes elements from array along a given axis. @@ -53,11 +52,10 @@ def take(x, indices, /, *, axis=None, mode="default"): Default: `None`. mode: How out-of-bounds indices will be handled. - "default" - clamps indices to (-n <= i < n), then wraps + "wrap" - clamps indices to (-n <= i < n), then wraps negative indices. "clip" - clips indices to (0 <= i < n) - "wrap" - wraps both negative and positive indices. - Default: `"default"`. + Default: `"wrap"`. Returns: out: usm_ndarray @@ -122,8 +120,8 @@ def take(x, indices, /, *, axis=None, mode="default"): return res -def put(x, indices, vals, /, *, axis=None, mode="default"): - """put(x, indices, vals, axis=None, mode="default") +def put(x, indices, vals, /, *, axis=None, mode="wrap"): + """put(x, indices, vals, axis=None, mode="wrap") Puts values of an array into another array along a given axis. @@ -142,11 +140,10 @@ def put(x, indices, vals, /, *, axis=None, mode="default"): Default: `None`. mode: How out-of-bounds indices will be handled. - "default" - clamps indices to (-n <= i < n), then wraps + "wrap" - clamps indices to (-n <= i < n), then wraps negative indices. "clip" - clips indices to (0 <= i < n) - "wrap" - wraps both negative and positive indices. - Default: `"default"`. + Default: `"wrap"`. """ if not isinstance(x, dpt.usm_ndarray): raise TypeError( diff --git a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp index 1c9fd39af2..40f457ce15 100644 --- a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp @@ -46,10 +46,10 @@ namespace py = pybind11; template class take_kernel; template class put_kernel; -class FancyIndex +class WrapIndex { public: - FancyIndex() = default; + WrapIndex() = default; void operator()(py::ssize_t max_item, py::ssize_t &ind) const { @@ -73,20 +73,6 @@ class ClipIndex } }; -class WrapIndex -{ -public: - WrapIndex() = default; - - void operator()(py::ssize_t max_item, py::ssize_t &ind) const - { - max_item = std::max(max_item, 1); - ind = (ind < 0) ? (ind + max_item * ((-ind / max_item) + 1)) % max_item - : ind % max_item; - return; - } -}; - template class TakeFunctor { private: @@ -361,22 +347,6 @@ sycl::event put_impl(sycl::queue q, return put_ev; } -template struct TakeFancyFactory -{ - fnT get() - { - if constexpr (std::is_integral::value && - !std::is_same::value) { - fnT fn = take_impl; - return fn; - } - else { - fnT fn = nullptr; - return fn; - } - } -}; - template struct TakeWrapFactory { fnT get() @@ -409,22 +379,6 @@ template struct TakeClipFactory } }; -template struct PutFancyFactory -{ - fnT get() - { - if constexpr (std::is_integral::value && - !std::is_same::value) { - fnT fn = put_impl; - return fn; - } - else { - fnT fn = nullptr; - return fn; - } - } -}; - template struct PutWrapFactory { fnT get() diff --git a/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp b/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp index 936a13e2bc..c2589836fd 100644 --- a/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp +++ b/dpctl/tensor/libtensor/source/integer_advanced_indexing.cpp @@ -39,10 +39,9 @@ #include "integer_advanced_indexing.hpp" -#define INDEXING_MODES 3 -#define FANCY_MODE 0 +#define INDEXING_MODES 2 +#define WRAP_MODE 0 #define CLIP_MODE 1 -#define WRAP_MODE 2 namespace dpctl { @@ -884,11 +883,6 @@ void init_advanced_indexing_dispatch_tables(void) { using namespace dpctl::tensor::detail; - using dpctl::tensor::kernels::indexing::TakeFancyFactory; - DispatchTableBuilder - dtb_takefancy; - dtb_takefancy.populate_dispatch_table(take_dispatch_table[FANCY_MODE]); - using dpctl::tensor::kernels::indexing::TakeClipFactory; DispatchTableBuilder dtb_takeclip; @@ -899,10 +893,6 @@ void init_advanced_indexing_dispatch_tables(void) dtb_takewrap; dtb_takewrap.populate_dispatch_table(take_dispatch_table[WRAP_MODE]); - using dpctl::tensor::kernels::indexing::PutFancyFactory; - DispatchTableBuilder dtb_putfancy; - dtb_putfancy.populate_dispatch_table(put_dispatch_table[FANCY_MODE]); - using dpctl::tensor::kernels::indexing::PutClipFactory; DispatchTableBuilder dtb_putclip; dtb_putclip.populate_dispatch_table(put_dispatch_table[CLIP_MODE]); diff --git a/dpctl/tests/test_usm_ndarray_indexing.py b/dpctl/tests/test_usm_ndarray_indexing.py index d2d26e6498..a57ea83cea 100644 --- a/dpctl/tests/test_usm_ndarray_indexing.py +++ b/dpctl/tests/test_usm_ndarray_indexing.py @@ -898,27 +898,19 @@ def test_integer_indexing_modes(): x = dpt.arange(5, sycl_queue=q) x_np = dpt.asnumpy(x) - ind = dpt.asarray([-6, -3, 0, 2, 6], dtype=np.intp, sycl_queue=q) - ind_np = dpt.asnumpy(ind) + # wrapping negative indices + ind = dpt.asarray([-4, -3, 0, 2, 4], dtype=np.intp, sycl_queue=q) - # wrapping res = dpt.take(x, ind, mode="wrap") - expected_arr = np.take(x_np, ind_np, mode="wrap") + expected_arr = np.take(x_np, dpt.asnumpy(ind), mode="raise") assert (dpt.asnumpy(res) == expected_arr).all() # clipping to 0 (disabling negative indices) - res = dpt.take(x, ind, mode="clip") - expected_arr = np.take(x_np, ind_np, mode="clip") - - assert (dpt.asnumpy(res) == expected_arr).all() - - # clipping to -n<=i