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

[antonwolfy]

dpnp.all and dpnp.any aborted the process when reducing over an empty axis, e.g.:

import dpnp
dpnp.all(dpnp.zeros((0, 3, 4)))          # reduces over 0 elements
dpnp.any(dpnp.zeros((0, 3, 4)), axis=0)  # zero-length reduced axis

_boolean_reduction always allocated a temporary and launched the reduction kernel. When the reduced extent is zero (a reduced dimension has length 0), the kernel is submitted with a zero-sized nd_range. That is a silent no-op on a SYCL runtime built with NDEBUG, but aborts on an assertions-enabled runtime, where adjustNDRangePerKernel asserts NDR.LocalSize[0] == 0 whenever GlobalSize[0] is 0:

Assertion `NDR.LocalSize[0] == 0' failed.
  .../sycl/source/detail/scheduler/commands.cpp
  void sycl::_V1::detail::adjustNDRangePerKernel(...)

This PR proposes to short-circuit in _boolean_reduction when the (permuted) input has no elements, building the result directly with the reduction identity instead of launching a kernel. all uses identity True, any uses False. A single x_tmp.size == 0 check covers both empty sub-cases correctly:

• empty output (a non-reduced dimension is zero) — the result is an empty array, so the fill value is irrelevant;
• zero-length reduced axis (the output is non-empty) — the reduction identity is the answer, e.g. all([]) is True, any([]) is False.

Note, the issue was identified when building with the nightly LLVM SYCL compiler and running dpnp tests.
And the tests are passing now without any crash.

• Have you provided a meaningful PR description?
• Have you added a test, reproducer or referred to an issue with a reproducer?
• Have you tested your changes locally for CPU and GPU devices?
• Have you made sure that new changes do not introduce compiler warnings?
• Have you checked performance impact of proposed changes?
• Have you added documentation for your changes, if necessary?
• Have you added your changes to the changelog?

[github-actions]

View rendered docs @ https://intelpython.github.io/dpnp/pull/3021/index.html

[coveralls]

coverage: 78.416% (+0.004%) from 78.412% — fix-empty-boolean-reduction into master

[github-actions]

Array API standard conformance tests for dpnp=0.21.0dev3=py314h509198e_42 ran successfully.
Passed: 1376
Failed: 1
Skipped: 5

[ndgrigorian]

@antonwolfy
We could also consider adding an early-exit path to the reduction pybind11 code. In fact, one should already exist, maybe it was missed for boolean reductions?

[antonwolfy]

I was thinking about that.. and it should be added to py_boolean_reduction, something like:

if (dst_nelems == 0) {
    // empty result: nothing to write
    return std::make_pair(sycl::event(), sycl::event());
}
if (red_nelems == 0) {
    // empty reduction extent: 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");
}

but I wondering why we already had in the code a hadnling depending on dst_nelems == 0:

    // TODO: should be dst_nelems == 0?
    if ((is_src_c_contig && is_dst_c_contig) ||
        (is_src_f_contig && dst_nelems == 0)) {

@ndgrigorian, if you remember why it was added this way and can confirm it's not needed, I'll add the above handling to py_boolean_reduction and will remove the TODO and obsolete path with dst_nelems == 0.