aboutsummaryrefslogtreecommitdiffhomepage
path: root/Eigen/src/Core/MathFunctions.h
diff options
context:
space:
mode:
authorGravatar Mehdi Goli <mehdi.goli@example.com>2016-12-01 13:02:27 +0000
committerGravatar Mehdi Goli <mehdi.goli@example.com>2016-12-01 13:02:27 +0000
commit79aa2b784ecc26d6a8ef6fb2b2b053f4ad81593b (patch)
tree626e91024c30ad3caa510ca2e06548dbd6ffadce /Eigen/src/Core/MathFunctions.h
parenta70393fd02fb56f432c6258ab1744e6d299797e3 (diff)
Adding sycl backend for TensorPadding.h; disbaling __unit128 for sycl in TensorIntDiv.h; disabling cashsize for sycl in tensorDeviceDefault.h; adding sycl backend for StrideSliceOP ; removing sycl compiler warning for creating an array of size 0 in CXX11Meta.h; cleaning up the sycl backend code.
Diffstat (limited to 'Eigen/src/Core/MathFunctions.h')
-rw-r--r--Eigen/src/Core/MathFunctions.h80
1 files changed, 79 insertions, 1 deletions
diff --git a/Eigen/src/Core/MathFunctions.h b/Eigen/src/Core/MathFunctions.h
index 7dfbc92d5..1ac0b2473 100644
--- a/Eigen/src/Core/MathFunctions.h
+++ b/Eigen/src/Core/MathFunctions.h
@@ -826,7 +826,7 @@ template<typename T> T generic_fast_tanh_float(const T& a_x);
namespace numext {
-#ifndef __CUDA_ARCH__
+#if !defined(__CUDA_ARCH__) && !defined(__SYCL_DEVICE_ONLY__)
template<typename T>
EIGEN_DEVICE_FUNC
EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y)
@@ -842,6 +842,84 @@ EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y)
EIGEN_USING_STD_MATH(max);
return max EIGEN_NOT_A_MACRO (x,y);
}
+
+
+#elif defined(__SYCL_DEVICE_ONLY__)
+template<typename T>
+EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y)
+{
+
+ return y < x ? y : x;
+}
+
+template<typename T>
+EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y)
+{
+
+ return x < y ? y : x;
+}
+
+EIGEN_ALWAYS_INLINE int mini(const int& x, const int& y)
+{
+ return cl::sycl::min(x,y);
+}
+
+EIGEN_ALWAYS_INLINE int maxi(const int& x, const int& y)
+{
+ return cl::sycl::max(x,y);
+}
+
+EIGEN_ALWAYS_INLINE unsigned int mini(const unsigned int& x, const unsigned int& y)
+{
+ return cl::sycl::min(x,y);
+}
+
+EIGEN_ALWAYS_INLINE unsigned int maxi(const unsigned int& x, const unsigned int& y)
+{
+ return cl::sycl::max(x,y);
+}
+
+EIGEN_ALWAYS_INLINE long mini(const long & x, const long & y)
+{
+ return cl::sycl::min(x,y);
+}
+
+EIGEN_ALWAYS_INLINE long maxi(const long & x, const long & y)
+{
+ return cl::sycl::max(x,y);
+}
+
+EIGEN_ALWAYS_INLINE unsigned long mini(const unsigned long& x, const unsigned long& y)
+{
+ return cl::sycl::min(x,y);
+}
+
+EIGEN_ALWAYS_INLINE unsigned long maxi(const unsigned long& x, const unsigned long& y)
+{
+ return cl::sycl::max(x,y);
+}
+
+
+EIGEN_ALWAYS_INLINE float mini(const float& x, const float& y)
+{
+ return cl::sycl::fmin(x,y);
+}
+
+EIGEN_ALWAYS_INLINE float maxi(const float& x, const float& y)
+{
+ return cl::sycl::fmax(x,y);
+}
+
+EIGEN_ALWAYS_INLINE double mini(const double& x, const double& y)
+{
+ return cl::sycl::fmin(x,y);
+}
+
+EIGEN_ALWAYS_INLINE double maxi(const double& x, const double& y)
+{
+ return cl::sycl::fmax(x,y);
+}
+
#else
template<typename T>
EIGEN_DEVICE_FUNC