[SYCL] Fix bitselect builtin for integer types

sycl/include/sycl/detail/builtins/relational_functions.inc

@@ -118,8 +118,12 @@ BUILTIN_REL(ONE_ARG, signbit, __spirv_SignBitSet)
 #undef BUILTIN_REL
 
 #ifdef __SYCL_DEVICE_ONLY__
-DEVICE_IMPL_TEMPLATE(THREE_ARGS, bitselect, builtin_enable_generic_t,
-                     __spirv_ocl_bitselect)
+DEVICE_IMPL_TEMPLATE(
+    THREE_ARGS, bitselect, builtin_enable_generic_t, [](auto... xs) {
+      using ret_ty = detail::builtin_enable_generic_t<THREE_ARGS_TEMPLATE_TYPE>;
+      using detail::builtins::convert_result;
+      return convert_result<ret_ty>(__spirv_ocl_bitselect(xs...));
+    })
 #else
 HOST_IMPL_TEMPLATE(THREE_ARGS, bitselect, builtin_enable_generic_t, rel,
                    default_ret_type)

sycl/test-e2e/Basic/built-ins/marray_relational.cpp

@@ -111,6 +111,61 @@ int main() {
   TEST2(sycl::any, int, EXPECTED(bool, false), 3, ma7);
   TEST(sycl::bitselect, float, EXPECTED(float, 1.0, 1.0), 2, ma8, ma9, ma10);
   TEST(sycl::select, float, EXPECTED(float, 1.0, 2.0, 8.0), 3, ma5, ma6, c);
+  {
+    // Extra tests for select/bitselect due to special handling required for
+    // integer return types.
+
+    auto Test = [&](auto F, auto Expected, auto... Args) {
+      std::tuple ArgsTuple{Args...};
+      auto Result = std::apply(F, ArgsTuple);
+      static_assert(std::is_same_v<decltype(Expected), decltype(Result)>);
+
+      // Note: operator==(vec, vec) return vec.
+      auto Equal = [](auto x, auto y) {
+        return std::equal(x.begin(), x.end(), y.begin());
+      };
+
+      assert(Equal(Result, Expected));
+
+      sycl::buffer<bool, 1> ResultBuf{1};
+      deviceQueue.submit([&](sycl::handler &cgh) {
+        sycl::accessor Result{ResultBuf, cgh};
+        cgh.single_task([=]() {
+          auto R = std::apply(F, ArgsTuple);
+          static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
+          Result[0] = Equal(R, Expected);
+        });
+      });
+      assert(sycl::host_accessor{ResultBuf}[0]);
+    };
+
+    sycl::marray<char, 2> a{0b1100, 0b0011};
+    sycl::marray<char, 2> b{0b0011, 0b1100};
+    sycl::marray<char, 2> c{0b1010, 0b1010};
+    sycl::marray<char, 2> r{0b0110, 0b1001};
+
+    auto BitSelect = [](auto... xs) { return sycl::bitselect(xs...); };
+    Test(BitSelect, r, a, b, c);
+    // Input values/results above are positive, so use the same values for
+    // signed/unsigned char tests.
+    [&](auto... xs) {
+      Test(BitSelect, sycl::marray<signed char, 2>{xs}...);
+    }(r, a, b, c);
+    [&](auto... xs) {
+      Test(BitSelect, sycl::marray<unsigned char, 2>{xs}...);
+    }(r, a, b, c);
+
+    auto Select = [](auto... xs) { return sycl::select(xs...); };
+    sycl::marray<bool, 2> c2{false, true};
+    sycl::marray<char, 2> r2{a[0], b[1]};
+    Test(Select, r2, a, b, c2);
+    [&](auto... xs) {
+      Test(Select, sycl::marray<signed char, 2>{xs}..., c2);
+    }(r2, a, b);
+    [&](auto... xs) {
+      Test(Select, sycl::marray<unsigned char, 2>{xs}..., c2);
+    }(r2, a, b);
+  }
 
   return 0;
 }

sycl/test-e2e/Basic/built-ins/scalar_relational.cpp

@@ -0,0 +1,59 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// RUN: %if preview-breaking-changes-supported %{ %{build} -fpreview-breaking-changes -o %t_preview.out %}
+// RUN: %if preview-breaking-changes-supported %{ %{run} %t_preview.out%}
+
+#include <sycl/sycl.hpp>
+
+template <typename... Ts, typename FuncTy> void TestTypes(FuncTy F) {
+  (F(Ts{}), ...);
+}
+
+int main() {
+  sycl::queue q;
+
+  auto Test = [&](auto F, auto Expected, auto... Args) {
+#if defined(__GNUC__) || defined(__clang__)
+    std::cout << __PRETTY_FUNCTION__ << std::endl;
+#endif
+    std::tuple ArgsTuple{Args...};
+    auto Result = std::apply(F, ArgsTuple);
+    static_assert(std::is_same_v<decltype(Expected), decltype(Result)>);
+    assert(Expected == Result);
+
+    sycl::buffer<bool, 1> ResultBuf{1};
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor Result{ResultBuf, cgh};
+      cgh.single_task([=]() {
+        auto R = std::apply(F, ArgsTuple);
+        static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
+        Result[0] = Expected == R;
+      });
+    });
+    assert(sycl::host_accessor{ResultBuf}[0]);
+  };
+
+  auto TestBitSelect = [&](auto type_val) {
+    using T = decltype(type_val);
+    auto BitSelect = [](auto... xs) { return sycl::bitselect(xs...); };
+
+    static_assert(std::is_integral_v<T>,
+                  "Only integer test is implemented here!");
+    Test(BitSelect, T{0b0110}, T{0b1100}, T{0b0011}, T{0b1010});
+  };
+
+  TestTypes<signed char, unsigned char, char, long, long long, unsigned long,
+            unsigned long long>(TestBitSelect);
+
+  auto TestSelect = [&](auto type_val) {
+    using T = decltype(type_val);
+    auto Select = [](auto... xs) { return sycl::select(xs...); };
+
+    Test(Select, T{0}, T{1}, T{0}, true);
+    Test(Select, T{1}, T{1}, T{0}, false);
+  };
+
+  TestTypes<signed char, unsigned char, char>(TestSelect);
+
+  return 0;
+}

sycl/test-e2e/Basic/built-ins/vec_relational.cpp

@@ -88,6 +88,70 @@ int main() {
   TEST2(sycl::any, int, EXPECTED(int32_t, 0), 3, va7);
   TEST(sycl::bitselect, float, EXPECTED(float, 1.0, 1.0), 2, va8, va9, va10);
   TEST(sycl::select, float, EXPECTED(float, 1.0, 2.0, 8.0), 3, va5, va6, c1);
+  {
+    // Extra tests for select/bitselect due to special handling required for
+    // integer return types.
+
+    auto Test = [&](auto F, auto Expected, auto... Args) {
+      std::tuple ArgsTuple{Args...};
+      auto Result = std::apply(F, ArgsTuple);
+      static_assert(std::is_same_v<decltype(Expected), decltype(Result)>);
+
+      // Note: operator==(vec, vec) return vec.
+      auto Equal = [](auto x, auto y) {
+        for (size_t i = 0; i < x.size(); ++i)
+          if (x[i] != y[i])
+            return false;
+
+        return true;
+      };
+
+      assert(Equal(Result, Expected));
+
+      sycl::buffer<bool, 1> ResultBuf{1};
+      deviceQueue.submit([&](sycl::handler &cgh) {
+        sycl::accessor Result{ResultBuf, cgh};
+        cgh.single_task([=]() {
+          auto R = std::apply(F, ArgsTuple);
+          static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
+          Result[0] = Equal(R, Expected);
+        });
+      });
+      assert(sycl::host_accessor{ResultBuf}[0]);
+    };
+
+    // Note that only int8_t/uint8_t are supported by the bitselect/select
+    // builtins and not all three char data types. Also, use positive numbers
+    // for the values below so that we could use the same for both
+    // signed/unsigned tests.
+    sycl::vec<uint8_t, 2> a{0b1100, 0b0011};
+    sycl::vec<uint8_t, 2> b{0b0011, 0b1100};
+    sycl::vec<uint8_t, 2> c{0b1010, 0b1010};
+    sycl::vec<uint8_t, 2> r{0b0110, 0b1001};
+
+    auto BitSelect = [](auto... xs) { return sycl::bitselect(xs...); };
+    Test(BitSelect, r, a, b, c);
+    [&](auto... xs) {
+      Test(BitSelect, xs.template as<sycl::vec<int8_t, 2>>()...);
+    }(r, a, b, c);
+
+    auto Select = [](auto... xs) { return sycl::select(xs...); };
+    sycl::vec<uint8_t, 2> c2{0x7F, 0xFF};
+    sycl::vec<uint8_t, 2> r2{a[0], b[1]};
+
+    Test(Select, r2, a, b, c2);
+    [&](auto... xs) {
+      Test(Select, xs.template as<sycl::vec<int8_t, 2>>()..., c2);
+    }(r2, a, b);
+
+    // Assume that MSB of a signed data type is the leftmost bit (signbit).
+    auto c3 = c2.template as<sycl::vec<int8_t, 2>>();
+
+    Test(Select, r2, a, b, c3);
+    [&](auto... xs) {
+      Test(Select, xs.template as<sycl::vec<int8_t, 2>>()..., c3);
+    }(r2, a, b);
+  }
 
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
   TEST(sycl::isequal, int32_t, EXPECTED(int32_t, 1, 1), 2, va11.swizzle<0, 1>(),