Fix abort in dpnp.all / dpnp.any over an empty axis

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
 

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):

dpnp/tensor/libtensor/source/reductions/reduction_over_axis.hpp

@@ -1099,11 +1099,22 @@ std::pair<sycl::event, sycl::event>
 
     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<std::size_t>(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<sycl::event, sycl::event>
     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;
 

dpnp/tests/test_logic.py

@@ -22,23 +22,22 @@
 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)
 
         expected = getattr(numpy, func)(np_array, axis=axis, keepdims=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,39 +48,43 @@ 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)
 
         result = getattr(dpnp, func)(dp_array, axis=axis)
         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)
 
         result = getattr(dp_array, 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)
 
         expected = getattr(numpy, func)(np_array, axis=axis, keepdims=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)