Skip to content

[SYCL] Fix marray ~operator for bool data type#22440

Open
uditagarwal97 wants to merge 2 commits into
syclfrom
private/udit/fix_marray_operator
Open

[SYCL] Fix marray ~operator for bool data type#22440
uditagarwal97 wants to merge 2 commits into
syclfrom
private/udit/fix_marray_operator

Conversation

@uditagarwal97

@uditagarwal97 uditagarwal97 commented Jun 25, 2026

Copy link
Copy Markdown
Contributor

We use uint8_t for bool when using ext_vector_type on device and
~uint8_t(1) = 0xFE, and static_cast<bool>(0xFE) is true, so the original storeVecResult path produced the wrong result for marray<bool>

Fixes CMPLRLLVM-76430

@uditagarwal97 uditagarwal97 self-assigned this Jun 25, 2026
@uditagarwal97 uditagarwal97 marked this pull request as ready for review June 25, 2026 21:38
@uditagarwal97 uditagarwal97 requested a review from a team as a code owner June 25, 2026 21:38
@uditagarwal97 uditagarwal97 requested a review from slawekptak June 25, 2026 21:38
@slawekptak

Copy link
Copy Markdown
Contributor

@uditagarwal97 could you please provide a bit of background info for this change? Thanks!

@uditagarwal97 uditagarwal97 changed the title [SYCL] Fix marray ~operator for bool data type [SYCL] Fix marray ~operator for bool data type Jun 26, 2026
@uditagarwal97 uditagarwal97 requested a review from Copilot June 27, 2026 07:20

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

This PR fixes the behavior of the bitwise NOT (~) operator for sycl::marray<bool, N>, ensuring it produces a proper boolean inversion instead of returning values that always convert to true after bitwise negation.

Changes:

  • Adjust marray’s operator~ implementation to correctly invert bool elements in both the ext-vector and scalar paths.
  • Add a new end-to-end test covering multiple marray<bool, N> sizes to exercise both implementation paths (including the NumElements == 3 non-ext-vector case).

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated no comments.

File Description
sycl/include/sycl/marray.hpp Fixes operator~ for marray<bool, N> by normalizing results to boolean inversion in both ext-vector and non-ext-vector implementations.
sycl/test-e2e/Basic/built-ins/marray_bitwise_not_bool.cpp Adds an E2E regression test validating ~ on marray<bool, N> across several sizes and value patterns.

@iclsrc iclsrc left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The fix is correct: ~uint8_t(1) = 0xFE, and static_cast<bool>(0xFE) is true, so the original storeVecResult path produced the wrong result for marray<bool>. Both the ext_vector and non-vector paths are now logically correct. One minor readability suggestion on the ext_vector path.

Comment on lines +460 to +462
if constexpr (std::is_same_v<DataT, bool>) {
for (size_t I = 0; I < NumElements; ++I)
Ret[I] = ResVec[I] & 1;

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 Suggestion: The ResVec[I] & 1 expression relies on a subtle property — that ~(0 or 1) always flips the LSB — and it also wastes the ResVec = ~LhsVec computation on line 459 (the compiler will likely optimize it away, but it's dead work). Consider !LhsVec[I] instead, which mirrors the non-vector path's !Lhs[I] pattern and is immediately readable without needing to reason about bitwise complement:

Suggested change
if constexpr (std::is_same_v<DataT, bool>) {
for (size_t I = 0; I < NumElements; ++I)
Ret[I] = ResVec[I] & 1;
if constexpr (std::is_same_v<DataT, bool>) {
for (size_t I = 0; I < NumElements; ++I)
Ret[I] = !LhsVec[I];

@uditagarwal97

Copy link
Copy Markdown
Contributor Author

@uditagarwal97 could you please provide a bit of background info for this change? Thanks!

@slawekptak I've updated the PR description.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants