diff options
| author |  Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-11-02 11:44:27 -0700 |
|---|---|---|
| committer | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-11-02 11:44:27 -0700 |
| commit | 0585b2965d06cc2c57be35844bd2d0d56e6334ac (patch) | |
| tree | 64a24b32a0d5a6cc882a78adb524b394b1c5ec0c | |
| parent | e44519744e6a788c12e572d654bb21de6bdf5684 (diff) | |
Disable vectorization on device only when compiling for sycl
| -rw-r--r-- | Eigen/Core | 10 | ||||
| -rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h | 21 |
2 files changed, 26 insertions, 5 deletions
diff --git a/Eigen/Core b/Eigen/Core index c82cddc5c..2d2616254 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -35,17 +35,17 @@ #undef EIGEN_INTERNAL_DEBUGGING #endif - // Do not try to vectorize on CUDA and SYCL! - #ifndef EIGEN_DONT_VECTORIZE - #define EIGEN_DONT_VECTORIZE - #endif - #ifdef EIGEN_EXCEPTIONS #undef EIGEN_EXCEPTIONS #endif // All functions callable from CUDA code must be qualified with __device__ #ifdef __CUDACC__ + // Do not try to vectorize on CUDA and SYCL! + #ifndef EIGEN_DONT_VECTORIZE + #define EIGEN_DONT_VECTORIZE + #endif + #define EIGEN_DEVICE_FUNC __host__ __device__ // We need math_functions.hpp to ensure that that EIGEN_USING_STD_MATH macro // works properly on the device side diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h b/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h index fdb5ee6b8..615559d44 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h @@ -83,6 +83,27 @@ struct PacketType<half, GpuDevice> { }; #endif +#if defined(EIGEN_USE_SYCL) +template <typename T> + struct PacketType<T, SyclDevice> { + typedef T type; + static const int size = 1; + enum { + HasAdd = 0, + HasSub = 0, + HasMul = 0, + HasNegate = 0, + HasAbs = 0, + HasArg = 0, + HasAbs2 = 0, + HasMin = 0, + HasMax = 0, + HasConj = 0, + HasSetLinear = 0, + HasBlend = 0 + }; +}; +#endif // Tuple mimics std::pair but works on e.g. nvcc. |