mirror/eigen
2
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2024-12-21 07:19:46 +08:00
Code Issues Packages Projects Releases Wiki Activity

Merged in codeplaysoftware/eigen-upstream-pure/SYCL-required-changes (pull request PR-454)

Browse Source 
SYCL required changes
This commit is contained in:
Benoit Steiner 2018-08-03 16:17:30 +00:00
commit cf17794ef4
6 changed files with 51 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
Eigen/Core
View File

@ -200,6 +200,12 @@ using std::ptrdiff_t;
#include "src/Core/arch/GPU/MathFunctions.h"
#endif
#if defined EIGEN_VECTORIZE_SYCL
#include "src/Core/arch/SYCL/InteropHeaders.h"
#include "src/Core/arch/SYCL/PacketMath.h"
#include "src/Core/arch/SYCL/MathFunctions.h"
#include "src/Core/arch/SYCL/TypeCasting.h"
#endif
#include "src/Core/arch/Default/Settings.h"
#include "src/Core/functors/TernaryFunctors.h"

7
Eigen/src/Core/arch/GPU/Half.h
View File

@ -83,7 +83,11 @@ struct __half_raw {
#if defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER < 90000
// In CUDA < 9.0, __half is the equivalent of CUDA 9's __half_raw
typedef __half __half_raw;
#endif
#endif // defined(EIGEN_HAS_CUDA_FP16)
#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
typedef cl::sycl::half __half_raw;
#endif
EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw raw_uint16_to_half(unsigned short x);
@ -200,6 +204,7 @@ struct half : public half_impl::half_base {
x = other.x;
return *this;
}
};
} // end namespace Eigen

1
unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h
View File

@ -21,6 +21,7 @@ namespace Eigen {
template<typename T> struct MakePointer {
typedef T* Type;
typedef T& RefType;
typedef T ScalarType;
};
namespace internal{

4
unsupported/Eigen/CXX11/src/Tensor/TensorTrace.h
View File

@ -273,11 +273,11 @@ struct TensorEvaluator<const TensorTraceOp<Dims, ArgType>, Device>
Dimensions m_dimensions;
TensorEvaluator<ArgType, Device> m_impl;
// Initialize the size of the trace dimension
Index m_traceDim;
const Device& m_device;
array<bool, NumInputDims> m_reduced;
array<Index, NumReducedDims> m_reducedDims;
// Initialize the size of the trace dimension
Index m_traceDim;
array<Index, NumOutputDims> m_outputStrides;
array<Index, NumReducedDims> m_reducedStrides;
array<Index, NumOutputDims> m_preservedStrides;

4
unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h
View File

@ -59,6 +59,7 @@ struct traits<Tensor<Scalar_, NumIndices_, Options_, IndexType_> >
template <typename T> struct MakePointer {
typedef T* Type;
typedef T& RefType;
typedef T ScalarType;
};
typedef typename MakePointer<Scalar>::Type PointerType;
@ -80,6 +81,7 @@ struct traits<TensorFixedSize<Scalar_, Dimensions, Options_, IndexType_> >
template <typename T> struct MakePointer {
typedef T* Type;
typedef T& RefType;
typedef T ScalarType;
};
typedef typename MakePointer<Scalar>::Type PointerType;
@ -105,6 +107,8 @@ struct traits<TensorMap<PlainObjectType, Options_, MakePointer_> >
typedef MakePointer_<T> MakePointerT;
typedef typename MakePointerT::Type Type;
typedef typename MakePointerT::RefType RefType;
typedef typename MakePointerT::ScalarType ScalarType;
};
typedef typename MakePointer<Scalar>::Type PointerType;

36
unsupported/Eigen/src/SpecialFunctions/SpecialFunctionsImpl.h
View File

@ -193,6 +193,8 @@ struct lgamma_impl<float> {
#if !defined(EIGEN_GPU_COMPILE_PHASE) && (defined(_BSD_SOURCE) || defined(_SVID_SOURCE)) && !defined(__APPLE__)
int dummy;
return ::lgammaf_r(x, &dummy);
#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
return cl::sycl::lgamma(x);
#else
return ::lgammaf(x);
#endif
@ -206,6 +208,8 @@ struct lgamma_impl<double> {
#if !defined(EIGEN_GPU_COMPILE_PHASE) && (defined(_BSD_SOURCE) || defined(_SVID_SOURCE)) && !defined(__APPLE__)
int dummy;
return ::lgamma_r(x, &dummy);
#elif defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
return cl::sycl::lgamma(x);
#else
return ::lgamma(x);
#endif
@ -423,13 +427,25 @@ struct erf_retval {
template <>
struct erf_impl<float> {
EIGEN_DEVICE_FUNC
static EIGEN_STRONG_INLINE float run(float x) { return ::erff(x); }
static EIGEN_STRONG_INLINE float run(float x) {
#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
return cl::sycl::erf(x);
#else
return ::erff(x);
#endif
}
};
template <>
struct erf_impl<double> {
EIGEN_DEVICE_FUNC
static EIGEN_STRONG_INLINE double run(double x) { return ::erf(x); }
static EIGEN_STRONG_INLINE double run(double x) {
#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
return cl::sycl::erf(x);
#else
return ::erf(x);
#endif
}
};
#endif // EIGEN_HAS_C99_MATH
@ -456,13 +472,25 @@ struct erfc_retval {
template <>
struct erfc_impl<float> {
EIGEN_DEVICE_FUNC
static EIGEN_STRONG_INLINE float run(const float x) { return ::erfcf(x); }
static EIGEN_STRONG_INLINE float run(const float x) {
#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
return cl::sycl::erfc(x);
#else
return ::erfcf(x);
#endif
}
};
template <>
struct erfc_impl<double> {
EIGEN_DEVICE_FUNC
static EIGEN_STRONG_INLINE double run(const double x) { return ::erfc(x); }
static EIGEN_STRONG_INLINE double run(const double x) {
#if defined(EIGEN_USE_SYCL) && defined(__SYCL_DEVICE_ONLY__)
return cl::sycl::erfc(x);
#else
return ::erfc(x);
#endif
}
};
#endif // EIGEN_HAS_C99_MATH