diff --git a/Eigen/Core b/Eigen/Core index 864bde551..080511fc9 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -200,6 +200,12 @@ using std::ptrdiff_t; #include "src/Core/arch/GPU/MathFunctions.h" #endif +#if defined EIGEN_VECTORIZE_SYCL + #include "src/Core/arch/SYCL/InteropHeaders.h" + #include "src/Core/arch/SYCL/PacketMath.h" + #include "src/Core/arch/SYCL/MathFunctions.h" + #include "src/Core/arch/SYCL/TypeCasting.h" +#endif #include "src/Core/arch/Default/Settings.h" #include "src/Core/functors/TernaryFunctors.h" diff --git a/Eigen/src/Core/arch/GPU/Half.h b/Eigen/src/Core/arch/GPU/Half.h index ab9d27591..aca56fa72 100644 --- a/Eigen/src/Core/arch/GPU/Half.h +++ b/Eigen/src/Core/arch/GPU/Half.h @@ -83,7 +83,11 @@ struct __half_raw { #if defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER < 90000 // In CUDA < 9.0, __half is the equivalent of CUDA 9's __half_raw typedef __half __half_raw; - #endif + #endif // defined(EIGEN_HAS_CUDA_FP16) + +#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) +typedef cl::sycl::half __half_raw; + #endif EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw raw_uint16_to_half(unsigned short x); @@ -200,6 +204,7 @@ struct half : public half_impl::half_base { x = other.x; return *this; } + }; } // end namespace Eigen diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h b/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h index 8ed1796df..ec1dc0fab 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h @@ -21,6 +21,7 @@ namespace Eigen { template struct MakePointer { typedef T* Type; typedef T& RefType; + typedef T ScalarType; }; namespace internal{ diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorTrace.h b/unsupported/Eigen/CXX11/src/Tensor/TensorTrace.h index c8b2fad1e..ea53bb04b 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorTrace.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorTrace.h @@ -273,11 +273,11 @@ struct TensorEvaluator, Device> Dimensions m_dimensions; TensorEvaluator m_impl; + // Initialize the size of the trace dimension + Index m_traceDim; const Device& m_device; array m_reduced; array m_reducedDims; - // Initialize the size of the trace dimension - Index m_traceDim; array m_outputStrides; array m_reducedStrides; array m_preservedStrides; diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h b/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h index 006b37921..0a394c88d 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h @@ -59,6 +59,7 @@ struct traits > template struct MakePointer { typedef T* Type; typedef T& RefType; + typedef T ScalarType; }; typedef typename MakePointer::Type PointerType; @@ -80,6 +81,7 @@ struct traits > template struct MakePointer { typedef T* Type; typedef T& RefType; + typedef T ScalarType; }; typedef typename MakePointer::Type PointerType; @@ -105,6 +107,8 @@ struct traits > typedef MakePointer_ MakePointerT; typedef typename MakePointerT::Type Type; typedef typename MakePointerT::RefType RefType; + typedef typename MakePointerT::ScalarType ScalarType; + }; typedef typename MakePointer::Type PointerType; diff --git a/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h b/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h index dbcc9d8ac..5784cbc86 100644 --- a/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h +++ b/unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h @@ -193,6 +193,8 @@ struct lgamma_impl { #if !defined(EIGEN_GPU_COMPILE_PHASE) && (defined(_BSD_SOURCE) || defined(_SVID_SOURCE)) && !defined(__APPLE__) int dummy; return ::lgammaf_r(x, &dummy); +#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::lgamma(x); #else return ::lgammaf(x); #endif @@ -206,6 +208,8 @@ struct lgamma_impl { #if !defined(EIGEN_GPU_COMPILE_PHASE) && (defined(_BSD_SOURCE) || defined(_SVID_SOURCE)) && !defined(__APPLE__) int dummy; return ::lgamma_r(x, &dummy); +#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::lgamma(x); #else return ::lgamma(x); #endif @@ -423,13 +427,25 @@ struct erf_retval { template <> struct erf_impl { EIGEN_DEVICE_FUNC - static EIGEN_STRONG_INLINE float run(float x) { return ::erff(x); } + static EIGEN_STRONG_INLINE float run(float x) { +#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::erf(x); +#else + return ::erff(x); +#endif + } }; template <> struct erf_impl { EIGEN_DEVICE_FUNC - static EIGEN_STRONG_INLINE double run(double x) { return ::erf(x); } + static EIGEN_STRONG_INLINE double run(double x) { +#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::erf(x); +#else + return ::erf(x); +#endif + } }; #endif // EIGEN_HAS_C99_MATH @@ -456,13 +472,25 @@ struct erfc_retval { template <> struct erfc_impl { EIGEN_DEVICE_FUNC - static EIGEN_STRONG_INLINE float run(const float x) { return ::erfcf(x); } + static EIGEN_STRONG_INLINE float run(const float x) { +#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::erfc(x); +#else + return ::erfcf(x); +#endif + } }; template <> struct erfc_impl { EIGEN_DEVICE_FUNC - static EIGEN_STRONG_INLINE double run(const double x) { return ::erfc(x); } + static EIGEN_STRONG_INLINE double run(const double x) { +#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::erfc(x); +#else + return ::erfc(x); +#endif + } }; #endif // EIGEN_HAS_C99_MATH