From 5159675c338ffef579fa7015fe5e05eb27bcbdb5 Mon Sep 17 00:00:00 2001 From: Luke Iwanski Date: Fri, 18 Nov 2016 16:01:48 +0000 Subject: [PATCH] Added isnan, isfinite and isinf for SYCL device. Plus test for that. --- Eigen/src/Core/MathFunctions.h | 9 ++++++ Eigen/src/Core/functors/UnaryFunctors.h | 24 +++++++++++++-- .../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 EIGEN_DEVICE_FUNC bool (isnan) (const T &x) { return inte template EIGEN_DEVICE_FUNC bool (isinf) (const T &x) { return internal::isinf_impl(x); } template 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 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 > template 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 struct functor_traits > @@ -696,7 +702,13 @@ struct functor_traits > template 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 struct functor_traits > @@ -714,7 +726,13 @@ struct functor_traits > template 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 struct functor_traits > 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 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 in(tensorRange); \ + Tensor out(tensorRange); \ + in = in.random() + static_cast(0.01); \ + SCALAR *gpu_data = static_cast( \ + sycl_device.allocate(in.size() * sizeof(SCALAR))); \ + bool *gpu_data_out = \ + static_cast(sycl_device.allocate(out.size() * sizeof(bool))); \ + TensorMap> gpu(gpu_data, tensorRange); \ + TensorMap> 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;