Update implementation of built-in functions

sycl/include/CL/sycl/detail/boolean.hpp

@@ -0,0 +1,144 @@
+//==----------- boolean.hpp - SYCL boolean type ----------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/detail/generic_type_traits.hpp>
+#include <CL/sycl/types.hpp>
+
+#include <initializer_list>
+#include <type_traits>
+
+namespace cl {
+namespace sycl {
+namespace detail {
+
+template <int Num> struct Assigner {
+  template <typename R, typename T> static void assign(R &r, const T x) {
+    Assigner<Num - 1>::assign(r, x);
+    r.template swizzle<Num>() = x.value[Num];
+  }
+
+  template <typename R, typename T, typename ET>
+  static void init(R &r, const T x) {
+    Assigner<Num - 1>::template init<R, T, ET>(r, x);
+    ET v = x.template swizzle<Num>();
+    r.value[Num] = msbIsSet(v);
+  }
+};
+
+template <> struct Assigner<0> {
+  template <typename R, typename T> static void assign(R &r, const T x) {
+    r.template swizzle<0>() = x.value[0];
+  }
+  template <typename R, typename T, typename ET>
+  static void init(R &r, const T x) {
+    ET v = x.template swizzle<0>();
+    r.value[0] = msbIsSet(v);
+  }
+};
+
+template <int N> struct alignas(N == 3 ? 4 : N) Boolean {
+  static_assert(((N == 2) || (N == 3) || (N == 4) || (N == 8) || (N == 16)),
+                "Invalid size");
+
+  using element_type = bool;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  using DataType =
+      element_type __attribute__((ext_vector_type(N == 3 ? 4 : N)));
+  using vector_t = DataType;
+#else
+  using DataType = element_type[N == 3 ? 4 : N];
+#endif
+
+  Boolean() : value{false} {}
+
+  Boolean(std::initializer_list<element_type> l) {
+    for (size_t I = 0; I < N; ++I) {
+      value[I] = *(l.begin() + I);
+    }
+  }
+
+  Boolean(const Boolean &rhs) {
+    for (size_t I = 0; I < N; ++I) {
+      value[I] = rhs.value[I];
+    }
+  }
+
+#ifdef __SYCL_DEVICE_ONLY__
+  // TODO change this to the vectors assignment when the assignment will be
+  // fixed on Intel GPU NEO OpenCL runtime
+  Boolean(const vector_t rhs) {
+    for (size_t I = 0; I < N; ++I) {
+      value[I] = rhs[I];
+    }
+  }
+#endif
+
+  template <typename T> Boolean(const T rhs) {
+    static_assert(is_vgeninteger<T>::value, "Invalid constructor");
+    Assigner<N - 1>::template init<Boolean<N>, T, typename T::element_type>(
+        *this, rhs);
+  }
+
+#ifdef __SYCL_DEVICE_ONLY__
+  operator vector_t() const { return value; }
+#endif
+
+  template <typename T> operator T() const {
+    static_assert(is_vgeninteger<T>::value, "Invalid conversion");
+    T r;
+    Assigner<N - 1>::assign(r, *this);
+    return r * -1;
+  }
+
+private:
+  template <int Num> friend struct Assigner;
+  DataType value;
+};
+
+template <> struct alignas(1) Boolean<1> {
+
+  using element_type = bool;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  using DataType = element_type;
+  using vector_t = DataType;
+#else
+  using DataType = element_type;
+#endif
+
+  Boolean() : value(false) {}
+
+  Boolean(const Boolean &rhs) : value(rhs.value) {}
+
+#ifdef __SYCL_DEVICE_ONLY__
+  Boolean(const vector_t rhs) : value(rhs) {}
+#endif
+
+  template <typename T> Boolean(T val) : value(val) {
+    static_assert(is_sgeninteger<T>::value, "Invalid constructor");
+  }
+
+#ifdef __SYCL_DEVICE_ONLY__
+  operator vector_t() const { return value; }
+#endif
+
+  template <typename T> operator T() const {
+    static_assert(is_sgeninteger<T>::value, "Invalid conversion");
+    return value;
+  }
+
+private:
+  DataType value;
+};
+
+} // namespace detail
+} // namespace sycl
+} // namespace cl

sycl/include/CL/sycl/detail/builtins.hpp

@@ -9,6 +9,7 @@
 #pragma once
 
 #include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/generic_type_traits.hpp>
 
 #include <type_traits>
 // TODO Delete this include after solving the problems in the test
@@ -18,84 +19,10 @@
 // TODO Decide whether to mark functions with this attribute.
 #define __NOEXC /*noexcept*/
 
-namespace cl {
-namespace sycl {
-namespace detail {
-
-// Try to get pointer_t, otherwise T
-template <typename T> class TryToGetPointerT {
-  static T check(...);
-  template <typename A> static typename A::pointer_t check(const A &);
-
-public:
-  using type = decltype(check(T()));
-  static constexpr bool value = !std::is_same<T, type>::value;
-};
-
-// Try to get element_type, otherwise T
-template <typename T> class TryToGetElementType {
-  static T check(...);
-  template <typename A> static typename A::element_type check(const A &);
-
-public:
-  using type = decltype(check(T()));
-  static constexpr bool value = !std::is_same<T, type>::value;
-};
-
-// Try to get vector_t, otherwise T
-template <typename T> class TryToGetVectorT {
-  static T check(...);
-  template <typename A> static typename A::vector_t check(const A &);
-
-public:
-  using type = decltype(check(T()));
-  static constexpr bool value = !std::is_same<T, type>::value;
-};
-
-// Try to get pointer_t (if pointer_t indicates on the type with vector_t
-// creates a pointer type on vector_t), otherwise T
-template <typename T> class TryToGetPointerVecT {
-  static T check(...);
-  template <typename A>
-  static typename PtrValueType<
-      typename TryToGetVectorT<typename TryToGetElementType<A>::type>::type,
-      A::address_space>::type *
-  check(const A &);
-
-public:
-  using type = decltype(check(T()));
-};
-
-template <typename T, typename = typename std::enable_if<
-                          TryToGetPointerT<T>::value, std::true_type>::type>
-typename TryToGetPointerVecT<T>::type TryToGetPointer(T &t) {
-  // TODO find the better way to get the pointer to underlying data from vec
-  // class
-  return reinterpret_cast<typename TryToGetPointerVecT<T>::type>(t.get());
-}
-
-template <typename T, typename = typename std::enable_if<
-                          !TryToGetPointerT<T>::value, std::false_type>::type>
-T TryToGetPointer(T &t) {
-  return t;
-}
-
-// Converts T to OpenCL friendly
-template <typename T>
-using ConvertToOpenCLType = std::conditional<
-    TryToGetVectorT<T>::value, typename TryToGetVectorT<T>::type,
-    typename std::conditional<TryToGetPointerT<T>::value,
-                              typename TryToGetPointerVecT<T>::type, T>::type>;
-
-} // namespace detail
-} // namespace sycl
-} // namespace cl
-
 #define MAKE_CALL_ARG1(call)                                                   \
   template <typename R, typename T1>                                           \
-  ALWAYS_INLINE                                                                \
-      typename cl::sycl::detail::ConvertToOpenCLType<R>::type __invoke_##call( \
-          T1 t1) __NOEXC {                                                     \
+  ALWAYS_INLINE typename cl::sycl::detail::ConvertToOpenCLType<R>::type        \
+  __invoke_##call(T1 t1) __NOEXC {                                             \
     using Ret = typename cl::sycl::detail::ConvertToOpenCLType<R>::type;       \
     using Arg1 = typename cl::sycl::detail::ConvertToOpenCLType<T1>::type;     \
     extern Ret call(Arg1);                                                     \
@@ -104,9 +31,8 @@ using ConvertToOpenCLType = std::conditional<
 
 #define MAKE_CALL_ARG2(call)                                                   \
   template <typename R, typename T1, typename T2>                              \
-  ALWAYS_INLINE                                                                \
-      typename cl::sycl::detail::ConvertToOpenCLType<R>::type __invoke_##call( \
-          T1 t1, T2 t2) __NOEXC {                                              \
+  ALWAYS_INLINE typename cl::sycl::detail::ConvertToOpenCLType<R>::type        \
+  __invoke_##call(T1 t1, T2 t2) __NOEXC {                                      \
     using Ret = typename cl::sycl::detail::ConvertToOpenCLType<R>::type;       \
     using Arg1 = typename cl::sycl::detail::ConvertToOpenCLType<T1>::type;     \
     using Arg2 = typename cl::sycl::detail::ConvertToOpenCLType<T2>::type;     \
@@ -117,9 +43,8 @@ using ConvertToOpenCLType = std::conditional<
 
 #define MAKE_CALL_ARG3(call)                                                   \
   template <typename R, typename T1, typename T2, typename T3>                 \
-  ALWAYS_INLINE                                                                \
-      typename cl::sycl::detail::ConvertToOpenCLType<R>::type __invoke_##call( \
-          T1 t1, T2 t2, T3 t3) __NOEXC {                                       \
+  ALWAYS_INLINE typename cl::sycl::detail::ConvertToOpenCLType<R>::type        \
+  __invoke_##call(T1 t1, T2 t2, T3 t3) __NOEXC {                               \
     using Ret = typename cl::sycl::detail::ConvertToOpenCLType<R>::type;       \
     using Arg1 = typename cl::sycl::detail::ConvertToOpenCLType<T1>::type;     \
     using Arg2 = typename cl::sycl::detail::ConvertToOpenCLType<T2>::type;     \
@@ -264,8 +189,8 @@ MAKE_CALL_ARG2(s_upsample)
 MAKE_CALL_ARG1(popcount)
 MAKE_CALL_ARG3(s_mad24)
 MAKE_CALL_ARG3(u_mad24)
-MAKE_CALL_ARG3(s_mul24)
-MAKE_CALL_ARG3(u_mul24)
+MAKE_CALL_ARG2(s_mul24)
+MAKE_CALL_ARG2(u_mul24)
 /* --------------- 4.13.5 Common functions. ---------------------------------*/
 MAKE_CALL_ARG3(fclamp)
 MAKE_CALL_ARG1(degrees)