From 13b81d1164a1a4da4fb8c5acd4973c8bd89a4e2d Mon Sep 17 00:00:00 2001
From: Michael Aziz <michael.aziz@intel.com>
Date: Fri, 28 Apr 2023 12:05:51 -0700
Subject: [PATCH] [SYCL] Fix `operator~` for `sycl::vec<bool, N>`

Signed-off-by: Michael Aziz <michael.aziz@intel.com>
---
 sycl/include/sycl/types.hpp      | 6 +++++-
 sycl/test-e2e/Basic/vec_bool.cpp | 2 +-
 2 files changed, 6 insertions(+), 2 deletions(-)

diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp
index b279a2ded7ad9..6af94bda04f03 100644
--- a/sycl/include/sycl/types.hpp
+++ b/sycl/include/sycl/types.hpp
@@ -1228,7 +1228,11 @@ template <typename Type, int NumElements> class vec {
 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
 // by SYCL device compiler only.
 #ifdef __SYCL_DEVICE_ONLY__
-    return vec{(typename vec::DataType) ~m_Data};
+    vec Ret{(typename vec::DataType) ~m_Data};
+    if constexpr (std::is_same<Type, bool>::value) {
+      Ret.ConvertToDataT();
+    }
+    return Ret;
 #else
     vec Ret;
     for (size_t I = 0; I < NumElements; ++I) {
diff --git a/sycl/test-e2e/Basic/vec_bool.cpp b/sycl/test-e2e/Basic/vec_bool.cpp
index 9d265772e128a..3bb33646fe26a 100644
--- a/sycl/test-e2e/Basic/vec_bool.cpp
+++ b/sycl/test-e2e/Basic/vec_bool.cpp
@@ -46,7 +46,7 @@ int main() {
 
   // Test negate (operator ~)
   {
-    init_arr(expected, false);
+    init_arr(expected, true);
 
     sycl::buffer<sycl::vec<bool, size>> bufVecTrue(&vec_true, 1);
     sycl::buffer<sycl::vec<bool, size>> bufResVec(&resVec, 1);