diff --git a/CHANGELOG.md b/CHANGELOG.md index ed4779cdb96d..6f7a1a60a939 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -84,6 +84,7 @@ This release is compatible with NumPy 2.5. * Fixed a crash in boolean-mask advanced indexing (`dpnp.ndarray` get/set item) when the selection is empty (e.g. a scalar `False` index that injects a length-0 axis) [#3019](https://github.com/IntelPython/dpnp/pull/3019) * Released the GIL before the remaining blocking OneMKL BLAS and LAPACK calls to prevent host tasks contention, completing the work started in [#2850](https://github.com/IntelPython/dpnp/pull/2850) [#3027](https://github.com/IntelPython/dpnp/pull/3027) * Fixed `dpnp.repeat` raising an unclear `TypeError` for a nested sequence of `repeats` [#3024](https://github.com/IntelPython/dpnp/pull/3024) +* 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 diff --git a/dpnp/tensor/_utility_functions.py b/dpnp/tensor/_utility_functions.py index 651ce0830266..8d0be2d85028 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,37 +77,48 @@ def _boolean_reduction(x, axis, keepdims, func): exec_q = x.sycl_queue res_usm_type = x.usm_type - _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 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, + ) + 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 @@ -142,7 +153,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 +182,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): 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; diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index e68ba8162442..a7696fe9852d 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -22,12 +22,12 @@ 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]) - 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) @@ -35,10 +35,9 @@ 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): + 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) @@ -49,10 +48,9 @@ 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): + def test_empty(self, func, axis, shape): dp_array = dpnp.empty(shape, dtype=dpnp.int64) np_array = dpnp.asnumpy(dp_array) @@ -60,8 +58,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_scalar(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) + + result = getattr(dpnp, func)(dp_array) + expected = getattr(numpy, func)(np_array) + assert_array_equal(result, expected) + + def test_scalar(self, func): dp_array = dpnp.array(0) np_array = dpnp.asnumpy(dp_array) @@ -69,10 +74,9 @@ 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): + 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) @@ -80,8 +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) - @pytest.mark.parametrize("func", ["all", "any"]) - 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)