From 9cfc7e218bd513e1a1ea685986a6953e82013b04 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Wed, 2 Sep 2026 10:35:20 -0500 Subject: [PATCH 1/5] fix: treat any non-zero bool byte as True in accumulators `NonZeroIndicator` compared against `inputT(0)` and the `Cumsum1D` factories used `NoOpTransformer`; for `bool` both fold into a raw byte load, so the scan summed byte values instead of 0/1. A mask stored as [0, 1, 2, 255, 0, 1] reported 259 non-zeros instead of 4, affecting `nonzero`, `where`, `extract`, `place` and `repeat`. Cast bool via `sycl::bit_cast` / `CastTransformer`, as `convert_impl` already does (gh-2121). --- CHANGELOG.md | 1 + .../include/kernels/accumulators.hpp | 24 ++++++-- .../tensor/test_usm_ndarray_manipulation.py | 12 ++++ dpnp/tests/test_indexing.py | 59 +++++++++++++++++++ 4 files changed, 90 insertions(+), 6 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index e37ae7055273..19a0501e2302 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -94,6 +94,7 @@ This release is compatible with NumPy 2.5. * Fixed `astype` casting an out-of-range floating point value to a signed narrow integer type saturating to the destination min/max instead of wrapping like NumPy, generalizing the earlier unsigned-only fix [#3033](https://github.com/IntelPython/dpnp/pull/3033) * Fixed `dpnp.insert` silently ignoring out-of-bounds negative indices in a multi-element `obj`, so a mix of in-bounds and out-of-bounds indices now consistently raises `IndexError` [#3041](https://github.com/IntelPython/dpnp/pull/3041) * Fixed a per-call `sycl::queue` leak in `usm_ndarray::get_queue()`/`get_device()` [#3042](https://github.com/IntelPython/dpnp/pull/3042) +* Fixed `dpnp.nonzero`, `dpnp.where`, `dpnp.extract`, `dpnp.place` and `dpnp.repeat` returning wrong results for a boolean array whose bytes are not `0x00`/`0x01`, e.g. a `view` of integer data [#3053](https://github.com/IntelPython/dpnp/pull/3053) ### Security diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp index 079fcf5e9c96..98a5b83d82eb 100644 --- a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp +++ b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp @@ -74,9 +74,17 @@ struct NonZeroIndicator { static constexpr outputT out_one(1); static constexpr outputT out_zero(0); - static constexpr inputT val_zero(0); - return (val == val_zero) ? out_zero : out_one; + if constexpr (std::is_same_v) { + // NumPy treats any non-zero byte as True, while a plain + // comparison may fold into a raw byte load, see gh-2121 + const std::uint8_t &u = sycl::bit_cast(val); + return (u == std::uint8_t{0}) ? out_zero : out_one; + } + else { + static constexpr inputT val_zero(0); + return (val == val_zero) ? out_zero : out_one; + } } }; @@ -1301,8 +1309,10 @@ struct Cumsum1DContigFactory { if constexpr (std::is_integral_v) { using cumsumT = std::int64_t; - fnT fn = - cumsum_val_contig_impl>; + // CastTransformer, not NoOpTransformer: an implicit bool + // conversion would read the raw byte, see gh-2121 + fnT fn = cumsum_val_contig_impl>; return fn; } else { @@ -1419,8 +1429,10 @@ struct Cumsum1DStridedFactory { if constexpr (std::is_integral_v) { using cumsumT = std::int64_t; - fnT fn = - cumsum_val_strided_impl>; + // CastTransformer, not NoOpTransformer: an implicit bool + // conversion would read the raw byte, see gh-2121 + fnT fn = cumsum_val_strided_impl>; return fn; } else { diff --git a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py index bf0bd3226309..88f8980a8c9d 100644 --- a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py +++ b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py @@ -1412,6 +1412,18 @@ def test_repeat_strided_repeats(): assert dpt.all(res == x) +def test_repeat_nonstandard_bool_bytes(): + # NumPy treats any non-zero byte of a bool as True, see gh-2121 + get_queue_or_skip() + + raw = dpt.asarray([0, 1, 2, 255, 0, 1], dtype="u1") + reps = dpt.usm_ndarray(raw.shape, dtype="?", buffer=raw.usm_data) + x = dpt.arange(reps.size, dtype="i4") + + res = dpt.repeat(x, reps) + assert_array_equal(dpt.asnumpy(res), np.array([1, 2, 3, 5], dtype="i4")) + + def test_repeat_size1_repeats(): get_queue_or_skip() diff --git a/dpnp/tests/test_indexing.py b/dpnp/tests/test_indexing.py index 84bf62d03562..378ca2fa10e6 100644 --- a/dpnp/tests/test_indexing.py +++ b/dpnp/tests/test_indexing.py @@ -277,6 +277,29 @@ def test_place_insert_from_empty_vals(self, xp): def test_place_wrong_array_type(self, xp): assert_raises(TypeError, xp.place, [1, 2, 3], [True, False], [0, 1]) + # NumPy treats any non-zero byte of a bool as True, see gh-2121 + def test_extract_nonstandard_bool_bytes(self): + raw = numpy.array([0, 1, 2, 255, 0, 1], dtype=numpy.uint8) + a = numpy.arange(raw.size, dtype=numpy.int32) + mask = raw.view(numpy.bool_) + ia = dpnp.asarray(a) + imask = dpnp.asarray(raw).view(dpnp.bool) + + result = dpnp.extract(imask, ia) + expected = numpy.extract(mask, a) + assert_array_equal(result, expected) + + def test_place_nonstandard_bool_bytes(self): + raw = numpy.array([0, 1, 2, 255, 0, 1], dtype=numpy.uint8) + a = numpy.arange(raw.size, dtype=numpy.int32) + mask = raw.view(numpy.bool_) + ia = dpnp.asarray(a) + imask = dpnp.asarray(raw).view(dpnp.bool) + + dpnp.place(ia, imask, [-1]) + numpy.place(a, mask, [-1]) + assert_array_equal(ia, a) + @pytest.mark.parametrize("dt", get_all_dtypes(no_none=True)) def test_both(self, dt): a = numpy.random.rand(10).astype(dt) @@ -616,6 +639,42 @@ def test_array_method(self, dtype): ia = dpnp.array(a) assert_array_equal(a.nonzero(), ia.nonzero()) + # NumPy treats any non-zero byte of a bool as True, see gh-2121 + @pytest.mark.parametrize( + "bytes_val", + [ + [0, 1, 2, 255, 0, 1], + [2] * 8, + [255], + [0] * 8, + [0, 128] * 64, + list(range(256)), + ], + ids=["mixed", "all_twos", "single_255", "all_zeros", "long", "range"], + ) + def test_nonstandard_bool_bytes(self, bytes_val): + a = numpy.array(bytes_val, dtype=numpy.uint8).view(numpy.bool_) + ia = dpnp.asarray(numpy.array(bytes_val, dtype=numpy.uint8)).view( + dpnp.bool + ) + + assert_array_equal(numpy.nonzero(a), dpnp.nonzero(ia)) + assert_array_equal(numpy.where(a), dpnp.where(ia)) + + def test_nonstandard_bool_bytes_strided(self): + raw = numpy.arange(24, dtype=numpy.uint8) + a = raw.view(numpy.bool_)[::3] + ia = dpnp.asarray(raw).view(dpnp.bool)[::3] + + assert_array_equal(numpy.nonzero(a), dpnp.nonzero(ia)) + + def test_nonstandard_bool_bytes_2d(self): + raw = numpy.array([[0, 1, 2], [255, 0, 7]], dtype=numpy.uint8) + a = raw.view(numpy.bool_) + ia = dpnp.asarray(raw).view(dpnp.bool) + + assert_array_equal(numpy.nonzero(a), dpnp.nonzero(ia)) + class TestPut: @pytest.mark.parametrize("a_dt", get_all_dtypes(no_none=True)) From 68e284bfd778ffbdb523ee053742003b11487659 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Wed, 2 Sep 2026 10:35:31 -0500 Subject: [PATCH 2/5] fix: treat any non-zero bool byte as True across kernels C++ may fold a bool comparison into a raw byte load, so a byte other than 0x00/0x01 compared unequal to a normalized True, ordered by its byte value, and leaked into computed bool results. Such a byte arises when a buffer is written through a raw pointer or viewed from integer data. Add `normalize_bool` and apply it where a bool is read from memory: elementwise operand loads (bool is excluded from the vector paths, which cannot normalize per element), the `convert_impl` same-type branch, the search-reduction loads, the `isin` equality test and the argsort projection. Add bool comparators for the merge-sort path. `sort` now orders False before True rather than reproducing NumPy's raw byte order, which would carry garbage bytes through a sort. --- CHANGELOG.md | 2 +- .../include/kernels/accumulators.hpp | 4 +- .../kernels/elementwise_functions/common.hpp | 50 +++++--- .../elementwise_functions/common_inplace.hpp | 20 ++-- .../libtensor/include/kernels/reductions.hpp | 8 +- .../include/kernels/sorting/isin.hpp | 11 +- .../include/kernels/sorting/radix_sort.hpp | 9 +- .../include/utils/rich_comparisons.hpp | 34 ++++++ .../libtensor/include/utils/type_utils.hpp | 19 ++- .../tensor/test_usm_ndarray_manipulation.py | 7 ++ dpnp/tests/test_logic.py | 111 ++++++++++++++++++ 11 files changed, 238 insertions(+), 37 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 19a0501e2302..66bb70d718a3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -94,7 +94,7 @@ This release is compatible with NumPy 2.5. * Fixed `astype` casting an out-of-range floating point value to a signed narrow integer type saturating to the destination min/max instead of wrapping like NumPy, generalizing the earlier unsigned-only fix [#3033](https://github.com/IntelPython/dpnp/pull/3033) * Fixed `dpnp.insert` silently ignoring out-of-bounds negative indices in a multi-element `obj`, so a mix of in-bounds and out-of-bounds indices now consistently raises `IndexError` [#3041](https://github.com/IntelPython/dpnp/pull/3041) * Fixed a per-call `sycl::queue` leak in `usm_ndarray::get_queue()`/`get_device()` [#3042](https://github.com/IntelPython/dpnp/pull/3042) -* Fixed `dpnp.nonzero`, `dpnp.where`, `dpnp.extract`, `dpnp.place` and `dpnp.repeat` returning wrong results for a boolean array whose bytes are not `0x00`/`0x01`, e.g. a `view` of integer data [#3053](https://github.com/IntelPython/dpnp/pull/3053) +* Fixed operations on a boolean array whose bytes are not `0x00`/`0x01`, e.g. a `view` of integer data, which returned wrong results or produced such bytes themselves. Affected comparison, logical, bitwise and arithmetic functions, `dpnp.max`/`dpnp.min`, `dpnp.argmax`/`dpnp.argmin`, `dpnp.sort`/`dpnp.argsort`, `dpnp.unique`, `dpnp.isin`, `dpnp.nonzero`, `dpnp.where`, `dpnp.extract`, `dpnp.place` and `dpnp.repeat` [#3053](https://github.com/IntelPython/dpnp/pull/3053) ### Security diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp index 98a5b83d82eb..584f913b8965 100644 --- a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp +++ b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp @@ -1311,8 +1311,8 @@ struct Cumsum1DContigFactory using cumsumT = std::int64_t; // CastTransformer, not NoOpTransformer: an implicit bool // conversion would read the raw byte, see gh-2121 - fnT fn = cumsum_val_contig_impl>; + fnT fn = + cumsum_val_contig_impl>; return fn; } else { diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp index bb310272e7a6..468d68b58579 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp @@ -46,6 +46,7 @@ #include "utils/offset_utils.hpp" #include "utils/sycl_alloc_utils.hpp" #include "utils/sycl_utils.hpp" +#include "utils/type_utils.hpp" #include "kernels/alignment.hpp" #include "kernels/dpnp_tensor_types.hpp" @@ -59,6 +60,7 @@ using dpnp::tensor::kernels::alignment_utils::required_alignment; using dpnp::tensor::sycl_utils::sub_group_load; using dpnp::tensor::sycl_utils::sub_group_store; +using dpnp::tensor::type_utils::normalize_bool; /*! @brief Functor for unary function evaluation on contiguous array */ template 1)) { + !std::is_same_v && (vec_sz > 1)) { auto sg = ndit.get_sub_group(); const std::uint16_t sgSize = sg.get_max_local_range()[0]; @@ -148,7 +152,7 @@ struct UnaryContigFunctor const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { // scalar call - out[k] = op(in[k]); + out[k] = op(normalize_bool(in[k])); } } } @@ -178,7 +182,7 @@ struct UnaryContigFunctor sub_group_load(sg, in_multi_ptr); #pragma unroll for (std::uint32_t k = 0; k < vec_sz; ++k) { - arg_vec[k] = op(arg_vec[k]); + arg_vec[k] = op(normalize_bool(arg_vec[k])); } sub_group_store(sg, arg_vec, out_multi_ptr); } @@ -186,7 +190,7 @@ struct UnaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in[k]); + out[k] = op(normalize_bool(in[k])); } } } @@ -216,7 +220,7 @@ struct UnaryContigFunctor sycl::vec res_vec; #pragma unroll for (std::uint8_t k = 0; k < vec_sz; ++k) { - res_vec[k] = op(arg_vec[k]); + res_vec[k] = op(normalize_bool(arg_vec[k])); } sub_group_store(sg, res_vec, out_multi_ptr); } @@ -224,7 +228,7 @@ struct UnaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in[k]); + out[k] = op(normalize_bool(in[k])); } } } @@ -238,7 +242,7 @@ struct UnaryContigFunctor (gid / sgSize) * (elems_per_sg - sgSize) + gid; const std::size_t end = std::min(nelems_, start + elems_per_sg); for (std::size_t offset = start; offset < end; offset += sgSize) { - out[offset] = op(in[offset]); + out[offset] = op(normalize_bool(in[offset])); } } } @@ -268,7 +272,7 @@ struct UnaryStridedFunctor UnaryOpT op{}; - res_[res_offset] = op(inp_[inp_offset]); + res_[res_offset] = op(normalize_bool(inp_[inp_offset])); } }; @@ -419,9 +423,13 @@ struct BinaryContigFunctor /* Each work-item processes vec_sz elements, contiguous in memory */ /* NOTE: work-group size must be divisible by sub-group size */ + // bool is excluded from the vector path: a byte other than 0x00/0x01 + // cannot be normalized element-wise there, see gh-2121 if constexpr (enable_sg_loadstore && BinaryOperatorT::supports_sg_loadstore::value && - BinaryOperatorT::supports_vec::value && (vec_sz > 1)) { + BinaryOperatorT::supports_vec::value && + !std::is_same_v && + !std::is_same_v && (vec_sz > 1)) { auto sg = ndit.get_sub_group(); std::uint16_t sgSize = sg.get_max_local_range()[0]; @@ -456,7 +464,7 @@ struct BinaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in1[k], in2[k]); + out[k] = op(normalize_bool(in1[k]), normalize_bool(in2[k])); } } } @@ -491,8 +499,8 @@ struct BinaryContigFunctor sycl::vec res_vec; #pragma unroll for (std::uint8_t vec_id = 0; vec_id < vec_sz; ++vec_id) { - res_vec[vec_id] = - op(arg1_vec[vec_id], arg2_vec[vec_id]); + res_vec[vec_id] = op(normalize_bool(arg1_vec[vec_id]), + normalize_bool(arg2_vec[vec_id])); } sub_group_store(sg, res_vec, out_multi_ptr); } @@ -500,7 +508,7 @@ struct BinaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in1[k], in2[k]); + out[k] = op(normalize_bool(in1[k]), normalize_bool(in2[k])); } } } @@ -514,7 +522,8 @@ struct BinaryContigFunctor (gid / sgSize) * (elems_per_sg - sgSize) + gid; const std::size_t end = std::min(nelems_, start + elems_per_sg); for (std::size_t offset = start; offset < end; offset += sgSize) { - out[offset] = op(in1[offset], in2[offset]); + out[offset] = op(normalize_bool(in1[offset]), + normalize_bool(in2[offset])); } } } @@ -553,7 +562,8 @@ struct BinaryStridedFunctor const auto &out_offset = three_offsets_.get_third_offset(); BinaryOperatorT op{}; - out[out_offset] = op(in1[inp1_offset], in2[inp2_offset]); + out[out_offset] = op(normalize_bool(in1[inp1_offset]), + normalize_bool(in2[inp2_offset])); } }; @@ -610,14 +620,15 @@ struct BinaryContigMatrixContigRowBroadcastingFunctor const argT1 mat_el = sub_group_load(sg, in1_multi_ptr); const argT2 vec_el = sub_group_load(sg, in2_multi_ptr); - resT res_el = op(mat_el, vec_el); + resT res_el = op(normalize_bool(mat_el), normalize_bool(vec_el)); sub_group_store(sg, res_el, out_multi_ptr); } else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < n_elems; k += sgSize) { - res[k] = op(mat[k], padded_vec[k % n1]); + res[k] = op(normalize_bool(mat[k]), + normalize_bool(padded_vec[k % n1])); } } } @@ -675,14 +686,15 @@ struct BinaryContigRowContigMatrixBroadcastingFunctor const argT2 mat_el = sub_group_load(sg, in2_multi_ptr); const argT1 vec_el = sub_group_load(sg, in1_multi_ptr); - resT res_el = op(vec_el, mat_el); + resT res_el = op(normalize_bool(vec_el), normalize_bool(mat_el)); sub_group_store(sg, res_el, out_multi_ptr); } else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < n_elems; k += sgSize) { - res[k] = op(padded_vec[k % n1], mat[k]); + res[k] = op(normalize_bool(padded_vec[k % n1]), + normalize_bool(mat[k])); } } } diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp index 9384ec603754..07d8deb2003e 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp @@ -44,6 +44,7 @@ #include "utils/offset_utils.hpp" #include "utils/sycl_alloc_utils.hpp" #include "utils/sycl_utils.hpp" +#include "utils/type_utils.hpp" #include "kernels/alignment.hpp" #include "kernels/dpnp_tensor_types.hpp" @@ -59,6 +60,7 @@ using dpnp::tensor::kernels::alignment_utils::required_alignment; using dpnp::tensor::sycl_utils::sub_group_load; using dpnp::tensor::sycl_utils::sub_group_store; +using dpnp::tensor::type_utils::normalize_bool; template 1)) { + !std::is_same_v && (vec_sz > 1)) { auto sg = ndit.get_sub_group(); std::uint16_t sgSize = sg.get_max_local_range()[0]; @@ -123,7 +127,7 @@ struct BinaryInplaceContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - op(lhs[k], rhs[k]); + op(lhs[k], normalize_bool(rhs[k])); } } } @@ -154,7 +158,7 @@ struct BinaryInplaceContigFunctor sub_group_load(sg, lhs_multi_ptr); #pragma unroll for (std::uint8_t vec_id = 0; vec_id < vec_sz; ++vec_id) { - op(res_vec[vec_id], arg_vec[vec_id]); + op(res_vec[vec_id], normalize_bool(arg_vec[vec_id])); } sub_group_store(sg, res_vec, lhs_multi_ptr); } @@ -162,7 +166,7 @@ struct BinaryInplaceContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - op(lhs[k], rhs[k]); + op(lhs[k], normalize_bool(rhs[k])); } } } @@ -176,7 +180,7 @@ struct BinaryInplaceContigFunctor (gid / sgSize) * (elems_per_sg - sgSize) + gid; const std::size_t end = std::min(nelems_, start + elems_per_sg); for (std::size_t offset = start; offset < end; offset += sgSize) { - op(lhs[offset], rhs[offset]); + op(lhs[offset], normalize_bool(rhs[offset])); } } } @@ -210,7 +214,7 @@ struct BinaryInplaceStridedFunctor const auto &lhs_offset = two_offsets_.get_second_offset(); BinaryInplaceOperatorT op{}; - op(lhs[lhs_offset], rhs[inp_offset]); + op(lhs[lhs_offset], normalize_bool(rhs[inp_offset])); } }; @@ -257,14 +261,14 @@ struct BinaryInplaceRowMatrixBroadcastingFunctor const argT vec_el = sub_group_load(sg, in_multi_ptr); resT mat_el = sub_group_load(sg, out_multi_ptr); - op(mat_el, vec_el); + op(mat_el, normalize_bool(vec_el)); sub_group_store(sg, mat_el, out_multi_ptr); } else { const std::size_t start = base + sg.get_local_id()[0]; for (std::size_t k = start; k < n_elems; k += sgSize) { - op(mat[k], padded_vec[k % n1]); + op(mat[k], normalize_bool(padded_vec[k % n1])); } } } diff --git a/dpnp/tensor/libtensor/include/kernels/reductions.hpp b/dpnp/tensor/libtensor/include/kernels/reductions.hpp index 42bea6f28126..9628e4028b26 100644 --- a/dpnp/tensor/libtensor/include/kernels/reductions.hpp +++ b/dpnp/tensor/libtensor/include/kernels/reductions.hpp @@ -56,6 +56,8 @@ namespace dpnp::tensor::kernels { +using dpnp::tensor::type_utils::normalize_bool; + using dpnp::tensor::ssize_t; namespace su_ns = dpnp::tensor::sycl_utils; @@ -1913,7 +1915,7 @@ struct SequentialSearchReduction const ssize_t inp_reduction_offset = inp_reduced_dims_indexer_(m); const ssize_t inp_offset = inp_iter_offset + inp_reduction_offset; - argT val = inp_[inp_offset]; + argT val = normalize_bool(inp_[inp_offset]); if (val == red_val) { idx_val = idx_reduction_op_(idx_val, static_cast(m)); } @@ -2058,7 +2060,7 @@ struct SearchReduction inp_reduced_dims_indexer_(arg_reduce_gid); auto inp_offset = inp_iter_offset + inp_reduction_offset; - argT val = inp_[inp_offset]; + argT val = normalize_bool(inp_[inp_offset]); if (val == local_red_val) { if constexpr (!First) { local_idx = @@ -2216,7 +2218,7 @@ struct CustomSearchReduction inp_reduced_dims_indexer_(arg_reduce_gid); auto inp_offset = inp_iter_offset + inp_reduction_offset; - argT val = inp_[inp_offset]; + argT val = normalize_bool(inp_[inp_offset]); if (val == local_red_val) { if constexpr (!First) { local_idx = diff --git a/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp b/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp index 2388f23fdf61..918eeb7e2f07 100644 --- a/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp +++ b/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp @@ -43,6 +43,7 @@ #include "kernels/sorting/search_sorted_detail.hpp" #include "utils/offset_utils.hpp" #include "utils/rich_comparisons.hpp" +#include "utils/type_utils.hpp" namespace dpnp::tensor::kernels { @@ -87,7 +88,10 @@ struct IsinFunctor static constexpr Compare comp{}; const std::size_t i = id[0]; - const T needle_v = needles_tp[needles_indexer(i)]; + // normalize: for bool a byte other than 0x00/0x01 would not compare + // equal to the normalized value in the hay array, see gh-2121 + const T needle_v = dpnp::tensor::type_utils::normalize_bool( + needles_tp[needles_indexer(i)]); // position of the needle_v in the hay array std::size_t pos{}; @@ -100,7 +104,10 @@ struct IsinFunctor // needle_v) is false, i.e. needle_v <= hay[pos] pos = search_sorted_detail::lower_bound_indexed_impl( hay_tp, zero, hay_nelems, needle_v, comp, hay_indexer); - bool out = (pos == hay_nelems ? false : hay_tp[pos] == needle_v); + bool out = + (pos == hay_nelems ? false + : dpnp::tensor::type_utils::normalize_bool( + hay_tp[pos]) == needle_v); out_tp[out_indexer(i)] = (invert) ? !out : out; } }; diff --git a/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp b/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp index 27eade9f3586..7c42ae3edbaa 100644 --- a/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp +++ b/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp @@ -50,6 +50,7 @@ #include "kernels/dpnp_tensor_types.hpp" #include "kernels/sorting/sort_utils.hpp" #include "utils/sycl_alloc_utils.hpp" +#include "utils/type_utils.hpp" namespace dpnp::tensor::kernels { @@ -1785,7 +1786,13 @@ struct IndexedProj { } - auto operator()(IndexT i) const { return value_projector(ptr[i]); } + auto operator()(IndexT i) const + { + // normalize the value read from memory: for bool a byte other than + // 0x00/0x01 would otherwise order by its raw value, see gh-2121 + return value_projector( + dpnp::tensor::type_utils::normalize_bool(ptr[i])); + } private: const ValueT *ptr; diff --git a/dpnp/tensor/libtensor/include/utils/rich_comparisons.hpp b/dpnp/tensor/libtensor/include/utils/rich_comparisons.hpp index 0db7168ab9f7..dcd6d740fab1 100644 --- a/dpnp/tensor/libtensor/include/utils/rich_comparisons.hpp +++ b/dpnp/tensor/libtensor/include/utils/rich_comparisons.hpp @@ -40,6 +40,8 @@ #include "sycl/sycl.hpp" +#include "utils/type_utils.hpp" + namespace dpnp::tensor::rich_comparisons { @@ -116,6 +118,26 @@ inline constexpr bool is_fp_v = (std::is_same_v || std::is_same_v || std::is_same_v); +// takes by reference: copying a bool first would let the compiler assume a +// 0/1 byte and fold the normalization away, see gh-2121 +struct BoolLess +{ + bool operator()(const bool &v1, const bool &v2) const + { + using dpnp::tensor::type_utils::normalize_bool; + return !normalize_bool(v1) && normalize_bool(v2); + } +}; + +struct BoolGreater +{ + bool operator()(const bool &v1, const bool &v2) const + { + using dpnp::tensor::type_utils::normalize_bool; + return normalize_bool(v1) && !normalize_bool(v2); + } +}; + } // namespace detail template @@ -126,6 +148,12 @@ struct AscendingSorter std::less>; }; +template <> +struct AscendingSorter +{ + using type = detail::BoolLess; +}; + template struct AscendingSorter> { @@ -140,6 +168,12 @@ struct DescendingSorter std::greater>; }; +template <> +struct DescendingSorter +{ + using type = detail::BoolGreater; +}; + template struct DescendingSorter> { diff --git a/dpnp/tensor/libtensor/include/utils/type_utils.hpp b/dpnp/tensor/libtensor/include/utils/type_utils.hpp index b62bfb38a2e6..1a1a19a86715 100644 --- a/dpnp/tensor/libtensor/include/utils/type_utils.hpp +++ b/dpnp/tensor/libtensor/include/utils/type_utils.hpp @@ -60,11 +60,28 @@ struct is_complex< template inline constexpr bool is_complex_v = is_complex::value; +// NumPy reads any non-zero byte of a bool as True, while a plain comparison +// may fold into a raw byte load, see gh-2121 +template +T normalize_bool(const T &v) +{ + if constexpr (std::is_same_v) { + // read the storage as a byte: a bool copy would let the compiler + // assume a 0/1 value and fold this away + const std::uint8_t u = *reinterpret_cast(&v); + return u != std::uint8_t{0}; + } + else { + return v; + } +} + template dstTy convert_impl(const srcTy &v) { if constexpr (std::is_same_v) { - return v; + // bool needs normalizing even here, the byte may not be 0x00/0x01 + return normalize_bool(v); } else if constexpr (std::is_same_v) { if constexpr (is_complex_v) { diff --git a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py index 88f8980a8c9d..cd5962615fb9 100644 --- a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py +++ b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py @@ -34,6 +34,7 @@ from numpy.testing import assert_, assert_array_equal, assert_raises_regex import dpnp.tensor as dpt +import dpnp.tensor._tensor_impl as ti from dpnp.tensor._numpy_helper import AxisError from .helper import get_queue_or_skip @@ -1423,6 +1424,12 @@ def test_repeat_nonstandard_bool_bytes(): res = dpt.repeat(x, reps) assert_array_equal(dpt.asnumpy(res), np.array([1, 2, 3, 5], dtype="i4")) + # `repeat` casts a non-int64 `reps` first, so drive the scan directly too + cumsum = dpt.empty(reps.size, dtype="i8") + total = ti._cumsum_1d(reps, cumsum, sycl_queue=reps.sycl_queue) + assert total == 4 + assert_array_equal(dpt.asnumpy(cumsum), np.array([0, 1, 2, 3, 3, 4])) + def test_repeat_size1_repeats(): get_queue_or_skip() diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index a7696fe9852d..8296d2609177 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -900,3 +900,114 @@ def test_isin_errors(self): with pytest.raises(ExecutionPlacementError): dpnp.isin(a, b) + + +# NumPy treats any non-zero byte of a bool as True, see gh-2121 +class TestNonstandardBoolBytes: + def _views(self, values): + raw = numpy.array(values, dtype=numpy.uint8) + return raw.view(numpy.bool_), dpnp.asarray(raw).view(dpnp.bool) + + @pytest.mark.parametrize( + "op", + [ + "equal", + "not_equal", + "less", + "less_equal", + "greater", + "greater_equal", + "logical_and", + "logical_or", + "logical_xor", + "bitwise_and", + "bitwise_or", + "bitwise_xor", + "maximum", + "minimum", + ], + ) + def test_binary(self, op): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + b, ib = self._views([1, 1, 0, 2, 0, 7]) + + result = getattr(dpnp, op)(ia, ib) + expected = getattr(numpy, op)(a, b) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("op", ["logical_not", "bitwise_invert"]) + def test_unary(self, op): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + result = getattr(dpnp, op)(ia) + expected = getattr(numpy, op)(a) + assert_array_equal(result, expected) + + @pytest.mark.parametrize( + "op", ["sum", "prod", "max", "min", "all", "any", "argmax", "argmin"] + ) + def test_reduction(self, op): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + result = getattr(dpnp, op)(ia) + expected = getattr(numpy, op)(a) + assert_array_equal(result, expected) + + def test_result_has_canonical_bytes(self): + # a computed bool result must only ever hold 0x00 or 0x01; `sort` and + # `unique` are excluded, they select elements and so carry the bytes + _, ia = self._views([0, 1, 2, 255, 3, 0]) + _, ib = self._views([1, 1, 0, 2, 0, 7]) + + for res in ( + dpnp.equal(ia, ib), + dpnp.logical_not(ia), + dpnp.maximum(ia, ib), + dpnp.isin(ia, ib), + dpnp.max(ia), + dpnp.min(ia), + dpnp.all(ia), + dpnp.any(ia), + ): + raw = dpnp.asnumpy(res).view(numpy.uint8) + assert numpy.all((raw == 0) | (raw == 1)) + + @pytest.mark.parametrize( + "values", + [[0, 1, 2, 255, 3, 0], [2] * 6, [255] * 6, [0, 128, 0, 127, 254, 1]], + ) + @pytest.mark.parametrize("kind", [None, "mergesort", "radixsort"]) + def test_sort_kinds_logically_ordered(self, values, kind): + a, ia = self._views(values) + kwargs = {} if kind is None else {"kind": kind} + + result = dpnp.asnumpy(dpnp.sort(ia, **kwargs)).view(numpy.uint8) != 0 + expected = numpy.sort(a).view(numpy.uint8) != 0 + assert_array_equal(result, expected) + + def test_unique(self): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + assert_array_equal(dpnp.unique(ia), numpy.unique(a)) + + def test_isin(self): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + b, ib = self._views([1, 1, 0, 2, 0, 7]) + + assert_array_equal(dpnp.isin(ia, ib), numpy.isin(a, b)) + + def test_sort_is_logically_ordered(self): + # dpnp normalizes, so False elements sort before True ones; NumPy + # leaves the raw bytes in place, so only the logical order matches + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + result = dpnp.asnumpy(dpnp.sort(ia)).view(numpy.uint8) != 0 + expected = numpy.sort(a).view(numpy.uint8) != 0 + assert_array_equal(result, expected) + + def test_argsort_is_logically_ordered(self): + raw = numpy.array([0, 1, 2, 255, 3, 0], dtype=numpy.uint8) + ia = dpnp.asarray(raw).view(dpnp.bool) + + gathered = raw[dpnp.asnumpy(dpnp.argsort(ia))] != 0 + assert_array_equal(gathered, numpy.sort(gathered)) From f5028a6a33e5caf954bdd54cc0da876707d5505f Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Wed, 2 Sep 2026 10:43:08 -0500 Subject: [PATCH 3/5] Update CHANGELOG.md --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 66bb70d718a3..ddec33786c7b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -94,7 +94,7 @@ This release is compatible with NumPy 2.5. * Fixed `astype` casting an out-of-range floating point value to a signed narrow integer type saturating to the destination min/max instead of wrapping like NumPy, generalizing the earlier unsigned-only fix [#3033](https://github.com/IntelPython/dpnp/pull/3033) * Fixed `dpnp.insert` silently ignoring out-of-bounds negative indices in a multi-element `obj`, so a mix of in-bounds and out-of-bounds indices now consistently raises `IndexError` [#3041](https://github.com/IntelPython/dpnp/pull/3041) * Fixed a per-call `sycl::queue` leak in `usm_ndarray::get_queue()`/`get_device()` [#3042](https://github.com/IntelPython/dpnp/pull/3042) -* Fixed operations on a boolean array whose bytes are not `0x00`/`0x01`, e.g. a `view` of integer data, which returned wrong results or produced such bytes themselves. Affected comparison, logical, bitwise and arithmetic functions, `dpnp.max`/`dpnp.min`, `dpnp.argmax`/`dpnp.argmin`, `dpnp.sort`/`dpnp.argsort`, `dpnp.unique`, `dpnp.isin`, `dpnp.nonzero`, `dpnp.where`, `dpnp.extract`, `dpnp.place` and `dpnp.repeat` [#3053](https://github.com/IntelPython/dpnp/pull/3053) +* Fixed operations on a boolean array whose bytes are not `0x00`/`0x01` [#3055](https://github.com/IntelPython/dpnp/pull/3055) ### Security From 202dfbc7f9964b6a3fdfcc95c1a9d027458a775f Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty Date: Wed, 2 Sep 2026 13:00:39 -0500 Subject: [PATCH 4/5] fix: normalize bool in radix sort's order-preserving cast `bool` takes a single radix pass, so the bucket index is `byte & 0xF`. An unnormalized byte whose low nibble is zero (0x10, 0x80, 0xF0) bucketed as False and sorted before True elements. Normalize in `order_preserving_cast`, where every radix path converges, and take the argument by reference so a copy cannot let the compiler assume a 0/1 byte. --- .../libtensor/include/kernels/sorting/radix_sort.hpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp b/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp index 7c42ae3edbaa..163f2ae64dcc 100644 --- a/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp +++ b/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp @@ -117,12 +117,15 @@ std::uint32_t ceil_log2(SizeT n) //---------------------------------------------------------- template -bool order_preserving_cast(bool val) +bool order_preserving_cast(const bool &val) { + // by reference: a bool copy lets the compiler assume a 0/1 byte, and the + // bucket index below reads only the low radix bits, see gh-2121 + const bool v = dpnp::tensor::type_utils::normalize_bool(val); if constexpr (is_ascending) - return val; + return v; else - return !val; + return !v; } template Date: Thu, 3 Sep 2026 13:42:11 -0500 Subject: [PATCH 5/5] fix: address PR review comments --- .../include/kernels/elementwise_functions/common.hpp | 4 ++-- .../kernels/elementwise_functions/common_inplace.hpp | 2 +- dpnp/tensor/libtensor/include/utils/type_utils.hpp | 2 +- dpnp/tests/test_logic.py | 11 ++--------- 4 files changed, 6 insertions(+), 13 deletions(-) diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp index 468d68b58579..39e472ddb727 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp @@ -152,7 +152,7 @@ struct UnaryContigFunctor const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { // scalar call - out[k] = op(normalize_bool(in[k])); + out[k] = op(in[k]); } } } @@ -464,7 +464,7 @@ struct BinaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(normalize_bool(in1[k]), normalize_bool(in2[k])); + out[k] = op(in1[k], in2[k]); } } } diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp index 07d8deb2003e..bd746ab182e4 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp @@ -127,7 +127,7 @@ struct BinaryInplaceContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - op(lhs[k], normalize_bool(rhs[k])); + op(lhs[k], rhs[k]); } } } diff --git a/dpnp/tensor/libtensor/include/utils/type_utils.hpp b/dpnp/tensor/libtensor/include/utils/type_utils.hpp index 1a1a19a86715..72ef6e903055 100644 --- a/dpnp/tensor/libtensor/include/utils/type_utils.hpp +++ b/dpnp/tensor/libtensor/include/utils/type_utils.hpp @@ -68,7 +68,7 @@ T normalize_bool(const T &v) if constexpr (std::is_same_v) { // read the storage as a byte: a bool copy would let the compiler // assume a 0/1 value and fold this away - const std::uint8_t u = *reinterpret_cast(&v); + const std::uint8_t u = sycl::bit_cast(v); return u != std::uint8_t{0}; } else { diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index 8296d2609177..b03144aa880f 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -978,6 +978,8 @@ def test_result_has_canonical_bytes(self): ) @pytest.mark.parametrize("kind", [None, "mergesort", "radixsort"]) def test_sort_kinds_logically_ordered(self, values, kind): + # dpnp normalizes, so False elements sort before True ones; NumPy + # leaves the raw bytes in place, so only the logical order matches a, ia = self._views(values) kwargs = {} if kind is None else {"kind": kind} @@ -996,15 +998,6 @@ def test_isin(self): assert_array_equal(dpnp.isin(ia, ib), numpy.isin(a, b)) - def test_sort_is_logically_ordered(self): - # dpnp normalizes, so False elements sort before True ones; NumPy - # leaves the raw bytes in place, so only the logical order matches - a, ia = self._views([0, 1, 2, 255, 3, 0]) - - result = dpnp.asnumpy(dpnp.sort(ia)).view(numpy.uint8) != 0 - expected = numpy.sort(a).view(numpy.uint8) != 0 - assert_array_equal(result, expected) - def test_argsort_is_logically_ordered(self): raw = numpy.array([0, 1, 2, 255, 3, 0], dtype=numpy.uint8) ia = dpnp.asarray(raw).view(dpnp.bool)