Fix nonzero bool nonstandard bytes - #3055
Conversation
`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<std::uint8_t>` / `CastTransformer`, as `convert_impl` already does (IntelPythongh-2121).
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.
|
Can one of the admins verify this patch? |
ndgrigorian
left a comment
ndgrigorian
left a comment
There was a problem hiding this comment.
I've looked over the changes, seems like a welcome fix for a previously unnoticed bug, LGTM
|
@abagusetty seems that the tests fail with the open-source compiler, interestingly enough. Possible bug in the nightly DPC++? |
|
coverage: 78.472% (-0.02%) from 78.491% — abagusetty:fix-nonzero-bool-nonstandard-bytes into IntelPython:master |
`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.
No bug in nightly. Interestingly, it is doing great. The issue was UB in the PR that the two compilers are treating it differently. A bool whose byte isnt |
| 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])); |
There was a problem hiding this comment.
It seems not needed here, since bool type was excluded above
There was a problem hiding this comment.
good catch! thanks
| 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); |
There was a problem hiding this comment.
Should we use sycl::bit_cast<std::uint8_t> as in convert_impl below?
There was a problem hiding this comment.
Thanks! I contemplated on using sycl::bit_cast<std::uint8_t> but thought about the support for CPU backend support (via opencl:cpu or native_cpu) and fell back to using std::uint8_t. Now that sycl::bit_cast support is robust, cleaner to have it
| 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])); |
There was a problem hiding this comment.
Seems no needed, as input array is never bool here
| 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])); |
There was a problem hiding this comment.
The same, not need to call normalize_bool here
There was a problem hiding this comment.
Perfect, removed it!
5 participants
Boolean arrays may contain non-zero bytes like
0x02or0xFFthat NumPy treats as True but dpnp does not.Fixes #3054