Specialised basic math functions for SYCL device.

2 files changed, 113 insertions, 8 deletions

diff --git a/Eigen/Core b/Eigen/Core
index 82558155e..55fc886b6 100644
--- a/Eigen/Core
+++ b/Eigen/Core
@@ -43,10 +43,12 @@
   #else
     #define EIGEN_DEVICE_FUNC
   #endif
-
 #else
   #define EIGEN_DEVICE_FUNC
+#endif
 
+#if defined(EIGEN_USE_SYCL)
+  #define EIGEN_DONT_VECTORIZE
 #endif
 
 // When compiling CUDA device code with NVCC, pull in math functions from the
@@ -283,6 +285,15 @@
   #include <intrin.h>
 #endif
 
+#if defined(__SYCL_DEVICE_ONLY__)
+  #undef min
+  #undef max
+  #undef isnan
+  #undef isinf
+  #undef isfinite
+  #include <SYCL/sycl.hpp>
+#endif
+
 /** \brief Namespace containing all symbols from the %Eigen library. */
 namespace Eigen {
 
diff --git a/Eigen/src/Core/MathFunctions.h b/Eigen/src/Core/MathFunctions.h
index 8d47fb8a4..142fec998 100644
--- a/Eigen/src/Core/MathFunctions.h
+++ b/Eigen/src/Core/MathFunctions.h
@@ -413,7 +413,7 @@ inline NewType cast(const OldType& x)
     static inline Scalar run(const Scalar& x)
     {
       EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL)
-      using std::round;
+      EIGEN_USING_STD_MATH(round);
       return round(x);
     }
   };
@@ -640,7 +640,7 @@ template<typename Scalar>
 struct random_default_impl<Scalar, false, true>
 {
   static inline Scalar run(const Scalar& x, const Scalar& y)
-  { 
+  {
     typedef typename conditional<NumTraits<Scalar>::IsSigned,std::ptrdiff_t,std::size_t>::type ScalarX;
     if(y<x)
       return x;
@@ -954,6 +954,11 @@ inline EIGEN_MATHFUNC_RETVAL(log1p, Scalar) log1p(const Scalar& x)
   return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   log1p(float x) { return cl::sycl::log1p(x); }
+EIGEN_ALWAYS_INLINE double  log1p(double x) { return cl::sycl::log1p(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float log1p(const float &x) { return ::log1pf(x); }
@@ -969,6 +974,11 @@ inline typename internal::pow_impl<ScalarX,ScalarY>::result_type pow(const Scala
   return internal::pow_impl<ScalarX,ScalarY>::run(x, y);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   pow(float x, float y) { return cl::sycl::pow(x, y); }
+EIGEN_ALWAYS_INLINE double  pow(double x, double y) { return cl::sycl::pow(x, y); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 template<typename T> EIGEN_DEVICE_FUNC bool (isnan)   (const T &x) { return internal::isnan_impl(x); }
 template<typename T> EIGEN_DEVICE_FUNC bool (isinf)   (const T &x) { return internal::isinf_impl(x); }
 template<typename T> EIGEN_DEVICE_FUNC bool (isfinite)(const T &x) { return internal::isfinite_impl(x); }
@@ -980,6 +990,11 @@ inline EIGEN_MATHFUNC_RETVAL(round, Scalar) round(const Scalar& x)
   return EIGEN_MATHFUNC_IMPL(round, Scalar)::run(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   round(float x) { return cl::sycl::round(x); }
+EIGEN_ALWAYS_INLINE double  round(double x) { return cl::sycl::round(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 template<typename T>
 EIGEN_DEVICE_FUNC
 T (floor)(const T& x)
@@ -988,6 +1003,11 @@ T (floor)(const T& x)
   return floor(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   floor(float x) { return cl::sycl::floor(x); }
+EIGEN_ALWAYS_INLINE double  floor(double x) { return cl::sycl::floor(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float floor(const float &x) { return ::floorf(x); }
@@ -1004,6 +1024,11 @@ T (ceil)(const T& x)
   return ceil(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   ceil(float x) { return cl::sycl::ceil(x); }
+EIGEN_ALWAYS_INLINE double  ceil(double x) { return cl::sycl::ceil(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float ceil(const float &x) { return ::ceilf(x); }
@@ -1044,6 +1069,11 @@ T sqrt(const T &x)
   return sqrt(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   sqrt(float x) { return cl::sycl::sqrt(x); }
+EIGEN_ALWAYS_INLINE double  sqrt(double x) { return cl::sycl::sqrt(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 template<typename T>
 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 T log(const T &x) {
@@ -1051,6 +1081,12 @@ T log(const T &x) {
   return log(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   log(float x) { return cl::sycl::log(x); }
+EIGEN_ALWAYS_INLINE double  log(double x) { return cl::sycl::log(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float log(const float &x) { return ::logf(x); }
@@ -1066,6 +1102,11 @@ typename NumTraits<T>::Real abs(const T &x) {
   return abs(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   abs(float x) { return cl::sycl::fabs(x); }
+EIGEN_ALWAYS_INLINE double  abs(double x) { return cl::sycl::fabs(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float abs(const float &x) { return ::fabsf(x); }
@@ -1091,6 +1132,11 @@ T exp(const T &x) {
   return exp(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   exp(float x) { return cl::sycl::exp(x); }
+EIGEN_ALWAYS_INLINE double  exp(double x) { return cl::sycl::exp(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float exp(const float &x) { return ::expf(x); }
@@ -1106,6 +1152,11 @@ T cos(const T &x) {
   return cos(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   cos(float x) { return cl::sycl::cos(x); }
+EIGEN_ALWAYS_INLINE double  cos(double x) { return cl::sycl::cos(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float cos(const float &x) { return ::cosf(x); }
@@ -1121,6 +1172,11 @@ T sin(const T &x) {
   return sin(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   sin(float x) { return cl::sycl::sin(x); }
+EIGEN_ALWAYS_INLINE double  sin(double x) { return cl::sycl::sin(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float sin(const float &x) { return ::sinf(x); }
@@ -1136,6 +1192,11 @@ T tan(const T &x) {
   return tan(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   tan(float x) { return cl::sycl::tan(x); }
+EIGEN_ALWAYS_INLINE double  tan(double x) { return cl::sycl::tan(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float tan(const float &x) { return ::tanf(x); }
@@ -1151,6 +1212,11 @@ T acos(const T &x) {
   return acos(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   acos(float x) { return cl::sycl::acos(x); }
+EIGEN_ALWAYS_INLINE double  acos(double x) { return cl::sycl::acos(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float acos(const float &x) { return ::acosf(x); }
@@ -1166,6 +1232,11 @@ T asin(const T &x) {
   return asin(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   asin(float x) { return cl::sycl::asin(x); }
+EIGEN_ALWAYS_INLINE double  asin(double x) { return cl::sycl::asin(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float asin(const float &x) { return ::asinf(x); }
@@ -1181,6 +1252,11 @@ T atan(const T &x) {
   return atan(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   atan(float x) { return cl::sycl::atan(x); }
+EIGEN_ALWAYS_INLINE double  atan(double x) { return cl::sycl::atan(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float atan(const float &x) { return ::atanf(x); }
@@ -1197,6 +1273,11 @@ T cosh(const T &x) {
   return cosh(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   cosh(float x) { return cl::sycl::cosh(x); }
+EIGEN_ALWAYS_INLINE double  cosh(double x) { return cl::sycl::cosh(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float cosh(const float &x) { return ::coshf(x); }
@@ -1212,6 +1293,11 @@ T sinh(const T &x) {
   return sinh(x);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   sinh(float x) { return cl::sycl::sinh(x); }
+EIGEN_ALWAYS_INLINE double  sinh(double x) { return cl::sycl::sinh(x); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float sinh(const float &x) { return ::sinhf(x); }
@@ -1227,7 +1313,10 @@ T tanh(const T &x) {
   return tanh(x);
 }
 
-#if (!defined(__CUDACC__)) && EIGEN_FAST_MATH
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   tanh(float x) { return cl::sycl::tanh(x); }
+EIGEN_ALWAYS_INLINE double  tanh(double x) { return cl::sycl::tanh(x); }
+#elif (!defined(__CUDACC__)) && EIGEN_FAST_MATH
 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
 float tanh(float x) { return internal::generic_fast_tanh_float(x); }
 #endif
@@ -1247,6 +1336,11 @@ T fmod(const T& a, const T& b) {
   return fmod(a, b);
 }
 
+#if defined(__SYCL_DEVICE_ONLY__)
+EIGEN_ALWAYS_INLINE float   fmod(float x, float y) { return cl::sycl::fmod(x, y); }
+EIGEN_ALWAYS_INLINE double  fmod(double x, double y) { return cl::sycl::fmod(x, y); }
+#endif // defined(__SYCL_DEVICE_ONLY__)
+
 #ifdef __CUDACC__
 template <>
 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
@@ -1389,13 +1483,13 @@ template<> struct random_impl<bool>
 template<> struct scalar_fuzzy_impl<bool>
 {
   typedef bool RealScalar;
-  
+
   template<typename OtherScalar> EIGEN_DEVICE_FUNC
   static inline bool isMuchSmallerThan(const bool& x, const bool&, const bool&)
   {
     return !x;
   }
-  
+
   EIGEN_DEVICE_FUNC
   static inline bool isApprox(bool x, bool y, bool)
   {
@@ -1407,10 +1501,10 @@ template<> struct scalar_fuzzy_impl<bool>
   {
     return (!x) || y;
   }
-  
+
 };
 
-  
+
 } // end namespace internal
 
 } // end namespace Eigen