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

Merged in benoitsteiner/opencl (pull request PR-323)

Browse Source 
Improved support for OpenCL
This commit is contained in:
Benoit Steiner 2017-07-07 16:27:33 +00:00
commit c5a241ab9b
34 changed files with 232 additions and 79 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1378,9 +1378,19 @@ T acos(const T &x) {
return acos(x); return acos(x);
} }
template<typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
T acosh(const T &x) {
EIGEN_USING_STD_MATH(acosh);
return acosh(x);
}
#if defined(__SYCL_DEVICE_ONLY__) #if defined(__SYCL_DEVICE_ONLY__)
EIGEN_ALWAYS_INLINE float acos(float x) { return cl::sycl::acos(x); } EIGEN_ALWAYS_INLINE float acos(float x) { return cl::sycl::acos(x); }
EIGEN_ALWAYS_INLINE double acos(double x) { return cl::sycl::acos(x); } EIGEN_ALWAYS_INLINE double acos(double x) { return cl::sycl::acos(x); }
EIGEN_ALWAYS_INLINE float acosh(float x) { return cl::sycl::acosh(x); }
EIGEN_ALWAYS_INLINE double acosh(double x) { return cl::sycl::acosh(x); }
#endif // defined(__SYCL_DEVICE_ONLY__) #endif // defined(__SYCL_DEVICE_ONLY__)
#ifdef __CUDACC__ #ifdef __CUDACC__
@ -1398,9 +1408,18 @@ T asin(const T &x) {
return asin(x); return asin(x);
} }
template<typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
T asinh(const T &x) {
EIGEN_USING_STD_MATH(asinh);
return asinh(x);
}
#if defined(__SYCL_DEVICE_ONLY__) #if defined(__SYCL_DEVICE_ONLY__)
EIGEN_ALWAYS_INLINE float asin(float x) { return cl::sycl::asin(x); } EIGEN_ALWAYS_INLINE float asin(float x) { return cl::sycl::asin(x); }
EIGEN_ALWAYS_INLINE double asin(double x) { return cl::sycl::asin(x); } EIGEN_ALWAYS_INLINE double asin(double x) { return cl::sycl::asin(x); }
EIGEN_ALWAYS_INLINE float asinh(float x) { return cl::sycl::asinh(x); }
EIGEN_ALWAYS_INLINE double asinh(double x) { return cl::sycl::asinh(x); }
#endif // defined(__SYCL_DEVICE_ONLY__) #endif // defined(__SYCL_DEVICE_ONLY__)
#ifdef __CUDACC__ #ifdef __CUDACC__
@ -1418,9 +1437,18 @@ T atan(const T &x) {
return atan(x); return atan(x);
} }
template<typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
T atanh(const T &x) {
EIGEN_USING_STD_MATH(atanh);
return atanh(x);
}
#if defined(__SYCL_DEVICE_ONLY__) #if defined(__SYCL_DEVICE_ONLY__)
EIGEN_ALWAYS_INLINE float atan(float x) { return cl::sycl::atan(x); } EIGEN_ALWAYS_INLINE float atan(float x) { return cl::sycl::atan(x); }
EIGEN_ALWAYS_INLINE double atan(double x) { return cl::sycl::atan(x); } EIGEN_ALWAYS_INLINE double atan(double x) { return cl::sycl::atan(x); }
EIGEN_ALWAYS_INLINE float atanh(float x) { return cl::sycl::atanh(x); }
EIGEN_ALWAYS_INLINE double atanh(double x) { return cl::sycl::atanh(x); }
#endif // defined(__SYCL_DEVICE_ONLY__) #endif // defined(__SYCL_DEVICE_ONLY__)
#ifdef __CUDACC__ #ifdef __CUDACC__

3
unsupported/Eigen/CXX11/src/Tensor/TensorAssign.h
View File

@ -34,6 +34,7 @@ struct traits<TensorAssignOp<LhsXprType, RhsXprType> >
typedef typename remove_reference<RhsNested>::type _RhsNested; typedef typename remove_reference<RhsNested>::type _RhsNested;
static const std::size_t NumDimensions = internal::traits<LhsXprType>::NumDimensions; static const std::size_t NumDimensions = internal::traits<LhsXprType>::NumDimensions;
static const int Layout = internal::traits<LhsXprType>::Layout; static const int Layout = internal::traits<LhsXprType>::Layout;
typedef typename traits<LhsXprType>::PointerType PointerType;
enum { enum {
Flags = 0 Flags = 0
@ -168,7 +169,7 @@ struct TensorEvaluator<const TensorAssignOp<LeftArgType, RightArgType>, Device>
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<RightArgType, Device>& right_impl() const { return m_rightImpl; } const TensorEvaluator<RightArgType, Device>& right_impl() const { return m_rightImpl; }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return m_leftImpl.data(); } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return m_leftImpl.data(); }
private: private:
TensorEvaluator<LeftArgType, Device> m_leftImpl; TensorEvaluator<LeftArgType, Device> m_leftImpl;

3
unsupported/Eigen/CXX11/src/Tensor/TensorBroadcasting.h
View File

@ -31,6 +31,7 @@ struct traits<TensorBroadcastingOp<Broadcast, XprType> > : public traits<XprType
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename Broadcast, typename XprType> template<typename Broadcast, typename XprType>
@ -372,7 +373,7 @@ struct TensorEvaluator<const TensorBroadcastingOp<Broadcast, ArgType>, Device>
TensorOpCost(0, 0, compute_cost, vectorized, PacketSize); TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

3
unsupported/Eigen/CXX11/src/Tensor/TensorChipping.h
View File

@ -32,6 +32,7 @@ struct traits<TensorChippingOp<DimId, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions - 1; static const int NumDimensions = XprTraits::NumDimensions - 1;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<DenseIndex DimId, typename XprType> template<DenseIndex DimId, typename XprType>
@ -264,7 +265,7 @@ struct TensorEvaluator<const TensorChippingOp<DimId, ArgType>, Device>
TensorOpCost(0, 0, cost, vectorized, PacketSize); TensorOpCost(0, 0, cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const {
CoeffReturnType* result = const_cast<CoeffReturnType*>(m_impl.data()); CoeffReturnType* result = const_cast<CoeffReturnType*>(m_impl.data());
if (((static_cast<int>(Layout) == static_cast<int>(ColMajor) && m_dim.actualDim() == NumDims) || if (((static_cast<int>(Layout) == static_cast<int>(ColMajor) && m_dim.actualDim() == NumDims) ||
(static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == 0)) && (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == 0)) &&

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

@ -37,6 +37,8 @@ struct traits<TensorConcatenationOp<Axis, LhsXprType, RhsXprType> >
static const int NumDimensions = traits<LhsXprType>::NumDimensions; static const int NumDimensions = traits<LhsXprType>::NumDimensions;
static const int Layout = traits<LhsXprType>::Layout; static const int Layout = traits<LhsXprType>::Layout;
enum { Flags = 0 }; enum { Flags = 0 };
typedef typename conditional<::Eigen::internal::Pointer_type_promotion<typename LhsXprType::Scalar, Scalar>::val,
typename traits<LhsXprType>::PointerType, typename traits<RhsXprType>::PointerType>::type PointerType;
}; };
template<typename Axis, typename LhsXprType, typename RhsXprType> template<typename Axis, typename LhsXprType, typename RhsXprType>
@ -275,7 +277,7 @@ struct TensorEvaluator<const TensorConcatenationOp<Axis, LeftArgType, RightArgTy
TensorOpCost(0, 0, compute_cost); TensorOpCost(0, 0, compute_cost);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; } const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor

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

@ -104,6 +104,8 @@ struct traits<TensorContractionOp<Dimensions, LhsXprType, RhsXprType> >
// From NumDims below. // From NumDims below.
static const int NumDimensions = traits<RhsXprType>::NumDimensions + traits<RhsXprType>::NumDimensions - 2 * array_size<Dimensions>::value; static const int NumDimensions = traits<RhsXprType>::NumDimensions + traits<RhsXprType>::NumDimensions - 2 * array_size<Dimensions>::value;
static const int Layout = traits<LhsXprType>::Layout; static const int Layout = traits<LhsXprType>::Layout;
typedef typename conditional<::Eigen::internal::Pointer_type_promotion<typename LhsXprType::Scalar, Scalar>::val,
typename traits<LhsXprType>::PointerType, typename traits<RhsXprType>::PointerType>::type PointerType;
enum { enum {
Flags = 0 Flags = 0
@ -609,7 +611,7 @@ struct TensorContractionEvaluatorBase
return internal::ploadt<PacketReturnType, LoadMode>(m_result + index); return internal::ploadt<PacketReturnType, LoadMode>(m_result + index);
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar* data() const { return m_result; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const { return m_result; }
protected: protected:
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void EnableXSMMIfPossible(const array<IndexPair<Index>, ContractDims>& eval_op_indices) { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void EnableXSMMIfPossible(const array<IndexPair<Index>, ContractDims>& eval_op_indices) {

3
unsupported/Eigen/CXX11/src/Tensor/TensorConversion.h
View File

@ -32,6 +32,7 @@ struct traits<TensorConversionOp<TargetType, XprType> >
static const int NumDimensions = traits<XprType>::NumDimensions; static const int NumDimensions = traits<XprType>::NumDimensions;
static const int Layout = traits<XprType>::Layout; static const int Layout = traits<XprType>::Layout;
enum { Flags = 0 }; enum { Flags = 0 };
typedef typename ::Eigen::internal::TypeConversion<Scalar, typename traits<XprType>::PointerType>::type PointerType;
}; };
template<typename TargetType, typename XprType> template<typename TargetType, typename XprType>
@ -244,7 +245,7 @@ struct TensorEvaluator<const TensorConversionOp<TargetType, ArgType>, Device>
} }
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the sycl accessor /// required by sycl in order to extract the sycl accessor
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

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

@ -231,6 +231,8 @@ struct traits<TensorConvolutionOp<Dimensions, InputXprType, KernelXprType> >
typedef typename remove_reference<RhsNested>::type _RhsNested; typedef typename remove_reference<RhsNested>::type _RhsNested;
static const int NumDimensions = traits<InputXprType>::NumDimensions; static const int NumDimensions = traits<InputXprType>::NumDimensions;
static const int Layout = traits<InputXprType>::Layout; static const int Layout = traits<InputXprType>::Layout;
typedef typename conditional<::Eigen::internal::Pointer_type_promotion<typename InputXprType::Scalar, Scalar>::val,
typename traits<InputXprType>::PointerType, typename traits<KernelXprType>::PointerType>::type PointerType;
enum { enum {
Flags = 0 Flags = 0
@ -465,7 +467,7 @@ struct TensorEvaluator<const TensorConvolutionOp<Indices, InputArgType, KernelAr
PacketSize)); PacketSize));
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
private: private:
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index firstInput(Index index) const { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index firstInput(Index index) const {

2
unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h
View File

@ -300,7 +300,7 @@ struct TensorEvaluator<const TensorConvolutionOp<Indices, InputArgType, KernelAr
/// used by sycl in order to build the sycl buffer /// used by sycl in order to build the sycl buffer
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const{return m_device;} EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const{return m_device;}
/// used by sycl in order to build the sycl buffer /// used by sycl in order to build the sycl buffer
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const { return m_buf; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const { return m_buf; }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void preloadKernel() { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void preloadKernel() {
// Don't make a local copy of the kernel unless we have to (i.e. it's an // Don't make a local copy of the kernel unless we have to (i.e. it's an

7
unsupported/Eigen/CXX11/src/Tensor/TensorCustomOp.h
View File

@ -30,6 +30,7 @@ struct traits<TensorCustomUnaryOp<CustomUnaryFunc, XprType> >
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = traits<XprType>::NumDimensions; static const int NumDimensions = traits<XprType>::NumDimensions;
static const int Layout = traits<XprType>::Layout; static const int Layout = traits<XprType>::Layout;
typedef typename traits<XprType>::PointerType PointerType;
}; };
template<typename CustomUnaryFunc, typename XprType> template<typename CustomUnaryFunc, typename XprType>
@ -138,7 +139,7 @@ struct TensorEvaluator<const TensorCustomUnaryOp<CustomUnaryFunc, XprType>, Devi
return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize); return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return m_result; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return m_result; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const { return m_device; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const { return m_device; }
@ -184,6 +185,8 @@ struct traits<TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType, RhsXprType> >
typedef typename remove_reference<RhsNested>::type _RhsNested; typedef typename remove_reference<RhsNested>::type _RhsNested;
static const int NumDimensions = traits<LhsXprType>::NumDimensions; static const int NumDimensions = traits<LhsXprType>::NumDimensions;
static const int Layout = traits<LhsXprType>::Layout; static const int Layout = traits<LhsXprType>::Layout;
typedef typename conditional<::Eigen::internal::Pointer_type_promotion<typename LhsXprType::Scalar, Scalar>::val,
typename traits<LhsXprType>::PointerType, typename traits<RhsXprType>::PointerType>::type PointerType;
}; };
template<typename CustomBinaryFunc, typename LhsXprType, typename RhsXprType> template<typename CustomBinaryFunc, typename LhsXprType, typename RhsXprType>
@ -297,7 +300,7 @@ struct TensorEvaluator<const TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType,
return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize); return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return m_result; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return m_result; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const { return m_device; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Device& device() const { return m_device; }

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

@ -32,6 +32,7 @@ struct traits<TensorEvalToOp<XprType, MakePointer_> >
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename MakePointer_<Scalar>::Type PointerType;
enum { enum {
Flags = 0 Flags = 0

10
unsupported/Eigen/CXX11/src/Tensor/TensorEvaluator.h
View File

@ -283,7 +283,7 @@ struct TensorEvaluator<const TensorCwiseNullaryOp<NullaryOp, ArgType>, Device>
internal::unpacket_traits<PacketReturnType>::size); internal::unpacket_traits<PacketReturnType>::size);
} }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<ArgType, Device>& impl() const { return m_argImpl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_argImpl; }
@ -353,7 +353,7 @@ struct TensorEvaluator<const TensorCwiseUnaryOp<UnaryOp, ArgType>, Device>
TensorOpCost(0, 0, functor_cost, vectorized, PacketSize); TensorOpCost(0, 0, functor_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<ArgType, Device> & impl() const { return m_argImpl; } const TensorEvaluator<ArgType, Device> & impl() const { return m_argImpl; }
@ -433,7 +433,7 @@ struct TensorEvaluator<const TensorCwiseBinaryOp<BinaryOp, LeftArgType, RightArg
TensorOpCost(0, 0, functor_cost, vectorized, PacketSize); TensorOpCost(0, 0, functor_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; } const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
@ -533,7 +533,7 @@ struct TensorEvaluator<const TensorCwiseTernaryOp<TernaryOp, Arg1Type, Arg2Type,
TensorOpCost(0, 0, functor_cost, vectorized, PacketSize); TensorOpCost(0, 0, functor_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<Arg1Type, Device> & arg1Impl() const { return m_arg1Impl; } const TensorEvaluator<Arg1Type, Device> & arg1Impl() const { return m_arg1Impl; }
@ -625,7 +625,7 @@ struct TensorEvaluator<const TensorSelectOp<IfArgType, ThenArgType, ElseArgType>
.cwiseMax(m_elseImpl.costPerCoeff(vectorized)); .cwiseMax(m_elseImpl.costPerCoeff(vectorized));
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const { return NULL; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<IfArgType, Device> & cond_impl() const { return m_condImpl; } const TensorEvaluator<IfArgType, Device> & cond_impl() const { return m_condImpl; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor

18
unsupported/Eigen/CXX11/src/Tensor/TensorExpr.h
View File

@ -38,7 +38,7 @@ struct traits<TensorCwiseNullaryOp<NullaryOp, XprType> >
typedef typename remove_reference<XprTypeNested>::type _XprTypeNested; typedef typename remove_reference<XprTypeNested>::type _XprTypeNested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
enum { enum {
Flags = 0 Flags = 0
}; };
@ -89,6 +89,7 @@ struct traits<TensorCwiseUnaryOp<UnaryOp, XprType> >
typedef typename remove_reference<XprTypeNested>::type _XprTypeNested; typedef typename remove_reference<XprTypeNested>::type _XprTypeNested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename ::Eigen::internal::TypeConversion<Scalar, typename XprTraits::PointerType>::type PointerType;
}; };
template<typename UnaryOp, typename XprType> template<typename UnaryOp, typename XprType>
@ -161,7 +162,11 @@ struct traits<TensorCwiseBinaryOp<BinaryOp, LhsXprType, RhsXprType> >
typedef typename remove_reference<RhsNested>::type _RhsNested; typedef typename remove_reference<RhsNested>::type _RhsNested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename ::Eigen::internal::TypeConversion<Scalar,
typename conditional<::Eigen::internal::Pointer_type_promotion<typename LhsXprType::Scalar, Scalar>::val,
typename traits<LhsXprType>::PointerType,
typename traits<RhsXprType>::PointerType>::type
>::type PointerType;
enum { enum {
Flags = 0 Flags = 0
}; };
@ -238,7 +243,11 @@ struct traits<TensorCwiseTernaryOp<TernaryOp, Arg1XprType, Arg2XprType, Arg3XprT
typedef typename remove_reference<Arg3Nested>::type _Arg3Nested; typedef typename remove_reference<Arg3Nested>::type _Arg3Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename ::Eigen::internal::TypeConversion<Scalar,
typename conditional<::Eigen::internal::Pointer_type_promotion<typename Arg2XprType::Scalar, Scalar>::val,
typename traits<Arg2XprType>::PointerType,
typename traits<Arg3XprType>::PointerType>::type
>::type PointerType;
enum { enum {
Flags = 0 Flags = 0
}; };
@ -314,6 +323,9 @@ struct traits<TensorSelectOp<IfXprType, ThenXprType, ElseXprType> >
typedef typename ElseXprType::Nested ElseNested; typedef typename ElseXprType::Nested ElseNested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename conditional<::Eigen::internal::Pointer_type_promotion<typename ThenXprType::Scalar, Scalar>::val,
typename traits<ThenXprType>::PointerType,
typename traits<ElseXprType>::PointerType>::type PointerType;
}; };
template<typename IfXprType, typename ThenXprType, typename ElseXprType> template<typename IfXprType, typename ThenXprType, typename ElseXprType>

3
unsupported/Eigen/CXX11/src/Tensor/TensorFFT.h
View File

@ -71,6 +71,7 @@ struct traits<TensorFFTOp<FFT, XprType, FFTResultType, FFTDir> > : public traits
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename traits<XprType>::PointerType PointerType;
}; };
template <typename FFT, typename XprType, int FFTResultType, int FFTDirection> template <typename FFT, typename XprType, int FFTResultType, int FFTDirection>
@ -234,7 +235,7 @@ struct TensorEvaluator<const TensorFFTOp<FFT, ArgType, FFTResultType, FFTDir>, D
if (line_len > 1) { if (line_len > 1) {
const RealScalar pi_over_len(EIGEN_PI / line_len); const RealScalar pi_over_len(EIGEN_PI / line_len);
const ComplexScalar pos_j_base = ComplexScalar( const ComplexScalar pos_j_base = ComplexScalar(
std::cos(pi_over_len), std::sin(pi_over_len)); std::cos(pi_over_len), std::sin(pi_over_len));
pos_j_base_powered[1] = pos_j_base; pos_j_base_powered[1] = pos_j_base;
if (line_len > 2) { if (line_len > 2) {
const ComplexScalar pos_j_base_sq = pos_j_base * pos_j_base; const ComplexScalar pos_j_base_sq = pos_j_base * pos_j_base;

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

@ -38,6 +38,7 @@ struct traits<TensorForcedEvalOp<XprType> >
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
enum { enum {
Flags = 0 Flags = 0
@ -143,7 +144,8 @@ struct TensorEvaluator<const TensorForcedEvalOp<ArgType>, Device>
return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize); return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const { return m_buffer; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename Eigen::internal::traits<XprType>::PointerType data() const { return m_buffer; }
/// required by sycl in order to extract the sycl accessor /// required by sycl in order to extract the sycl accessor
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() { return m_impl; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() { return m_impl; }

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

@ -22,6 +22,22 @@ template<typename T> struct MakePointer {
typedef T* Type; typedef T* Type;
typedef T& RefType; typedef T& RefType;
}; };
namespace internal{
template<typename A, typename B> struct Pointer_type_promotion {
static const bool val=false;
};
template<typename A> struct Pointer_type_promotion<A, A> {
static const bool val = true;
};
template<typename A, typename B> struct TypeConversion;
#ifndef __SYCL_DEVICE_ONLY__
template<typename A, typename B> struct TypeConversion{
typedef A* type;
};
#endif
}
#if defined(EIGEN_USE_SYCL) #if defined(EIGEN_USE_SYCL)
namespace TensorSycl { namespace TensorSycl {
namespace internal{ namespace internal{

3
unsupported/Eigen/CXX11/src/Tensor/TensorGenerator.h
View File

@ -31,6 +31,7 @@ struct traits<TensorGeneratorOp<Generator, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename Generator, typename XprType> template<typename Generator, typename XprType>
@ -156,7 +157,7 @@ struct TensorEvaluator<const TensorGeneratorOp<Generator, ArgType>, Device>
TensorOpCost::MulCost<Scalar>()); TensorOpCost::MulCost<Scalar>());
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_argImpl; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_argImpl; }

63
unsupported/Eigen/CXX11/src/Tensor/TensorImagePatch.h
View File

@ -27,6 +27,7 @@ namespace Eigen {
* patch_cols, and 1 for all the additional dimensions. * patch_cols, and 1 for all the additional dimensions.
*/ */
namespace internal { namespace internal {
template<DenseIndex Rows, DenseIndex Cols, typename XprType> template<DenseIndex Rows, DenseIndex Cols, typename XprType>
struct traits<TensorImagePatchOp<Rows, Cols, XprType> > : public traits<XprType> struct traits<TensorImagePatchOp<Rows, Cols, XprType> > : public traits<XprType>
{ {
@ -38,6 +39,7 @@ struct traits<TensorImagePatchOp<Rows, Cols, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions + 1; static const int NumDimensions = XprTraits::NumDimensions + 1;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<DenseIndex Rows, DenseIndex Cols, typename XprType> template<DenseIndex Rows, DenseIndex Cols, typename XprType>
@ -70,12 +72,12 @@ class TensorImagePatchOp : public TensorBase<TensorImagePatchOp<Rows, Cols, XprT
DenseIndex in_row_strides, DenseIndex in_col_strides, DenseIndex in_row_strides, DenseIndex in_col_strides,
DenseIndex row_inflate_strides, DenseIndex col_inflate_strides, DenseIndex row_inflate_strides, DenseIndex col_inflate_strides,
PaddingType padding_type, Scalar padding_value) PaddingType padding_type, Scalar padding_value)
: m_xpr(expr), m_patch_rows(patch_rows), m_patch_cols(patch_cols), : m_xpr(expr), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_row_strides(row_strides), m_col_strides(col_strides), m_row_strides(row_strides), m_col_strides(col_strides),
m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides), m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides),
m_padding_explicit(false), m_padding_top(0), m_padding_bottom(0), m_padding_left(0), m_padding_right(0), m_padding_explicit(false), m_padding_top(0), m_padding_bottom(0), m_padding_left(0), m_padding_right(0),
m_padding_type(padding_type), m_padding_value(padding_value) {} m_padding_type(padding_type), m_padding_value(padding_value) {}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorImagePatchOp(const XprType& expr, DenseIndex patch_rows, DenseIndex patch_cols, EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorImagePatchOp(const XprType& expr, DenseIndex patch_rows, DenseIndex patch_cols,
DenseIndex row_strides, DenseIndex col_strides, DenseIndex row_strides, DenseIndex col_strides,
@ -84,13 +86,31 @@ class TensorImagePatchOp : public TensorBase<TensorImagePatchOp<Rows, Cols, XprT
DenseIndex padding_top, DenseIndex padding_bottom, DenseIndex padding_top, DenseIndex padding_bottom,
DenseIndex padding_left, DenseIndex padding_right, DenseIndex padding_left, DenseIndex padding_right,
Scalar padding_value) Scalar padding_value)
: m_xpr(expr), m_patch_rows(patch_rows), m_patch_cols(patch_cols), : m_xpr(expr), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_row_strides(row_strides), m_col_strides(col_strides), m_row_strides(row_strides), m_col_strides(col_strides),
m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides), m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides),
m_padding_explicit(true), m_padding_top(padding_top), m_padding_bottom(padding_bottom), m_padding_explicit(true), m_padding_top(padding_top), m_padding_bottom(padding_bottom),
m_padding_left(padding_left), m_padding_right(padding_right), m_padding_left(padding_left), m_padding_right(padding_right),
m_padding_type(PADDING_VALID), m_padding_value(padding_value) {} m_padding_type(PADDING_VALID), m_padding_value(padding_value) {}
#ifdef EIGEN_USE_SYCL // this is work around for sycl as Eigen could not use c++11 deligate constructor feature
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorImagePatchOp(const XprType& expr, DenseIndex patch_rows, DenseIndex patch_cols,
DenseIndex row_strides, DenseIndex col_strides,
DenseIndex in_row_strides, DenseIndex in_col_strides,
DenseIndex row_inflate_strides, DenseIndex col_inflate_strides,
bool padding_explicit, DenseIndex padding_top, DenseIndex padding_bottom,
DenseIndex padding_left, DenseIndex padding_right, PaddingType padding_type,
Scalar padding_value)
: m_xpr(expr), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_row_strides(row_strides), m_col_strides(col_strides),
m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides),
m_padding_explicit(padding_explicit), m_padding_top(padding_top), m_padding_bottom(padding_bottom),
m_padding_left(padding_left), m_padding_right(padding_right),
m_padding_type(padding_type), m_padding_value(padding_value) {}
#endif
EIGEN_DEVICE_FUNC EIGEN_DEVICE_FUNC
DenseIndex patch_rows() const { return m_patch_rows; } DenseIndex patch_rows() const { return m_patch_rows; }
@ -171,7 +191,11 @@ struct TensorEvaluator<const TensorImagePatchOp<Rows, Cols, ArgType>, Device>
RawAccess = false RawAccess = false
}; };
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device) #ifdef __SYCL_DEVICE_ONLY__
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator( const XprType op, const Device& device)
#else
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator( const XprType& op, const Device& device)
#endif
: m_impl(op.expression(), device) : m_impl(op.expression(), device)
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
, m_op(op) , m_op(op)
@ -423,11 +447,12 @@ struct TensorEvaluator<const TensorImagePatchOp<Rows, Cols, ArgType>, Device>
return packetWithPossibleZero(index); return packetWithPossibleZero(index);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
// Required by SYCL in order to construct the expression tree on the device
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& xpr() const { return m_op; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& xpr() const { return m_op; }
#endif #endif
@ -510,10 +535,10 @@ struct TensorEvaluator<const TensorImagePatchOp<Rows, Cols, ArgType>, Device>
Scalar m_paddingValue; Scalar m_paddingValue;
TensorEvaluator<ArgType, Device> m_impl; TensorEvaluator<ArgType, Device> m_impl;
#ifdef EIGEN_USE_SYCL
#ifdef EIGEN_USE_SYCL // Required for SYCL in order to construct the expression tree on the device
const XprType& m_op; XprType m_op;
#endif #endif
}; };

3
unsupported/Eigen/CXX11/src/Tensor/TensorInflation.h
View File

@ -31,6 +31,7 @@ struct traits<TensorInflationOp<Strides, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename Strides, typename XprType> template<typename Strides, typename XprType>
@ -213,7 +214,7 @@ struct TensorEvaluator<const TensorInflationOp<Strides, ArgType>, Device>
compute_cost, vectorized, PacketSize); compute_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

3
unsupported/Eigen/CXX11/src/Tensor/TensorLayoutSwap.h
View File

@ -46,6 +46,7 @@ struct traits<TensorLayoutSwapOp<XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = traits<XprType>::NumDimensions; static const int NumDimensions = traits<XprType>::NumDimensions;
static const int Layout = (traits<XprType>::Layout == ColMajor) ? RowMajor : ColMajor; static const int Layout = (traits<XprType>::Layout == ColMajor) ? RowMajor : ColMajor;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename XprType> template<typename XprType>
@ -159,7 +160,7 @@ struct TensorEvaluator<const TensorLayoutSwapOp<ArgType>, Device>
return m_impl.costPerCoeff(vectorized); return m_impl.costPerCoeff(vectorized);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return m_impl.data(); } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return m_impl.data(); }
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

9
unsupported/Eigen/CXX11/src/Tensor/TensorMorphing.h
View File

@ -31,6 +31,7 @@ struct traits<TensorReshapingOp<NewDimensions, XprType> > : public traits<XprTyp
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = array_size<NewDimensions>::value; static const int NumDimensions = array_size<NewDimensions>::value;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename NewDimensions, typename XprType> template<typename NewDimensions, typename XprType>
@ -146,7 +147,7 @@ struct TensorEvaluator<const TensorReshapingOp<NewDimensions, ArgType>, Device>
return m_impl.costPerCoeff(vectorized); return m_impl.costPerCoeff(vectorized);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return const_cast<Scalar*>(m_impl.data()); } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return const_cast<Scalar*>(m_impl.data()); }
EIGEN_DEVICE_FUNC const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } EIGEN_DEVICE_FUNC const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
@ -214,6 +215,7 @@ struct traits<TensorSlicingOp<StartIndices, Sizes, XprType> > : public traits<Xp
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = array_size<StartIndices>::value; static const int NumDimensions = array_size<StartIndices>::value;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename StartIndices, typename Sizes, typename XprType> template<typename StartIndices, typename Sizes, typename XprType>
@ -468,7 +470,7 @@ struct TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, ArgType>, Devi
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar* data() const { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const {
Scalar* result = m_impl.data(); Scalar* result = m_impl.data();
if (result) { if (result) {
Index offset = 0; Index offset = 0;
@ -633,6 +635,7 @@ struct traits<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprTyp
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = array_size<StartIndices>::value; static const int NumDimensions = array_size<StartIndices>::value;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename StartIndices, typename StopIndices, typename Strides, typename XprType> template<typename StartIndices, typename StopIndices, typename Strides, typename XprType>
@ -823,7 +826,7 @@ struct TensorEvaluator<const TensorStridingSlicingOp<StartIndices, StopIndices,
return m_impl.costPerCoeff(vectorized) + TensorOpCost(0, 0, NumDims); return m_impl.costPerCoeff(vectorized) + TensorOpCost(0, 0, NumDims);
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar* data() const { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const {
return NULL; return NULL;
} }

3
unsupported/Eigen/CXX11/src/Tensor/TensorPadding.h
View File

@ -31,6 +31,7 @@ struct traits<TensorPaddingOp<PaddingDimensions, XprType> > : public traits<XprT
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename PaddingDimensions, typename XprType> template<typename PaddingDimensions, typename XprType>
@ -198,7 +199,7 @@ struct TensorEvaluator<const TensorPaddingOp<PaddingDimensions, ArgType>, Device
return cost; return cost;
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// used by sycl /// used by sycl
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const PaddingDimensions& padding() const { return m_padding; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const PaddingDimensions& padding() const { return m_padding; }

3
unsupported/Eigen/CXX11/src/Tensor/TensorPatch.h
View File

@ -31,6 +31,7 @@ struct traits<TensorPatchOp<PatchDim, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions + 1; static const int NumDimensions = XprTraits::NumDimensions + 1;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename PatchDim, typename XprType> template<typename PatchDim, typename XprType>
@ -256,7 +257,7 @@ struct TensorEvaluator<const TensorPatchOp<PatchDim, ArgType>, Device>
TensorOpCost(0, 0, compute_cost, vectorized, PacketSize); TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

3
unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h
View File

@ -44,6 +44,7 @@ namespace internal {
typedef typename XprType::Nested Nested; typedef typename XprType::Nested Nested;
static const int NumDimensions = XprTraits::NumDimensions - array_size<Dims>::value; static const int NumDimensions = XprTraits::NumDimensions - array_size<Dims>::value;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
template <class T> struct MakePointer { template <class T> struct MakePointer {
// Intermediate typedef to workaround MSVC issue. // Intermediate typedef to workaround MSVC issue.
@ -677,7 +678,7 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType, MakePointer_>,
} }
} }
EIGEN_DEVICE_FUNC typename MakePointer_<Scalar>::Type data() const { return m_result; } EIGEN_DEVICE_FUNC typename MakePointer_<CoeffReturnType>::Type data() const { return m_result; }
#if defined(EIGEN_USE_SYCL) #if defined(EIGEN_USE_SYCL)
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

3
unsupported/Eigen/CXX11/src/Tensor/TensorReverse.h
View File

@ -31,6 +31,7 @@ struct traits<TensorReverseOp<ReverseDimensions,
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename ReverseDimensions, typename XprType> template<typename ReverseDimensions, typename XprType>
@ -222,7 +223,7 @@ struct TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device
TensorOpCost(0, 0, compute_cost, false /* vectorized */, PacketSize); TensorOpCost(0, 0, compute_cost, false /* vectorized */, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<ArgType, Device> & impl() const { return m_impl; } const TensorEvaluator<ArgType, Device> & impl() const { return m_impl; }

3
unsupported/Eigen/CXX11/src/Tensor/TensorScan.h
View File

@ -24,6 +24,7 @@ struct traits<TensorScanOp<Op, XprType> >
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename Op, typename XprType> template<typename Op, typename XprType>
@ -175,7 +176,7 @@ struct TensorEvaluator<const TensorScanOp<Op, ArgType>, Device> {
return internal::ploadt<PacketReturnType, LoadMode>(m_output + index); return internal::ploadt<PacketReturnType, LoadMode>(m_output + index);
} }
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Eigen::internal::traits<XprType>::PointerType data() const
{ {
return m_output; return m_output;
} }

3
unsupported/Eigen/CXX11/src/Tensor/TensorShuffling.h
View File

@ -31,6 +31,7 @@ struct traits<TensorShufflingOp<Shuffle, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename Shuffle, typename XprType> template<typename Shuffle, typename XprType>
@ -185,7 +186,7 @@ struct TensorEvaluator<const TensorShufflingOp<Shuffle, ArgType>, Device>
TensorOpCost(0, 0, compute_cost, false /* vectorized */, PacketSize); TensorOpCost(0, 0, compute_cost, false /* vectorized */, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
// required by sycl // required by sycl
EIGEN_STRONG_INLINE const Shuffle& shufflePermutation() const {return m_shuffle;} EIGEN_STRONG_INLINE const Shuffle& shufflePermutation() const {return m_shuffle;}

3
unsupported/Eigen/CXX11/src/Tensor/TensorStriding.h
View File

@ -31,6 +31,7 @@ struct traits<TensorStridingOp<Strides, XprType> > : public traits<XprType>
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions; static const int NumDimensions = XprTraits::NumDimensions;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<typename Strides, typename XprType> template<typename Strides, typename XprType>
@ -222,7 +223,7 @@ struct TensorEvaluator<const TensorStridingOp<Strides, ArgType>, Device>
TensorOpCost(0, 0, compute_cost, vectorized, PacketSize); TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
/// required by sycl in order to extract the accessor /// required by sycl in order to extract the accessor
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }

24
unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h
View File

@ -34,12 +34,26 @@ struct MakeLocalPointer {
namespace Eigen { namespace Eigen {
template<typename StrideDims, typename XprType> class TensorTupleReducerDeviceOp; template<typename StrideDims, typename XprType> class TensorTupleReducerDeviceOp;
template<typename StrideDims, typename ArgType> struct TensorEvaluator<const TensorTupleReducerDeviceOp<StrideDims, ArgType>, SyclKernelDevice>; template<typename StrideDims, typename ArgType> struct TensorEvaluator<const TensorTupleReducerDeviceOp<StrideDims, ArgType>, SyclKernelDevice>;
namespace internal {
#ifdef __SYCL_DEVICE_ONLY__
template<typename A, typename B> struct TypeConversion {
template<typename T>
static typename MakeGlobalPointer<A>::Type get_address_space_pointer(typename MakeGlobalPointer<T>::Type p);
template<typename T>
static typename MakeLocalPointer<A>::Type get_address_space_pointer(typename MakeLocalPointer<T>::Type p);
template<typename T>
static A* get_address_space_pointer(T* p);
typedef decltype(get_address_space_pointer(B())) type;
};
#endif
}
namespace TensorSycl { namespace TensorSycl {
namespace internal { namespace internal {
template<typename CoeffReturnType, typename OP, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer; template<typename CoeffReturnType, typename OP, typename OutputAccessor, typename InputAccessor, typename LocalAccessor> struct GenericKernelReducer;
/// This struct is used for special expression nodes with no operations (for example assign and selectOP). /// This struct is used for special expression nodes with no operations (for example assign and selectOP).
struct NoOP; struct NoOP;
@ -51,10 +65,10 @@ template<typename T> struct GetType<false, T>{
}; };
template <bool Conds, size_t X , size_t Y > struct ValueCondition { template <bool Conds, size_t X , size_t Y > struct ValueCondition {
static const size_t Res =X; static constexpr size_t Res =X;
}; };
template<size_t X, size_t Y> struct ValueCondition<false, X , Y> { template<size_t X, size_t Y> struct ValueCondition<false, X, Y> {
static const size_t Res =Y; static constexpr size_t Res =Y;
}; };
} }

8
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h
View File

@ -448,8 +448,8 @@ struct ExprConstructor<CVQual TensorImagePatchOp<Rows, Cols, OrigXprType>, CVQu
template <typename FuncDetector>\ template <typename FuncDetector>\
ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.m_patch_rows, funcD.m_patch_cols, funcD.m_row_strides, funcD.m_col_strides,\ : xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.m_patch_rows, funcD.m_patch_cols, funcD.m_row_strides, funcD.m_col_strides,\
funcD.m_in_row_strides, funcD.m_in_col_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, \ funcD.m_in_row_strides, funcD.m_in_col_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, funcD.m_padding_explicit, \
funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, funcD.m_padding_value, funcD.m_padding_type, funcD.m_padding_explicit){}\ funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, funcD.m_padding_type, funcD.m_padding_value){}\
}; };
SYCLTENSORIMAGEPATCHOPEXPR(const) SYCLTENSORIMAGEPATCHOPEXPR(const)
@ -468,8 +468,8 @@ struct ExprConstructor<CVQual TensorVolumePatchOp<Planes, Rows, Cols, OrigXprTy
ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.m_patch_planes, funcD.m_patch_rows, funcD.m_patch_cols, funcD.m_plane_strides, funcD.m_row_strides, funcD.m_col_strides,\ : xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.m_patch_planes, funcD.m_patch_rows, funcD.m_patch_cols, funcD.m_plane_strides, funcD.m_row_strides, funcD.m_col_strides,\
funcD.m_in_plane_strides, funcD.m_in_row_strides, funcD.m_in_col_strides,funcD.m_plane_inflate_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, \ funcD.m_in_plane_strides, funcD.m_in_row_strides, funcD.m_in_col_strides,funcD.m_plane_inflate_strides, funcD.m_row_inflate_strides, funcD.m_col_inflate_strides, \
funcD.m_padding_top_z, funcD.m_padding_bottom_z, funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, funcD.m_padding_value,\ funcD.m_padding_explicit, funcD.m_padding_top_z, funcD.m_padding_bottom_z, funcD.m_padding_top, funcD.m_padding_bottom, funcD.m_padding_left, funcD.m_padding_right, \
funcD.m_padding_type, funcD.m_padding_explicit){\ funcD.m_padding_type, funcD.m_padding_value ){\
}\ }\
}; };

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

@ -61,6 +61,7 @@ struct traits<Tensor<Scalar_, NumIndices_, Options_, IndexType_> >
typedef T& RefType; typedef T& RefType;
}; };
typedef typename MakePointer<Scalar>::Type PointerType;
}; };
@ -81,6 +82,7 @@ struct traits<TensorFixedSize<Scalar_, Dimensions, Options_, IndexType_> >
typedef T& RefType; typedef T& RefType;
}; };
typedef typename MakePointer<Scalar>::Type PointerType;
}; };
@ -105,6 +107,7 @@ struct traits<TensorMap<PlainObjectType, Options_, MakePointer_> >
typedef typename MakePointerT::RefType RefType; typedef typename MakePointerT::RefType RefType;
}; };
typedef typename MakePointer<Scalar>::Type PointerType;
}; };
template<typename PlainObjectType> template<typename PlainObjectType>
@ -121,6 +124,7 @@ struct traits<TensorRef<PlainObjectType> >
Options = BaseTraits::Options, Options = BaseTraits::Options,
Flags = BaseTraits::Flags Flags = BaseTraits::Flags
}; };
typedef typename BaseTraits::PointerType PointerType;
}; };

52
unsupported/Eigen/CXX11/src/Tensor/TensorVolumePatch.h
View File

@ -22,6 +22,7 @@ namespace Eigen {
* dimensions. * dimensions.
*/ */
namespace internal { namespace internal {
template<DenseIndex Planes, DenseIndex Rows, DenseIndex Cols, typename XprType> template<DenseIndex Planes, DenseIndex Rows, DenseIndex Cols, typename XprType>
struct traits<TensorVolumePatchOp<Planes, Rows, Cols, XprType> > : public traits<XprType> struct traits<TensorVolumePatchOp<Planes, Rows, Cols, XprType> > : public traits<XprType>
{ {
@ -33,6 +34,8 @@ struct traits<TensorVolumePatchOp<Planes, Rows, Cols, XprType> > : public traits
typedef typename remove_reference<Nested>::type _Nested; typedef typename remove_reference<Nested>::type _Nested;
static const int NumDimensions = XprTraits::NumDimensions + 1; static const int NumDimensions = XprTraits::NumDimensions + 1;
static const int Layout = XprTraits::Layout; static const int Layout = XprTraits::Layout;
typedef typename XprTraits::PointerType PointerType;
}; };
template<DenseIndex Planes, DenseIndex Rows, DenseIndex Cols, typename XprType> template<DenseIndex Planes, DenseIndex Rows, DenseIndex Cols, typename XprType>
@ -65,12 +68,12 @@ class TensorVolumePatchOp : public TensorBase<TensorVolumePatchOp<Planes, Rows,
DenseIndex in_plane_strides, DenseIndex in_row_strides, DenseIndex in_col_strides, DenseIndex in_plane_strides, DenseIndex in_row_strides, DenseIndex in_col_strides,
DenseIndex plane_inflate_strides, DenseIndex row_inflate_strides, DenseIndex col_inflate_strides, DenseIndex plane_inflate_strides, DenseIndex row_inflate_strides, DenseIndex col_inflate_strides,
PaddingType padding_type, Scalar padding_value) PaddingType padding_type, Scalar padding_value)
: m_xpr(expr), m_patch_planes(patch_planes), m_patch_rows(patch_rows), m_patch_cols(patch_cols), : m_xpr(expr), m_patch_planes(patch_planes), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_plane_strides(plane_strides), m_row_strides(row_strides), m_col_strides(col_strides), m_plane_strides(plane_strides), m_row_strides(row_strides), m_col_strides(col_strides),
m_in_plane_strides(in_plane_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides), m_in_plane_strides(in_plane_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_plane_inflate_strides(plane_inflate_strides), m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides), m_plane_inflate_strides(plane_inflate_strides), m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides),
m_padding_explicit(false), m_padding_top_z(0), m_padding_bottom_z(0), m_padding_top(0), m_padding_bottom(0), m_padding_left(0), m_padding_right(0), m_padding_explicit(false), m_padding_top_z(0), m_padding_bottom_z(0), m_padding_top(0), m_padding_bottom(0), m_padding_left(0), m_padding_right(0),
m_padding_type(padding_type), m_padding_value(padding_value) {} m_padding_type(padding_type), m_padding_value(padding_value) {}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorVolumePatchOp(const XprType& expr, DenseIndex patch_planes, DenseIndex patch_rows, DenseIndex patch_cols, EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorVolumePatchOp(const XprType& expr, DenseIndex patch_planes, DenseIndex patch_rows, DenseIndex patch_cols,
DenseIndex plane_strides, DenseIndex row_strides, DenseIndex col_strides, DenseIndex plane_strides, DenseIndex row_strides, DenseIndex col_strides,
@ -80,13 +83,31 @@ class TensorVolumePatchOp : public TensorBase<TensorVolumePatchOp<Planes, Rows,
DenseIndex padding_top, DenseIndex padding_bottom, DenseIndex padding_top, DenseIndex padding_bottom,
DenseIndex padding_left, DenseIndex padding_right, DenseIndex padding_left, DenseIndex padding_right,
Scalar padding_value) Scalar padding_value)
: m_xpr(expr), m_patch_planes(patch_planes), m_patch_rows(patch_rows), m_patch_cols(patch_cols), : m_xpr(expr), m_patch_planes(patch_planes), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_plane_strides(plane_strides), m_row_strides(row_strides), m_col_strides(col_strides), m_plane_strides(plane_strides), m_row_strides(row_strides), m_col_strides(col_strides),
m_in_plane_strides(in_plane_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides), m_in_plane_strides(in_plane_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_plane_inflate_strides(plane_inflate_strides), m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides), m_plane_inflate_strides(plane_inflate_strides), m_row_inflate_strides(row_inflate_strides), m_col_inflate_strides(col_inflate_strides),
m_padding_explicit(true), m_padding_top_z(padding_top_z), m_padding_bottom_z(padding_bottom_z), m_padding_top(padding_top), m_padding_bottom(padding_bottom), m_padding_explicit(true), m_padding_top_z(padding_top_z), m_padding_bottom_z(padding_bottom_z), m_padding_top(padding_top), m_padding_bottom(padding_bottom),
m_padding_left(padding_left), m_padding_right(padding_right), m_padding_left(padding_left), m_padding_right(padding_right),
m_padding_type(PADDING_VALID), m_padding_value(padding_value) {} m_padding_type(PADDING_VALID), m_padding_value(padding_value) {}
#ifdef EIGEN_USE_SYCL // this is work around for sycl as Eigen could not use c++11 deligate constructor feature
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorVolumePatchOp(const XprType& expr, DenseIndex patch_planes, DenseIndex patch_rows, DenseIndex patch_cols,
DenseIndex plane_strides, DenseIndex row_strides, DenseIndex col_strides,
DenseIndex in_plane_strides, DenseIndex in_row_strides, DenseIndex in_col_strides,
DenseIndex plane_inflate_strides, DenseIndex row_inflate_strides, DenseIndex col_inflate_strides,
bool padding_explicit, DenseIndex padding_top_z, DenseIndex padding_bottom_z,
DenseIndex padding_top, DenseIndex padding_bottom, DenseIndex padding_left,
DenseIndex padding_right, PaddingType padding_type, Scalar padding_value)
: m_xpr(expr), m_patch_planes(patch_planes), m_patch_rows(patch_rows), m_patch_cols(patch_cols),
m_plane_strides(plane_strides), m_row_strides(row_strides), m_col_strides(col_strides),
m_in_plane_strides(in_plane_strides), m_in_row_strides(in_row_strides), m_in_col_strides(in_col_strides),
m_plane_inflate_strides(plane_inflate_strides), m_row_inflate_strides(row_inflate_strides),
m_col_inflate_strides(col_inflate_strides), m_padding_explicit(padding_explicit), m_padding_top_z(padding_top_z),
m_padding_bottom_z(padding_bottom_z), m_padding_top(padding_top), m_padding_bottom(padding_bottom), m_padding_left(padding_left),
m_padding_right(padding_right), m_padding_type(padding_type), m_padding_value(padding_value) {}
#endif
EIGEN_DEVICE_FUNC EIGEN_DEVICE_FUNC
DenseIndex patch_planes() const { return m_patch_planes; } DenseIndex patch_planes() const { return m_patch_planes; }
@ -509,11 +530,12 @@ struct TensorEvaluator<const TensorVolumePatchOp<Planes, Rows, Cols, ArgType>, D
return TensorOpCost(0, 0, compute_cost, vectorized, PacketSize); return TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
} }
EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; } EIGEN_DEVICE_FUNC typename Eigen::internal::traits<XprType>::PointerType data() const { return NULL; }
const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
// Required by SYCL in order to construct the expression on the device
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& xpr() const { return m_op; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& xpr() const { return m_op; }
#endif #endif
@ -613,8 +635,10 @@ struct TensorEvaluator<const TensorVolumePatchOp<Planes, Rows, Cols, ArgType>, D
TensorEvaluator<ArgType, Device> m_impl; TensorEvaluator<ArgType, Device> m_impl;
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
// Required by SYCL in order to construct the expression on the device
XprType m_op; XprType m_op;
#endif #endif
}; };

4
unsupported/test/cxx11_tensor_image_patch_sycl.cpp
View File

@ -13,7 +13,7 @@
#define EIGEN_TEST_NO_LONGDOUBLE #define EIGEN_TEST_NO_LONGDOUBLE
#define EIGEN_TEST_NO_COMPLEX #define EIGEN_TEST_NO_COMPLEX
#define EIGEN_TEST_FUNC cxx11_tensor_image_patchOP_sycl #define EIGEN_TEST_FUNC cxx11_tensor_image_patch_sycl
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t #define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t
#define EIGEN_USE_SYCL #define EIGEN_USE_SYCL
@ -1084,7 +1084,7 @@ test_patch_padding_same_sycl<DataType, int64_t>(sycl_device);
test_patch_no_extra_dim_sycl<DataType, int64_t>(sycl_device); test_patch_no_extra_dim_sycl<DataType, int64_t>(sycl_device);
test_imagenet_patches_sycl<DataType, int64_t>(sycl_device); test_imagenet_patches_sycl<DataType, int64_t>(sycl_device);
} }
void test_cxx11_tensor_image_patchOP_sycl() void test_cxx11_tensor_image_patch_sycl()
{ {
for (const auto& device :Eigen::get_sycl_supported_devices()) { for (const auto& device :Eigen::get_sycl_supported_devices()) {
CALL_SUBTEST(sycl_tensor_image_patch_test_per_device<float>(device)); CALL_SUBTEST(sycl_tensor_image_patch_test_per_device<float>(device));

4
unsupported/test/cxx11_tensor_volume_patch_sycl.cpp
View File

@ -13,7 +13,7 @@
#define EIGEN_TEST_NO_LONGDOUBLE #define EIGEN_TEST_NO_LONGDOUBLE
#define EIGEN_TEST_NO_COMPLEX #define EIGEN_TEST_NO_COMPLEX
#define EIGEN_TEST_FUNC cxx11_tensor_volume_patchOP_sycl #define EIGEN_TEST_FUNC cxx11_tensor_volume_patch_sycl
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t #define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t
#define EIGEN_USE_SYCL #define EIGEN_USE_SYCL
@ -214,7 +214,7 @@ std::cout << "Running on " << s.template get_info<cl::sycl::info::device::name>(
test_single_voxel_patch_sycl<DataType, int64_t>(sycl_device); test_single_voxel_patch_sycl<DataType, int64_t>(sycl_device);
test_entire_volume_patch_sycl<DataType, int64_t>(sycl_device); test_entire_volume_patch_sycl<DataType, int64_t>(sycl_device);
} }
void test_cxx11_tensor_volume_patchOP_sycl() void test_cxx11_tensor_volume_patch_sycl()
{ {
for (const auto& device :Eigen::get_sycl_supported_devices()) { for (const auto& device :Eigen::get_sycl_supported_devices()) {
CALL_SUBTEST(sycl_tensor_volume_patch_test_per_device<float>(device)); CALL_SUBTEST(sycl_tensor_volume_patch_test_per_device<float>(device));