Update Eigen to commit:0cef325b07b18fff9da30f9a4b838ae2f617e2f6
CHANGELOG
=========
0cef325b0 - Fix another UB access.
5527e78a6 - Add missing x86 pcasts
24d15e086 - [SYCL-2020] Add test to validate SYCL in Eigen core.
PiperOrigin-RevId: 552576025
Change-Id: I3d9e2dab0c5a1526594b03de5e00a14a79b40863
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..d4aa4236
--- /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)
