diff options
author | Mehdi Goli <mehdi.goli@example.com> | 2016-12-01 13:02:27 +0000 |
---|---|---|
committer | Mehdi Goli <mehdi.goli@example.com> | 2016-12-01 13:02:27 +0000 |
commit | 79aa2b784ecc26d6a8ef6fb2b2b053f4ad81593b (patch) | |
tree | 626e91024c30ad3caa510ca2e06548dbd6ffadce /Eigen/src/Core/MathFunctions.h | |
parent | a70393fd02fb56f432c6258ab1744e6d299797e3 (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.h | 80 |
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 |