diff --git a/Eigen/src/Core/GenericPacketMath.h b/Eigen/src/Core/GenericPacketMath.h
index ceb7a0a..f3d607a 100644
--- a/Eigen/src/Core/GenericPacketMath.h
+++ b/Eigen/src/Core/GenericPacketMath.h
@@ -206,6 +206,17 @@
   };
 };
 
+// provides a succint template to define vectorized casting traits with respect to the largest accessible packet types
+template <typename Src, typename Tgt>
+struct vectorized_type_casting_traits {
+  enum : int {
+    DefaultSrcPacketSize = packet_traits<Src>::size,
+    DefaultTgtPacketSize = packet_traits<Tgt>::size,
+    VectorizedCast = 1,
+    SrcCoeffRatio = plain_enum_max(DefaultTgtPacketSize / DefaultSrcPacketSize, 1),
+    TgtCoeffRatio = plain_enum_max(DefaultSrcPacketSize / DefaultTgtPacketSize, 1)
+  };
+};
 
 /** \internal Wrapper to ensure that multiple packet types can map to the same
     same underlying vector type. */
diff --git a/Eigen/src/Core/SolveTriangular.h b/Eigen/src/Core/SolveTriangular.h
index 71d6f85..23df508 100644
--- a/Eigen/src/Core/SolveTriangular.h
+++ b/Eigen/src/Core/SolveTriangular.h
@@ -97,6 +97,11 @@
     typedef internal::gemm_blocking_space<(Rhs::Flags&RowMajorBit) ? RowMajor : ColMajor,Scalar,Scalar,
               Rhs::MaxRowsAtCompileTime, Rhs::MaxColsAtCompileTime, Lhs::MaxRowsAtCompileTime,4> BlockingType;
 
+    // Nothing to solve.
+    if (actualLhs.size() == 0 || rhs.size() == 0) {
+      return;
+    }
+
     BlockingType blocking(rhs.rows(), rhs.cols(), size, 1, false);
 
     triangular_solve_matrix<Scalar,Index,Side,Mode,LhsProductTraits::NeedToConjugate,(int(Lhs::Flags) & RowMajorBit) ? RowMajor : ColMajor,
diff --git a/Eigen/src/Core/arch/AVX/TypeCasting.h b/Eigen/src/Core/arch/AVX/TypeCasting.h
index 461f3a6..9853347 100644
--- a/Eigen/src/Core/arch/AVX/TypeCasting.h
+++ b/Eigen/src/Core/arch/AVX/TypeCasting.h
@@ -17,76 +17,24 @@
 namespace internal {
 
 #ifndef EIGEN_VECTORIZE_AVX512
-template <>
-struct type_casting_traits<Eigen::half, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<float, bool> : vectorized_type_casting_traits<float, bool> {};
+template<> struct type_casting_traits<bool, float> : vectorized_type_casting_traits<bool, float> {};
 
+template<> struct type_casting_traits<float, int> : vectorized_type_casting_traits<float, int> {};
+template<> struct type_casting_traits<int, float> : vectorized_type_casting_traits<int, float> {};
 
-template <>
-struct type_casting_traits<float, Eigen::half> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<float, double> : vectorized_type_casting_traits<float, double> {};
+template<> struct type_casting_traits<double, float> : vectorized_type_casting_traits<double, float> {};
 
-template <>
-struct type_casting_traits<bfloat16, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<double, int> : vectorized_type_casting_traits<double, int> {};
+template<> struct type_casting_traits<int, double> : vectorized_type_casting_traits<int, double> {};
 
-template <>
-struct type_casting_traits<float, bfloat16> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<half, float> : vectorized_type_casting_traits<half, float> {};
+template<> struct type_casting_traits<float, half> : vectorized_type_casting_traits<float, half> {};
 
-template <>
-struct type_casting_traits<float, bool> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 2,
-    TgtCoeffRatio = 1
-  };
-};
-#endif  // EIGEN_VECTORIZE_AVX512
-
-template<> EIGEN_STRONG_INLINE Packet8i pcast<Packet8f, Packet8i>(const Packet8f& a) {
-  return _mm256_cvttps_epi32(a);
-}
-
-template<> EIGEN_STRONG_INLINE Packet8f pcast<Packet8i, Packet8f>(const Packet8i& a) {
-  return _mm256_cvtepi32_ps(a);
-}
-
-template<> EIGEN_STRONG_INLINE Packet8f pcast<Packet4d, Packet8f>(const Packet4d& a, const Packet4d& b) {
-  return _mm256_set_m128(_mm256_cvtpd_ps(b), _mm256_cvtpd_ps(a));
-}
-
-template<> EIGEN_STRONG_INLINE Packet8i pcast<Packet4d, Packet8i>(const Packet4d& a, const Packet4d& b) {
-  return _mm256_set_m128i(_mm256_cvttpd_epi32(b), _mm256_cvttpd_epi32(a));
-}
-
-template <> EIGEN_STRONG_INLINE Packet4f pcast<Packet4d, Packet4f>(const Packet4d& a) {
-  return _mm256_cvtpd_ps(a);
-}
-
-template <> EIGEN_STRONG_INLINE Packet4i pcast<Packet4d, Packet4i>(const Packet4d& a) {
-  return _mm256_cvttpd_epi32(a);
-}
+template<> struct type_casting_traits<bfloat16, float> : vectorized_type_casting_traits<bfloat16, float> {};
+template<> struct type_casting_traits<float, bfloat16> : vectorized_type_casting_traits<float, bfloat16> {};
+#endif
 
 template <>
 EIGEN_STRONG_INLINE Packet16b pcast<Packet8f, Packet16b>(const Packet8f& a,
@@ -118,6 +66,63 @@
 #endif
 }
 
+template <>
+EIGEN_STRONG_INLINE Packet8f pcast<Packet16b, Packet8f>(const Packet16b& a) {
+  const __m256 cst_one = _mm256_set1_ps(1.0f);
+  #ifdef EIGEN_VECTORIZE_AVX2
+  __m256i a_extended = _mm256_cvtepi8_epi32(a);
+  __m256i abcd_efgh = _mm256_cmpeq_epi32(a_extended, _mm256_setzero_si256());
+  #else
+  __m128i abcd_efhg_ijkl_mnop = _mm_cmpeq_epi8(a, _mm_setzero_si128());
+  __m128i aabb_ccdd_eeff_gghh = _mm_unpacklo_epi8(abcd_efhg_ijkl_mnop, abcd_efhg_ijkl_mnop);
+  __m128i aaaa_bbbb_cccc_dddd = _mm_unpacklo_epi8(aabb_ccdd_eeff_gghh, aabb_ccdd_eeff_gghh);
+  __m128i eeee_ffff_gggg_hhhh = _mm_unpackhi_epi8(aabb_ccdd_eeff_gghh, aabb_ccdd_eeff_gghh);
+  __m256i abcd_efgh = _mm256_setr_m128i(aaaa_bbbb_cccc_dddd, eeee_ffff_gggg_hhhh);
+  #endif
+  __m256 result = _mm256_andnot_ps(_mm256_castsi256_ps(abcd_efgh), cst_one);
+  return result;
+}
+
+template<> EIGEN_STRONG_INLINE Packet8i pcast<Packet8f, Packet8i>(const Packet8f& a) {
+  return _mm256_cvttps_epi32(a);
+}
+
+template<> EIGEN_STRONG_INLINE Packet8i pcast<Packet4d, Packet8i>(const Packet4d& a, const Packet4d& b) {
+  return _mm256_set_m128i(_mm256_cvttpd_epi32(b), _mm256_cvttpd_epi32(a));
+}
+
+template <> EIGEN_STRONG_INLINE Packet4i pcast<Packet4d, Packet4i>(const Packet4d& a) {
+  return _mm256_cvttpd_epi32(a);
+}
+
+template<> EIGEN_STRONG_INLINE Packet8f pcast<Packet8i, Packet8f>(const Packet8i& a) {
+  return _mm256_cvtepi32_ps(a);
+}
+
+template<> EIGEN_STRONG_INLINE Packet8f pcast<Packet4d, Packet8f>(const Packet4d& a, const Packet4d& b) {
+  return _mm256_set_m128(_mm256_cvtpd_ps(b), _mm256_cvtpd_ps(a));
+}
+
+template <> EIGEN_STRONG_INLINE Packet4f pcast<Packet4d, Packet4f>(const Packet4d& a) {
+  return _mm256_cvtpd_ps(a);
+}
+
+template <> EIGEN_STRONG_INLINE Packet4d pcast<Packet8i, Packet4d>(const Packet8i& a) {
+  return _mm256_cvtepi32_pd(_mm256_castsi256_si128(a));
+}
+
+template <> EIGEN_STRONG_INLINE Packet4d pcast<Packet4i, Packet4d>(const Packet4i& a) {
+  return _mm256_cvtepi32_pd(a);
+}
+
+template <> EIGEN_STRONG_INLINE Packet4d pcast<Packet8f, Packet4d>(const Packet8f& a) {
+  return _mm256_cvtps_pd(_mm256_castps256_ps128(a));
+}
+
+template <> EIGEN_STRONG_INLINE Packet4d pcast<Packet4f, Packet4d>(const Packet4f& a) {
+  return _mm256_cvtps_pd(a);
+}
+
 template<> EIGEN_STRONG_INLINE Packet8i preinterpret<Packet8i,Packet8f>(const Packet8f& a) {
   return _mm256_castps_si256(a);
 }
diff --git a/Eigen/src/Core/arch/AVX512/TypeCasting.h b/Eigen/src/Core/arch/AVX512/TypeCasting.h
index 2f38d7f..02c5628 100644
--- a/Eigen/src/Core/arch/AVX512/TypeCasting.h
+++ b/Eigen/src/Core/arch/AVX512/TypeCasting.h
@@ -16,23 +16,23 @@
 
 namespace internal {
 
-template <>
-struct type_casting_traits<float, bool> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<float, bool> : vectorized_type_casting_traits<float, bool> {};
+template<> struct type_casting_traits<bool, float> : vectorized_type_casting_traits<bool, float> {};
 
-template <>
-struct type_casting_traits<bool, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<float, int> : vectorized_type_casting_traits<float, int> {};
+template<> struct type_casting_traits<int, float> : vectorized_type_casting_traits<int, float> {};
+
+template<> struct type_casting_traits<float, double> : vectorized_type_casting_traits<float, double> {};
+template<> struct type_casting_traits<double, float> : vectorized_type_casting_traits<double, float> {};
+
+template<> struct type_casting_traits<double, int> : vectorized_type_casting_traits<double, int> {};
+template<> struct type_casting_traits<int, double> : vectorized_type_casting_traits<int, double> {};
+
+template<> struct type_casting_traits<half, float> : vectorized_type_casting_traits<half, float> {};
+template<> struct type_casting_traits<float, half> : vectorized_type_casting_traits<float, half> {};
+
+template<> struct type_casting_traits<bfloat16, float> : vectorized_type_casting_traits<bfloat16, float> {};
+template<> struct type_casting_traits<float, bfloat16> : vectorized_type_casting_traits<float, bfloat16> {};
 
 template<> EIGEN_STRONG_INLINE Packet16b pcast<Packet16f, Packet16b>(const Packet16f& a) {
   __mmask16 mask = _mm512_cmpneq_ps_mask(a, pzero(a));
@@ -47,10 +47,26 @@
   return _mm512_cvttps_epi32(a);
 }
 
+template<> EIGEN_STRONG_INLINE Packet8d pcast<Packet16f, Packet8d>(const Packet16f& a) {
+  return _mm512_cvtps_pd(_mm512_castps512_ps256(a));
+}
+
+template<> EIGEN_STRONG_INLINE Packet8d pcast<Packet8f, Packet8d>(const Packet8f& a) {
+  return _mm512_cvtps_pd(a);
+}
+
 template<> EIGEN_STRONG_INLINE Packet16f pcast<Packet16i, Packet16f>(const Packet16i& a) {
   return _mm512_cvtepi32_ps(a);
 }
 
+template<> EIGEN_STRONG_INLINE Packet8d pcast<Packet16i, Packet8d>(const Packet16i& a) {
+  return _mm512_cvtepi32_pd(_mm512_castsi512_si256(a));
+}
+
+template<> EIGEN_STRONG_INLINE Packet8d pcast<Packet8i, Packet8d>(const Packet8i& a) {
+  return _mm512_cvtepi32_pd(a);
+}
+
 template<> EIGEN_STRONG_INLINE Packet16f pcast<Packet8d, Packet16f>(const Packet8d& a, const Packet8d& b) {
   return  cat256(_mm512_cvtpd_ps(a), _mm512_cvtpd_ps(b));
 }
@@ -131,80 +147,26 @@
 
 #ifndef EIGEN_VECTORIZE_AVX512FP16
 
-template <>
-struct type_casting_traits<half, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
-
 template<> EIGEN_STRONG_INLINE Packet16f pcast<Packet16h, Packet16f>(const Packet16h& a) {
   return half2float(a);
 }
 
-template <>
-struct type_casting_traits<float, half> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
-
 template<> EIGEN_STRONG_INLINE Packet16h pcast<Packet16f, Packet16h>(const Packet16f& a) {
   return float2half(a);
 }
 
 #endif
 
-template <>
-struct type_casting_traits<bfloat16, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
-
 template<> EIGEN_STRONG_INLINE Packet16f pcast<Packet16bf, Packet16f>(const Packet16bf& a) {
   return Bf16ToF32(a);
 }
 
-template <>
-struct type_casting_traits<float, bfloat16> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
-
 template<> EIGEN_STRONG_INLINE Packet16bf pcast<Packet16f, Packet16bf>(const Packet16f& a) {
   return F32ToBf16(a);
 }
 
 #ifdef EIGEN_VECTORIZE_AVX512FP16
 
-template <>
-struct type_casting_traits<half, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 2
-  };
-};
-
-template <>
-struct type_casting_traits<float, half> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 2,
-    TgtCoeffRatio = 1
-  };
-};
-
 template<> EIGEN_STRONG_INLINE Packet16h preinterpret<Packet16h, Packet32h>(const Packet32h& a) {
   return _mm256_castpd_si256(_mm512_extractf64x4_pd(_mm512_castph_pd(a), 0));
 }
@@ -257,7 +219,7 @@
   __m256 result = _mm256_undefined_ps();
   result = _mm256_insertf128_ps(result, a, 0);
   result = _mm256_insertf128_ps(result, b, 1);
-  return _mm256_cvtps_ph(result, _MM_FROUND_TO_NEAREST_INT|_MM_FROUND_NO_EXC);
+  return _mm256_cvtps_ph(result, _MM_FROUND_TO_NEAREST_INT);
 }
 
 
diff --git a/Eigen/src/Core/arch/SSE/TypeCasting.h b/Eigen/src/Core/arch/SSE/TypeCasting.h
index 0b5aa1c..bb28170 100644
--- a/Eigen/src/Core/arch/SSE/TypeCasting.h
+++ b/Eigen/src/Core/arch/SSE/TypeCasting.h
@@ -17,62 +17,20 @@
 namespace internal {
 
 #ifndef EIGEN_VECTORIZE_AVX
-template <>
-struct type_casting_traits<float, bool> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 4,
-    TgtCoeffRatio = 1
-  };
-};
+template<> struct type_casting_traits<float, bool> : vectorized_type_casting_traits<float, bool> {};
+template<> struct type_casting_traits<bool, float> : vectorized_type_casting_traits<bool, float> {};
 
-template <>
-struct type_casting_traits<float, double> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 2
-  };
-};
+template<> struct type_casting_traits<float, int> : vectorized_type_casting_traits<float, int> {};
+template<> struct type_casting_traits<int, float> : vectorized_type_casting_traits<int, float> {};
+
+template<> struct type_casting_traits<float, double> : vectorized_type_casting_traits<float, double> {};
+template<> struct type_casting_traits<double, float> : vectorized_type_casting_traits<double, float> {};
+
+template<> struct type_casting_traits<double, int> : vectorized_type_casting_traits<double, int> {};
+template<> struct type_casting_traits<int, double> : vectorized_type_casting_traits<int, double> {};
 #endif
 
 template <>
-struct type_casting_traits<int, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
-
-template <>
-struct type_casting_traits<float, int> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 1,
-    TgtCoeffRatio = 1
-  };
-};
-
-template <>
-struct type_casting_traits<double, int> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 2,
-    TgtCoeffRatio = 1
-  };
-};
-
-template <>
-struct type_casting_traits<double, float> {
-  enum {
-    VectorizedCast = 1,
-    SrcCoeffRatio = 2,
-    TgtCoeffRatio = 1
-  };
-};
-
-template <>
 EIGEN_STRONG_INLINE Packet16b pcast<Packet4f, Packet16b>(const Packet4f& a,
                                                          const Packet4f& b,
                                                          const Packet4f& c,
@@ -88,10 +46,31 @@
   return _mm_and_si128(merged, _mm_set1_epi8(1));
 }
 
+template <>
+EIGEN_STRONG_INLINE Packet4f pcast<Packet16b, Packet4f>(const Packet16b& a) {
+  const __m128 cst_one = _mm_set_ps1(1.0f);
+  #ifdef EIGEN_VECTORIZE_SSE4_1
+  __m128i a_extended = _mm_cvtepi8_epi32(a);
+  __m128i abcd = _mm_cmpeq_epi32(a_extended, _mm_setzero_si128());
+  #else
+  __m128i abcd_efhg_ijkl_mnop = _mm_cmpeq_epi8(a, _mm_setzero_si128());
+  __m128i aabb_ccdd_eeff_gghh = _mm_unpacklo_epi8(abcd_efhg_ijkl_mnop, abcd_efhg_ijkl_mnop);
+  __m128i abcd = _mm_unpacklo_epi8(aabb_ccdd_eeff_gghh, aabb_ccdd_eeff_gghh);
+  #endif
+  __m128 result = _mm_andnot_ps(_mm_castsi128_ps(abcd), cst_one);
+  return result;
+}
+
 template<> EIGEN_STRONG_INLINE Packet4i pcast<Packet4f, Packet4i>(const Packet4f& a) {
   return _mm_cvttps_epi32(a);
 }
 
+template<> EIGEN_STRONG_INLINE Packet4i pcast<Packet2d, Packet4i>(const Packet2d& a, const Packet2d& b) {
+  return _mm_castps_si128(_mm_shuffle_ps(_mm_castsi128_ps(_mm_cvttpd_epi32(a)),
+                                         _mm_castsi128_ps(_mm_cvttpd_epi32(b)),
+                                         (1 << 2) | (1 << 6)));
+}
+
 template<> EIGEN_STRONG_INLINE Packet4f pcast<Packet4i, Packet4f>(const Packet4i& a) {
   return _mm_cvtepi32_ps(a);
 }
@@ -100,10 +79,9 @@
   return _mm_shuffle_ps(_mm_cvtpd_ps(a), _mm_cvtpd_ps(b), (1 << 2) | (1 << 6));
 }
 
-template<> EIGEN_STRONG_INLINE Packet4i pcast<Packet2d, Packet4i>(const Packet2d& a, const Packet2d& b) {
-  return _mm_castps_si128(_mm_shuffle_ps(_mm_castsi128_ps(_mm_cvttpd_epi32(a)),
-                                         _mm_castsi128_ps(_mm_cvttpd_epi32(b)),
-                                         (1 << 2) | (1 << 6)));
+template<> EIGEN_STRONG_INLINE Packet2d pcast<Packet4i, Packet2d>(const Packet4i& a) {
+  // Simply discard the second half of the input
+  return _mm_cvtepi32_pd(a);
 }
 
 template<> EIGEN_STRONG_INLINE Packet2d pcast<Packet4f, Packet2d>(const Packet4f& a) {
diff --git a/Eigen/src/Core/products/TriangularMatrixMatrix.h b/Eigen/src/Core/products/TriangularMatrixMatrix.h
index 94eabdc..80c98dd 100644
--- a/Eigen/src/Core/products/TriangularMatrixMatrix.h
+++ b/Eigen/src/Core/products/TriangularMatrixMatrix.h
@@ -424,7 +424,7 @@
 
     // Empty product, return early.  Otherwise, we get `nullptr` use errors below when we try to access
     // coeffRef(0,0).
-    if (a_lhs.size() == 0 || a_rhs.size() == 0) {
+    if (lhs.size() == 0 || rhs.size() == 0) {
       return;
     }
 
diff --git a/cmake/EigenTesting.cmake b/cmake/EigenTesting.cmake
index 639790c..2022cf0 100644
--- a/cmake/EigenTesting.cmake
+++ b/cmake/EigenTesting.cmake
@@ -368,8 +368,10 @@
     if(EIGEN_TEST_SYCL)
       if(EIGEN_SYCL_TRISYCL)
         message(STATUS "SYCL:              ON (using triSYCL)")
-      else()
+      elseif(EIGEN_SYCL_ComputeCpp)
         message(STATUS "SYCL:              ON (using computeCPP)")
+      elseif(EIGEN_SYCL_DPCPP)
+        message(STATUS "SYCL:              ON (using DPCPP)")
       endif()
     else()
       message(STATUS "SYCL:              OFF")
diff --git a/cmake/SyclConfigureTesting.cmake b/cmake/SyclConfigureTesting.cmake
new file mode 100644
index 0000000..d4aa423
--- /dev/null
+++ b/cmake/SyclConfigureTesting.cmake
@@ -0,0 +1,64 @@
+set(CMAKE_CXX_STANDARD 17)
+# Forward CMake options as preprocessor definitions
+if(EIGEN_SYCL_USE_DEFAULT_SELECTOR)
+    add_definitions(-DEIGEN_SYCL_USE_DEFAULT_SELECTOR=${EIGEN_SYCL_USE_DEFAULT_SELECTOR})
+endif()
+if(EIGEN_SYCL_NO_LOCAL_MEM)
+    add_definitions(-DEIGEN_SYCL_NO_LOCAL_MEM=${EIGEN_SYCL_NO_LOCAL_MEM})
+endif()
+if(EIGEN_SYCL_LOCAL_MEM)
+    add_definitions(-DEIGEN_SYCL_LOCAL_MEM=${EIGEN_SYCL_LOCAL_MEM})
+endif()
+if(EIGEN_SYCL_MAX_GLOBAL_RANGE)
+    add_definitions(-DEIGEN_SYCL_MAX_GLOBAL_RANGE=${EIGEN_SYCL_MAX_GLOBAL_RANGE})
+endif()
+if(EIGEN_SYCL_LOCAL_THREAD_DIM0)
+    add_definitions(-DEIGEN_SYCL_LOCAL_THREAD_DIM0=${EIGEN_SYCL_LOCAL_THREAD_DIM0})
+endif()
+if(EIGEN_SYCL_LOCAL_THREAD_DIM1)
+    add_definitions(-DEIGEN_SYCL_LOCAL_THREAD_DIM1=${EIGEN_SYCL_LOCAL_THREAD_DIM1})
+endif()
+if(EIGEN_SYCL_REG_M)
+    add_definitions(-DEIGEN_SYCL_REG_M=${EIGEN_SYCL_REG_M})
+endif()
+if(EIGEN_SYCL_REG_N)
+    add_definitions(-DEIGEN_SYCL_REG_N=${EIGEN_SYCL_REG_N})
+endif()
+if(EIGEN_SYCL_ASYNC_EXECUTION)
+    add_definitions(-DEIGEN_SYCL_ASYNC_EXECUTION=${EIGEN_SYCL_ASYNC_EXECUTION})
+endif()
+if(EIGEN_SYCL_DISABLE_SKINNY)
+    add_definitions(-DEIGEN_SYCL_DISABLE_SKINNY=${EIGEN_SYCL_DISABLE_SKINNY})
+endif()
+if(EIGEN_SYCL_DISABLE_DOUBLE_BUFFER)
+    add_definitions(-DEIGEN_SYCL_DISABLE_DOUBLE_BUFFER=${EIGEN_SYCL_DISABLE_DOUBLE_BUFFER})
+endif()
+if(EIGEN_SYCL_DISABLE_SCALAR)
+    add_definitions(-DEIGEN_SYCL_DISABLE_SCALAR=${EIGEN_SYCL_DISABLE_SCALAR})
+endif()
+if(EIGEN_SYCL_DISABLE_GEMV)
+    add_definitions(-DEIGEN_SYCL_DISABLE_GEMV=${EIGEN_SYCL_DISABLE_GEMV})
+endif()
+if(EIGEN_SYCL_DISABLE_ARM_GPU_CACHE_OPTIMISATION)
+    add_definitions(-DEIGEN_SYCL_DISABLE_ARM_GPU_CACHE_OPTIMISATION=${EIGEN_SYCL_DISABLE_ARM_GPU_CACHE_OPTIMISATION})
+endif()
+
+if(EIGEN_SYCL_ComputeCpp)
+    if(MSVC)
+        list(APPEND COMPUTECPP_USER_FLAGS -DWIN32)
+    else()
+        list(APPEND COMPUTECPP_USER_FLAGS -Wall)
+    endif()
+    # The following flags are not supported by Clang and can cause warnings
+    # if used with -Werror so they are removed here.
+    if(COMPUTECPP_USE_COMPILER_DRIVER)
+        set(CMAKE_CXX_COMPILER ${ComputeCpp_DEVICE_COMPILER_EXECUTABLE})
+        string(REPLACE "-Wlogical-op" "" CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS})
+        string(REPLACE "-Wno-psabi" "" CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS})
+    endif()
+    list(APPEND COMPUTECPP_USER_FLAGS
+            -DEIGEN_NO_ASSERTION_CHECKING=1
+            -no-serial-memop
+            -Xclang
+            -cl-mad-enable)
+endif(EIGEN_SYCL_ComputeCpp)
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 98d1bad..e1a056f 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -477,6 +477,14 @@
   endif()
 endif()
 
+if(EIGEN_TEST_SYCL)
+  set(EIGEN_SYCL ON)
+  include(SyclConfigureTesting)
+
+  ei_add_test(sycl_basic)
+  set(EIGEN_SYCL OFF)
+endif()
+
 cmake_dependent_option(EIGEN_TEST_BUILD_DOCUMENTATION "Test building the doxygen documentation" OFF "EIGEN_BUILD_DOC" OFF)
 if(EIGEN_TEST_BUILD_DOCUMENTATION)
   add_dependencies(buildtests doc)
diff --git a/test/array_cwise.cpp b/test/array_cwise.cpp
index d06fa2c..49e6672 100644
--- a/test/array_cwise.cpp
+++ b/test/array_cwise.cpp
@@ -1211,6 +1211,22 @@
     typed_logicals_test_impl<ArrayType>::run(m);
 }
 
+// print non-mangled typenames
+template<typename T> std::string printTypeInfo(const T&) { return typeid(T).name(); }
+template<> std::string printTypeInfo(const int8_t&) { return "int8_t"; }
+template<> std::string printTypeInfo(const int16_t&) { return "int16_t"; }
+template<> std::string printTypeInfo(const int32_t&) { return "int32_t"; }
+template<> std::string printTypeInfo(const int64_t&) { return "int64_t"; }
+template<> std::string printTypeInfo(const uint8_t&) { return "uint8_t"; }
+template<> std::string printTypeInfo(const uint16_t&) { return "uint16_t"; }
+template<> std::string printTypeInfo(const uint32_t&) { return "uint32_t"; }
+template<> std::string printTypeInfo(const uint64_t&) { return "uint64_t"; }
+template<> std::string printTypeInfo(const float&) { return "float"; }
+template<> std::string printTypeInfo(const double&) { return "double"; }
+//template<> std::string printTypeInfo(const long double&) { return "long double"; }
+template<> std::string printTypeInfo(const half&) { return "half"; }
+template<> std::string printTypeInfo(const bfloat16&) { return "bfloat16"; }
+
 template <typename SrcType, typename DstType, int RowsAtCompileTime, int ColsAtCompileTime>
 struct cast_test_impl {
   using SrcArray = Array<SrcType, RowsAtCompileTime, ColsAtCompileTime>;
@@ -1225,63 +1241,30 @@
   static constexpr int DstPacketSize = internal::packet_traits<DstType>::size;
   static constexpr int MaxPacketSize = internal::plain_enum_max(SrcPacketSize, DstPacketSize);
 
-  // print non-mangled typenames
-  template <typename T>
-  static std::string printTypeInfo(const T&) {
-    if (internal::is_same<bool, T>::value)
-      return "bool";
-    else if (internal::is_same<int8_t, T>::value)
-      return "int8_t";
-    else if (internal::is_same<int16_t, T>::value)
-      return "int16_t";
-    else if (internal::is_same<int32_t, T>::value)
-      return "int32_t";
-    else if (internal::is_same<int64_t, T>::value)
-      return "int64_t";
-    else if (internal::is_same<uint8_t, T>::value)
-      return "uint8_t";
-    else if (internal::is_same<uint16_t, T>::value)
-      return "uint16_t";
-    else if (internal::is_same<uint32_t, T>::value)
-      return "uint32_t";
-    else if (internal::is_same<uint64_t, T>::value)
-      return "uint64_t";
-    else if (internal::is_same<float, T>::value)
-      return "float";
-    else if (internal::is_same<double, T>::value)
-      return "double";
-    //else if (internal::is_same<long double, T>::value)
-    //  return "long double";
-    else if (internal::is_same<half, T>::value)
-      return "half";
-    else if (internal::is_same<bfloat16, T>::value)
-      return "bfloat16";
-    else
-      return typeid(T).name();
-  }
-
   static void run() {
     const Index testRows = RowsAtCompileTime == Dynamic ? ((10 * MaxPacketSize) + 1) : RowsAtCompileTime;
     const Index testCols = ColsAtCompileTime == Dynamic ? ((10 * MaxPacketSize) + 1) : ColsAtCompileTime;
     const Index testSize = testRows * testCols;
     const Index minTestSize = 100;
     const Index repeats = numext::div_ceil(minTestSize, testSize);
+
     SrcArray src(testRows, testCols);
     DstArray dst(testRows, testCols);
+
     for (Index repeat = 0; repeat < repeats; repeat++) {
       src = src.unaryExpr(RandomOp());
       dst = src.template cast<DstType>();
-      for (Index i = 0; i < testRows; i++)
-        for (Index j = 0; j < testCols; j++) {
-          DstType ref = internal::cast_impl<SrcType, DstType>::run(src(i, j));
-          bool all_nan = ((numext::isnan)(src(i, j)) && (numext::isnan)(ref) && (numext::isnan)(dst(i, j)));
-          bool is_equal = ref == dst(i, j);
-          bool pass = all_nan || is_equal;
-          if (!pass) {
-            std::cout << printTypeInfo(SrcType()) << ": [" << +src(i, j) << "] to " << printTypeInfo(DstType()) << ": ["
-                      << +dst(i, j) << "] != [" << +ref << "]\n";
-          }
-          VERIFY(pass);
+
+      for (Index j = 0; j < testCols; j++)
+        for (Index i = 0; i < testRows; i++) {
+          SrcType srcVal = src(i, j);
+          DstType refVal = internal::cast_impl<SrcType, DstType>::run(srcVal);
+          DstType dstVal = dst(i, j);
+          bool isApprox = verifyIsApprox(dstVal, refVal);
+          if (!isApprox)
+            std::cout << printTypeInfo(srcVal) << ": [" << +srcVal << "] to " << printTypeInfo(dstVal) << ": ["
+                      << +dstVal << "] != [" << +refVal << "]\n";
+          VERIFY(isApprox);
         }
     }
   }
diff --git a/test/sycl_basic.cpp b/test/sycl_basic.cpp
new file mode 100644
index 0000000..06f03c4
--- /dev/null
+++ b/test/sycl_basic.cpp
@@ -0,0 +1,382 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Copyright (C) 2023
+// Alejandro Acosta    Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+// Copyright (C) 2015-2016 Gael Guennebaud <gael.guennebaud@inria.fr>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+#define EIGEN_TEST_NO_LONGDOUBLE
+#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int
+
+#define EIGEN_USE_SYCL
+#include "main.h"
+
+#include <Eigen/Dense>
+
+template <bool verifyNan = false, bool singleTask = false, typename Operation, typename Input, typename Output>
+void run_and_verify(Operation& ope, size_t num_elements, const Input& in, Output& out) {
+  Output out_gpu, out_cpu;
+  out_gpu = out_cpu = out;
+  auto queue = sycl::queue{sycl::default_selector_v};
+
+  auto in_size_bytes = sizeof(typename Input::Scalar) * in.size();
+  auto out_size_bytes = sizeof(typename Output::Scalar) * out.size();
+  auto in_d = sycl::malloc_device<typename Input::Scalar>(in.size(), queue);
+  auto out_d = sycl::malloc_device<typename Output::Scalar>(out.size(), queue);
+
+  queue.memcpy(in_d, in.data(), in_size_bytes).wait();
+  queue.memcpy(out_d, out.data(), out_size_bytes).wait();
+
+  if constexpr (singleTask) {
+    queue.single_task([=]() { ope(in_d, out_d); }).wait();
+  } else {
+    queue
+        .parallel_for(sycl::range{num_elements},
+                      [=](sycl::id<1> idx) {
+                        auto id = idx[0];
+                        ope(id, in_d, out_d);
+                      })
+        .wait();
+  }
+
+  queue.memcpy(out_gpu.data(), out_d, out_size_bytes).wait();
+
+  sycl::free(in_d, queue);
+  sycl::free(out_d, queue);
+
+  queue.throw_asynchronous();
+
+  // Run on CPU and compare the output
+  if constexpr (singleTask == 1) {
+    ope(in.data(), out_cpu.data());
+  } else {
+    for (size_t i = 0; i < num_elements; ++i) {
+      ope(i, in.data(), out_cpu.data());
+    }
+  }
+  if constexpr (verifyNan) {
+    VERIFY_IS_CWISE_APPROX(out_gpu, out_cpu);
+  } else {
+    VERIFY_IS_APPROX(out_gpu, out_cpu);
+  }
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_coeff_wise(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    DataType x1(in + i);
+    DataType x2(in + i + 1);
+    DataType x3(in + i + 2);
+    Map<DataType> res(out + i * DataType::MaxSizeAtCompileTime);
+
+    res.array() += (in[0] * x1 + x2).array() * x3.array();
+  };
+
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_complex_sqrt(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    using namespace Eigen;
+    typedef typename DataType::Scalar ComplexType;
+    typedef typename DataType::Scalar::value_type ValueType;
+    const int num_special_inputs = 18;
+
+    if (i == 0) {
+      const ValueType nan = std::numeric_limits<ValueType>::quiet_NaN();
+      typedef Eigen::Vector<ComplexType, num_special_inputs> SpecialInputs;
+      SpecialInputs special_in;
+      special_in.setZero();
+      int idx = 0;
+      special_in[idx++] = ComplexType(0, 0);
+      special_in[idx++] = ComplexType(-0, 0);
+      special_in[idx++] = ComplexType(0, -0);
+      special_in[idx++] = ComplexType(-0, -0);
+      const ValueType inf = std::numeric_limits<ValueType>::infinity();
+      special_in[idx++] = ComplexType(1.0, inf);
+      special_in[idx++] = ComplexType(nan, inf);
+      special_in[idx++] = ComplexType(1.0, -inf);
+      special_in[idx++] = ComplexType(nan, -inf);
+      special_in[idx++] = ComplexType(-inf, 1.0);
+      special_in[idx++] = ComplexType(inf, 1.0);
+      special_in[idx++] = ComplexType(-inf, -1.0);
+      special_in[idx++] = ComplexType(inf, -1.0);
+      special_in[idx++] = ComplexType(-inf, nan);
+      special_in[idx++] = ComplexType(inf, nan);
+      special_in[idx++] = ComplexType(1.0, nan);
+      special_in[idx++] = ComplexType(nan, 1.0);
+      special_in[idx++] = ComplexType(nan, -1.0);
+      special_in[idx++] = ComplexType(nan, nan);
+
+      Map<SpecialInputs> special_out(out);
+      special_out = special_in.cwiseSqrt();
+    }
+
+    DataType x1(in + i);
+    Map<DataType> res(out + num_special_inputs + i * DataType::MaxSizeAtCompileTime);
+    res = x1.cwiseSqrt();
+  };
+  run_and_verify<true>(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_complex_operators(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    using namespace Eigen;
+    typedef typename DataType::Scalar ComplexType;
+    typedef typename DataType::Scalar::value_type ValueType;
+    const int num_scalar_operators = 24;
+    const int num_vector_operators = 23;  // no unary + operator.
+    size_t out_idx = i * (num_scalar_operators + num_vector_operators * DataType::MaxSizeAtCompileTime);
+
+    // Scalar operators.
+    const ComplexType a = in[i];
+    const ComplexType b = in[i + 1];
+
+    out[out_idx++] = +a;
+    out[out_idx++] = -a;
+
+    out[out_idx++] = a + b;
+    out[out_idx++] = a + numext::real(b);
+    out[out_idx++] = numext::real(a) + b;
+    out[out_idx++] = a - b;
+    out[out_idx++] = a - numext::real(b);
+    out[out_idx++] = numext::real(a) - b;
+    out[out_idx++] = a * b;
+    out[out_idx++] = a * numext::real(b);
+    out[out_idx++] = numext::real(a) * b;
+    out[out_idx++] = a / b;
+    out[out_idx++] = a / numext::real(b);
+    out[out_idx++] = numext::real(a) / b;
+
+    out[out_idx] = a;
+    out[out_idx++] += b;
+    out[out_idx] = a;
+    out[out_idx++] -= b;
+    out[out_idx] = a;
+    out[out_idx++] *= b;
+    out[out_idx] = a;
+    out[out_idx++] /= b;
+
+    const ComplexType true_value = ComplexType(ValueType(1), ValueType(0));
+    const ComplexType false_value = ComplexType(ValueType(0), ValueType(0));
+    out[out_idx++] = (a == b ? true_value : false_value);
+    out[out_idx++] = (a == numext::real(b) ? true_value : false_value);
+    out[out_idx++] = (numext::real(a) == b ? true_value : false_value);
+    out[out_idx++] = (a != b ? true_value : false_value);
+    out[out_idx++] = (a != numext::real(b) ? true_value : false_value);
+    out[out_idx++] = (numext::real(a) != b ? true_value : false_value);
+
+    // Vector versions.
+    DataType x1(in + i);
+    DataType x2(in + i + 1);
+    const int res_size = DataType::MaxSizeAtCompileTime * num_scalar_operators;
+    const int size = DataType::MaxSizeAtCompileTime;
+    int block_idx = 0;
+
+    Map<VectorX<ComplexType>> res(out + out_idx, res_size);
+    res.segment(block_idx, size) = -x1;
+    block_idx += size;
+
+    res.segment(block_idx, size) = x1 + x2;
+    block_idx += size;
+    res.segment(block_idx, size) = x1 + x2.real();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.real() + x2;
+    block_idx += size;
+    res.segment(block_idx, size) = x1 - x2;
+    block_idx += size;
+    res.segment(block_idx, size) = x1 - x2.real();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.real() - x2;
+    block_idx += size;
+    res.segment(block_idx, size) = x1.array() * x2.array();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.array() * x2.real().array();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.real().array() * x2.array();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.array() / x2.array();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.array() / x2.real().array();
+    block_idx += size;
+    res.segment(block_idx, size) = x1.real().array() / x2.array();
+    block_idx += size;
+
+    res.segment(block_idx, size) = x1;
+    res.segment(block_idx, size) += x2;
+    block_idx += size;
+    res.segment(block_idx, size) = x1;
+    res.segment(block_idx, size) -= x2;
+    block_idx += size;
+    res.segment(block_idx, size) = x1;
+    res.segment(block_idx, size).array() *= x2.array();
+    block_idx += size;
+    res.segment(block_idx, size) = x1;
+    res.segment(block_idx, size).array() /= x2.array();
+    block_idx += size;
+
+    const DataType true_vector = DataType::Constant(true_value);
+    const DataType false_vector = DataType::Constant(false_value);
+    res.segment(block_idx, size) = (x1 == x2 ? true_vector : false_vector);
+    block_idx += size;
+    res.segment(block_idx, size) = (x1 == x2.real() ? true_vector : false_vector);
+    block_idx += size;
+    //        res.segment(block_idx, size) = (x1.real() == x2) ? true_vector : false_vector;
+    //        block_idx += size;
+    res.segment(block_idx, size) = (x1 != x2 ? true_vector : false_vector);
+    block_idx += size;
+    res.segment(block_idx, size) = (x1 != x2.real() ? true_vector : false_vector);
+    block_idx += size;
+    //        res.segment(block_idx, size) = (x1.real() != x2 ? true_vector : false_vector);
+    //        block_idx += size;
+  };
+  run_and_verify<true>(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_redux(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    using namespace Eigen;
+    int N = 10;
+    DataType x1(in + i);
+    out[i * N + 0] = x1.minCoeff();
+    out[i * N + 1] = x1.maxCoeff();
+    out[i * N + 2] = x1.sum();
+    out[i * N + 3] = x1.prod();
+    out[i * N + 4] = x1.matrix().squaredNorm();
+    out[i * N + 5] = x1.matrix().norm();
+    out[i * N + 6] = x1.colwise().sum().maxCoeff();
+    out[i * N + 7] = x1.rowwise().maxCoeff().sum();
+    out[i * N + 8] = x1.matrix().colwise().squaredNorm().sum();
+  };
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_replicate(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    using namespace Eigen;
+    DataType x1(in + i);
+    int step = x1.size() * 4;
+    int stride = 3 * step;
+
+    typedef Map<Array<typename DataType::Scalar, Dynamic, Dynamic>> MapType;
+    MapType(out + i * stride + 0 * step, x1.rows() * 2, x1.cols() * 2) = x1.replicate(2, 2);
+    MapType(out + i * stride + 1 * step, x1.rows() * 3, x1.cols()) = in[i] * x1.colwise().replicate(3);
+    MapType(out + i * stride + 2 * step, x1.rows(), x1.cols() * 3) = in[i] * x1.rowwise().replicate(3);
+  };
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType1, typename DataType2, typename Input, typename Output>
+void test_product(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType1::Scalar* in, typename DataType1::Scalar* out) {
+    using namespace Eigen;
+    typedef Matrix<typename DataType1::Scalar, DataType1::RowsAtCompileTime, DataType2::ColsAtCompileTime> DataType3;
+    DataType1 x1(in + i);
+    DataType2 x2(in + i + 1);
+    Map<DataType3> res(out + i * DataType3::MaxSizeAtCompileTime);
+    res += in[i] * x1 * x2;
+  };
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType1, typename DataType2, typename Input, typename Output>
+void test_diagonal(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType1::Scalar* in, typename DataType1::Scalar* out) {
+    using namespace Eigen;
+    DataType1 x1(in + i);
+    Map<DataType2> res(out + i * DataType2::MaxSizeAtCompileTime);
+    res += x1.diagonal();
+  };
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_eigenvalues_direct(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    using namespace Eigen;
+    typedef Matrix<typename DataType::Scalar, DataType::RowsAtCompileTime, 1> Vec;
+    DataType M(in + i);
+    Map<Vec> res(out + i * Vec::MaxSizeAtCompileTime);
+    DataType A = M * M.adjoint();
+    SelfAdjointEigenSolver<DataType> eig;
+    eig.computeDirect(A);
+    res = eig.eigenvalues();
+  };
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_matrix_inverse(size_t num_elements, const Input& in, Output& out) {
+  auto operation = [](size_t i, const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    using namespace Eigen;
+    DataType M(in + i);
+    Map<DataType> res(out + i * DataType::MaxSizeAtCompileTime);
+    res = M.inverse();
+  };
+  run_and_verify(operation, num_elements, in, out);
+}
+
+template <typename DataType, typename Input, typename Output>
+void test_numeric_limits(const Input& in, Output& out) {
+  auto operation = [](const typename DataType::Scalar* in, typename DataType::Scalar* out) {
+    EIGEN_UNUSED_VARIABLE(in)
+    out[0] = numext::numeric_limits<float>::epsilon();
+    out[1] = (numext::numeric_limits<float>::max)();
+    out[2] = (numext::numeric_limits<float>::min)();
+    out[3] = numext::numeric_limits<float>::infinity();
+    out[4] = numext::numeric_limits<float>::quiet_NaN();
+  };
+  run_and_verify<true, true>(operation, 1, in, out);
+}
+
+EIGEN_DECLARE_TEST(sycl_basic) {
+  Eigen::VectorXf in, out;
+  Eigen::VectorXcf cfin, cfout;
+
+  constexpr size_t num_elements = 100;
+  constexpr size_t data_size = num_elements * 512;
+  in.setRandom(data_size);
+  out.setConstant(data_size, -1);
+  cfin.setRandom(data_size);
+  cfout.setConstant(data_size, -1);
+
+  CALL_SUBTEST(test_coeff_wise<Vector3f>(num_elements, in, out));
+  CALL_SUBTEST(test_coeff_wise<Array44f>(num_elements, in, out));
+
+  CALL_SUBTEST(test_complex_operators<Vector3cf>(num_elements, cfin, cfout));
+  CALL_SUBTEST(test_complex_sqrt<Vector3cf>(num_elements, cfin, cfout));
+
+  CALL_SUBTEST(test_redux<Array4f>(num_elements, in, out));
+  CALL_SUBTEST(test_redux<Matrix3f>(num_elements, in, out));
+
+  CALL_SUBTEST(test_replicate<Array4f>(num_elements, in, out));
+  CALL_SUBTEST(test_replicate<Array33f>(num_elements, in, out));
+
+  auto test_prod_mm = [&]() { test_product<Matrix3f, Matrix3f>(num_elements, in, out); };
+  auto test_prod_mv = [&]() { test_product<Matrix4f, Vector4f>(num_elements, in, out); };
+  CALL_SUBTEST(test_prod_mm());
+  CALL_SUBTEST(test_prod_mv());
+
+  auto test_diagonal_mv3f = [&]() { test_diagonal<Matrix3f, Vector3f>(num_elements, in, out); };
+  auto test_diagonal_mv4f = [&]() { test_diagonal<Matrix4f, Vector4f>(num_elements, in, out); };
+  CALL_SUBTEST(test_diagonal_mv3f());
+  CALL_SUBTEST(test_diagonal_mv4f());
+
+  CALL_SUBTEST(test_eigenvalues_direct<Matrix3f>(num_elements, in, out));
+  CALL_SUBTEST(test_eigenvalues_direct<Matrix2f>(num_elements, in, out));
+
+  CALL_SUBTEST(test_matrix_inverse<Matrix2f>(num_elements, in, out));
+  CALL_SUBTEST(test_matrix_inverse<Matrix3f>(num_elements, in, out));
+  CALL_SUBTEST(test_matrix_inverse<Matrix4f>(num_elements, in, out));
+
+  CALL_SUBTEST(test_numeric_limits<Vector3f>(in, out));
+}
diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt
index 2bb5518..1d40ae5 100644
--- a/unsupported/test/CMakeLists.txt
+++ b/unsupported/test/CMakeLists.txt
@@ -122,73 +122,7 @@
 
 if(EIGEN_TEST_SYCL)
   set(EIGEN_SYCL ON)
-  set(CMAKE_CXX_STANDARD 17)
-  # Forward CMake options as preprocessor definitions
-  if(EIGEN_SYCL_USE_DEFAULT_SELECTOR)
-    add_definitions(-DEIGEN_SYCL_USE_DEFAULT_SELECTOR=${EIGEN_SYCL_USE_DEFAULT_SELECTOR})
-  endif()
-  if(EIGEN_SYCL_NO_LOCAL_MEM)
-    add_definitions(-DEIGEN_SYCL_NO_LOCAL_MEM=${EIGEN_SYCL_NO_LOCAL_MEM})
-  endif()
-  if(EIGEN_SYCL_LOCAL_MEM)
-    add_definitions(-DEIGEN_SYCL_LOCAL_MEM=${EIGEN_SYCL_LOCAL_MEM})
-  endif()
-  if(EIGEN_SYCL_MAX_GLOBAL_RANGE)
-    add_definitions(-DEIGEN_SYCL_MAX_GLOBAL_RANGE=${EIGEN_SYCL_MAX_GLOBAL_RANGE})
-  endif()
-  if(EIGEN_SYCL_LOCAL_THREAD_DIM0)
-    add_definitions(-DEIGEN_SYCL_LOCAL_THREAD_DIM0=${EIGEN_SYCL_LOCAL_THREAD_DIM0})
-  endif()
-  if(EIGEN_SYCL_LOCAL_THREAD_DIM1)
-    add_definitions(-DEIGEN_SYCL_LOCAL_THREAD_DIM1=${EIGEN_SYCL_LOCAL_THREAD_DIM1})
-  endif()
-  if(EIGEN_SYCL_REG_M)
-    add_definitions(-DEIGEN_SYCL_REG_M=${EIGEN_SYCL_REG_M})
-  endif()
-  if(EIGEN_SYCL_REG_N)
-    add_definitions(-DEIGEN_SYCL_REG_N=${EIGEN_SYCL_REG_N})
-  endif()
-  if(EIGEN_SYCL_ASYNC_EXECUTION)
-    add_definitions(-DEIGEN_SYCL_ASYNC_EXECUTION=${EIGEN_SYCL_ASYNC_EXECUTION})
-  endif()
-  if(EIGEN_SYCL_DISABLE_SKINNY)
-    add_definitions(-DEIGEN_SYCL_DISABLE_SKINNY=${EIGEN_SYCL_DISABLE_SKINNY})
-  endif()
-  if(EIGEN_SYCL_DISABLE_DOUBLE_BUFFER)
-    add_definitions(-DEIGEN_SYCL_DISABLE_DOUBLE_BUFFER=${EIGEN_SYCL_DISABLE_DOUBLE_BUFFER})
-  endif()
-  if(EIGEN_SYCL_DISABLE_RANK1)
-    add_definitions(-DEIGEN_SYCL_DISABLE_RANK1=${EIGEN_SYCL_DISABLE_RANK1})
-  endif()
-  if(EIGEN_SYCL_DISABLE_SCALAR)
-    add_definitions(-DEIGEN_SYCL_DISABLE_SCALAR=${EIGEN_SYCL_DISABLE_SCALAR})
-  endif()
-  if(EIGEN_SYCL_DISABLE_GEMV)
-    add_definitions(-DEIGEN_SYCL_DISABLE_GEMV=${EIGEN_SYCL_DISABLE_GEMV})
-  endif()
-  if(EIGEN_SYCL_DISABLE_ARM_GPU_CACHE_OPTIMISATION)
-    add_definitions(-DEIGEN_SYCL_DISABLE_ARM_GPU_CACHE_OPTIMISATION=${EIGEN_SYCL_DISABLE_ARM_GPU_CACHE_OPTIMISATION})
-  endif()
-
-  if(EIGEN_SYCL_ComputeCpp)
-    if(MSVC)
-      list(APPEND COMPUTECPP_USER_FLAGS -DWIN32)
-    else()
-      list(APPEND COMPUTECPP_USER_FLAGS -Wall)
-    endif()
-    # The following flags are not supported by Clang and can cause warnings
-    # if used with -Werror so they are removed here.
-    if(COMPUTECPP_USE_COMPILER_DRIVER)
-      set(CMAKE_CXX_COMPILER ${ComputeCpp_DEVICE_COMPILER_EXECUTABLE})
-      string(REPLACE "-Wlogical-op" "" CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS})
-      string(REPLACE "-Wno-psabi" "" CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS})
-    endif()
-    list(APPEND COMPUTECPP_USER_FLAGS
-        -DEIGEN_NO_ASSERTION_CHECKING=1
-        -no-serial-memop
-        -Xclang
-        -cl-mad-enable)
-  endif(EIGEN_SYCL_ComputeCpp)
+  include(SyclConfigureTesting)
 
   ei_add_test(cxx11_tensor_sycl)
   ei_add_test(cxx11_tensor_image_op_sycl)