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)