diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
index a5866d0a391e..7182e9f79f22 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 955d7b8a9cfc..ba3e4d6110bc 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,36 @@ 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));
+};
+
+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>
+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,9 +175,29 @@ 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));
 };
 
+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>
 struct known_identity_impl<
     BinaryOperation, half,
@@ -165,48 +215,163 @@ 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));
 };
 
+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.
 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)());
 };
 
+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.
 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());
+};
+
+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.
 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));
 };
 
+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.
 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));
+};
+
+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
 
 // ---- has_known_identity
diff --git a/sycl/test/basic_tests/known_identity.cpp b/sycl/test/basic_tests/known_identity.cpp
index 06908d1ca247..2b4f9b92c915 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,285 @@ 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]);
+  }
+  return res;
+}
+
+template <typename T, int Num>
+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));
+  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_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_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_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));
+  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);
+}
+
+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)
+  checkVecTypeKnownIdentity<std::byte>();
+#endif
+  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>();
+}
+
+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]);
+  }
+  return res;
+}
+
+template <typename T, size_t Num>
+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));
+  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_v<T, std::byte>) {
+    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_v<T, std::byte>) {
+    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_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));
+  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);
+}
+
+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)
+  checkMarrayTypeKnownIdentity<std::byte>();
+#endif
+  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() {
   checkIntKnownIdentity<int8_t>();
   checkIntKnownIdentity<char>();
@@ -143,6 +424,9 @@ 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);
   static_assert(!has_known_identity<sycl::bit_or<>, float>::value);