From eaadfb3432c499abbb292fe489ce2895c29a5813 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Thu, 25 Jun 2026 20:16:17 +0200 Subject: [PATCH 1/2] [SYCL] Fix marray ~ operator for bool data type --- sycl/include/sycl/marray.hpp | 68 +++++++------ .../built-ins/marray_bitwise_not_bool.cpp | 99 +++++++++++++++++++ 2 files changed, 138 insertions(+), 29 deletions(-) create mode 100644 sycl/test-e2e/Basic/built-ins/marray_bitwise_not_bool.cpp diff --git a/sycl/include/sycl/marray.hpp b/sycl/include/sycl/marray.hpp index 37fba95b1e538..8e430990a4117 100644 --- a/sycl/include/sycl/marray.hpp +++ b/sycl/include/sycl/marray.hpp @@ -266,7 +266,8 @@ template class marray { std::is_convertible_v && \ (std::is_fundamental_v || \ std::is_same_v::type, half>), \ - marray> operator BINOP(const marray & Lhs, const T & Rhs) { \ + marray> \ + operator BINOP(const marray &Lhs, const T &Rhs) { \ return Lhs BINOP marray(static_cast(Rhs)); \ } \ template \ @@ -274,16 +275,17 @@ template class marray { std::is_convertible_v && \ (std::is_fundamental_v || \ std::is_same_v::type, half>), \ - marray> operator BINOP(const T & Lhs, const marray & Rhs) { \ + marray> \ + operator BINOP(const T &Lhs, const marray &Rhs) { \ return marray(static_cast(Lhs)) BINOP Rhs; \ } \ - friend marray &operator OPASSIGN(marray & Lhs, const marray & Rhs) { \ + friend marray &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \ Lhs = Lhs BINOP Rhs; \ return Lhs; \ } \ template \ friend typename std::enable_if_t operator OPASSIGN( \ - marray & Lhs, const DataT & Rhs) { \ + marray &Lhs, const DataT &Rhs) { \ Lhs = Lhs BINOP marray(Rhs); \ return Lhs; \ } @@ -309,29 +311,28 @@ template class marray { friend typename std::enable_if_t && \ std::is_integral_v && \ std::is_integral_v, \ - marray> operator BINOP(const marray & Lhs, \ - const T & Rhs) { \ + marray> \ + operator BINOP(const marray &Lhs, const T &Rhs) { \ return Lhs BINOP marray(static_cast(Rhs)); \ } \ template \ - friend \ - typename std::enable_if_t && \ - std::is_integral_v && \ - std::is_integral_v, \ - marray> operator BINOP(const T & Lhs, \ - const marray & Rhs) { \ + friend typename std::enable_if_t && \ + std::is_integral_v && \ + std::is_integral_v, \ + marray> \ + operator BINOP(const T &Lhs, const marray &Rhs) { \ return marray(static_cast(Lhs)) BINOP Rhs; \ } \ template , marray>> \ - friend marray &operator OPASSIGN(marray & Lhs, const marray & Rhs) { \ + friend marray &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \ Lhs = Lhs BINOP Rhs; \ return Lhs; \ } \ template \ friend \ typename std::enable_if_t, marray &> \ - operator OPASSIGN(marray & Lhs, const DataT & Rhs) { \ + operator OPASSIGN(marray &Lhs, const DataT &Rhs) { \ Lhs = Lhs BINOP marray(Rhs); \ return Lhs; \ } @@ -359,8 +360,8 @@ template class marray { #endif #define __SYCL_RELLOGOP(RELLOGOP) \ - friend marray operator RELLOGOP(const marray & Lhs, \ - const marray & Rhs) { \ + friend marray operator RELLOGOP(const marray &Lhs, \ + const marray &Rhs) { \ marray Ret; \ if constexpr (use_ext_vector_type) { \ ext_vector_t LhsVec = sycl::bit_cast(Lhs.MData); \ @@ -377,19 +378,19 @@ template class marray { return Ret; \ } \ template \ - friend typename std::enable_if_t< \ - std::is_convertible_v && \ - (std::is_fundamental_v || std::is_same_v), \ - marray> operator RELLOGOP(const marray & Lhs, \ - const T & Rhs) { \ + friend typename std::enable_if_t && \ + (std::is_fundamental_v || \ + std::is_same_v), \ + marray> \ + operator RELLOGOP(const marray &Lhs, const T &Rhs) { \ return Lhs RELLOGOP marray(static_cast(Rhs)); \ } \ template \ - friend typename std::enable_if_t< \ - std::is_convertible_v && \ - (std::is_fundamental_v || std::is_same_v), \ - marray> operator RELLOGOP(const T & Lhs, \ - const marray & Rhs) { \ + friend typename std::enable_if_t && \ + (std::is_fundamental_v || \ + std::is_same_v), \ + marray> \ + operator RELLOGOP(const T &Lhs, const marray &Rhs) { \ return marray(static_cast(Lhs)) RELLOGOP Rhs; \ } @@ -457,10 +458,19 @@ template class marray { if constexpr (use_ext_vector_type) { ext_vector_t LhsVec = sycl::bit_cast(Lhs.MData); ext_vector_t ResVec = ~LhsVec; - sycl::detail::memcpy_no_adl(Ret.MData, &ResVec, sizeof(ResVec)); + if constexpr (std::is_same_v) { + for (size_t I = 0; I < NumElements; ++I) + Ret[I] = ResVec[I] & 1; + } else { + storeVecResult(Ret.MData, ResVec); + } } else { - for (size_t I = 0; I < NumElements; ++I) { - Ret[I] = ~Lhs[I]; + if constexpr (std::is_same_v) { + for (size_t I = 0; I < NumElements; ++I) + Ret[I] = !Lhs[I]; + } else { + for (size_t I = 0; I < NumElements; ++I) + Ret[I] = ~Lhs[I]; } } return Ret; diff --git a/sycl/test-e2e/Basic/built-ins/marray_bitwise_not_bool.cpp b/sycl/test-e2e/Basic/built-ins/marray_bitwise_not_bool.cpp new file mode 100644 index 0000000000000..f012b06182623 --- /dev/null +++ b/sycl/test-e2e/Basic/built-ins/marray_bitwise_not_bool.cpp @@ -0,0 +1,99 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Test bitwise NOT operator (~) on marray + +#include +#include + +#include +#include + +using namespace sycl; + +template +bool test_bitwise_not_bool(queue &q, const marray &input, + const marray &expected) { + bool result[N]; + { + buffer out_buf(result, N); + q.submit([&](handler &h) { + accessor out_acc(out_buf, h, write_only); + h.single_task([=]() { + marray res = ~input; + for (size_t i = 0; i < N; ++i) { + out_acc[i] = res[i]; + } + }); + }).wait(); + } + + // Verify results + for (size_t i = 0; i < N; ++i) { + if (result[i] != expected[i]) { + std::cout << "FAILED at index " << i << ": input=" << input[i] + << ", expected=" << expected[i] << ", got=" << result[i] + << std::endl; + return false; + } + } + return true; +} + +int main() { + queue q; + + std::cout << "Testing bitwise NOT (~) on marray\n"; + + // Test case 1: Size 2 + { + marray input{true, false}; + marray expected{false, true}; + assert(test_bitwise_not_bool(q, input, expected) && + "Test failed for size 2"); + } + + // Test case 2: Size 4 + { + marray input{true, false, true, false}; + marray expected{false, true, false, true}; + assert(test_bitwise_not_bool(q, input, expected) && + "Test failed for size 4"); + } + + // Test case 3: Size 3 (no padding, different code path) + { + marray input{false, true, false}; + marray expected{true, false, true}; + assert(test_bitwise_not_bool(q, input, expected) && + "Test failed for size 3"); + } + + // Test case 4: Size 8 + { + marray input{true, true, false, false, true, false, true, false}; + marray expected{false, false, true, true, + false, true, false, true}; + assert(test_bitwise_not_bool(q, input, expected) && + "Test failed for size 8"); + } + + // Test case 5: All true + { + marray input{true, true, true, true}; + marray expected{false, false, false, false}; + assert(test_bitwise_not_bool(q, input, expected) && + "Test failed for all true"); + } + + // Test case 6: All false + { + marray input{false, false, false, false}; + marray expected{true, true, true, true}; + assert(test_bitwise_not_bool(q, input, expected) && + "Test failed for all false"); + } + + std::cout << "All tests passed!\n"; + return 0; +} From e42b1f5a98df476780c3e2dd4139aad6f00635fa Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Thu, 25 Jun 2026 22:53:39 +0200 Subject: [PATCH 2/2] Fix formatting --- sycl/include/sycl/marray.hpp | 53 ++++++++++++++++++------------------ 1 file changed, 26 insertions(+), 27 deletions(-) diff --git a/sycl/include/sycl/marray.hpp b/sycl/include/sycl/marray.hpp index 8e430990a4117..502b57ff3f7fc 100644 --- a/sycl/include/sycl/marray.hpp +++ b/sycl/include/sycl/marray.hpp @@ -266,8 +266,7 @@ template class marray { std::is_convertible_v && \ (std::is_fundamental_v || \ std::is_same_v::type, half>), \ - marray> \ - operator BINOP(const marray &Lhs, const T &Rhs) { \ + marray> operator BINOP(const marray & Lhs, const T & Rhs) { \ return Lhs BINOP marray(static_cast(Rhs)); \ } \ template \ @@ -275,17 +274,16 @@ template class marray { std::is_convertible_v && \ (std::is_fundamental_v || \ std::is_same_v::type, half>), \ - marray> \ - operator BINOP(const T &Lhs, const marray &Rhs) { \ + marray> operator BINOP(const T & Lhs, const marray & Rhs) { \ return marray(static_cast(Lhs)) BINOP Rhs; \ } \ - friend marray &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \ + friend marray &operator OPASSIGN(marray & Lhs, const marray & Rhs) { \ Lhs = Lhs BINOP Rhs; \ return Lhs; \ } \ template \ friend typename std::enable_if_t operator OPASSIGN( \ - marray &Lhs, const DataT &Rhs) { \ + marray & Lhs, const DataT & Rhs) { \ Lhs = Lhs BINOP marray(Rhs); \ return Lhs; \ } @@ -311,28 +309,29 @@ template class marray { friend typename std::enable_if_t && \ std::is_integral_v && \ std::is_integral_v, \ - marray> \ - operator BINOP(const marray &Lhs, const T &Rhs) { \ + marray> operator BINOP(const marray & Lhs, \ + const T & Rhs) { \ return Lhs BINOP marray(static_cast(Rhs)); \ } \ template \ - friend typename std::enable_if_t && \ - std::is_integral_v && \ - std::is_integral_v, \ - marray> \ - operator BINOP(const T &Lhs, const marray &Rhs) { \ + friend \ + typename std::enable_if_t && \ + std::is_integral_v && \ + std::is_integral_v, \ + marray> operator BINOP(const T & Lhs, \ + const marray & Rhs) { \ return marray(static_cast(Lhs)) BINOP Rhs; \ } \ template , marray>> \ - friend marray &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \ + friend marray &operator OPASSIGN(marray & Lhs, const marray & Rhs) { \ Lhs = Lhs BINOP Rhs; \ return Lhs; \ } \ template \ friend \ typename std::enable_if_t, marray &> \ - operator OPASSIGN(marray &Lhs, const DataT &Rhs) { \ + operator OPASSIGN(marray & Lhs, const DataT & Rhs) { \ Lhs = Lhs BINOP marray(Rhs); \ return Lhs; \ } @@ -360,8 +359,8 @@ template class marray { #endif #define __SYCL_RELLOGOP(RELLOGOP) \ - friend marray operator RELLOGOP(const marray &Lhs, \ - const marray &Rhs) { \ + friend marray operator RELLOGOP(const marray & Lhs, \ + const marray & Rhs) { \ marray Ret; \ if constexpr (use_ext_vector_type) { \ ext_vector_t LhsVec = sycl::bit_cast(Lhs.MData); \ @@ -378,19 +377,19 @@ template class marray { return Ret; \ } \ template \ - friend typename std::enable_if_t && \ - (std::is_fundamental_v || \ - std::is_same_v), \ - marray> \ - operator RELLOGOP(const marray &Lhs, const T &Rhs) { \ + friend typename std::enable_if_t< \ + std::is_convertible_v && \ + (std::is_fundamental_v || std::is_same_v), \ + marray> operator RELLOGOP(const marray & Lhs, \ + const T & Rhs) { \ return Lhs RELLOGOP marray(static_cast(Rhs)); \ } \ template \ - friend typename std::enable_if_t && \ - (std::is_fundamental_v || \ - std::is_same_v), \ - marray> \ - operator RELLOGOP(const T &Lhs, const marray &Rhs) { \ + friend typename std::enable_if_t< \ + std::is_convertible_v && \ + (std::is_fundamental_v || std::is_same_v), \ + marray> operator RELLOGOP(const T & Lhs, \ + const marray & Rhs) { \ return marray(static_cast(Lhs)) RELLOGOP Rhs; \ }