GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-09-23 23:03:15 +08:00
Code Issues Packages Projects Releases Wiki Activity

Added isnan, isfinite and isinf for SYCL device. Plus test for that.

Browse Source
This commit is contained in:
Luke Iwanski 2016-11-18 16:01:48 +00:00
parent 927bd62d2a
commit 5159675c33
3 changed files with 58 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
Eigen/src/Core/MathFunctions.h
View File

@ -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 (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); } 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> template<typename Scalar>
EIGEN_DEVICE_FUNC EIGEN_DEVICE_FUNC
inline EIGEN_MATHFUNC_RETVAL(round, Scalar) round(const Scalar& x) inline EIGEN_MATHFUNC_RETVAL(round, Scalar) round(const Scalar& x)

24
Eigen/src/Core/functors/UnaryFunctors.h
View File

@ -678,7 +678,13 @@ struct functor_traits<scalar_ceil_op<Scalar> >
template<typename Scalar> struct scalar_isnan_op { template<typename Scalar> struct scalar_isnan_op {
EIGEN_EMPTY_STRUCT_CTOR(scalar_isnan_op) EIGEN_EMPTY_STRUCT_CTOR(scalar_isnan_op)
typedef bool result_type; 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> template<typename Scalar>
struct functor_traits<scalar_isnan_op<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 { template<typename Scalar> struct scalar_isinf_op {
EIGEN_EMPTY_STRUCT_CTOR(scalar_isinf_op) EIGEN_EMPTY_STRUCT_CTOR(scalar_isinf_op)
typedef bool result_type; 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> template<typename Scalar>
struct functor_traits<scalar_isinf_op<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 { template<typename Scalar> struct scalar_isfinite_op {
EIGEN_EMPTY_STRUCT_CTOR(scalar_isfinite_op) EIGEN_EMPTY_STRUCT_CTOR(scalar_isfinite_op)
typedef bool result_type; 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> template<typename Scalar>
struct functor_traits<scalar_isfinite_op<Scalar> > struct functor_traits<scalar_isfinite_op<Scalar> >

29
unsupported/test/cxx11_tensor_builtins_sycl.cpp
View File

@ -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(round, SCALAR, OPERATOR) \
TEST_UNARY_BUILTINS_FOR_SCALAR(log1p, 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) \ #define TEST_UNARY_BUILTINS(SCALAR) \
TEST_UNARY_BUILTINS_OPERATOR(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) { static void test_builtin_unary_sycl(const Eigen::SyclDevice &sycl_device) {
int sizeDim1 = 10; int sizeDim1 = 10;