diff options
| -rw-r--r-- | Eigen/src/Core/MathFunctions.h | 9 | ||||
| -rw-r--r-- | Eigen/src/Core/functors/UnaryFunctors.h | 24 | ||||
| -rw-r--r-- | unsupported/test/cxx11_tensor_builtins_sycl.cpp | 29 |
3 files changed, 58 insertions, 4 deletions
diff --git a/Eigen/src/Core/MathFunctions.h b/Eigen/src/Core/MathFunctions.h index 142fec998..7dfbc92d5 100644 --- a/Eigen/src/Core/MathFunctions.h +++ b/Eigen/src/Core/MathFunctions.h @@ -983,6 +983,15 @@ template<typename T> EIGEN_DEVICE_FUNC bool (isnan) (const T &x) { return inte 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); } +#if defined(__SYCL_DEVICE_ONLY__) +EIGEN_ALWAYS_INLINE float isnan(float x) { return cl::sycl::isnan(x); } +EIGEN_ALWAYS_INLINE double isnan(double x) { return cl::sycl::isnan(x); } +EIGEN_ALWAYS_INLINE float isinf(float x) { return cl::sycl::isinf(x); } +EIGEN_ALWAYS_INLINE double isinf(double x) { return cl::sycl::isinf(x); } +EIGEN_ALWAYS_INLINE float isfinite(float x) { return cl::sycl::isfinite(x); } +EIGEN_ALWAYS_INLINE double isfinite(double x) { return cl::sycl::isfinite(x); } +#endif // defined(__SYCL_DEVICE_ONLY__) + template<typename Scalar> EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(round, Scalar) round(const Scalar& x) diff --git a/Eigen/src/Core/functors/UnaryFunctors.h b/Eigen/src/Core/functors/UnaryFunctors.h index 2e6a00ffd..9d4d3eece 100644 --- a/Eigen/src/Core/functors/UnaryFunctors.h +++ b/Eigen/src/Core/functors/UnaryFunctors.h @@ -678,7 +678,13 @@ struct functor_traits<scalar_ceil_op<Scalar> > template<typename Scalar> struct scalar_isnan_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_isnan_op) typedef bool result_type; - EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return (numext::isnan)(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { +#if defined(__SYCL_DEVICE_ONLY__) + return numext::isnan(a); +#else + return (numext::isnan)(a); +#endif + } }; template<typename Scalar> struct functor_traits<scalar_isnan_op<Scalar> > @@ -696,7 +702,13 @@ struct functor_traits<scalar_isnan_op<Scalar> > template<typename Scalar> struct scalar_isinf_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_isinf_op) typedef bool result_type; - EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return (numext::isinf)(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { +#if defined(__SYCL_DEVICE_ONLY__) + return numext::isinf(a); +#else + return (numext::isinf)(a); +#endif + } }; template<typename Scalar> struct functor_traits<scalar_isinf_op<Scalar> > @@ -714,7 +726,13 @@ struct functor_traits<scalar_isinf_op<Scalar> > template<typename Scalar> struct scalar_isfinite_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_isfinite_op) typedef bool result_type; - EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { return (numext::isfinite)(a); } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE result_type operator() (const Scalar& a) const { +#if defined(__SYCL_DEVICE_ONLY__) + return numext::isfinite(a); +#else + return (numext::isfinite)(a); +#endif + } }; template<typename Scalar> struct functor_traits<scalar_isfinite_op<Scalar> > diff --git a/unsupported/test/cxx11_tensor_builtins_sycl.cpp b/unsupported/test/cxx11_tensor_builtins_sycl.cpp index 0a284e95c..d57d502ca 100644 --- a/unsupported/test/cxx11_tensor_builtins_sycl.cpp +++ b/unsupported/test/cxx11_tensor_builtins_sycl.cpp @@ -98,9 +98,36 @@ template <typename T> T inverse(T x) { return 1 / x; } TEST_UNARY_BUILTINS_FOR_SCALAR(round, SCALAR, OPERATOR) \ TEST_UNARY_BUILTINS_FOR_SCALAR(log1p, SCALAR, OPERATOR) +#define TEST_IS_THAT_RETURNS_BOOL(SCALAR, FUNC) \ + { \ + /* out OPERATOR in.FUNC() */ \ + Tensor<SCALAR, 3> in(tensorRange); \ + Tensor<bool, 3> out(tensorRange); \ + in = in.random() + static_cast<SCALAR>(0.01); \ + SCALAR *gpu_data = static_cast<SCALAR *>( \ + sycl_device.allocate(in.size() * sizeof(SCALAR))); \ + bool *gpu_data_out = \ + static_cast<bool *>(sycl_device.allocate(out.size() * sizeof(bool))); \ + TensorMap<Tensor<SCALAR, 3>> gpu(gpu_data, tensorRange); \ + TensorMap<Tensor<bool, 3>> gpu_out(gpu_data_out, tensorRange); \ + sycl_device.memcpyHostToDevice(gpu_data, in.data(), \ + (in.size()) * sizeof(SCALAR)); \ + gpu_out.device(sycl_device) = gpu.FUNC(); \ + sycl_device.memcpyDeviceToHost(out.data(), gpu_data_out, \ + (out.size()) * sizeof(bool)); \ + for (int i = 0; i < out.size(); ++i) { \ + VERIFY_IS_EQUAL(out(i), std::FUNC(in(i))); \ + } \ + sycl_device.deallocate(gpu_data); \ + sycl_device.deallocate(gpu_data_out); \ + } + #define TEST_UNARY_BUILTINS(SCALAR) \ TEST_UNARY_BUILTINS_OPERATOR(SCALAR, += ) \ - TEST_UNARY_BUILTINS_OPERATOR(SCALAR, = ) + TEST_UNARY_BUILTINS_OPERATOR(SCALAR, = ) \ + TEST_IS_THAT_RETURNS_BOOL(SCALAR, isnan) \ + TEST_IS_THAT_RETURNS_BOOL(SCALAR, isfinite) \ + TEST_IS_THAT_RETURNS_BOOL(SCALAR, isinf) static void test_builtin_unary_sycl(const Eigen::SyclDevice &sycl_device) { int sizeDim1 = 10; |