From 5ae497c4cbe8d4a35f5a1f13e349a2eb8d22a827 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 12 Aug 2026 13:32:27 +0200 Subject: [PATCH 1/7] Handle empty boolean reductions without launching a kernel dpnp.all / dpnp.any submitted the reduction kernel unconditionally. When the input has no elements to reduce over - e.g. all(zeros((0, 3, 4))) or a reduction along a zero-length axis - the reduction extent is zero and the kernel is launched with a zero-sized nd_range. That is a silent no-op on runtimes built with NDEBUG, but aborts on an assertions-enabled SYCL runtime (adjustNDRangePerKernel asserts NDR.LocalSize[0] == 0 when GlobalSize is 0). Short-circuit in _boolean_reduction when the (permuted) input is empty: build the result directly with the reduction identity (True for all, False for any). This is correct for both empty sub-cases - an empty output (fill value is irrelevant) and a zero-length reduced axis (identity is the answer) - and submits no kernel. --- dpnp/tensor/_utility_functions.py | 28 ++++++++++++++++++++++------ 1 file changed, 22 insertions(+), 6 deletions(-) diff --git a/dpnp/tensor/_utility_functions.py b/dpnp/tensor/_utility_functions.py index 651ce0830266..3b4158abf56c 100644 --- a/dpnp/tensor/_utility_functions.py +++ b/dpnp/tensor/_utility_functions.py @@ -48,7 +48,7 @@ ) -def _boolean_reduction(x, axis, keepdims, func): +def _boolean_reduction(x, axis, keepdims, func, identity): if not isinstance(x, dpt.usm_ndarray): raise TypeError(f"Expected dpnp.tensor.usm_ndarray, got {type(x)}") @@ -77,11 +77,27 @@ def _boolean_reduction(x, axis, keepdims, func): exec_q = x.sycl_queue res_usm_type = x.usm_type + if x_tmp.size == 0: + # nothing to reduce over: the result is either empty (a non-reduced + # dimension is zero) or filled with the reduction identity (a reduced + # dimension is zero, e.g. all([]) is True and any([]) is False) + res = dpt.full( + res_shape, + identity, + dtype=dpt.bool, + usm_type=res_usm_type, + sycl_queue=exec_q, + ) + if keepdims: + res_shape = res_shape + (1,) * red_nd + inv_perm = sorted(range(nd), key=lambda d: perm[d]) + res = dpt.permute_dims(dpt.reshape(res, res_shape), inv_perm) + return res + _manager = du.SequentialOrderManager[exec_q] dep_evs = _manager.submitted_events - # always allocate the temporary as - # int32 and usm-device to ensure that atomic updates - # are supported + # always allocate the temporary as int32 and usm-device to ensure + # that atomic updates are supported res_tmp = dpt.empty( res_shape, dtype=dpt.int32, @@ -142,7 +158,7 @@ def all(x, /, *, axis=None, keepdims=False): An array with a data type of `bool` containing the results of the logical AND reduction. """ - return _boolean_reduction(x, axis, keepdims, tri._all) + return _boolean_reduction(x, axis, keepdims, tri._all, True) def any(x, /, *, axis=None, keepdims=False): @@ -171,7 +187,7 @@ def any(x, /, *, axis=None, keepdims=False): An array with a data type of `bool` containing the results of the logical OR reduction. """ - return _boolean_reduction(x, axis, keepdims, tri._any) + return _boolean_reduction(x, axis, keepdims, tri._any, False) def _validate_diff_shape(sh1, sh2, axis): From 99bda422ea6a80714964421df2fe654a391d6700 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 12 Aug 2026 14:05:33 +0200 Subject: [PATCH 2/7] Update changelog for empty boolean reduction fix --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 8d01978e60df..8defd2a707a5 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -78,6 +78,7 @@ This release is compatible with NumPy 2.5. * Fixed `dpnp.interp` returning `nan` when querying at an exact knot point whose adjacent `fp` value is `inf` [#2986](https://github.com/IntelPython/dpnp/pull/2986) * Fixed missing strides validation in `dpnp.tensor.usm_ndarray` constructor when allocating new memory [#2927](https://github.com/IntelPython/dpnp/pull/2927) * Fixed `dpnp.bincount` raising a `ValueError` on an empty input array instead of returning an empty `intp` array [#3018](https://github.com/IntelPython/dpnp/pull/3018) +* Fixed `dpnp.all` and `dpnp.any` aborting when reducing over an empty axis (e.g. an array with a zero-length dimension) [#3021](https://github.com/IntelPython/dpnp/pull/3021) ### Security From a90dafd635146eac3d01d9bccfd8944ce331be55 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 12 Aug 2026 14:13:32 +0200 Subject: [PATCH 3/7] Deduplicate keepdims handling in _boolean_reduction Use an if/else so the empty-input and reduction paths share the single trailing keepdims block instead of repeating it. --- dpnp/tensor/_utility_functions.py | 67 ++++++++++++++----------------- 1 file changed, 31 insertions(+), 36 deletions(-) diff --git a/dpnp/tensor/_utility_functions.py b/dpnp/tensor/_utility_functions.py index 3b4158abf56c..8d0be2d85028 100644 --- a/dpnp/tensor/_utility_functions.py +++ b/dpnp/tensor/_utility_functions.py @@ -88,42 +88,37 @@ def _boolean_reduction(x, axis, keepdims, func, identity): usm_type=res_usm_type, sycl_queue=exec_q, ) - if keepdims: - res_shape = res_shape + (1,) * red_nd - inv_perm = sorted(range(nd), key=lambda d: perm[d]) - res = dpt.permute_dims(dpt.reshape(res, res_shape), inv_perm) - return res - - _manager = du.SequentialOrderManager[exec_q] - dep_evs = _manager.submitted_events - # always allocate the temporary as int32 and usm-device to ensure - # that atomic updates are supported - res_tmp = dpt.empty( - res_shape, - dtype=dpt.int32, - usm_type="device", - sycl_queue=exec_q, - ) - hev0, ev0 = func( - src=x_tmp, - trailing_dims_to_reduce=red_nd, - dst=res_tmp, - sycl_queue=exec_q, - depends=dep_evs, - ) - _manager.add_event_pair(hev0, ev0) - - # copy to boolean result array - res = dpt.empty( - res_shape, - dtype=dpt.bool, - usm_type=res_usm_type, - sycl_queue=exec_q, - ) - hev1, ev1 = ti._copy_usm_ndarray_into_usm_ndarray( - src=res_tmp, dst=res, sycl_queue=exec_q, depends=[ev0] - ) - _manager.add_event_pair(hev1, ev1) + else: + _manager = du.SequentialOrderManager[exec_q] + dep_evs = _manager.submitted_events + # always allocate the temporary as int32 and usm-device to ensure + # that atomic updates are supported + res_tmp = dpt.empty( + res_shape, + dtype=dpt.int32, + usm_type="device", + sycl_queue=exec_q, + ) + hev0, ev0 = func( + src=x_tmp, + trailing_dims_to_reduce=red_nd, + dst=res_tmp, + sycl_queue=exec_q, + depends=dep_evs, + ) + _manager.add_event_pair(hev0, ev0) + + # copy to boolean result array + res = dpt.empty( + res_shape, + dtype=dpt.bool, + usm_type=res_usm_type, + sycl_queue=exec_q, + ) + hev1, ev1 = ti._copy_usm_ndarray_into_usm_ndarray( + src=res_tmp, dst=res, sycl_queue=exec_q, depends=[ev0] + ) + _manager.add_event_pair(hev1, ev1) if keepdims: res_shape = res_shape + (1,) * red_nd From 3a894115288ebff89f40c9d7ebdf8ec2ede3561e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 17 Aug 2026 21:25:20 +0200 Subject: [PATCH 4/7] Adding an early-exit path to the reduction pybind11 code --- .../source/reductions/reduction_over_axis.hpp | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/dpnp/tensor/libtensor/source/reductions/reduction_over_axis.hpp b/dpnp/tensor/libtensor/source/reductions/reduction_over_axis.hpp index ee65cd8b45ac..6be8646e4056 100644 --- a/dpnp/tensor/libtensor/source/reductions/reduction_over_axis.hpp +++ b/dpnp/tensor/libtensor/source/reductions/reduction_over_axis.hpp @@ -1099,11 +1099,22 @@ std::pair std::size_t dst_nelems = dst.get_size(); + if (dst_nelems == 0) { + // empty result: nothing to write + return std::make_pair(sycl::event(), sycl::event()); + } + std::size_t red_nelems(1); for (int i = dst_nd; i < src_nd; ++i) { red_nelems *= static_cast(src_shape_ptr[i]); } + if (red_nelems == 0) { + // empty reduction extent: the result is the op identity, which this + // kernel cannot produce; the caller must handle it + throw py::value_error("Reduction over an empty axis is not supported"); + } + auto const &overlap = dpnp::tensor::overlap::MemoryOverlap(); if (overlap(dst, src)) { throw py::value_error("Arrays are expected to have no memory overlap"); @@ -1142,9 +1153,8 @@ std::pair bool is_src_f_contig = src.is_f_contiguous(); bool is_dst_c_contig = dst.is_c_contiguous(); - // TODO: should be dst_nelems == 0? if ((is_src_c_contig && is_dst_c_contig) || - (is_src_f_contig && dst_nelems == 0)) { + (is_src_f_contig && dst_nelems == 1)) { auto fn = axis1_contig_dispatch_vector[src_typeid]; static constexpr py::ssize_t zero_offset = 0; From ac258fbf2304ab9dde2ec6d88e7267b3b4b9013a Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 17 Aug 2026 21:26:12 +0200 Subject: [PATCH 5/7] Add test to cover new fast reduction path --- dpnp/tests/test_logic.py | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index e68ba8162442..3bb16575fcf8 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -60,6 +60,15 @@ def test_all_any_empty(self, func, axis, shape): expected = getattr(numpy, func)(np_array, axis=axis) assert_allclose(result, expected) + @pytest.mark.parametrize("func", ["all", "any"]) + def test_all_any_f_contig_full(self, func): + dp_array = dpnp.array([[0, 1, 2], [3, 4, 0]], order="F") + np_array = dpnp.asnumpy(dp_array) + + result = getattr(dpnp, func)(dp_array) + expected = getattr(numpy, func)(np_array) + assert_array_equal(result, expected) + @pytest.mark.parametrize("func", ["all", "any"]) def test_all_any_scalar(self, func): dp_array = dpnp.array(0) From 8c82912d9e5acf8b25853522e9eb0f6191ed4008 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 17 Aug 2026 21:28:55 +0200 Subject: [PATCH 6/7] Switched to decorating the class directly to reduce duplication --- dpnp/tests/test_logic.py | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index 3bb16575fcf8..796fb37083cf 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -22,8 +22,8 @@ from .third_party.cupy import testing +@pytest.mark.parametrize("func", ["all", "any"]) class TestAllAny: - @pytest.mark.parametrize("func", ["all", "any"]) @pytest.mark.parametrize("dtype", get_all_dtypes()) @pytest.mark.parametrize("axis", [None, 0, 1, (0, 1)]) @pytest.mark.parametrize("keepdims", [True, False]) @@ -35,7 +35,6 @@ def test_all_any(self, func, dtype, axis, keepdims): result = getattr(dpnp, func)(dp_array, axis=axis, keepdims=keepdims) assert_allclose(result, expected) - @pytest.mark.parametrize("func", ["all", "any"]) @pytest.mark.parametrize("a_dtype", get_all_dtypes(no_none=True)) @pytest.mark.parametrize("out_dtype", get_all_dtypes(no_none=True)) def test_all_any_out(self, func, a_dtype, out_dtype): @@ -49,7 +48,6 @@ def test_all_any_out(self, func, a_dtype, out_dtype): # out kwarg is not used with NumPy, dtype may differ assert_array_equal(result, expected, strict=False) - @pytest.mark.parametrize("func", ["all", "any"]) @pytest.mark.parametrize("axis", [None, 0, 1, (0, 1)]) @pytest.mark.parametrize("shape", [(2, 3), (2, 0), (0, 3)]) def test_all_any_empty(self, func, axis, shape): @@ -60,7 +58,6 @@ def test_all_any_empty(self, func, axis, shape): expected = getattr(numpy, func)(np_array, axis=axis) assert_allclose(result, expected) - @pytest.mark.parametrize("func", ["all", "any"]) def test_all_any_f_contig_full(self, func): dp_array = dpnp.array([[0, 1, 2], [3, 4, 0]], order="F") np_array = dpnp.asnumpy(dp_array) @@ -69,7 +66,6 @@ def test_all_any_f_contig_full(self, func): expected = getattr(numpy, func)(np_array) assert_array_equal(result, expected) - @pytest.mark.parametrize("func", ["all", "any"]) def test_all_any_scalar(self, func): dp_array = dpnp.array(0) np_array = dpnp.asnumpy(dp_array) @@ -78,7 +74,6 @@ def test_all_any_scalar(self, func): expected = getattr(np_array, func)() assert_allclose(result, expected) - @pytest.mark.parametrize("func", ["all", "any"]) @pytest.mark.parametrize("axis", [None, 0, 1]) @pytest.mark.parametrize("keepdims", [True, False]) def test_all_any_nan_inf(self, func, axis, keepdims): @@ -89,7 +84,6 @@ def test_all_any_nan_inf(self, func, axis, keepdims): result = getattr(dpnp, func)(dp_array, axis=axis, keepdims=keepdims) assert_allclose(result, expected) - @pytest.mark.parametrize("func", ["all", "any"]) def test_all_any_error(self, func): def check_raises(func_name, exception, *args, **kwargs): assert_raises( From 6a7ef40c0c7a7a1da4cdae3285ea3a3e9b54720f Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 17 Aug 2026 21:29:47 +0200 Subject: [PATCH 7/7] Rename the tests --- dpnp/tests/test_logic.py | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index 796fb37083cf..a7696fe9852d 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -27,7 +27,7 @@ class TestAllAny: @pytest.mark.parametrize("dtype", get_all_dtypes()) @pytest.mark.parametrize("axis", [None, 0, 1, (0, 1)]) @pytest.mark.parametrize("keepdims", [True, False]) - def test_all_any(self, func, dtype, axis, keepdims): + def test_basic(self, func, dtype, axis, keepdims): dp_array = dpnp.array([[0, 1, 2], [3, 4, 0]], dtype=dtype) np_array = dpnp.asnumpy(dp_array) @@ -37,7 +37,7 @@ def test_all_any(self, func, dtype, axis, keepdims): @pytest.mark.parametrize("a_dtype", get_all_dtypes(no_none=True)) @pytest.mark.parametrize("out_dtype", get_all_dtypes(no_none=True)) - def test_all_any_out(self, func, a_dtype, out_dtype): + def test_out(self, func, a_dtype, out_dtype): dp_array = dpnp.array([[0, 1, 2], [3, 4, 0]], dtype=a_dtype) np_array = dpnp.asnumpy(dp_array) @@ -50,7 +50,7 @@ def test_all_any_out(self, func, a_dtype, out_dtype): @pytest.mark.parametrize("axis", [None, 0, 1, (0, 1)]) @pytest.mark.parametrize("shape", [(2, 3), (2, 0), (0, 3)]) - def test_all_any_empty(self, func, axis, shape): + def test_empty(self, func, axis, shape): dp_array = dpnp.empty(shape, dtype=dpnp.int64) np_array = dpnp.asnumpy(dp_array) @@ -58,7 +58,7 @@ def test_all_any_empty(self, func, axis, shape): expected = getattr(numpy, func)(np_array, axis=axis) assert_allclose(result, expected) - def test_all_any_f_contig_full(self, func): + def test_f_contig_full(self, func): dp_array = dpnp.array([[0, 1, 2], [3, 4, 0]], order="F") np_array = dpnp.asnumpy(dp_array) @@ -66,7 +66,7 @@ def test_all_any_f_contig_full(self, func): expected = getattr(numpy, func)(np_array) assert_array_equal(result, expected) - def test_all_any_scalar(self, func): + def test_scalar(self, func): dp_array = dpnp.array(0) np_array = dpnp.asnumpy(dp_array) @@ -76,7 +76,7 @@ def test_all_any_scalar(self, func): @pytest.mark.parametrize("axis", [None, 0, 1]) @pytest.mark.parametrize("keepdims", [True, False]) - def test_all_any_nan_inf(self, func, axis, keepdims): + def test_nan_inf(self, func, axis, keepdims): dp_array = dpnp.array([[dpnp.nan, 1, 2], [dpnp.inf, -dpnp.inf, 0]]) np_array = dpnp.asnumpy(dp_array) @@ -84,7 +84,7 @@ def test_all_any_nan_inf(self, func, axis, keepdims): result = getattr(dpnp, func)(dp_array, axis=axis, keepdims=keepdims) assert_allclose(result, expected) - def test_all_any_error(self, func): + def test_error(self, func): def check_raises(func_name, exception, *args, **kwargs): assert_raises( exception, lambda: getattr(dpnp, func_name)(*args, **kwargs)