Fix nonzero bool nonstandard bytes

CHANGELOG.md

@@ -98,6 +98,7 @@ This release is compatible with NumPy 2.5.
 * Fixed `dpnp.linspace` returning `nan` for equal infinite endpoints [#3043](https://github.com/IntelPython/dpnp/pull/3043)
 * Fixed `dpnp.cumsum`, `dpnp.cumprod`, and their `nan`/`cumulative_*` variants (including `dpnp.tensor.cumulative_sum`/`cumulative_prod`) silently returning incorrect results when accumulating along an axis of an array with more than one row [#3063](https://github.com/IntelPython/dpnp/pull/3063)
 * Fixed `dpnp.einsum` returning a result whose memory layout differs from NumPy for the default `order="K"`, and ignoring `out` and `order` for a contraction over a size-0 dimension [#3058](https://github.com/IntelPython/dpnp/pull/3058)
+* 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; read the raw byte
+            // rather than the bool value, 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;
+        }
     }
 };
 
@@ -85,7 +93,12 @@ struct NoOpTransformer
 {
     constexpr NoOpTransformer() {}
 
-    T operator()(const T &val) const { return val; }
+    // bool is normalized: a byte other than 0x00/0x01 would otherwise reach
+    // the scan operation and leak into the result, see gh-2121
+    T operator()(const T &val) const
+    {
+        return dpnp::tensor::type_utils::normalize_bool(val);
+    }
 };
 
 template <typename srcTy, typename dstTy>
@@ -1303,8 +1316,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 {
@@ -1421,8 +1436,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,22 @@ 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;
+
+// a bool lhs is read as a value by the functors, so normalize it too and write
+// the result back as a canonical byte, see gh-2121
+template <typename OpT, typename resT, typename argT>
+void apply_inplace(OpT &op, resT &res, const argT &rhs)
+{
+    if constexpr (std::is_same_v<resT, bool>) {
+        resT tmp = normalize_bool(res);
+        op(tmp, normalize_bool(rhs));
+        res = tmp;
+    }
+    else {
+        op(res, normalize_bool(rhs));
+    }
+}
 
 template <typename argT,
           typename resT,
@@ -88,10 +105,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 +173,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]);
+                        apply_inplace(op, res_vec[vec_id], 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]);
+                    apply_inplace(op, lhs[k], rhs[k]);
                 }
             }
         }
@@ -176,7 +195,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]);
+                apply_inplace(op, lhs[offset], rhs[offset]);
             }
         }
     }
@@ -210,7 +229,7 @@ struct BinaryInplaceStridedFunctor
         const auto &lhs_offset = two_offsets_.get_second_offset();
 
         BinaryInplaceOperatorT op{};
-        op(lhs[lhs_offset], rhs[inp_offset]);
+        apply_inplace(op, lhs[lhs_offset], rhs[inp_offset]);
     }
 };
 
@@ -257,14 +276,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);
+            apply_inplace(op, mat_el, 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]);
+                apply_inplace(op, mat[k], 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 =