diff --git a/sycl/include/sycl/marray.hpp b/sycl/include/sycl/marray.hpp index 37fba95b1e538..502b57ff3f7fc 100644 --- a/sycl/include/sycl/marray.hpp +++ b/sycl/include/sycl/marray.hpp @@ -457,10 +457,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; +}