From e00afe4b6e3fa33d5fba1b90974f82c352c4cd04 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy@intel.com>
Date: Thu, 16 Dec 2021 15:39:38 +0300
Subject: [PATCH 1/5] Add sycl::known_identity for sycl::vec<type, N>

---
 sycl/include/CL/sycl/known_identity.hpp  | 130 ++++++++++++++++---
 sycl/test/basic_tests/known_identity.cpp | 152 ++++++++++++++++++++++-
 2 files changed, 263 insertions(+), 19 deletions(-)

diff --git a/sycl/include/CL/sycl/known_identity.hpp b/sycl/include/CL/sycl/known_identity.hpp
index 955d7b8a9cfca..55511bf6e44f2 100644
--- a/sycl/include/CL/sycl/known_identity.hpp
+++ b/sycl/include/CL/sycl/known_identity.hpp
@@ -65,33 +65,34 @@ using IsLogicalOR =
 
 // Identity = 0
 template <typename T, class BinaryOperation>
-using IsZeroIdentityOp = bool_constant<
-    (is_sgeninteger<T>::value &&
-     (IsPlus<T, BinaryOperation>::value || IsBitOR<T, BinaryOperation>::value ||
-      IsBitXOR<T, BinaryOperation>::value)) ||
-    (is_sgenfloat<T>::value && IsPlus<T, BinaryOperation>::value)>;
+using IsZeroIdentityOp =
+    bool_constant<(is_geninteger<T>::value &&
+                   (IsPlus<T, BinaryOperation>::value ||
+                    IsBitOR<T, BinaryOperation>::value ||
+                    IsBitXOR<T, BinaryOperation>::value)) ||
+                  (is_genfloat<T>::value && IsPlus<T, BinaryOperation>::value)>;
 
 // Identity = 1
 template <typename T, class BinaryOperation>
 using IsOneIdentityOp =
-    bool_constant<(is_sgeninteger<T>::value || is_sgenfloat<T>::value) &&
+    bool_constant<(is_geninteger<T>::value || is_genfloat<T>::value) &&
                   IsMultiplies<T, BinaryOperation>::value>;
 
 // Identity = ~0
 template <typename T, class BinaryOperation>
-using IsOnesIdentityOp = bool_constant<is_sgeninteger<T>::value &&
+using IsOnesIdentityOp = bool_constant<is_geninteger<T>::value &&
                                        IsBitAND<T, BinaryOperation>::value>;
 
 // Identity = <max possible value>
 template <typename T, class BinaryOperation>
 using IsMinimumIdentityOp =
-    bool_constant<(is_sgeninteger<T>::value || is_sgenfloat<T>::value) &&
+    bool_constant<(is_geninteger<T>::value || is_genfloat<T>::value) &&
                   IsMinimum<T, BinaryOperation>::value>;
 
 // Identity = <min possible value>
 template <typename T, class BinaryOperation>
 using IsMaximumIdentityOp =
-    bool_constant<(is_sgeninteger<T>::value || is_sgenfloat<T>::value) &&
+    bool_constant<(is_geninteger<T>::value || is_genfloat<T>::value) &&
                   IsMaximum<T, BinaryOperation>::value>;
 
 // Identity = false
@@ -125,7 +126,27 @@ template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<
     BinaryOperation, AccumulatorT,
     std::enable_if_t<IsZeroIdentityOp<AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value = 0;
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(0);
+};
+
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<IsZeroIdentityOp<vec<std::byte, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value =
+      vec<std::byte, NumElements>(std::byte(0));
+};
+#endif
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<sycl::half, NumElements>,
+    std::enable_if_t<IsZeroIdentityOp<vec<sycl::half, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr vec<sycl::half, NumElements> value =
+      vec<sycl::half, NumElements>(sycl::half());
 };
 
 template <typename BinaryOperation>
@@ -145,8 +166,19 @@ template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<
     BinaryOperation, AccumulatorT,
     std::enable_if_t<IsOneIdentityOp<AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value = 1;
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(1);
+};
+
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<
+        IsOneIdentityOp<vec<std::byte, NumElements>, BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value =
+      vec<std::byte, NumElements>(std::byte(1));
 };
+#endif
 
 template <typename BinaryOperation>
 struct known_identity_impl<
@@ -165,48 +197,110 @@ template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<
     BinaryOperation, AccumulatorT,
     std::enable_if_t<IsOnesIdentityOp<AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value = ~static_cast<AccumulatorT>(0);
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(-1LL);
+};
+
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<IsOnesIdentityOp<vec<std::byte, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value =
+      vec<std::byte, NumElements>(std::byte(-1LL));
 };
+#endif
 
 /// Returns maximal possible value as identity for MIN operations.
 template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<BinaryOperation, AccumulatorT,
                            std::enable_if_t<IsMinimumIdentityOp<
                                AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value =
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(
       std::numeric_limits<AccumulatorT>::has_infinity
           ? std::numeric_limits<AccumulatorT>::infinity()
-          : (std::numeric_limits<AccumulatorT>::max)();
+          : (std::numeric_limits<AccumulatorT>::max)());
 };
 
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<IsMinimumIdentityOp<vec<std::byte, NumElements>,
+                                         BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value =
+      static_cast<vec<std::byte, NumElements>>(
+          std::numeric_limits<vec<std::byte, NumElements>>::has_infinity
+              ? std::numeric_limits<vec<std::byte, NumElements>>::infinity()
+              : (std::numeric_limits<vec<std::byte, NumElements>>::max)());
+};
+#endif
+
 /// Returns minimal possible value as identity for MAX operations.
 template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<BinaryOperation, AccumulatorT,
                            std::enable_if_t<IsMaximumIdentityOp<
                                AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value =
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(
       std::numeric_limits<AccumulatorT>::has_infinity
           ? static_cast<AccumulatorT>(
                 -std::numeric_limits<AccumulatorT>::infinity())
-          : std::numeric_limits<AccumulatorT>::lowest();
+          : std::numeric_limits<AccumulatorT>::lowest());
+};
+
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<IsMaximumIdentityOp<vec<std::byte, NumElements>,
+                                         BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value = static_cast<
+      vec<std::byte, NumElements>>(
+      std::numeric_limits<vec<std::byte, NumElements>>::has_infinity
+          ? static_cast<vec<std::byte, NumElements>>(
+                -std::numeric_limits<vec<std::byte, NumElements>>::infinity())
+          : std::numeric_limits<vec<std::byte, NumElements>>::lowest());
 };
+#endif
 
 /// Returns false as identity for LOGICAL OR operations.
 template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<
     BinaryOperation, AccumulatorT,
     std::enable_if_t<IsFalseIdentityOp<AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value = false;
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(false);
 };
 
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<IsFalseIdentityOp<vec<std::byte, NumElements>,
+                                       BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value =
+      vec<std::byte, NumElements>(std::byte(false));
+};
+#endif
+
 /// Returns true as identity for LOGICAL AND operations.
 template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<
     BinaryOperation, AccumulatorT,
     std::enable_if_t<IsTrueIdentityOp<AccumulatorT, BinaryOperation>::value>> {
-  static constexpr AccumulatorT value = true;
+  static constexpr AccumulatorT value = static_cast<AccumulatorT>(true);
 };
 
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, vec<std::byte, NumElements>,
+    std::enable_if_t<IsTrueIdentityOp<vec<std::byte, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr vec<std::byte, NumElements> value =
+      vec<std::byte, NumElements>(std::byte(true));
+};
+#endif
+
 } // namespace detail
 
 // ---- has_known_identity
diff --git a/sycl/test/basic_tests/known_identity.cpp b/sycl/test/basic_tests/known_identity.cpp
index 06908d1ca247d..a758e7dd7f583 100644
--- a/sycl/test/basic_tests/known_identity.cpp
+++ b/sycl/test/basic_tests/known_identity.cpp
@@ -1,4 +1,5 @@
-// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
+// RUN: %clangxx -fsycl -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning -o %t.out -std=c++17
+// RUN: %RUN_ON_HOST %t.out
 // expected-no-diagnostics
 
 // This test performs basic checks of has_known_identity and known_identity
@@ -6,6 +7,7 @@
 
 #include <CL/sycl.hpp>
 #include <cassert>
+#include <cstddef>
 
 using namespace cl::sycl;
 
@@ -93,6 +95,152 @@ template <typename T> void checkBoolKnownIdentity() {
   static_assert(known_identity<sycl::logical_or<T>, T>::value == false);
 }
 
+template <typename T, int Num>
+bool compareVectors(const vec<T, Num> a, const vec<T, Num> b) {
+  bool res = true;
+  for (int i = 0; i < Num; ++i) {
+    res &= (a[i] == b[i]);
+  }
+  if (!res) {
+    for (int i = 0; i < Num; ++i) {
+      std::cout << "(" << (int)a[i] << " == " << (int)b[i] << ")" << std::endl;
+    }
+  }
+  return res;
+}
+
+template <typename T, int Num>
+typename std::enable_if<!std::is_same<T, half>::value &&
+                            !std::is_same<T, float>::value &&
+                            !std::is_same<T, double>::value,
+                        void>::type
+checkVecKnownIdentity() {
+  constexpr vec<T, Num> zeros(T(0));
+  constexpr vec<T, Num> ones(T(1));
+  constexpr vec<T, Num> bit_ones(~T(0));
+
+  static_assert(has_known_identity<plus<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<plus<vec<T, Num>>, vec<T, Num>>::value);
+  assert(compareVectors(known_identity<plus<>, vec<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<bit_or<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<bit_or<vec<T, Num>>, vec<T, Num>>::value);
+  assert(compareVectors(known_identity<bit_or<>, vec<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<bit_xor<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<bit_xor<vec<T, Num>>, vec<T, Num>>::value);
+  assert(compareVectors(known_identity<bit_xor<>, vec<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<bit_and<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<bit_and<vec<T, Num>>, vec<T, Num>>::value);
+  assert(
+      compareVectors(known_identity<bit_and<>, vec<T, Num>>::value, bit_ones));
+
+  static_assert(has_known_identity<logical_or<>, vec<T, Num>>::value);
+  static_assert(
+      has_known_identity<logical_or<vec<T, Num>>, vec<T, Num>>::value);
+  assert(
+      compareVectors(known_identity<logical_or<>, vec<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<logical_and<>, vec<T, Num>>::value);
+  static_assert(
+      has_known_identity<logical_and<vec<T, Num>>, vec<T, Num>>::value);
+  assert(
+      compareVectors(known_identity<logical_and<>, vec<T, Num>>::value, ones));
+
+  static_assert(has_known_identity<multiplies<>, vec<T, Num>>::value);
+  static_assert(
+      has_known_identity<multiplies<vec<T, Num>>, vec<T, Num>>::value);
+  assert(
+      compareVectors(known_identity<multiplies<>, vec<T, Num>>::value, ones));
+
+  static_assert(has_known_identity<minimum<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<minimum<vec<T, Num>>, vec<T, Num>>::value);
+  if constexpr (!std::is_same<T, std::byte>::value) {
+    constexpr vec<T, Num> maxs(-std::numeric_limits<T>::infinity());
+    assert(compareVectors(known_identity<minimum<>, vec<T, Num>>::value, maxs));
+  }
+
+  static_assert(has_known_identity<maximum<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<maximum<vec<T, Num>>, vec<T, Num>>::value);
+  if constexpr (!std::is_same<T, std::byte>::value) {
+    constexpr vec<T, Num> mins(std::numeric_limits<T>::infinity());
+    assert(compareVectors(known_identity<maximum<>, vec<T, Num>>::value, mins));
+  }
+}
+
+template <typename T, int Num>
+typename std::enable_if<std::is_same<T, sycl::half>::value ||
+                            std::is_same<T, float>::value ||
+                            std::is_same<T, double>::value,
+                        void>::type
+checkVecKnownIdentity() {
+  constexpr vec<T, Num> zeros(T(0.0f));
+  constexpr vec<T, Num> ones(T(1.0f));
+
+  static_assert(has_known_identity<plus<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<plus<vec<T, Num>>, vec<T, Num>>::value);
+  assert(compareVectors(known_identity<plus<>, vec<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<multiplies<>, vec<T, Num>>::value);
+  static_assert(
+      has_known_identity<multiplies<vec<T, Num>>, vec<T, Num>>::value);
+  assert(
+      compareVectors(known_identity<multiplies<>, vec<T, Num>>::value, ones));
+
+  static_assert(has_known_identity<minimum<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<minimum<vec<T, Num>>, vec<T, Num>>::value);
+
+  static_assert(has_known_identity<maximum<>, vec<T, Num>>::value);
+  static_assert(has_known_identity<maximum<vec<T, Num>>, vec<T, Num>>::value);
+}
+
+void checkVecTypesKnownIdentity() {
+
+#define CHECK_VEC(type)                                                        \
+  do {                                                                         \
+    checkVecKnownIdentity<type, 1>();                                          \
+    checkVecKnownIdentity<type, 2>();                                          \
+    checkVecKnownIdentity<type, 3>();                                          \
+    checkVecKnownIdentity<type, 4>();                                          \
+    checkVecKnownIdentity<type, 8>();                                          \
+    checkVecKnownIdentity<type, 16>();                                         \
+  } while (0)
+
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+  CHECK_VEC(std::byte);
+#endif
+  CHECK_VEC(int8_t);
+  CHECK_VEC(int16_t);
+  CHECK_VEC(int32_t);
+  CHECK_VEC(int64_t);
+  CHECK_VEC(uint8_t);
+  CHECK_VEC(uint16_t);
+  CHECK_VEC(uint32_t);
+  CHECK_VEC(uint64_t);
+
+  CHECK_VEC(char);
+  CHECK_VEC(short int);
+  CHECK_VEC(int);
+  CHECK_VEC(long);
+  CHECK_VEC(long long);
+  CHECK_VEC(unsigned char);
+  CHECK_VEC(unsigned short int);
+  CHECK_VEC(unsigned int);
+  CHECK_VEC(unsigned long);
+  CHECK_VEC(unsigned long long);
+  CHECK_VEC(float);
+  CHECK_VEC(double);
+
+  checkVecKnownIdentity<half, 2>();
+  checkVecKnownIdentity<half, 3>();
+  checkVecKnownIdentity<half, 4>();
+  checkVecKnownIdentity<half, 8>();
+  checkVecKnownIdentity<half, 16>();
+
+#undef CHECK_VEC
+}
+
 int main() {
   checkIntKnownIdentity<int8_t>();
   checkIntKnownIdentity<char>();
@@ -143,6 +291,8 @@ int main() {
 
   checkBoolKnownIdentity<bool>();
 
+  checkVecTypesKnownIdentity();
+
   // Few negative tests just to check that it does not always return true.
   static_assert(!has_known_identity<std::minus<>, int>::value);
   static_assert(!has_known_identity<sycl::bit_or<>, float>::value);

From 78e77a69d5b1392dc805531d506f2742e407e3db Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy@intel.com>
Date: Thu, 16 Dec 2021 15:53:18 +0300
Subject: [PATCH 2/5] Add sycl::known_identity for sycl::marray<type, N>

---
 .../CL/sycl/detail/generic_type_lists.hpp     | 210 +++++++++++++++---
 sycl/include/CL/sycl/known_identity.hpp       |  71 ++++++
 sycl/test/basic_tests/known_identity.cpp      | 155 +++++++++++++
 3 files changed, 408 insertions(+), 28 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
index a5866d0a391e5..7182e9f79f22e 100644
--- a/sycl/include/CL/sycl/detail/generic_type_lists.hpp
+++ b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
@@ -21,6 +21,7 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 template <typename T, int N> class vec;
+template <typename Type, std::size_t NumElements> class marray;
 namespace detail {
 namespace half_impl {
 class half;
@@ -40,7 +41,12 @@ using scalar_half_list = type_list<half>;
 using vector_half_list = type_list<vec<half, 1>, vec<half, 2>, vec<half, 3>,
                                    vec<half, 4>, vec<half, 8>, vec<half, 16>>;
 
-using half_list = type_list<scalar_half_list, vector_half_list>;
+using marray_half_list =
+    type_list<marray<half, 1>, marray<half, 2>, marray<half, 3>,
+              marray<half, 4>, marray<half, 8>, marray<half, 16>>;
+
+using half_list =
+    type_list<scalar_half_list, vector_half_list, marray_half_list>;
 
 using scalar_float_list = type_list<float>;
 
@@ -48,7 +54,12 @@ using vector_float_list =
     type_list<vec<float, 1>, vec<float, 2>, vec<float, 3>, vec<float, 4>,
               vec<float, 8>, vec<float, 16>>;
 
-using float_list = type_list<scalar_float_list, vector_float_list>;
+using marray_float_list =
+    type_list<marray<float, 1>, marray<float, 2>, marray<float, 3>,
+              marray<float, 4>, marray<float, 8>, marray<float, 16>>;
+
+using float_list =
+    type_list<scalar_float_list, vector_float_list, marray_float_list>;
 
 using scalar_double_list = type_list<double>;
 
@@ -56,7 +67,12 @@ using vector_double_list =
     type_list<vec<double, 1>, vec<double, 2>, vec<double, 3>, vec<double, 4>,
               vec<double, 8>, vec<double, 16>>;
 
-using double_list = type_list<scalar_double_list, vector_double_list>;
+using marray_double_list =
+    type_list<marray<double, 1>, marray<double, 2>, marray<double, 3>,
+              marray<double, 4>, marray<double, 8>, marray<double, 16>>;
+
+using double_list =
+    type_list<scalar_double_list, vector_double_list, marray_double_list>;
 
 using scalar_floating_list =
     type_list<scalar_float_list, scalar_double_list, scalar_half_list>;
@@ -64,7 +80,11 @@ using scalar_floating_list =
 using vector_floating_list =
     type_list<vector_float_list, vector_double_list, vector_half_list>;
 
-using floating_list = type_list<scalar_floating_list, vector_floating_list>;
+using marray_floating_list =
+    type_list<marray_float_list, marray_double_list, marray_half_list>;
+
+using floating_list =
+    type_list<scalar_floating_list, vector_floating_list, marray_floating_list>;
 
 // geometric floating point types
 using scalar_geo_half_list = type_list<half>;
@@ -113,8 +133,13 @@ using vector_default_char_list =
     type_list<vec<char, 1>, vec<char, 2>, vec<char, 3>, vec<char, 4>,
               vec<char, 8>, vec<char, 16>>;
 
+using marray_default_char_list =
+    type_list<marray<char, 1>, marray<char, 2>, marray<char, 3>,
+              marray<char, 4>, marray<char, 8>, marray<char, 16>>;
+
 using default_char_list =
-    type_list<scalar_default_char_list, vector_default_char_list>;
+    type_list<scalar_default_char_list, vector_default_char_list,
+              marray_default_char_list>;
 
 using scalar_signed_char_list = type_list<signed char>;
 
@@ -122,8 +147,14 @@ using vector_signed_char_list =
     type_list<vec<signed char, 1>, vec<signed char, 2>, vec<signed char, 3>,
               vec<signed char, 4>, vec<signed char, 8>, vec<signed char, 16>>;
 
+using marray_signed_char_list =
+    type_list<marray<signed char, 1>, marray<signed char, 2>,
+              marray<signed char, 3>, marray<signed char, 4>,
+              marray<signed char, 8>, marray<signed char, 16>>;
+
 using signed_char_list =
-    type_list<scalar_signed_char_list, vector_signed_char_list>;
+    type_list<scalar_signed_char_list, vector_signed_char_list,
+              marray_signed_char_list>;
 
 using scalar_unsigned_char_list = type_list<unsigned char>;
 
@@ -132,8 +163,14 @@ using vector_unsigned_char_list =
               vec<unsigned char, 3>, vec<unsigned char, 4>,
               vec<unsigned char, 8>, vec<unsigned char, 16>>;
 
+using marray_unsigned_char_list =
+    type_list<marray<unsigned char, 1>, marray<unsigned char, 2>,
+              marray<unsigned char, 3>, marray<unsigned char, 4>,
+              marray<unsigned char, 8>, marray<unsigned char, 16>>;
+
 using unsigned_char_list =
-    type_list<scalar_unsigned_char_list, vector_unsigned_char_list>;
+    type_list<scalar_unsigned_char_list, vector_unsigned_char_list,
+              marray_unsigned_char_list>;
 
 using scalar_char_list =
     type_list<scalar_default_char_list, scalar_signed_char_list,
@@ -143,6 +180,10 @@ using vector_char_list =
     type_list<vector_default_char_list, vector_signed_char_list,
               vector_unsigned_char_list>;
 
+using marray_char_list =
+    type_list<marray_default_char_list, marray_signed_char_list,
+              marray_unsigned_char_list>;
+
 using char_list = type_list<scalar_char_list, vector_char_list>;
 
 // short int types
@@ -153,8 +194,14 @@ using vector_signed_short_list =
               vec<signed short, 4>, vec<signed short, 8>,
               vec<signed short, 16>>;
 
+using marray_signed_short_list =
+    type_list<marray<signed short, 1>, marray<signed short, 2>,
+              marray<signed short, 3>, marray<signed short, 4>,
+              marray<signed short, 8>, marray<signed short, 16>>;
+
 using signed_short_list =
-    type_list<scalar_signed_short_list, vector_signed_short_list>;
+    type_list<scalar_signed_short_list, vector_signed_short_list,
+              marray_signed_short_list>;
 
 using scalar_unsigned_short_list = type_list<unsigned short>;
 
@@ -163,14 +210,21 @@ using vector_unsigned_short_list =
               vec<unsigned short, 3>, vec<unsigned short, 4>,
               vec<unsigned short, 8>, vec<unsigned short, 16>>;
 
+using marray_unsigned_short_list =
+    type_list<marray<unsigned short, 1>, marray<unsigned short, 2>,
+              marray<unsigned short, 3>, marray<unsigned short, 4>,
+              marray<unsigned short, 8>, marray<unsigned short, 16>>;
+
 using unsigned_short_list =
-    type_list<scalar_unsigned_short_list, vector_unsigned_short_list>;
+    type_list<scalar_unsigned_short_list, vector_unsigned_short_list,
+              marray_unsigned_short_list>;
 
 using scalar_short_list =
     type_list<scalar_signed_short_list, scalar_unsigned_short_list>;
 
 using vector_short_list =
-    type_list<vector_signed_short_list, vector_unsigned_short_list>;
+    type_list<vector_signed_short_list, vector_unsigned_short_list,
+              marray_unsigned_short_list>;
 
 using short_list = type_list<scalar_short_list, vector_short_list>;
 
@@ -181,8 +235,14 @@ using vector_signed_int_list =
     type_list<vec<signed int, 1>, vec<signed int, 2>, vec<signed int, 3>,
               vec<signed int, 4>, vec<signed int, 8>, vec<signed int, 16>>;
 
+using marray_signed_int_list =
+    type_list<marray<signed int, 1>, marray<signed int, 2>,
+              marray<signed int, 3>, marray<signed int, 4>,
+              marray<signed int, 8>, marray<signed int, 16>>;
+
 using signed_int_list =
-    type_list<scalar_signed_int_list, vector_signed_int_list>;
+    type_list<scalar_signed_int_list, vector_signed_int_list,
+              marray_signed_int_list>;
 
 using scalar_unsigned_int_list = type_list<unsigned int>;
 
@@ -191,8 +251,14 @@ using vector_unsigned_int_list =
               vec<unsigned int, 4>, vec<unsigned int, 8>,
               vec<unsigned int, 16>>;
 
+using marray_unsigned_int_list =
+    type_list<marray<unsigned int, 1>, marray<unsigned int, 2>,
+              marray<unsigned int, 3>, marray<unsigned int, 4>,
+              marray<unsigned int, 8>, marray<unsigned int, 16>>;
+
 using unsigned_int_list =
-    type_list<scalar_unsigned_int_list, vector_unsigned_int_list>;
+    type_list<scalar_unsigned_int_list, vector_unsigned_int_list,
+              marray_unsigned_int_list>;
 
 using scalar_int_list =
     type_list<scalar_signed_int_list, scalar_unsigned_int_list>;
@@ -200,7 +266,10 @@ using scalar_int_list =
 using vector_int_list =
     type_list<vector_signed_int_list, vector_unsigned_int_list>;
 
-using int_list = type_list<scalar_int_list, vector_int_list>;
+using marray_int_list =
+    type_list<marray_signed_int_list, marray_unsigned_int_list>;
+
+using int_list = type_list<scalar_int_list, vector_int_list, marray_int_list>;
 
 // long types
 using scalar_signed_long_list = type_list<signed long>;
@@ -209,8 +278,14 @@ using vector_signed_long_list =
     type_list<vec<signed long, 1>, vec<signed long, 2>, vec<signed long, 3>,
               vec<signed long, 4>, vec<signed long, 8>, vec<signed long, 16>>;
 
+using marray_signed_long_list =
+    type_list<marray<signed long, 1>, marray<signed long, 2>,
+              marray<signed long, 3>, marray<signed long, 4>,
+              marray<signed long, 8>, marray<signed long, 16>>;
+
 using signed_long_list =
-    type_list<scalar_signed_long_list, vector_signed_long_list>;
+    type_list<scalar_signed_long_list, vector_signed_long_list,
+              marray_signed_long_list>;
 
 using scalar_unsigned_long_list = type_list<unsigned long>;
 
@@ -219,8 +294,14 @@ using vector_unsigned_long_list =
               vec<unsigned long, 3>, vec<unsigned long, 4>,
               vec<unsigned long, 8>, vec<unsigned long, 16>>;
 
+using marray_unsigned_long_list =
+    type_list<marray<unsigned long, 1>, marray<unsigned long, 2>,
+              marray<unsigned long, 3>, marray<unsigned long, 4>,
+              marray<unsigned long, 8>, marray<unsigned long, 16>>;
+
 using unsigned_long_list =
-    type_list<scalar_unsigned_long_list, vector_unsigned_long_list>;
+    type_list<scalar_unsigned_long_list, vector_unsigned_long_list,
+              marray_unsigned_long_list>;
 
 using scalar_long_list =
     type_list<scalar_signed_long_list, scalar_unsigned_long_list>;
@@ -228,7 +309,11 @@ using scalar_long_list =
 using vector_long_list =
     type_list<vector_signed_long_list, vector_unsigned_long_list>;
 
-using long_list = type_list<scalar_long_list, vector_long_list>;
+using marray_long_list =
+    type_list<marray_signed_long_list, marray_unsigned_long_list>;
+
+using long_list =
+    type_list<scalar_long_list, vector_long_list, marray_long_list>;
 
 // long long types
 using scalar_signed_longlong_list = type_list<signed long long>;
@@ -238,8 +323,14 @@ using vector_signed_longlong_list =
               vec<signed long long, 3>, vec<signed long long, 4>,
               vec<signed long long, 8>, vec<signed long long, 16>>;
 
+using marray_signed_longlong_list =
+    type_list<marray<signed long long, 1>, marray<signed long long, 2>,
+              marray<signed long long, 3>, marray<signed long long, 4>,
+              marray<signed long long, 8>, marray<signed long long, 16>>;
+
 using signed_longlong_list =
-    type_list<scalar_signed_longlong_list, vector_signed_longlong_list>;
+    type_list<scalar_signed_longlong_list, vector_signed_longlong_list,
+              marray_signed_longlong_list>;
 
 using scalar_unsigned_longlong_list = type_list<unsigned long long>;
 
@@ -248,8 +339,14 @@ using vector_unsigned_longlong_list =
               vec<unsigned long long, 3>, vec<unsigned long long, 4>,
               vec<unsigned long long, 8>, vec<unsigned long long, 16>>;
 
+using marray_unsigned_longlong_list =
+    type_list<marray<unsigned long long, 1>, marray<unsigned long long, 2>,
+              marray<unsigned long long, 3>, marray<unsigned long long, 4>,
+              marray<unsigned long long, 8>, marray<unsigned long long, 16>>;
+
 using unsigned_longlong_list =
-    type_list<scalar_unsigned_longlong_list, vector_unsigned_longlong_list>;
+    type_list<scalar_unsigned_longlong_list, vector_unsigned_longlong_list,
+              marray_unsigned_longlong_list>;
 
 using scalar_longlong_list =
     type_list<scalar_signed_longlong_list, scalar_unsigned_longlong_list>;
@@ -257,7 +354,11 @@ using scalar_longlong_list =
 using vector_longlong_list =
     type_list<vector_signed_longlong_list, vector_unsigned_longlong_list>;
 
-using longlong_list = type_list<scalar_longlong_list, vector_longlong_list>;
+using marray_longlong_list =
+    type_list<marray_signed_longlong_list, marray_unsigned_longlong_list>;
+
+using longlong_list =
+    type_list<scalar_longlong_list, vector_longlong_list, marray_longlong_list>;
 
 // long integer types
 using scalar_signed_long_integer_list =
@@ -266,8 +367,12 @@ using scalar_signed_long_integer_list =
 using vector_signed_long_integer_list =
     type_list<vector_signed_long_list, vector_signed_longlong_list>;
 
+using marray_signed_long_integer_list =
+    type_list<marray_signed_long_list, marray_signed_longlong_list>;
+
 using signed_long_integer_list =
-    type_list<scalar_signed_long_integer_list, vector_signed_long_integer_list>;
+    type_list<scalar_signed_long_integer_list, vector_signed_long_integer_list,
+              marray_signed_long_integer_list>;
 
 using scalar_unsigned_long_integer_list =
     type_list<scalar_unsigned_long_list, scalar_unsigned_longlong_list>;
@@ -275,8 +380,12 @@ using scalar_unsigned_long_integer_list =
 using vector_unsigned_long_integer_list =
     type_list<vector_unsigned_long_list, vector_unsigned_longlong_list>;
 
+using marray_unsigned_long_integer_list =
+    type_list<marray_unsigned_long_list, marray_unsigned_longlong_list>;
+
 using unsigned_long_integer_list = type_list<scalar_unsigned_long_integer_list,
-                                             vector_unsigned_long_integer_list>;
+                                             vector_unsigned_long_integer_list,
+                                             marray_unsigned_long_integer_list>;
 
 using scalar_long_integer_list = type_list<scalar_signed_long_integer_list,
                                            scalar_unsigned_long_integer_list>;
@@ -284,8 +393,12 @@ using scalar_long_integer_list = type_list<scalar_signed_long_integer_list,
 using vector_long_integer_list = type_list<vector_signed_long_integer_list,
                                            vector_unsigned_long_integer_list>;
 
+using marray_long_integer_list = type_list<marray_signed_long_integer_list,
+                                           marray_unsigned_long_integer_list>;
+
 using long_integer_list =
-    type_list<scalar_long_integer_list, vector_long_integer_list>;
+    type_list<scalar_long_integer_list, vector_long_integer_list,
+              marray_long_integer_list>;
 
 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
 // std::byte
@@ -294,6 +407,10 @@ using scalar_byte_list = type_list<std::byte>;
 using vector_byte_list =
     type_list<vec<std::byte, 1>, vec<std::byte, 2>, vec<std::byte, 3>,
               vec<std::byte, 4>, vec<std::byte, 8>, vec<std::byte, 16>>;
+
+using marray_byte_list = type_list<marray<std::byte, 1>, marray<std::byte, 2>,
+                                   marray<std::byte, 3>, marray<std::byte, 4>,
+                                   marray<std::byte, 8>, marray<std::byte, 16>>;
 #endif
 
 // integer types
@@ -311,8 +428,16 @@ using vector_signed_integer_list = type_list<
     vector_signed_short_list, vector_signed_int_list, vector_signed_long_list,
     vector_signed_longlong_list>;
 
+using marray_signed_integer_list = type_list<
+    conditional_t<std::is_signed<char>::value,
+                  type_list<marray_default_char_list, marray_signed_char_list>,
+                  marray_signed_char_list>,
+    marray_signed_short_list, marray_signed_int_list, marray_signed_long_list,
+    marray_signed_longlong_list>;
+
 using signed_integer_list =
-    type_list<scalar_signed_integer_list, vector_signed_integer_list>;
+    type_list<scalar_signed_integer_list, vector_signed_integer_list,
+              marray_signed_integer_list>;
 
 using scalar_unsigned_integer_list =
     type_list<conditional_t<std::is_unsigned<char>::value,
@@ -340,8 +465,22 @@ using vector_unsigned_integer_list =
 #endif
               >;
 
+using marray_unsigned_integer_list =
+    type_list<conditional_t<std::is_unsigned<char>::value,
+                            type_list<marray_default_char_list,
+                                      marray_unsigned_char_list>,
+                            marray_unsigned_char_list>,
+              marray_unsigned_short_list, marray_unsigned_int_list,
+              marray_unsigned_long_list, marray_unsigned_longlong_list
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+              ,
+              marray_byte_list
+#endif
+              >;
+
 using unsigned_integer_list =
-    type_list<scalar_unsigned_integer_list, vector_unsigned_integer_list>;
+    type_list<scalar_unsigned_integer_list, vector_unsigned_integer_list,
+              marray_unsigned_integer_list>;
 
 using scalar_integer_list =
     type_list<scalar_signed_integer_list, scalar_unsigned_integer_list>;
@@ -349,7 +488,11 @@ using scalar_integer_list =
 using vector_integer_list =
     type_list<vector_signed_integer_list, vector_unsigned_integer_list>;
 
-using integer_list = type_list<scalar_integer_list, vector_integer_list>;
+using marray_integer_list =
+    type_list<marray_signed_integer_list, marray_unsigned_integer_list>;
+
+using integer_list =
+    type_list<scalar_integer_list, vector_integer_list, marray_integer_list>;
 
 // basic types
 using scalar_signed_basic_list =
@@ -358,15 +501,22 @@ using scalar_signed_basic_list =
 using vector_signed_basic_list =
     type_list<vector_floating_list, vector_signed_integer_list>;
 
+using marray_signed_basic_list =
+    type_list<marray_floating_list, marray_signed_integer_list>;
+
 using signed_basic_list =
-    type_list<scalar_signed_basic_list, vector_signed_basic_list>;
+    type_list<scalar_signed_basic_list, vector_signed_basic_list,
+              marray_signed_basic_list>;
 
 using scalar_unsigned_basic_list = type_list<scalar_unsigned_integer_list>;
 
 using vector_unsigned_basic_list = type_list<vector_unsigned_integer_list>;
 
+using marray_unsigned_basic_list = type_list<marray_unsigned_integer_list>;
+
 using unsigned_basic_list =
-    type_list<scalar_unsigned_basic_list, vector_unsigned_basic_list>;
+    type_list<scalar_unsigned_basic_list, vector_unsigned_basic_list,
+              marray_unsigned_basic_list>;
 
 using scalar_basic_list =
     type_list<scalar_signed_basic_list, scalar_unsigned_basic_list>;
@@ -374,7 +524,11 @@ using scalar_basic_list =
 using vector_basic_list =
     type_list<vector_signed_basic_list, vector_unsigned_basic_list>;
 
-using basic_list = type_list<scalar_basic_list, vector_basic_list>;
+using marray_basic_list =
+    type_list<marray_signed_basic_list, marray_unsigned_basic_list>;
+
+using basic_list =
+    type_list<scalar_basic_list, vector_basic_list, marray_basic_list>;
 
 // nan builtin types
 using nan_list = type_list<gtl::unsigned_short_list, gtl::unsigned_int_list,
diff --git a/sycl/include/CL/sycl/known_identity.hpp b/sycl/include/CL/sycl/known_identity.hpp
index 55511bf6e44f2..ba3e4d6110bc7 100644
--- a/sycl/include/CL/sycl/known_identity.hpp
+++ b/sycl/include/CL/sycl/known_identity.hpp
@@ -138,6 +138,15 @@ struct known_identity_impl<
   static constexpr vec<std::byte, NumElements> value =
       vec<std::byte, NumElements>(std::byte(0));
 };
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsZeroIdentityOp<marray<std::byte, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      marray<std::byte, NumElements>(std::byte(0));
+};
 #endif
 
 template <typename BinaryOperation, int NumElements>
@@ -178,6 +187,15 @@ struct known_identity_impl<
   static constexpr vec<std::byte, NumElements> value =
       vec<std::byte, NumElements>(std::byte(1));
 };
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsOneIdentityOp<marray<std::byte, NumElements>,
+                                     BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      marray<std::byte, NumElements>(std::byte(1));
+};
 #endif
 
 template <typename BinaryOperation>
@@ -209,6 +227,15 @@ struct known_identity_impl<
   static constexpr vec<std::byte, NumElements> value =
       vec<std::byte, NumElements>(std::byte(-1LL));
 };
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsOnesIdentityOp<marray<std::byte, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      marray<std::byte, NumElements>(std::byte(-1LL));
+};
 #endif
 
 /// Returns maximal possible value as identity for MIN operations.
@@ -234,6 +261,18 @@ struct known_identity_impl<
               ? std::numeric_limits<vec<std::byte, NumElements>>::infinity()
               : (std::numeric_limits<vec<std::byte, NumElements>>::max)());
 };
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsMinimumIdentityOp<marray<std::byte, NumElements>,
+                                         BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      static_cast<marray<std::byte, NumElements>>(
+          std::numeric_limits<marray<std::byte, NumElements>>::has_infinity
+              ? std::numeric_limits<marray<std::byte, NumElements>>::infinity()
+              : (std::numeric_limits<marray<std::byte, NumElements>>::max)());
+};
 #endif
 
 /// Returns minimal possible value as identity for MAX operations.
@@ -261,6 +300,20 @@ struct known_identity_impl<
                 -std::numeric_limits<vec<std::byte, NumElements>>::infinity())
           : std::numeric_limits<vec<std::byte, NumElements>>::lowest());
 };
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsMaximumIdentityOp<marray<std::byte, NumElements>,
+                                         BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      static_cast<marray<std::byte, NumElements>>(
+          std::numeric_limits<marray<std::byte, NumElements>>::has_infinity
+              ? static_cast<marray<std::byte, NumElements>>(
+                    -std::numeric_limits<
+                        marray<std::byte, NumElements>>::infinity())
+              : std::numeric_limits<marray<std::byte, NumElements>>::lowest());
+};
 #endif
 
 /// Returns false as identity for LOGICAL OR operations.
@@ -280,6 +333,15 @@ struct known_identity_impl<
   static constexpr vec<std::byte, NumElements> value =
       vec<std::byte, NumElements>(std::byte(false));
 };
+
+template <typename BinaryOperation, size_t NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsFalseIdentityOp<marray<std::byte, NumElements>,
+                                       BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      marray<std::byte, NumElements>(std::byte(false));
+};
 #endif
 
 /// Returns true as identity for LOGICAL AND operations.
@@ -299,6 +361,15 @@ struct known_identity_impl<
   static constexpr vec<std::byte, NumElements> value =
       vec<std::byte, NumElements>(std::byte(true));
 };
+
+template <typename BinaryOperation, int NumElements>
+struct known_identity_impl<
+    BinaryOperation, marray<std::byte, NumElements>,
+    std::enable_if_t<IsTrueIdentityOp<marray<std::byte, NumElements>,
+                                      BinaryOperation>::value>> {
+  static constexpr marray<std::byte, NumElements> value =
+      marray<std::byte, NumElements>(std::byte(true));
+};
 #endif
 
 } // namespace detail
diff --git a/sycl/test/basic_tests/known_identity.cpp b/sycl/test/basic_tests/known_identity.cpp
index a758e7dd7f583..c6e85b26f9407 100644
--- a/sycl/test/basic_tests/known_identity.cpp
+++ b/sycl/test/basic_tests/known_identity.cpp
@@ -241,6 +241,160 @@ void checkVecTypesKnownIdentity() {
 #undef CHECK_VEC
 }
 
+template <typename T, size_t Num>
+bool compareMarrays(const marray<T, Num> a, const marray<T, Num> b) {
+  bool res = true;
+  for (int i = 0; i < Num; ++i) {
+    res &= (a[i] == b[i]);
+  }
+  if (!res) {
+    for (int i = 0; i < Num; ++i) {
+      std::cout << "(" << (int)a[i] << " == " << (int)b[i] << ")" << std::endl;
+    }
+  }
+  return res;
+}
+
+template <typename T, size_t Num>
+typename std::enable_if<!std::is_same<T, half>::value &&
+                            !std::is_same<T, float>::value &&
+                            !std::is_same<T, double>::value,
+                        void>::type
+checkMarrayKnownIdentity() {
+  constexpr marray<T, Num> zeros(T(0));
+  constexpr marray<T, Num> ones(T(1));
+  constexpr marray<T, Num> bit_ones(~T(0));
+
+  static_assert(has_known_identity<plus<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<plus<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<plus<>, marray<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<bit_or<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<bit_or<marray<T, Num>>, marray<T, Num>>::value);
+  assert(
+      compareMarrays(known_identity<bit_or<>, marray<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<bit_xor<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<bit_xor<marray<T, Num>>, marray<T, Num>>::value);
+  assert(
+      compareMarrays(known_identity<bit_xor<>, marray<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<bit_and<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<bit_and<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<bit_and<>, marray<T, Num>>::value,
+                        bit_ones));
+
+  static_assert(has_known_identity<logical_or<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<logical_or<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<logical_or<>, marray<T, Num>>::value,
+                        zeros));
+
+  static_assert(has_known_identity<logical_and<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<logical_and<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<logical_and<>, marray<T, Num>>::value,
+                        ones));
+
+  static_assert(has_known_identity<multiplies<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<multiplies<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<multiplies<>, marray<T, Num>>::value,
+                        ones));
+
+  static_assert(has_known_identity<minimum<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<minimum<marray<T, Num>>, marray<T, Num>>::value);
+  if constexpr (!std::is_same<T, std::byte>::value) {
+    constexpr marray<T, Num> maxs(-std::numeric_limits<T>::infinity());
+    assert(
+        compareMarrays(known_identity<minimum<>, marray<T, Num>>::value, maxs));
+  }
+
+  static_assert(has_known_identity<maximum<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<maximum<marray<T, Num>>, marray<T, Num>>::value);
+  if constexpr (!std::is_same<T, std::byte>::value) {
+    constexpr marray<T, Num> mins(std::numeric_limits<T>::infinity());
+    assert(
+        compareMarrays(known_identity<maximum<>, marray<T, Num>>::value, mins));
+  }
+}
+
+template <typename T, int Num>
+typename std::enable_if<std::is_same<T, sycl::half>::value ||
+                            std::is_same<T, float>::value ||
+                            std::is_same<T, double>::value,
+                        void>::type
+checkMarrayKnownIdentity() {
+  constexpr marray<T, Num> zeros(T(0.0f));
+  constexpr marray<T, Num> ones(T(1.0f));
+
+  static_assert(has_known_identity<plus<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<plus<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<plus<>, marray<T, Num>>::value, zeros));
+
+  static_assert(has_known_identity<multiplies<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<multiplies<marray<T, Num>>, marray<T, Num>>::value);
+  assert(compareMarrays(known_identity<multiplies<>, marray<T, Num>>::value,
+                        ones));
+
+  static_assert(has_known_identity<minimum<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<minimum<marray<T, Num>>, marray<T, Num>>::value);
+
+  static_assert(has_known_identity<maximum<>, marray<T, Num>>::value);
+  static_assert(
+      has_known_identity<maximum<marray<T, Num>>, marray<T, Num>>::value);
+}
+
+void checkMarrayTypesKnownIdentity() {
+
+#define CHECK_MARRAY(type)                                                     \
+  do {                                                                         \
+    checkMarrayKnownIdentity<type, 1>();                                       \
+    checkMarrayKnownIdentity<type, 2>();                                       \
+    checkMarrayKnownIdentity<type, 3>();                                       \
+    checkMarrayKnownIdentity<type, 4>();                                       \
+    checkMarrayKnownIdentity<type, 8>();                                       \
+    checkMarrayKnownIdentity<type, 16>();                                      \
+  } while (0)
+
+#if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+  CHECK_MARRAY(std::byte);
+#endif
+  CHECK_MARRAY(int8_t);
+  CHECK_MARRAY(int16_t);
+  CHECK_MARRAY(int32_t);
+  CHECK_MARRAY(int64_t);
+  CHECK_MARRAY(uint8_t);
+  CHECK_MARRAY(uint16_t);
+  CHECK_MARRAY(uint32_t);
+  CHECK_MARRAY(uint64_t);
+
+  CHECK_MARRAY(char);
+  CHECK_MARRAY(short int);
+  CHECK_MARRAY(int);
+  CHECK_MARRAY(long);
+  CHECK_MARRAY(long long);
+  CHECK_MARRAY(unsigned char);
+  CHECK_MARRAY(unsigned short int);
+  CHECK_MARRAY(unsigned int);
+  CHECK_MARRAY(unsigned long);
+  CHECK_MARRAY(unsigned long long);
+  CHECK_MARRAY(half);
+  CHECK_MARRAY(float);
+  CHECK_MARRAY(double);
+
+#undef CHECK_MARRAY
+}
+
 int main() {
   checkIntKnownIdentity<int8_t>();
   checkIntKnownIdentity<char>();
@@ -292,6 +446,7 @@ int main() {
   checkBoolKnownIdentity<bool>();
 
   checkVecTypesKnownIdentity();
+  checkMarrayTypesKnownIdentity();
 
   // Few negative tests just to check that it does not always return true.
   static_assert(!has_known_identity<std::minus<>, int>::value);

From d7b908d2ca71e5b02de032b66bd4314e770a2601 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy@intel.com>
Date: Thu, 16 Dec 2021 16:05:27 +0300
Subject: [PATCH 3/5] Test cleanup

---
 sycl/test/basic_tests/known_identity.cpp | 144 ++++++++++-------------
 1 file changed, 63 insertions(+), 81 deletions(-)

diff --git a/sycl/test/basic_tests/known_identity.cpp b/sycl/test/basic_tests/known_identity.cpp
index c6e85b26f9407..950d2873bdf74 100644
--- a/sycl/test/basic_tests/known_identity.cpp
+++ b/sycl/test/basic_tests/known_identity.cpp
@@ -101,11 +101,6 @@ bool compareVectors(const vec<T, Num> a, const vec<T, Num> b) {
   for (int i = 0; i < Num; ++i) {
     res &= (a[i] == b[i]);
   }
-  if (!res) {
-    for (int i = 0; i < Num; ++i) {
-      std::cout << "(" << (int)a[i] << " == " << (int)b[i] << ")" << std::endl;
-    }
-  }
   return res;
 }
 
@@ -195,50 +190,46 @@ checkVecKnownIdentity() {
   static_assert(has_known_identity<maximum<vec<T, Num>>, vec<T, Num>>::value);
 }
 
-void checkVecTypesKnownIdentity() {
-
-#define CHECK_VEC(type)                                                        \
-  do {                                                                         \
-    checkVecKnownIdentity<type, 1>();                                          \
-    checkVecKnownIdentity<type, 2>();                                          \
-    checkVecKnownIdentity<type, 3>();                                          \
-    checkVecKnownIdentity<type, 4>();                                          \
-    checkVecKnownIdentity<type, 8>();                                          \
-    checkVecKnownIdentity<type, 16>();                                         \
-  } while (0)
+template <typename T> void checkVecTypeKnownIdentity() {
+  checkVecKnownIdentity<T, 1>();
+  checkVecKnownIdentity<T, 2>();
+  checkVecKnownIdentity<T, 3>();
+  checkVecKnownIdentity<T, 4>();
+  checkVecKnownIdentity<T, 8>();
+  checkVecKnownIdentity<T, 16>();
+}
 
+void checkVecTypesKnownIdentity() {
 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-  CHECK_VEC(std::byte);
+  checkVecTypeKnownIdentity<std::byte>();
 #endif
-  CHECK_VEC(int8_t);
-  CHECK_VEC(int16_t);
-  CHECK_VEC(int32_t);
-  CHECK_VEC(int64_t);
-  CHECK_VEC(uint8_t);
-  CHECK_VEC(uint16_t);
-  CHECK_VEC(uint32_t);
-  CHECK_VEC(uint64_t);
-
-  CHECK_VEC(char);
-  CHECK_VEC(short int);
-  CHECK_VEC(int);
-  CHECK_VEC(long);
-  CHECK_VEC(long long);
-  CHECK_VEC(unsigned char);
-  CHECK_VEC(unsigned short int);
-  CHECK_VEC(unsigned int);
-  CHECK_VEC(unsigned long);
-  CHECK_VEC(unsigned long long);
-  CHECK_VEC(float);
-  CHECK_VEC(double);
+  checkVecTypeKnownIdentity<int8_t>();
+  checkVecTypeKnownIdentity<int16_t>();
+  checkVecTypeKnownIdentity<int32_t>();
+  checkVecTypeKnownIdentity<int64_t>();
+  checkVecTypeKnownIdentity<uint8_t>();
+  checkVecTypeKnownIdentity<uint16_t>();
+  checkVecTypeKnownIdentity<uint32_t>();
+  checkVecTypeKnownIdentity<uint64_t>();
+
+  checkVecTypeKnownIdentity<char>();
+  checkVecTypeKnownIdentity<short int>();
+  checkVecTypeKnownIdentity<int>();
+  checkVecTypeKnownIdentity<long>();
+  checkVecTypeKnownIdentity<long long>();
+  checkVecTypeKnownIdentity<unsigned char>();
+  checkVecTypeKnownIdentity<unsigned short int>();
+  checkVecTypeKnownIdentity<unsigned int>();
+  checkVecTypeKnownIdentity<unsigned long>();
+  checkVecTypeKnownIdentity<unsigned long long>();
+  checkVecTypeKnownIdentity<float>();
+  checkVecTypeKnownIdentity<double>();
 
   checkVecKnownIdentity<half, 2>();
   checkVecKnownIdentity<half, 3>();
   checkVecKnownIdentity<half, 4>();
   checkVecKnownIdentity<half, 8>();
   checkVecKnownIdentity<half, 16>();
-
-#undef CHECK_VEC
 }
 
 template <typename T, size_t Num>
@@ -247,11 +238,6 @@ bool compareMarrays(const marray<T, Num> a, const marray<T, Num> b) {
   for (int i = 0; i < Num; ++i) {
     res &= (a[i] == b[i]);
   }
-  if (!res) {
-    for (int i = 0; i < Num; ++i) {
-      std::cout << "(" << (int)a[i] << " == " << (int)b[i] << ")" << std::endl;
-    }
-  }
   return res;
 }
 
@@ -354,45 +340,41 @@ checkMarrayKnownIdentity() {
       has_known_identity<maximum<marray<T, Num>>, marray<T, Num>>::value);
 }
 
-void checkMarrayTypesKnownIdentity() {
-
-#define CHECK_MARRAY(type)                                                     \
-  do {                                                                         \
-    checkMarrayKnownIdentity<type, 1>();                                       \
-    checkMarrayKnownIdentity<type, 2>();                                       \
-    checkMarrayKnownIdentity<type, 3>();                                       \
-    checkMarrayKnownIdentity<type, 4>();                                       \
-    checkMarrayKnownIdentity<type, 8>();                                       \
-    checkMarrayKnownIdentity<type, 16>();                                      \
-  } while (0)
+template <typename T> void checkMarrayTypeKnownIdentity() {
+  checkMarrayKnownIdentity<T, 1>();
+  checkMarrayKnownIdentity<T, 2>();
+  checkMarrayKnownIdentity<T, 3>();
+  checkMarrayKnownIdentity<T, 4>();
+  checkMarrayKnownIdentity<T, 8>();
+  checkMarrayKnownIdentity<T, 16>();
+}
 
+void checkMarrayTypesKnownIdentity() {
 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-  CHECK_MARRAY(std::byte);
+  checkMarrayTypeKnownIdentity<std::byte>();
 #endif
-  CHECK_MARRAY(int8_t);
-  CHECK_MARRAY(int16_t);
-  CHECK_MARRAY(int32_t);
-  CHECK_MARRAY(int64_t);
-  CHECK_MARRAY(uint8_t);
-  CHECK_MARRAY(uint16_t);
-  CHECK_MARRAY(uint32_t);
-  CHECK_MARRAY(uint64_t);
-
-  CHECK_MARRAY(char);
-  CHECK_MARRAY(short int);
-  CHECK_MARRAY(int);
-  CHECK_MARRAY(long);
-  CHECK_MARRAY(long long);
-  CHECK_MARRAY(unsigned char);
-  CHECK_MARRAY(unsigned short int);
-  CHECK_MARRAY(unsigned int);
-  CHECK_MARRAY(unsigned long);
-  CHECK_MARRAY(unsigned long long);
-  CHECK_MARRAY(half);
-  CHECK_MARRAY(float);
-  CHECK_MARRAY(double);
-
-#undef CHECK_MARRAY
+  checkMarrayTypeKnownIdentity<int8_t>();
+  checkMarrayTypeKnownIdentity<int16_t>();
+  checkMarrayTypeKnownIdentity<int32_t>();
+  checkMarrayTypeKnownIdentity<int64_t>();
+  checkMarrayTypeKnownIdentity<uint8_t>();
+  checkMarrayTypeKnownIdentity<uint16_t>();
+  checkMarrayTypeKnownIdentity<uint32_t>();
+  checkMarrayTypeKnownIdentity<uint64_t>();
+
+  checkMarrayTypeKnownIdentity<char>();
+  checkMarrayTypeKnownIdentity<short int>();
+  checkMarrayTypeKnownIdentity<int>();
+  checkMarrayTypeKnownIdentity<long>();
+  checkMarrayTypeKnownIdentity<long long>();
+  checkMarrayTypeKnownIdentity<unsigned char>();
+  checkMarrayTypeKnownIdentity<unsigned short int>();
+  checkMarrayTypeKnownIdentity<unsigned int>();
+  checkMarrayTypeKnownIdentity<unsigned long>();
+  checkMarrayTypeKnownIdentity<unsigned long long>();
+  checkMarrayTypeKnownIdentity<half>();
+  checkMarrayTypeKnownIdentity<float>();
+  checkMarrayTypeKnownIdentity<double>();
 }
 
 int main() {

From 124f9476454f0042b23bd11ad18494762edd13ce Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <gooddoog@student.su>
Date: Fri, 17 Dec 2021 18:29:14 +0300
Subject: [PATCH 4/5] Update sycl/test/basic_tests/known_identity.cpp

Co-authored-by: Steffen Larsen <steffen.larsen@intel.com>
---
 sycl/test/basic_tests/known_identity.cpp | 6 ++----
 1 file changed, 2 insertions(+), 4 deletions(-)

diff --git a/sycl/test/basic_tests/known_identity.cpp b/sycl/test/basic_tests/known_identity.cpp
index 950d2873bdf74..75a3d35464877 100644
--- a/sycl/test/basic_tests/known_identity.cpp
+++ b/sycl/test/basic_tests/known_identity.cpp
@@ -105,10 +105,8 @@ bool compareVectors(const vec<T, Num> a, const vec<T, Num> b) {
 }
 
 template <typename T, int Num>
-typename std::enable_if<!std::is_same<T, half>::value &&
-                            !std::is_same<T, float>::value &&
-                            !std::is_same<T, double>::value,
-                        void>::type
+std::enable_if_t<!std::is_same_v<T, half> && !std::is_same_v<T, float> &&
+                 !std::is_same_v<T, double>>
 checkVecKnownIdentity() {
   constexpr vec<T, Num> zeros(T(0));
   constexpr vec<T, Num> ones(T(1));

From 8f73ea6f63b0c38f5e1e37cb30eb13e55b049b36 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy@intel.com>
Date: Fri, 17 Dec 2021 19:15:45 +0300
Subject: [PATCH 5/5] Replace std::is_same with std::is_same_v

---
 sycl/test/basic_tests/known_identity.cpp | 25 ++++++++++++------------
 1 file changed, 12 insertions(+), 13 deletions(-)

diff --git a/sycl/test/basic_tests/known_identity.cpp b/sycl/test/basic_tests/known_identity.cpp
index 75a3d35464877..2b4f9b92c9158 100644
--- a/sycl/test/basic_tests/known_identity.cpp
+++ b/sycl/test/basic_tests/known_identity.cpp
@@ -149,23 +149,23 @@ checkVecKnownIdentity() {
 
   static_assert(has_known_identity<minimum<>, vec<T, Num>>::value);
   static_assert(has_known_identity<minimum<vec<T, Num>>, vec<T, Num>>::value);
-  if constexpr (!std::is_same<T, std::byte>::value) {
+  if constexpr (!std::is_same_v<T, std::byte>) {
     constexpr vec<T, Num> maxs(-std::numeric_limits<T>::infinity());
     assert(compareVectors(known_identity<minimum<>, vec<T, Num>>::value, maxs));
   }
 
   static_assert(has_known_identity<maximum<>, vec<T, Num>>::value);
   static_assert(has_known_identity<maximum<vec<T, Num>>, vec<T, Num>>::value);
-  if constexpr (!std::is_same<T, std::byte>::value) {
+  if constexpr (!std::is_same_v<T, std::byte>) {
     constexpr vec<T, Num> mins(std::numeric_limits<T>::infinity());
     assert(compareVectors(known_identity<maximum<>, vec<T, Num>>::value, mins));
   }
 }
 
 template <typename T, int Num>
-typename std::enable_if<std::is_same<T, sycl::half>::value ||
-                            std::is_same<T, float>::value ||
-                            std::is_same<T, double>::value,
+typename std::enable_if<std::is_same_v<T, sycl::half> ||
+                            std::is_same_v<T, float> ||
+                            std::is_same_v<T, double>,
                         void>::type
 checkVecKnownIdentity() {
   constexpr vec<T, Num> zeros(T(0.0f));
@@ -240,9 +240,8 @@ bool compareMarrays(const marray<T, Num> a, const marray<T, Num> b) {
 }
 
 template <typename T, size_t Num>
-typename std::enable_if<!std::is_same<T, half>::value &&
-                            !std::is_same<T, float>::value &&
-                            !std::is_same<T, double>::value,
+typename std::enable_if<!std::is_same_v<T, half> && !std::is_same_v<T, float> &&
+                            !std::is_same_v<T, double>,
                         void>::type
 checkMarrayKnownIdentity() {
   constexpr marray<T, Num> zeros(T(0));
@@ -293,7 +292,7 @@ checkMarrayKnownIdentity() {
   static_assert(has_known_identity<minimum<>, marray<T, Num>>::value);
   static_assert(
       has_known_identity<minimum<marray<T, Num>>, marray<T, Num>>::value);
-  if constexpr (!std::is_same<T, std::byte>::value) {
+  if constexpr (!std::is_same_v<T, std::byte>) {
     constexpr marray<T, Num> maxs(-std::numeric_limits<T>::infinity());
     assert(
         compareMarrays(known_identity<minimum<>, marray<T, Num>>::value, maxs));
@@ -302,7 +301,7 @@ checkMarrayKnownIdentity() {
   static_assert(has_known_identity<maximum<>, marray<T, Num>>::value);
   static_assert(
       has_known_identity<maximum<marray<T, Num>>, marray<T, Num>>::value);
-  if constexpr (!std::is_same<T, std::byte>::value) {
+  if constexpr (!std::is_same_v<T, std::byte>) {
     constexpr marray<T, Num> mins(std::numeric_limits<T>::infinity());
     assert(
         compareMarrays(known_identity<maximum<>, marray<T, Num>>::value, mins));
@@ -310,9 +309,9 @@ checkMarrayKnownIdentity() {
 }
 
 template <typename T, int Num>
-typename std::enable_if<std::is_same<T, sycl::half>::value ||
-                            std::is_same<T, float>::value ||
-                            std::is_same<T, double>::value,
+typename std::enable_if<std::is_same_v<T, sycl::half> ||
+                            std::is_same_v<T, float> ||
+                            std::is_same_v<T, double>,
                         void>::type
 checkMarrayKnownIdentity() {
   constexpr marray<T, Num> zeros(T(0.0f));