Fix nonzero bool nonstandard bytes

CHANGELOG.md

@@ -95,6 +95,7 @@ This release is compatible with NumPy 2.5.
 * 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.linspace` returning `nan` for equal infinite endpoints [#3043](https://github.com/IntelPython/dpnp/pull/3043)
+* Fixed operations on a boolean array whose bytes are not `0x00`/`0x01` [#3055](https://github.com/IntelPython/dpnp/pull/3055)
 
 ### Security
 

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<inputT, bool>) {
+            // 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<std::uint8_t>(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<T>) {
             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<T, cumsumT, NoOpTransformer<cumsumT>>;
+                cumsum_val_contig_impl<T, cumsumT, CastTransformer<T, cumsumT>>;
             return fn;
         }
         else {
@@ -1419,8 +1429,10 @@ struct Cumsum1DStridedFactory
     {
         if constexpr (std::is_integral_v<T>) {
             using cumsumT = std::int64_t;
-            fnT fn =
-                cumsum_val_strided_impl<T, cumsumT, NoOpTransformer<cumsumT>>;
+            // CastTransformer, not NoOpTransformer: an implicit bool
+            // conversion would read the raw byte, see gh-2121
+            fnT fn = cumsum_val_strided_impl<T, cumsumT,
+                                             CastTransformer<T, cumsumT>>;
             return fn;
         }
         else {

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 <typename argT,
@@ -117,10 +119,12 @@ struct UnaryContigFunctor
                 }
             }
         }
+        // bool is excluded from the vector path: a byte other than 0x00/0x01
+        // cannot be normalized element-wise there, see gh-2121
         else if constexpr (enable_sg_loadstore &&
                            UnaryOperatorT::supports_sg_loadstore::value &&
                            UnaryOperatorT::supports_vec::value &&
-                           (vec_sz > 1)) {
+                           !std::is_same_v<argT, bool> && (vec_sz > 1)) {
             auto sg = ndit.get_sub_group();
             const std::uint16_t sgSize = sg.get_max_local_range()[0];
 
@@ -178,15 +182,15 @@ struct UnaryContigFunctor
                         sub_group_load<vec_sz>(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<vec_sz>(sg, arg_vec, out_multi_ptr);
                 }
             }
             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,15 +220,15 @@ struct UnaryContigFunctor
                     sycl::vec<resT, vec_sz> 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<vec_sz>(sg, res_vec, out_multi_ptr);
                 }
             }
             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<argT1, bool> &&
+                      !std::is_same_v<argT2, bool> && (vec_sz > 1)) {
             auto sg = ndit.get_sub_group();
             std::uint16_t sgSize = sg.get_max_local_range()[0];
 
@@ -491,16 +499,16 @@ struct BinaryContigFunctor
                     sycl::vec<resT, vec_sz> 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<vec_sz>(sg, res_vec, out_multi_ptr);
                 }
             }
             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]));
             }
         }
     }

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 <typename argT,
           typename resT,
@@ -88,10 +90,12 @@ struct BinaryInplaceContigFunctor
         /* Each work-item processes vec_sz elements, contiguous in memory */
         /* NB: Workgroup 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 &&
                       BinaryInplaceOperatorT::supports_sg_loadstore::value &&
                       BinaryInplaceOperatorT::supports_vec::value &&
-                      (vec_sz > 1)) {
+                      !std::is_same_v<argT, bool> && (vec_sz > 1)) {
             auto sg = ndit.get_sub_group();
             std::uint16_t sgSize = sg.get_max_local_range()[0];
 
@@ -154,15 +158,15 @@ struct BinaryInplaceContigFunctor
                         sub_group_load<vec_sz>(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<vec_sz>(sg, res_vec, lhs_multi_ptr);
                 }
             }
             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]));
             }
         }
     }

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<outT>(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 =

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;
     }
 };