[SYCL] Add support of sycl::marray to relational functions (#8427)

This patch adds support of relational functions which take mgentype as
an argument as it is stated in Table 182 of SYCL 2020 specification.

E2E tests: intel/llvm-test-suite#1617

Author: dm-vodopyanov

sycl/include/sycl/builtins.hpp

@@ -1377,6 +1377,45 @@ detail::common_rel_ret_t<T> signbit(T x) __NOEXC {
       __sycl_std::__invoke_SignBitSet<detail::internal_rel_ret_t<T>>(x));
 }
 
+// marray relational functions
+
+#define __SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(NAME)                 \
+  template <typename T,                                                        \
+            typename = std::enable_if_t<detail::is_mgenfloat<T>::value>>       \
+  sycl::marray<bool, T::size()> NAME(T x, T y) __NOEXC {                       \
+    sycl::marray<bool, T::size()> res;                                         \
+    for (int i = 0; i < x.size(); i++) {                                       \
+      res[i] = NAME(x[i], y[i]);                                               \
+    }                                                                          \
+    return res;                                                                \
+  }
+
+#define __SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(NAME)                  \
+  template <typename T,                                                        \
+            typename = std::enable_if_t<detail::is_mgenfloat<T>::value>>       \
+  sycl::marray<bool, T::size()> NAME(T x) __NOEXC {                            \
+    sycl::marray<bool, T::size()> res;                                         \
+    for (int i = 0; i < x.size(); i++) {                                       \
+      res[i] = NAME(x[i]);                                                     \
+    }                                                                          \
+    return res;                                                                \
+  }
+
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isequal)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isnotequal)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isgreater)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isgreaterequal)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isless)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(islessequal)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(islessgreater)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(isfinite)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(isinf)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(isnan)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(isnormal)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isordered)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isunordered)
+__SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(signbit)
+
 namespace detail {
 #if defined(SYCL2020_CONFORMANT_APIS) && SYCL_LANGUAGE_VERSION >= 202001
 using anyall_ret_t = bool;
@@ -1431,6 +1470,18 @@ detail::enable_if_t<detail::is_sgentype<T>::value, T> select(T a, T b,
   return __sycl_std::__invoke_select<T>(a, b, static_cast<int>(c));
 }
 
+// mgentype select (mgentype a, mgentype b, marray<bool, { N }> c)
+template <typename T,
+          typename = std::enable_if_t<detail::is_mgenfloat<T>::value>>
+sycl::marray<detail::marray_element_type<T>, T::size()>
+select(T a, T b, sycl::marray<bool, T::size()> c) __NOEXC {
+  sycl::marray<detail::marray_element_type<T>, T::size()> res;
+  for (int i = 0; i < a.size(); i++) {
+    res[i] = select(a[i], b[i], c[i]);
+  }
+  return res;
+}
+
 // geninteger select (geninteger a, geninteger b, igeninteger c)
 template <typename T, typename T2>
 detail::enable_if_t<

sycl/include/sycl/detail/generic_type_traits.hpp

@@ -57,6 +57,13 @@ using is_vgenfloat = is_contained<T, gtl::vector_floating_list>;
 template <typename T>
 using is_svgenfloat = is_contained<T, gtl::scalar_vector_floating_list>;
 
+template <typename T> using marray_element_type = typename T::value_type;
+
+template <typename T>
+using is_mgenfloat = bool_constant<
+    std::is_same<T, sycl::marray<marray_element_type<T>, T::size()>>::value &&
+    is_svgenfloat<marray_element_type<T>>::value>;
+
 template <typename T>
 using is_gengeofloat = is_contained<T, gtl::geo_float_list>;