Fix abort on empty boolean-mask advanced indexing - #3019
Merged
Conversation
py_place and py_extract computed the orthogonal and masked element counts and dispatched to the strided kernels without checking for zero. When a boolean index produces an empty selection - e.g. a leading scalar `False` that injects a length-0 axis - ortho_nelems becomes 0, and the some-slices kernel builds sycl::nd_range<2> with GlobalSize[0] == 0 and a non-zero local size. That is a degenerate launch: it 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). Return early with empty events when ortho_nelems or the masked element count is zero, so no kernel is submitted for an empty selection.
Contributor
|
View rendered docs @ https://intelpython.github.io/dpnp/pull/3019/index.html |
antonwolfy
marked this pull request as ready for review
August 12, 2026 10:54
antonwolfy
requested review from
ndgrigorian and
vlad-perevezentsev
as code owners
August 12, 2026 10:54
Contributor
|
Array API standard conformance tests for dpnp=0.21.0dev4=py314h509198e_7 ran successfully. |
Collaborator
|
No base build to compare — fix-empty-boolean-adv-indexing into master |
github-actions Bot
added a commit
that referenced
this pull request
Aug 14, 2026
Boolean-mask advanced indexing (`dpnp.ndarray.__getitem__`/`__setitem__`) aborted when the selection was empty. For example: ```python import dpnp a = dpnp.zeros((2, 3, 4)) a[False, [True, False]] = 1 # scalar `False` injects a length-0 axis ``` `py_place` and `py_extract` compute the orthogonal and masked element counts and dispatch to the strided masked place/extract kernels without checking for zero. When the boolean index produces an empty selection — e.g. a leading scalar `False` that injects a length-0 axis — `ortho_nelems` becomes `0`, and the some-slices kernel builds `sycl::nd_range<2>` with `GlobalSize[0] == 0` and a non-zero local size. That is a degenerate launch. It is a silent no-op on a SYCL runtime built with `NDEBUG`, so it goes unnoticed with the released oneAPI runtime. On an assertions-enabled SYCL runtime it aborts, because `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 return early with empty events when the orthogonal or masked element count is zero, before any kernel is submitted, in both `py_place` and `py_extract`. An empty selection is a no-op, so this is semantically correct and mirrors the existing `cumsum_sz == 0` guard in the same file. Note, the issue was identified when building with the nightly LLVM SYCL compiler and [running](https://github.com/antonwolfy/dpnp/actions/runs/31576568012/job/94049919151) dpnp tests. And the test is [passing](https://github.com/antonwolfy/dpnp/actions/runs/31585418522/job/94078011200) with the fix: > dpnp/tests/third_party/cupy/core_tests/test_ndarray_adv_indexing.py::TestArrayAdvancedIndexingSetitemScalarValue::test_adv_setitem[42] PASSED [ 77%] 12b436f
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Boolean-mask advanced indexing (
dpnp.ndarray.__getitem__/__setitem__) aborted when the selection was empty. For example:py_placeandpy_extractcompute the orthogonal and masked element counts and dispatch to the strided masked place/extract kernels without checking for zero. When the boolean index produces an empty selection — e.g. a leading scalarFalsethat injects a length-0 axis —ortho_nelemsbecomes0, and the some-slices kernel buildssycl::nd_range<2>withGlobalSize[0] == 0and a non-zero local size.That is a degenerate launch. It is a silent no-op on a SYCL runtime built with
NDEBUG, so it goes unnoticed with the released oneAPI runtime. On an assertions-enabled SYCL runtime it aborts, becauseadjustNDRangePerKernelassertsNDR.LocalSize[0] == 0wheneverGlobalSize[0]is0:This PR proposes to return early with empty events when the orthogonal or masked element count is zero, before any kernel is submitted, in both
py_placeandpy_extract. An empty selection is a no-op, so this is semantically correct and mirrors the existingcumsum_sz == 0guard in the same file.Note, the issue was identified when building with the nightly LLVM SYCL compiler and running dpnp tests.
And the test is passing with the fix: