diff --git a/CHANGELOG.md b/CHANGELOG.md
index e37ae7055273..ddec33786c7b 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 operations on a boolean array whose bytes are not `0x00`/`0x01` [#3055](https://github.com/IntelPython/dpnp/pull/3055)
 
 ### Security
 
diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp
index 079fcf5e9c96..584f913b8965 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<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 {
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 <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];
 
@@ -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<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);
                 }
@@ -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<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);
                 }
@@ -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<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];
 
@@ -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<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);
                 }
@@ -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 <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];
 
@@ -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<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);
                 }
@@ -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<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 =
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..163f2ae64dcc 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
 {
@@ -116,12 +117,15 @@ std::uint32_t ceil_log2(SizeT n)
 //----------------------------------------------------------
 
 template <bool is_ascending>
-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 <bool is_ascending,
@@ -1785,7 +1789,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<T, sycl::half> || std::is_same_v<T, float> ||
      std::is_same_v<T, double>);
 
+// 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 <typename argTy>
@@ -126,6 +148,12 @@ struct AscendingSorter
                                     std::less<argTy>>;
 };
 
+template <>
+struct AscendingSorter<bool>
+{
+    using type = detail::BoolLess;
+};
+
 template <typename T>
 struct AscendingSorter<std::complex<T>>
 {
@@ -140,6 +168,12 @@ struct DescendingSorter
                                     std::greater<argTy>>;
 };
 
+template <>
+struct DescendingSorter<bool>
+{
+    using type = detail::BoolGreater;
+};
+
 template <typename T>
 struct DescendingSorter<std::complex<T>>
 {
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 <typename T>
 inline constexpr bool is_complex_v = is_complex<T>::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 <typename T>
+T normalize_bool(const T &v)
+{
+    if constexpr (std::is_same_v<T, bool>) {
+        // 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<const std::uint8_t *>(&v);
+        return u != std::uint8_t{0};
+    }
+    else {
+        return v;
+    }
+}
+
 template <typename dstTy, typename srcTy>
 dstTy convert_impl(const srcTy &v)
 {
     if constexpr (std::is_same_v<dstTy, srcTy>) {
-        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<dstTy, bool>) {
         if constexpr (is_complex_v<srcTy>) {
diff --git a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py
index bf0bd3226309..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
@@ -1412,6 +1413,24 @@ 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"))
+
+    # `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_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))
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))