diff options
17 files changed, 727 insertions, 1883 deletions
diff --git a/Eigen/Core b/Eigen/Core index ed007dfa9..53cb72c95 100644 --- a/Eigen/Core +++ b/Eigen/Core @@ -11,46 +11,8 @@ #ifndef EIGEN_CORE_H #define EIGEN_CORE_H -// first thing Eigen does: stop the compiler from committing suicide -#include "src/Core/util/DisableStupidWarnings.h" - -// Handle NVCC/CUDA -#ifdef __CUDACC__ - // Do not try asserts on CUDA! - #ifndef EIGEN_NO_DEBUG - #define EIGEN_NO_DEBUG - #endif - - #ifdef EIGEN_INTERNAL_DEBUGGING - #undef EIGEN_INTERNAL_DEBUGGING - #endif - - // Do not try to vectorize on CUDA! - #ifndef EIGEN_DONT_VECTORIZE - #define EIGEN_DONT_VECTORIZE - #endif - - #ifdef EIGEN_EXCEPTIONS - #undef EIGEN_EXCEPTIONS - #endif - - // All functions callable from CUDA code must be qualified with __device__ - #define EIGEN_DEVICE_FUNC __host__ __device__ - -#else - #define EIGEN_DEVICE_FUNC - -#endif - -// When compiling CUDA device code with NVCC, pull in math functions from the -// global namespace. In host mode, and when device doee with clang, use the -// std versions. -#if defined(__CUDA_ARCH__) && defined(__NVCC__) - #define EIGEN_USING_STD_MATH(FUNC) using ::FUNC; -#else - #define EIGEN_USING_STD_MATH(FUNC) using std::FUNC; -#endif +/// This will no longer be needed after the next release of the computecppCE #ifdef EIGEN_USE_SYCL #undef min #undef max @@ -60,9 +22,9 @@ #include <SYCL/sycl.hpp> #endif -// We need these predefines to determine if asserts need to be disabled for the device compiler -#if defined(__SYCL_DEVICE_ONLY__) - // Do not try asserts on SYCL! +// Handle NVCC/CUDA/SYCL +#if defined(__CUDACC__) || defined(__SYCL_DEVICE_ONLY__) + // Do not try asserts on CUDA and SYCL! #ifndef EIGEN_NO_DEBUG #define EIGEN_NO_DEBUG #endif @@ -71,7 +33,7 @@ #undef EIGEN_INTERNAL_DEBUGGING #endif - // Do not try to vectorize on SYCL! + // Do not try to vectorize on CUDA and SYCL! #ifndef EIGEN_DONT_VECTORIZE #define EIGEN_DONT_VECTORIZE #endif @@ -80,12 +42,25 @@ #undef EIGEN_EXCEPTIONS #endif + // All functions callable from CUDA code must be qualified with __device__ + #ifdef __CUDACC__ + #define EIGEN_DEVICE_FUNC __host__ __device__ + // For Sycl we dont need that + #else + #define EIGEN_DEVICE_FUNC + #endif +#else #define EIGEN_DEVICE_FUNC #endif -#if (defined(_CPPUNWIND) || defined(__EXCEPTIONS)) && !defined(__CUDA_ARCH__) && !defined(EIGEN_EXCEPTIONS) && !defined(EIGEN_USE_SYCL) - #define EIGEN_EXCEPTIONS +// When compiling CUDA device code with NVCC, pull in math functions from the +// global namespace. In host mode, and when device doee with clang, use the +// std versions. +#if defined(__CUDA_ARCH__) && defined(__NVCC__) + #define EIGEN_USING_STD_MATH(FUNC) using ::FUNC; +#else + #define EIGEN_USING_STD_MATH(FUNC) using std::FUNC; #endif #ifdef EIGEN_EXCEPTIONS diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h index a87777b22..2cbcc87c6 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h @@ -423,15 +423,15 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device> // Precompute output strides. if (NumOutputDims > 0) { if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) { - m_outputStrides[0] = 1; - for (int i = 1; i < NumOutputDims; ++i) { - m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1]; - } + m_outputStrides[0] = 1; + for (int i = 1; i < NumOutputDims; ++i) { + m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1]; + } } else { - m_outputStrides.back() = 1; - for (int i = NumOutputDims - 2; i >= 0; --i) { - m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1]; - } + m_outputStrides.back() = 1; + for (int i = NumOutputDims - 2; i >= 0; --i) { + m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1]; + } } } @@ -439,27 +439,27 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device> if (NumInputDims > 0) { array<Index, NumInputDims> input_strides; if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) { - input_strides[0] = 1; - for (int i = 1; i < NumInputDims; ++i) { - input_strides[i] = input_strides[i-1] * input_dims[i-1]; - } + input_strides[0] = 1; + for (int i = 1; i < NumInputDims; ++i) { + input_strides[i] = input_strides[i-1] * input_dims[i-1]; + } } else { - input_strides.back() = 1; - for (int i = NumInputDims - 2; i >= 0; --i) { - input_strides[i] = input_strides[i + 1] * input_dims[i + 1]; - } + input_strides.back() = 1; + for (int i = NumInputDims - 2; i >= 0; --i) { + input_strides[i] = input_strides[i + 1] * input_dims[i + 1]; + } } int outputIndex = 0; int reduceIndex = 0; for (int i = 0; i < NumInputDims; ++i) { - if (m_reduced[i]) { - m_reducedStrides[reduceIndex] = input_strides[i]; - ++reduceIndex; - } else { - m_preservedStrides[outputIndex] = input_strides[i]; - ++outputIndex; - } + if (m_reduced[i]) { + m_reducedStrides[reduceIndex] = input_strides[i]; + ++reduceIndex; + } else { + m_preservedStrides[outputIndex] = input_strides[i]; + ++outputIndex; + } } } @@ -578,7 +578,7 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device> Op reducer(m_reducer); if (ReducingInnerMostDims || RunningFullReduction) { const Index num_values_to_reduce = - (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1]; + (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1]; return internal::InnerMostDimReducer<Self, Op>::reduce(*this, firstInput(index), num_values_to_reduce, reducer); } else { @@ -602,7 +602,7 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device> EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[PacketSize]; if (ReducingInnerMostDims) { const Index num_values_to_reduce = - (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1]; + (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1]; const Index firstIndex = firstInput(index); for (Index i = 0; i < PacketSize; ++i) { Op reducer(m_reducer); diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h index 277dd739c..da15f7942 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h @@ -11,20 +11,35 @@ // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. // General include header of SYCL target for Tensor Module -#ifndef TENSORSYCL_H -#define TENSORSYCL_H +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H #ifdef EIGEN_USE_SYCL -// trait class to extract different attribute contents -template <typename T> -struct Trait; // global pointer to set different attribute state for a class template <class T> struct MakeGlobalPointer { typedef typename cl::sycl::global_ptr<T>::pointer_t Type; }; +namespace Eigen { +namespace TensorSycl { +namespace internal { + +/// This struct is used for special expression nodes with no operations (for example assign and selectOP). + struct NoOP; + +template<bool IsConst, typename T> struct GetType{ + typedef const T Type; +}; +template<typename T> struct GetType<false, T>{ + typedef T Type; +}; + +} +} +} + // tuple construction #include "TensorSyclTuple.h" @@ -59,4 +74,4 @@ struct MakeGlobalPointer { #include "TensorSyclRun.h" #endif // end of EIGEN_USE_SYCL -#endif // TENSORSYCL_H +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h index b3748131b..a94c30426 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h @@ -19,12 +19,13 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP namespace Eigen { namespace TensorSycl { namespace internal { + /// \struct ConvertToDeviceExpression /// \brief This struct is used to convert the MakePointer in the host expression /// to the MakeGlobalPointer for the device expression. For the leafNodes @@ -33,204 +34,74 @@ namespace internal { template <typename Expr> struct ConvertToDeviceExpression; -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorMap -template <typename Scalar_, int Options_, int Options2_, int NumIndices_, - typename IndexType_, template <class> class MakePointer_> -struct ConvertToDeviceExpression< - TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, - MakePointer_>> { - using Type = TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakeGlobalPointer>; +template<template<class...> class NonOpCategory, bool IsConst, typename... Args> +struct NonOpConversion{ + typedef typename GetType<IsConst, NonOpCategory<typename ConvertToDeviceExpression<Args>::Type...> >::Type Type; }; -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorMap -template <typename Scalar_, int Options_, int Options2_, int NumIndices_, - typename IndexType_, template <class> class MakePointer_> -struct ConvertToDeviceExpression< - const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakePointer_>> { - using Type = - const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakeGlobalPointer>; -}; -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorCwiseNullaryOp -template <typename OP, typename RHSExpr> -struct ConvertToDeviceExpression<const TensorCwiseNullaryOp<OP, RHSExpr>> { - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = const TensorCwiseNullaryOp<OP, RHSPlaceHolderType>; +template<template<class, template <class> class > class NonOpCategory, bool IsConst, typename Args> +struct DeviceConvertor{ + typedef typename GetType<IsConst, NonOpCategory<typename ConvertToDeviceExpression<Args>::Type, MakeGlobalPointer> >::Type Type; }; /// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorCwiseNullaryOp -template <typename OP, typename RHSExpr> -struct ConvertToDeviceExpression<TensorCwiseNullaryOp<OP, RHSExpr>> { - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = TensorCwiseNullaryOp<OP, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorBroadcastingOp -template <typename OP, typename RHSExpr> -struct ConvertToDeviceExpression<const TensorBroadcastingOp<OP, RHSExpr>> { - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = const TensorBroadcastingOp<OP, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorBroadcastingOp -template <typename OP, typename RHSExpr> -struct ConvertToDeviceExpression<TensorBroadcastingOp<OP, RHSExpr>> { - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = TensorBroadcastingOp<OP, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorCwiseUnaryOp -template <typename OP, typename RHSExpr> -struct ConvertToDeviceExpression<const TensorCwiseUnaryOp<OP, RHSExpr>> { - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = const TensorCwiseUnaryOp<OP, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorCwiseUnaryOp -template <typename OP, typename RHSExpr> -struct ConvertToDeviceExpression<TensorCwiseUnaryOp<OP, RHSExpr>> { - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = TensorCwiseUnaryOp<OP, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr> -struct ConvertToDeviceExpression< - const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> { - using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type; - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = - const TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr> -struct ConvertToDeviceExpression<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> { - using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type; - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>; +/// type is TensorMap +#define TENSORMAPCONVERT(CVQual)\ +template <typename Scalar_, int Options_, int Options2_, int NumIndices_, typename IndexType_, template <class> class MakePointer_>\ +struct ConvertToDeviceExpression<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_> > {\ + typedef CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer> Type;\ }; -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorCwiseTernaryOp -template <typename OP, typename Arg1Impl, typename Arg2Impl, typename Arg3Impl> -struct ConvertToDeviceExpression< - const TensorCwiseTernaryOp<OP, Arg1Impl, Arg2Impl, Arg3Impl>> { - using Arg1PlaceHolderType = - typename ConvertToDeviceExpression<Arg1Impl>::Type; - using Arg2PlaceHolderType = - typename ConvertToDeviceExpression<Arg2Impl>::Type; - using Arg3PlaceHolderType = - typename ConvertToDeviceExpression<Arg3Impl>::Type; - using Type = - const TensorCwiseTernaryOp<OP, Arg1PlaceHolderType, Arg2PlaceHolderType, - Arg3PlaceHolderType>; -}; +TENSORMAPCONVERT(const) +TENSORMAPCONVERT() +#undef TENSORMAPCONVERT /// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorCwiseTernaryOp -template <typename OP, typename Arg1Impl, typename Arg2Impl, typename Arg3Impl> -struct ConvertToDeviceExpression< - TensorCwiseTernaryOp<OP, Arg1Impl, Arg2Impl, Arg3Impl>> { - using Arg1PlaceHolderType = - typename ConvertToDeviceExpression<Arg1Impl>::Type; - using Arg2PlaceHolderType = - typename ConvertToDeviceExpression<Arg2Impl>::Type; - using Arg3PlaceHolderType = - typename ConvertToDeviceExpression<Arg3Impl>::Type; - using Type = TensorCwiseTernaryOp<OP, Arg1PlaceHolderType, - Arg2PlaceHolderType, Arg3PlaceHolderType>; +/// type is TensorCwiseNullaryOp, TensorCwiseUnaryOp, TensorCwiseBinaryOp, TensorCwiseTernaryOp, TensorBroadcastingOp +#define CATEGORYCONVERT(CVQual)\ +template <template<class, class...> class Category, typename OP, typename... subExprs>\ +struct ConvertToDeviceExpression<CVQual Category<OP, subExprs...> > {\ + typedef CVQual Category<OP, typename ConvertToDeviceExpression<subExprs>::Type... > Type;\ }; +CATEGORYCONVERT(const) +CATEGORYCONVERT() +#undef CATEGORYCONVERT -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorCwiseSelectOp -template <typename IfExpr, typename ThenExpr, typename ElseExpr> -struct ConvertToDeviceExpression< - const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> { - using IfPlaceHolderType = typename ConvertToDeviceExpression<IfExpr>::Type; - using ThenPlaceHolderType = - typename ConvertToDeviceExpression<ThenExpr>::Type; - using ElsePlaceHolderType = - typename ConvertToDeviceExpression<ElseExpr>::Type; - using Type = const TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType, - ElsePlaceHolderType>; -}; /// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorCwiseSelectOp -template <typename IfExpr, typename ThenExpr, typename ElseExpr> -struct ConvertToDeviceExpression<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> { - using IfPlaceHolderType = typename ConvertToDeviceExpression<IfExpr>::Type; - using ThenPlaceHolderType = - typename ConvertToDeviceExpression<ThenExpr>::Type; - using ElsePlaceHolderType = - typename ConvertToDeviceExpression<ElseExpr>::Type; - using Type = TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType, - ElsePlaceHolderType>; -}; +/// type is TensorCwiseSelectOp +#define SELECTOPCONVERT(CVQual, Res)\ +template <typename IfExpr, typename ThenExpr, typename ElseExpr>\ +struct ConvertToDeviceExpression<CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr> >\ +: NonOpConversion<TensorSelectOp, Res, IfExpr, ThenExpr, ElseExpr> {}; +SELECTOPCONVERT(const, true) +SELECTOPCONVERT(, false) +#undef SELECTOPCONVERT /// specialisation of the \ref ConvertToDeviceExpression struct when the node /// type is const AssingOP -template <typename LHSExpr, typename RHSExpr> -struct ConvertToDeviceExpression<const TensorAssignOp<LHSExpr, RHSExpr>> { - using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type; - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = const TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is AssingOP -template <typename LHSExpr, typename RHSExpr> -struct ConvertToDeviceExpression<TensorAssignOp<LHSExpr, RHSExpr>> { - using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type; - using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; - using Type = TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorForcedEvalOp -template <typename Expr> -struct ConvertToDeviceExpression<const TensorForcedEvalOp<Expr>> { - using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type; - using Type = const TensorForcedEvalOp<PlaceHolderType, MakeGlobalPointer>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorForcedEvalOp -template <typename Expr> -struct ConvertToDeviceExpression<TensorForcedEvalOp<Expr>> { - using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type; - using Type = TensorForcedEvalOp<PlaceHolderType, MakeGlobalPointer>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is const TensorEvalToOp -template <typename Expr> -struct ConvertToDeviceExpression<const TensorEvalToOp<Expr>> { - using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type; - using Type = const TensorEvalToOp<PlaceHolderType, MakeGlobalPointer>; -}; - -/// specialisation of the \ref ConvertToDeviceExpression struct when the node -/// type is TensorEvalToOp -template <typename Expr> -struct ConvertToDeviceExpression<TensorEvalToOp<Expr>> { - using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type; - using Type = TensorEvalToOp<PlaceHolderType, MakeGlobalPointer>; -}; +#define ASSIGNCONVERT(CVQual, Res)\ +template <typename LHSExpr, typename RHSExpr>\ +struct ConvertToDeviceExpression<CVQual TensorAssignOp<LHSExpr, RHSExpr> >\ +: NonOpConversion<TensorAssignOp, Res, LHSExpr, RHSExpr>{}; + +ASSIGNCONVERT(const, true) +ASSIGNCONVERT(, false) +#undef ASSIGNCONVERT + +/// specialisation of the \ref ConvertToDeviceExpression struct when the node +/// type is either TensorForcedEvalOp or TensorEvalToOp +#define KERNELBROKERCONVERT(CVQual, Res, ExprNode)\ +template <typename Expr>\ +struct ConvertToDeviceExpression<CVQual ExprNode<Expr> > \ +: DeviceConvertor<ExprNode, Res, Expr>{}; + +KERNELBROKERCONVERT(const, true, TensorForcedEvalOp) +KERNELBROKERCONVERT(, false, TensorForcedEvalOp) +KERNELBROKERCONVERT(const, true, TensorEvalToOp) +KERNELBROKERCONVERT(, false, TensorEvalToOp) +#undef KERNELBROKERCONVERT } // namespace internal } // namespace TensorSycl } // namespace Eigen diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h index fe3994175..833d5e271 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h @@ -19,8 +19,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXPR_CONSTRUCTOR_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXPR_CONSTRUCTOR_HPP namespace Eigen { namespace TensorSycl { @@ -30,8 +30,7 @@ namespace internal { template <typename PtrType, size_t N, typename... Params> struct EvalToLHSConstructor { PtrType expr; - EvalToLHSConstructor(const utility::tuple::Tuple<Params...> &t) - : expr((&(*(utility::tuple::get<N>(t).get_pointer())))) {} + EvalToLHSConstructor(const utility::tuple::Tuple<Params...> &t): expr((&(*(utility::tuple::get<N>(t).get_pointer())))) {} }; /// \struct ExprConstructor is used to reconstruct the expression on the device @@ -45,446 +44,165 @@ template <typename OrigExpr, typename IndexExpr, typename... Params> struct ExprConstructor; /// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorMap -template <typename Scalar_, int Options_, int Options2_, int Options3_, - int NumIndices_, typename IndexType_, - template <class> class MakePointer_, size_t N, typename... Params> -struct ExprConstructor< - const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakeGlobalPointer>, - const Eigen::internal::PlaceHolder< - const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options3_, MakePointer_>, - N>, - Params...> { - using Type = - const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakeGlobalPointer>; - - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t) - : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), - fd.dimensions())) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is /// TensorMap -template <typename Scalar_, int Options_, int Options2_, int Options3_, - int NumIndices_, typename IndexType_, - template <class> class MakePointer_, size_t N, typename... Params> -struct ExprConstructor< - TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, - MakeGlobalPointer>, - Eigen::internal::PlaceHolder< - TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options3_, - MakePointer_>, - N>, - Params...> { - using Type = TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakeGlobalPointer>; - - Type expr; - template <typename FuncDetector> - ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t) - : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), - fd.dimensions())) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// TensorCwiseNullaryOp -template <typename OP, typename OrigRHSExpr, typename RHSExpr, - typename... Params> -struct ExprConstructor<TensorCwiseNullaryOp<OP, OrigRHSExpr>, - TensorCwiseNullaryOp<OP, RHSExpr>, Params...> { - using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - my_type rhsExpr; - using Type = TensorCwiseNullaryOp<OP, typename my_type::Type>; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorCwiseNullaryOp -template <typename OP, typename OrigRHSExpr, typename RHSExpr, - typename... Params> -struct ExprConstructor<const TensorCwiseNullaryOp<OP, OrigRHSExpr>, - const TensorCwiseNullaryOp<OP, RHSExpr>, Params...> { - using my_type = const ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - my_type rhsExpr; - using Type = const TensorCwiseNullaryOp<OP, typename my_type::Type>; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// TensorBroadcastingOp -template <typename OP, typename OrigRHSExpr, typename RHSExpr, - typename... Params> -struct ExprConstructor<TensorBroadcastingOp<OP, OrigRHSExpr>, - TensorBroadcastingOp<OP, RHSExpr>, Params...> { - using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - my_type rhsExpr; - using Type = TensorBroadcastingOp<OP, typename my_type::Type>; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorBroadcastingOp -template <typename OP, typename OrigRHSExpr, typename RHSExpr, - typename... Params> -struct ExprConstructor<const TensorBroadcastingOp<OP, OrigRHSExpr>, - const TensorBroadcastingOp<OP, RHSExpr>, Params...> { - using my_type = const ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - my_type rhsExpr; - using Type = const TensorBroadcastingOp<OP, typename my_type::Type>; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// TensorCwiseUnaryOp -template <typename OP, typename OrigRHSExpr, typename RHSExpr, - typename... Params> -struct ExprConstructor<TensorCwiseUnaryOp<OP, OrigRHSExpr>, - TensorCwiseUnaryOp<OP, RHSExpr>, Params...> { - using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - using Type = TensorCwiseUnaryOp<OP, typename my_type::Type>; - my_type rhsExpr; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, utility::tuple::Tuple<Params...> &t) - : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorCwiseUnaryOp -template <typename OP, typename OrigRHSExpr, typename RHSExpr, - typename... Params> -struct ExprConstructor<const TensorCwiseUnaryOp<OP, OrigRHSExpr>, - const TensorCwiseUnaryOp<OP, RHSExpr>, Params...> { - using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - using Type = const TensorCwiseUnaryOp<OP, typename my_type::Type>; - my_type rhsExpr; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// TensorCwiseBinaryOp -template <typename OP, typename OrigLHSExpr, typename OrigRHSExpr, - typename LHSExpr, typename RHSExpr, typename... Params> -struct ExprConstructor<TensorCwiseBinaryOp<OP, OrigLHSExpr, OrigRHSExpr>, - TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Params...> { - using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>; - using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - using Type = TensorCwiseBinaryOp<OP, typename my_left_type::Type, - typename my_right_type::Type>; - - my_left_type lhsExpr; - my_right_type rhsExpr; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : lhsExpr(funcD.lhsExpr, t), - rhsExpr(funcD.rhsExpr, t), - expr(lhsExpr.expr, rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorCwiseBinaryOp -template <typename OP, typename OrigLHSExpr, typename OrigRHSExpr, - typename LHSExpr, typename RHSExpr, typename... Params> -struct ExprConstructor<const TensorCwiseBinaryOp<OP, OrigLHSExpr, OrigRHSExpr>, - const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, - Params...> { - using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>; - using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - using Type = const TensorCwiseBinaryOp<OP, typename my_left_type::Type, - typename my_right_type::Type>; - - my_left_type lhsExpr; - my_right_type rhsExpr; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : lhsExpr(funcD.lhsExpr, t), - rhsExpr(funcD.rhsExpr, t), - expr(lhsExpr.expr, rhsExpr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorCwiseTernaryOp -template <typename OP, typename OrigArg1Expr, typename OrigArg2Expr, - typename OrigArg3Expr, typename Arg1Expr, typename Arg2Expr, - typename Arg3Expr, typename... Params> -struct ExprConstructor< - const TensorCwiseTernaryOp<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>, - const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Params...> { - using my_arg1_type = ExprConstructor<OrigArg1Expr, Arg1Expr, Params...>; - using my_arg2_type = ExprConstructor<OrigArg2Expr, Arg2Expr, Params...>; - using my_arg3_type = ExprConstructor<OrigArg3Expr, Arg3Expr, Params...>; - using Type = const TensorCwiseTernaryOp<OP, typename my_arg1_type::Type, - typename my_arg2_type::Type, - typename my_arg3_type::Type>; - - my_arg1_type arg1Expr; - my_arg2_type arg2Expr; - my_arg3_type arg3Expr; - Type expr; - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : arg1Expr(funcD.arg1Expr, t), - arg2Expr(funcD.arg2Expr, t), - arg3Expr(funcD.arg3Expr, t), - expr(arg1Expr.expr, arg2Expr.expr, arg3Expr.expr, funcD.func) {} -}; +#define TENSORMAP(CVQual)\ +template <typename Scalar_, int Options_, int Options2_, int Options3_, int NumIndices_, typename IndexType_,\ +template <class> class MakePointer_, size_t N, typename... Params>\ +struct ExprConstructor< CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer>,\ +CVQual Eigen::internal::PlaceHolder<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options3_, MakePointer_>, N>, Params...>{\ + typedef CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer> Type;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)\ + : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), fd.dimensions())) {}\ +}; + +TENSORMAP(const) +TENSORMAP() +#undef TENSORMAP + +#define UNARYCATEGORY(CVQual)\ +template <template<class, class> class UnaryCategory, typename OP, typename OrigRHSExpr, typename RHSExpr, typename... Params>\ +struct ExprConstructor<CVQual UnaryCategory<OP, OrigRHSExpr>, CVQual UnaryCategory<OP, RHSExpr>, Params...> {\ + typedef ExprConstructor<OrigRHSExpr, RHSExpr, Params...> my_type;\ + my_type rhsExpr;\ + typedef CVQual UnaryCategory<OP, typename my_type::Type> Type;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}\ +}; + +UNARYCATEGORY(const) +UNARYCATEGORY() +#undef UNARYCATEGORY + +/// specialisation of the \ref ExprConstructor struct when the node type is +/// TensorBinaryOp +#define BINARYCATEGORY(CVQual)\ +template <template<class, class, class> class BinaryCategory, typename OP, typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr,\ +typename RHSExpr, typename... Params>\ +struct ExprConstructor<CVQual BinaryCategory<OP, OrigLHSExpr, OrigRHSExpr>, CVQual BinaryCategory<OP, LHSExpr, RHSExpr>, Params...> {\ + typedef ExprConstructor<OrigLHSExpr, LHSExpr, Params...> my_left_type;\ + typedef ExprConstructor<OrigRHSExpr, RHSExpr, Params...> my_right_type;\ + typedef CVQual BinaryCategory<OP, typename my_left_type::Type, typename my_right_type::Type> Type;\ + my_left_type lhsExpr;\ + my_right_type rhsExpr;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : lhsExpr(funcD.lhsExpr, t),rhsExpr(funcD.rhsExpr, t), expr(lhsExpr.expr, rhsExpr.expr, funcD.func) {}\ +}; + +BINARYCATEGORY(const) +BINARYCATEGORY() +#undef BINARYCATEGORY /// specialisation of the \ref ExprConstructor struct when the node type is /// TensorCwiseTernaryOp -template <typename OP, typename OrigArg1Expr, typename OrigArg2Expr, - typename OrigArg3Expr, typename Arg1Expr, typename Arg2Expr, - typename Arg3Expr, typename... Params> -struct ExprConstructor< - TensorCwiseTernaryOp<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>, - TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Params...> { - using my_arg1_type = ExprConstructor<OrigArg1Expr, Arg1Expr, Params...>; - using my_arg2_type = ExprConstructor<OrigArg2Expr, Arg2Expr, Params...>; - using my_arg3_type = ExprConstructor<OrigArg3Expr, Arg3Expr, Params...>; - using Type = TensorCwiseTernaryOp<OP, typename my_arg1_type::Type, - typename my_arg2_type::Type, - typename my_arg3_type::Type>; - - my_arg1_type arg1Expr; - my_arg2_type arg2Expr; - my_arg3_type arg3Expr; - Type expr; - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : arg1Expr(funcD.arg1Expr, t), - arg2Expr(funcD.arg2Expr, t), - arg3Expr(funcD.arg3Expr, t), - expr(arg1Expr.expr, arg2Expr.expr, arg3Expr.expr, funcD.func) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorCwiseSelectOp -template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr, - typename IfExpr, typename ThenExpr, typename ElseExpr, - typename... Params> -struct ExprConstructor< - const TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>, - const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Params...> { - using my_if_type = ExprConstructor<OrigIfExpr, IfExpr, Params...>; - using my_then_type = ExprConstructor<OrigThenExpr, ThenExpr, Params...>; - using my_else_type = ExprConstructor<OrigElseExpr, ElseExpr, Params...>; - using Type = const TensorSelectOp<typename my_if_type::Type, - typename my_then_type::Type, - typename my_else_type::Type>; - - my_if_type ifExpr; - my_then_type thenExpr; - my_else_type elseExpr; - Type expr; - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : ifExpr(funcD.ifExpr, t), - thenExpr(funcD.thenExpr, t), - elseExpr(funcD.elseExpr, t), - expr(ifExpr.expr, thenExpr.expr, elseExpr.expr) {} -}; +#define TERNARYCATEGORY(CVQual)\ +template <template <class, class, class, class> class TernaryCategory, typename OP, typename OrigArg1Expr, typename OrigArg2Expr,typename OrigArg3Expr,\ +typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename... Params>\ +struct ExprConstructor<CVQual TernaryCategory<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>, CVQual TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Params...> {\ + typedef ExprConstructor<OrigArg1Expr, Arg1Expr, Params...> my_arg1_type;\ + typedef ExprConstructor<OrigArg2Expr, Arg2Expr, Params...> my_arg2_type;\ + typedef ExprConstructor<OrigArg3Expr, Arg3Expr, Params...> my_arg3_type;\ + typedef CVQual TernaryCategory<OP, typename my_arg1_type::Type, typename my_arg2_type::Type, typename my_arg3_type::Type> Type;\ + my_arg1_type arg1Expr;\ + my_arg2_type arg2Expr;\ + my_arg3_type arg3Expr;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD,const utility::tuple::Tuple<Params...> &t)\ + : arg1Expr(funcD.arg1Expr, t), arg2Expr(funcD.arg2Expr, t), arg3Expr(funcD.arg3Expr, t), expr(arg1Expr.expr, arg2Expr.expr, arg3Expr.expr, funcD.func) {}\ +}; + +TERNARYCATEGORY(const) +TERNARYCATEGORY() +#undef TERNARYCATEGORY /// specialisation of the \ref ExprConstructor struct when the node type is /// TensorCwiseSelectOp -template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr, - typename IfExpr, typename ThenExpr, typename ElseExpr, - typename... Params> -struct ExprConstructor<TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>, - TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Params...> { - using my_if_type = ExprConstructor<OrigIfExpr, IfExpr, Params...>; - using my_then_type = ExprConstructor<OrigThenExpr, ThenExpr, Params...>; - using my_else_type = ExprConstructor<OrigElseExpr, ElseExpr, Params...>; - using Type = - TensorSelectOp<typename my_if_type::Type, typename my_then_type::Type, - typename my_else_type::Type>; - - my_if_type ifExpr; - my_then_type thenExpr; - my_else_type elseExpr; - Type expr; - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : ifExpr(funcD.ifExpr, t), - thenExpr(funcD.thenExpr, t), - elseExpr(funcD.elseExpr, t), - expr(ifExpr.expr, thenExpr.expr, elseExpr.expr) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// TensorAssignOp -template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr, - typename RHSExpr, typename... Params> -struct ExprConstructor<TensorAssignOp<OrigLHSExpr, OrigRHSExpr>, - TensorAssignOp<LHSExpr, RHSExpr>, Params...> { - using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>; - using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - using Type = - TensorAssignOp<typename my_left_type::Type, typename my_right_type::Type>; - - my_left_type lhsExpr; - my_right_type rhsExpr; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : lhsExpr(funcD.lhsExpr, t), - rhsExpr(funcD.rhsExpr, t), - expr(lhsExpr.expr, rhsExpr.expr) {} -}; +#define SELECTOP(CVQual)\ +template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr, typename IfExpr, typename ThenExpr, typename ElseExpr, typename... Params>\ +struct ExprConstructor< CVQual TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>, CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Params...> {\ + typedef ExprConstructor<OrigIfExpr, IfExpr, Params...> my_if_type;\ + typedef ExprConstructor<OrigThenExpr, ThenExpr, Params...> my_then_type;\ + typedef ExprConstructor<OrigElseExpr, ElseExpr, Params...> my_else_type;\ + typedef CVQual TensorSelectOp<typename my_if_type::Type, typename my_then_type::Type, typename my_else_type::Type> Type;\ + my_if_type ifExpr;\ + my_then_type thenExpr;\ + my_else_type elseExpr;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : ifExpr(funcD.ifExpr, t), thenExpr(funcD.thenExpr, t), elseExpr(funcD.elseExpr, t), expr(ifExpr.expr, thenExpr.expr, elseExpr.expr) {}\ +}; + +SELECTOP(const) +SELECTOP() +#undef SELECTOP /// specialisation of the \ref ExprConstructor struct when the node type is /// const TensorAssignOp -template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr, - typename RHSExpr, typename... Params> -struct ExprConstructor<const TensorAssignOp<OrigLHSExpr, OrigRHSExpr>, - const TensorAssignOp<LHSExpr, RHSExpr>, Params...> { - using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>; - using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; - using Type = const TensorAssignOp<typename my_left_type::Type, - typename my_right_type::Type>; - - my_left_type lhsExpr; - my_right_type rhsExpr; - Type expr; - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : lhsExpr(funcD.lhsExpr, t), - rhsExpr(funcD.rhsExpr, t), - expr(lhsExpr.expr, rhsExpr.expr) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorEvalToOp -template <typename OrigExpr, typename Expr, typename... Params> -struct ExprConstructor<const TensorEvalToOp<OrigExpr, MakeGlobalPointer>, - const TensorEvalToOp<Expr>, Params...> { - using my_expr_type = ExprConstructor<OrigExpr, Expr, Params...>; - using my_buffer_type = - typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType; - using Type = - const TensorEvalToOp<typename my_expr_type::Type, MakeGlobalPointer>; - my_expr_type nestedExpression; - EvalToLHSConstructor<my_buffer_type, 0, Params...> buffer; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : nestedExpression(funcD.rhsExpr, t), - buffer(t), - expr(buffer.expr, nestedExpression.expr) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// TensorEvalToOp -template <typename OrigExpr, typename Expr, typename... Params> -struct ExprConstructor<TensorEvalToOp<OrigExpr, MakeGlobalPointer>, - TensorEvalToOp<Expr>, Params...> { - using my_expr_type = ExprConstructor<OrigExpr, Expr, Params...>; - using my_buffer_type = - typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType; - using Type = TensorEvalToOp<typename my_expr_type::Type>; - my_expr_type nestedExpression; - EvalToLHSConstructor<my_buffer_type, 0, Params...> buffer; - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &funcD, - const utility::tuple::Tuple<Params...> &t) - : nestedExpression(funcD.rhsExpr, t), - buffer(t), - expr(buffer.expr, nestedExpression.expr) {} -}; - -/// specialisation of the \ref ExprConstructor struct when the node type is -/// const TensorForcedEvalOp -template <typename OrigExpr, typename DevExpr, size_t N, typename... Params> -struct ExprConstructor< - const TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>, - const Eigen::internal::PlaceHolder<const TensorForcedEvalOp<DevExpr>, N>, - Params...> { - using Type = const TensorMap< - Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar, - TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::NumDimensions, 0, - typename TensorForcedEvalOp<DevExpr>::Index>, - 0, MakeGlobalPointer>; - - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t) - : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), - fd.dimensions())) {} -}; +#define ASSIGN(CVQual)\ +template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr, typename RHSExpr, typename... Params>\ +struct ExprConstructor<CVQual TensorAssignOp<OrigLHSExpr, OrigRHSExpr>, CVQual TensorAssignOp<LHSExpr, RHSExpr>, Params...> {\ + typedef ExprConstructor<OrigLHSExpr, LHSExpr, Params...> my_left_type;\ + typedef ExprConstructor<OrigRHSExpr, RHSExpr, Params...> my_right_type;\ + typedef CVQual TensorAssignOp<typename my_left_type::Type, typename my_right_type::Type> Type;\ + my_left_type lhsExpr;\ + my_right_type rhsExpr;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : lhsExpr(funcD.lhsExpr, t), rhsExpr(funcD.rhsExpr, t), expr(lhsExpr.expr, rhsExpr.expr) {}\ + }; + + ASSIGN(const) + ASSIGN() + #undef ASSIGN +/// specialisation of the \ref ExprConstructor struct when the node type is +/// TensorEvalToOp +#define EVALTO(CVQual)\ +template <typename OrigExpr, typename Expr, typename... Params>\ +struct ExprConstructor<CVQual TensorEvalToOp<OrigExpr, MakeGlobalPointer>, CVQual TensorEvalToOp<Expr>, Params...> {\ + typedef ExprConstructor<OrigExpr, Expr, Params...> my_expr_type;\ + typedef typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType my_buffer_type;\ + typedef CVQual TensorEvalToOp<typename my_expr_type::Type, MakeGlobalPointer> Type;\ + my_expr_type nestedExpression;\ + EvalToLHSConstructor<my_buffer_type, 0, Params...> buffer;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : nestedExpression(funcD.rhsExpr, t), buffer(t), expr(buffer.expr, nestedExpression.expr) {}\ +}; + +EVALTO(const) +EVALTO() +#undef EVALTO /// specialisation of the \ref ExprConstructor struct when the node type is /// TensorForcedEvalOp -template <typename OrigExpr, typename DevExpr, size_t N, typename... Params> -struct ExprConstructor< - const TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>, - const Eigen::internal::PlaceHolder<TensorForcedEvalOp<DevExpr>, N>, - Params...> { - using Type = TensorMap< - Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar, 1, - 0, typename TensorForcedEvalOp<DevExpr>::Index>, - 0, MakeGlobalPointer>; - - Type expr; - - template <typename FuncDetector> - ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t) - : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), - fd.dimensions())) {} -}; +#define FORCEDEVAL(CVQual)\ +template <typename OrigExpr, typename DevExpr, size_t N, typename... Params>\ +struct ExprConstructor<CVQual TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>,\ +CVQual Eigen::internal::PlaceHolder<CVQual TensorForcedEvalOp<DevExpr>, N>, Params...> {\ + typedef CVQual TensorMap<Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar,\ + TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::NumDimensions, 0, typename TensorForcedEvalOp<DevExpr>::Index>, 0, MakeGlobalPointer> Type;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)\ + : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), fd.dimensions())) {}\ +}; + +FORCEDEVAL(const) +FORCEDEVAL() +#undef FORCEDEVAL /// template deduction for \ref ExprConstructor struct -template <typename OrigExpr, typename IndexExpr, typename FuncD, - typename... Params> -auto createDeviceExpression(FuncD &funcD, - const utility::tuple::Tuple<Params...> &t) +template <typename OrigExpr, typename IndexExpr, typename FuncD, typename... Params> +auto createDeviceExpression(FuncD &funcD, const utility::tuple::Tuple<Params...> &t) -> decltype(ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t)) { return ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t); } @@ -492,4 +210,4 @@ auto createDeviceExpression(FuncD &funcD, } } // namespace Eigen -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXPR_CONSTRUCTOR_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h index cb0ac131d..ceec528ea 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h @@ -29,8 +29,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP namespace Eigen { namespace TensorSycl { @@ -42,425 +42,160 @@ namespace internal { template <typename Evaluator> struct ExtractAccessor; -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// const TensorMap -template <typename PlainObjectType, int Options_, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>> { - using actual_type = typename Eigen::internal::remove_all< - typename Eigen::internal::traits<PlainObjectType>::Scalar>::type; - static inline auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev> - eval) - -> decltype(utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::read, true, - actual_type>( - eval.dimensions().TotalSize(), cgh, - eval.derived().data())))) { - return utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::read, true, - actual_type>( - eval.dimensions().TotalSize(), cgh, eval.derived().data()))); +struct AccessorConstructor{ + template<typename Arg> static inline auto getTuple(cl::sycl::handler& cgh, Arg eval) + -> decltype(ExtractAccessor<Arg>::getTuple(cgh, eval)) { + return ExtractAccessor<Arg>::getTuple(cgh, eval); } -}; - -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// TensorMap -template <typename PlainObjectType, int Options_, typename Dev> -struct ExtractAccessor< - TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>> { - using actual_type = typename Eigen::internal::remove_all< - typename Eigen::internal::traits<PlainObjectType>::Scalar>::type; - static inline auto getTuple( - cl::sycl::handler& cgh, - TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev> eval) - -> decltype(utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::read_write, - true, actual_type>( - eval.dimensions().TotalSize(), cgh, - eval.derived().data())))) { - return utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::read_write, - true, actual_type>( - eval.dimensions().TotalSize(), cgh, eval.derived().data()))); + template<typename Arg1, typename Arg2> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1, Arg2 eval2) + -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2))) { + return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2)); } -}; - -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// const TensorCwiseNullaryOp -template <typename OP, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev> eval) - -> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl())) { - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl()); - return RHSTuple; + template<typename Arg1, typename Arg2, typename Arg3> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1 , Arg2 eval2 , Arg3 eval3) + -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3)))) { + return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3))); } -}; - -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev> eval) - -> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl())) { - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl()); - return RHSTuple; - } -}; - -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// const TensorBroadcastingOp -template <typename OP, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev> eval) - -> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl())) { - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl()); - return RHSTuple; - } -}; - -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// TensorBroadcastingOp -template <typename OP, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev> eval) - -> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl())) { - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl()); - return RHSTuple; + template< cl::sycl::access::mode AcM, typename Arg> static inline auto getAccessor(cl::sycl::handler& cgh, Arg eval) + -> decltype(utility::tuple::make_tuple( eval.device().template get_sycl_accessor<AcM, true, + typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data()))){ + return utility::tuple::make_tuple(eval.device().template get_sycl_accessor<AcM, true, typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data())); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is -/// const TenosorCwiseUnary -template <typename OP, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev> eval) - -> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl())) { - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl()); - return RHSTuple; +/// const TensorCwiseNullaryOp, const TensorCwiseUnaryOp and const TensorBroadcastingOp +template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> +struct ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> eval) + -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){ + return AccessorConstructor::getTuple(cgh, eval.impl()); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is -/// TenosorCwiseUnary -template <typename OP, typename RHSExpr, typename Dev> -struct ExtractAccessor<TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev> eval) - -> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl())) { - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.impl()); - return RHSTuple; - } -}; +/// TensorCwiseNullaryOp, TensorCwiseUnaryOp and TensorBroadcastingOp +template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> +struct ExtractAccessor<TensorEvaluator<UnaryCategory<OP, RHSExpr>, Dev> > +: ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > {}; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { - static auto getTuple(cl::sycl::handler& cgh, - const TensorEvaluator< - const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev> - eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()), - ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()))) { - auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()); - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()); - return utility::tuple::append(LHSTuple, RHSTuple); +template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> +struct ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> eval) + -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){ + return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl()); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is /// TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev> - eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()), - ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()))) { - auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()); - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()); - return utility::tuple::append(LHSTuple, RHSTuple); - } -}; +template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> +struct ExtractAccessor<TensorEvaluator<BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > +: ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorCwiseTernaryOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, - typename Dev> -struct ExtractAccessor<TensorEvaluator< - const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator< - const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> - eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple( - cgh, eval.arg1Impl()), - utility::tuple::append( - ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple( - cgh, eval.arg2Impl()), - ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple( - cgh, eval.arg3Impl())))) { - auto Arg1Tuple = ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple( - cgh, eval.arg1Impl()); - auto Arg2Tuple = ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple( - cgh, eval.arg2Impl()); - auto Arg3Tuple = ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple( - cgh, eval.arg3Impl()); - return utility::tuple::append(Arg1Tuple, - utility::tuple::append(Arg2Tuple, Arg3Tuple)); +template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> +struct ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> eval) + -> decltype(AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl())){ + return AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl()); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is /// TensorCwiseTernaryOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, - typename Dev> -struct ExtractAccessor<TensorEvaluator< - TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator< - TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> - eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple( - cgh, eval.arg1Impl()), - utility::tuple::append( - ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple( - cgh, eval.arg2Impl()), - ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple( - cgh, eval.arg3Impl())))) { - auto Arg1Tuple = ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple( - cgh, eval.arg1Impl()); - auto Arg2Tuple = ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple( - cgh, eval.arg2Impl()); - auto Arg3Tuple = ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple( - cgh, eval.arg3Impl()); - return utility::tuple::append(Arg1Tuple, - utility::tuple::append(Arg2Tuple, Arg3Tuple)); - } -}; +template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> +struct ExtractAccessor<TensorEvaluator<TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > +: ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is -/// const TensorCwiseSelectOp +/// const TensorCwiseSelectOp. This is a special case where there is no OP template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, - Dev> - eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple( - cgh, eval.cond_impl()), - utility::tuple::append( - ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple( - cgh, eval.then_impl()), - ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple( - cgh, eval.else_impl())))) { - auto IfTuple = ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple( - cgh, eval.cond_impl()); - auto ThenTuple = ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple( - cgh, eval.then_impl()); - auto ElseTuple = ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple( - cgh, eval.else_impl()); - return utility::tuple::append(IfTuple, - utility::tuple::append(ThenTuple, ElseTuple)); +struct ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> eval) + -> decltype(AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl())){ + return AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl()); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is -/// TensorCwiseSelectOp +/// TensorCwiseSelectOp. This is a special case where there is no OP template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> - eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple( - cgh, eval.cond_impl()), - utility::tuple::append( - ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple( - cgh, eval.then_impl()), - ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple( - cgh, eval.else_impl())))) { - auto IfTuple = ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple( - cgh, eval.cond_impl()); - auto ThenTuple = ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple( - cgh, eval.then_impl()); - auto ElseTuple = ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple( - cgh, eval.else_impl()); - return utility::tuple::append(IfTuple, - utility::tuple::append(ThenTuple, ElseTuple)); - } -}; +struct ExtractAccessor<TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > +: ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorAssignOp template <typename LHSExpr, typename RHSExpr, typename Dev> -struct ExtractAccessor< - TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()), - ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()))) { - auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()); - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()); - return utility::tuple::append(LHSTuple, RHSTuple); - } +struct ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) + -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){ + return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl()); + } }; /// specialisation of the \ref ExtractAccessor struct when the node type is /// TensorAssignOp template <typename LHSExpr, typename RHSExpr, typename Dev> -struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) - -> decltype(utility::tuple::append( - ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - eval.left_impl()), - ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - eval.right_impl()))) { - auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( - cgh, eval.left_impl()); - auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( - cgh, eval.right_impl()); - return utility::tuple::append(LHSTuple, RHSTuple); - } +struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> > +: ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{}; + +/// specialisation of the \ref ExtractAccessor struct when the node type is +/// const TensorMap +#define TENSORMAPEXPR(CVQual, ACCType)\ +template <typename PlainObjectType, int Options_, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> eval)\ + -> decltype(AccessorConstructor::template getAccessor<ACCType>(cgh, eval)){\ + return AccessorConstructor::template getAccessor<ACCType>(cgh, eval);\ + }\ }; +TENSORMAPEXPR(const, cl::sycl::access::mode::read) +TENSORMAPEXPR(, cl::sycl::access::mode::read_write) +#undef TENSORMAPEXPR /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorForcedEvalOp template <typename Expr, typename Dev> -struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> { - using actual_type = - typename Eigen::internal::remove_all<typename TensorEvaluator< - const TensorForcedEvalOp<Expr>, Dev>::CoeffReturnType>::type; - static auto getTuple( - cl::sycl::handler& cgh, - const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> eval) - -> decltype(utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::read, false, - actual_type>( - eval.dimensions().TotalSize(), cgh, eval.data())))) { - return utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::read, false, - actual_type>( - eval.dimensions().TotalSize(), cgh, eval.data()))); +struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> eval) + -> decltype(AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval)){ + return AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is /// TensorForcedEvalOp template <typename Expr, typename Dev> -struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>> - : ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> {}; +struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev> > +: ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorEvalToOp template <typename Expr, typename Dev> -struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev>> { - using actual_type = - typename Eigen::internal::remove_all<typename TensorEvaluator< - const TensorEvalToOp<Expr>, Dev>::CoeffReturnType>::type; - - static auto getTuple(cl::sycl::handler& cgh, - TensorEvaluator<const TensorEvalToOp<Expr>, Dev> eval) - -> decltype(utility::tuple::append( - utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::write, - false, actual_type>( - eval.dimensions().TotalSize(), cgh, eval.data()))), - ExtractAccessor<TensorEvaluator<Expr, Dev>>::getTuple(cgh, - eval.impl()))) { - auto LHSTuple = utility::tuple::make_tuple( - (eval.device() - .template get_sycl_accessor<cl::sycl::access::mode::write, false, - actual_type>( - eval.dimensions().TotalSize(), cgh, eval.data()))); - - auto RHSTuple = - ExtractAccessor<TensorEvaluator<Expr, Dev>>::getTuple(cgh, eval.impl()); - return utility::tuple::append(LHSTuple, RHSTuple); +struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> > { + static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<const TensorEvalToOp<Expr>, Dev> eval) + -> decltype(utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl()))){ + return utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl())); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is /// TensorEvalToOp template <typename Expr, typename Dev> -struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev>> - : ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev>> {}; +struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev> > +: ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> >{}; /// template deduction for \ref ExtractAccessor template <typename Evaluator> auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) - -> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) { +-> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) { return ExtractAccessor<Evaluator>::getTuple(cgh, expr); } } } } -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h index f69c5afcb..801b4f5d7 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h @@ -19,8 +19,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_FUNCTORS_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_FUNCTORS_HPP namespace Eigen { namespace TensorSycl { @@ -31,283 +31,124 @@ namespace internal { /// expression on the device. /// We have to do that as in Eigen the functors are not stateless so we cannot /// re-instantiate them on the device. -/// We have to pass whatever instantiated to the device. -template <typename Evaluator> -struct FunctorExtractor; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorMap: -template <typename PlainObjectType, int Options_, typename Dev> -struct FunctorExtractor< - TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>> { - using Dimensions = typename PlainObjectType::Dimensions; - const Dimensions m_dimensions; - const Dimensions& dimensions() const { return m_dimensions; } - FunctorExtractor( - const TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>& expr) - : m_dimensions(expr.dimensions()) {} -}; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorMap -template <typename PlainObjectType, int Options_, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>> { - using Dimensions = typename PlainObjectType::Dimensions; +/// We have to pass instantiated functors to the device. +// This struct is used for leafNode (TensorMap) and nodes behaving like leafNode (TensorForcedEval). +template <typename Evaluator> struct FunctorExtractor{ + typedef typename Evaluator::Dimensions Dimensions; const Dimensions m_dimensions; const Dimensions& dimensions() const { return m_dimensions; } - FunctorExtractor( - const TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>& - expr) - : m_dimensions(expr.dimensions()) {} -}; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorForcedEvalOp -template <typename Expr, typename Dev> -struct FunctorExtractor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>> { - using Dimensions = typename Expr::Dimensions; - const Dimensions m_dimensions; - const Dimensions& dimensions() const { return m_dimensions; } - FunctorExtractor(const TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>& expr) - : m_dimensions(expr.dimensions()) {} -}; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorForcedEvalOp -template <typename Expr, typename Dev> -struct FunctorExtractor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> { - using Dimensions = - typename TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>::Dimensions; - const Dimensions m_dimensions; - const Dimensions& dimensions() const { return m_dimensions; } - FunctorExtractor( - const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>& expr) - : m_dimensions(expr.dimensions()) {} -}; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - OP func; - FunctorExtractor( - TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()), func(expr.functor()) {} -}; + FunctorExtractor(const Evaluator& expr) + : m_dimensions(expr.dimensions()) {} -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorCwiseNullaryOp -template <typename OP, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - OP func; - FunctorExtractor( - const TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()), func(expr.functor()) {} -}; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorBroadcastingOp -template <typename OP, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - OP func; - FunctorExtractor( - const TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()), func(expr.functor()) {} }; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorBroadcastingOp -template <typename OP, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; +/// const TensorCwiseNullaryOp, const TensorCwiseUnaryOp, and const TensorBroadcastingOp +template <template <class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> +struct FunctorExtractor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > { + FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr; OP func; - FunctorExtractor( - const TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()), func(expr.functor()) {} + FunctorExtractor(const TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev>& expr) + : rhsExpr(expr.impl()), func(expr.functor()) {} }; - /// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorCwiseUnaryOp -template <typename OP, typename RHSExpr, typename Dev> -struct FunctorExtractor<TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - OP func; - FunctorExtractor( - const TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()), func(expr.functor()) {} -}; +/// TensorCwiseNullaryOp, TensorCwiseUnaryOp, and TensorBroadcastingOp +template <template <class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> +struct FunctorExtractor<TensorEvaluator<UnaryCategory<OP, RHSExpr>, Dev> > +: FunctorExtractor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> >{}; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorCwiseUnaryOp -template <typename OP, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - OP func; - FunctorExtractor( - const TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()), func(expr.functor()) {} -}; - -/// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr; - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; +/// const TensorCwiseBinaryOp +template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> +struct FunctorExtractor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > { + FunctorExtractor<TensorEvaluator<LHSExpr, Dev> > lhsExpr; + FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr; OP func; - FunctorExtractor( - const TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>& - expr) - : lhsExpr(expr.left_impl()), - rhsExpr(expr.right_impl()), - func(expr.functor()) {} + FunctorExtractor(const TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev>& expr) + : lhsExpr(expr.left_impl()),rhsExpr(expr.right_impl()),func(expr.functor()) {} }; /// specialisation of the \ref FunctorExtractor struct when the node type is /// const TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr; - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - OP func; - FunctorExtractor(const TensorEvaluator< - const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>& expr) - : lhsExpr(expr.left_impl()), - rhsExpr(expr.right_impl()), - func(expr.functor()) {} -}; +template <template <class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> +struct FunctorExtractor<TensorEvaluator<BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > +: FunctorExtractor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >{}; /// specialisation of the \ref FunctorExtractor struct when the node type is /// const TensorCwiseTernaryOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, - typename Dev> -struct FunctorExtractor<TensorEvaluator< - const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> { - FunctorExtractor<TensorEvaluator<Arg1Expr, Dev>> arg1Expr; - FunctorExtractor<TensorEvaluator<Arg2Expr, Dev>> arg2Expr; - FunctorExtractor<TensorEvaluator<Arg3Expr, Dev>> arg3Expr; +template <template <class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,typename Dev> +struct FunctorExtractor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > { + FunctorExtractor<TensorEvaluator<Arg1Expr, Dev> > arg1Expr; + FunctorExtractor<TensorEvaluator<Arg2Expr, Dev> > arg2Expr; + FunctorExtractor<TensorEvaluator<Arg3Expr, Dev> > arg3Expr; OP func; - FunctorExtractor(const TensorEvaluator< - const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, - Dev>& expr) - : arg1Expr(expr.arg1Impl()), - arg2Expr(expr.arg2Impl()), - arg3Expr(expr.arg3Impl()), - func(expr.functor()) {} + FunctorExtractor(const TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>& expr) + : arg1Expr(expr.arg1Impl()), arg2Expr(expr.arg2Impl()), arg3Expr(expr.arg3Impl()), func(expr.functor()) {} }; /// specialisation of the \ref FunctorExtractor struct when the node type is /// TensorCwiseTernaryOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, - typename Dev> -struct FunctorExtractor<TensorEvaluator< - TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> { - FunctorExtractor<TensorEvaluator<Arg1Expr, Dev>> arg1Expr; - FunctorExtractor<TensorEvaluator<Arg2Expr, Dev>> arg2Expr; - FunctorExtractor<TensorEvaluator<Arg3Expr, Dev>> arg3Expr; - OP func; - FunctorExtractor( - const TensorEvaluator< - TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>& expr) - : arg1Expr(expr.arg1Impl()), - arg2Expr(expr.arg2Impl()), - arg3Expr(expr.arg3Impl()), - func(expr.functor()) {} -}; +template <template <class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> +struct FunctorExtractor<TensorEvaluator< TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > +:FunctorExtractor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >{}; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorCwiseSelectOp +/// const TensorCwiseSelectOp. This is an specialisation without OP so it has to be separated. template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<IfExpr, Dev>> ifExpr; - FunctorExtractor<TensorEvaluator<ThenExpr, Dev>> thenExpr; - FunctorExtractor<TensorEvaluator<ElseExpr, Dev>> elseExpr; - FunctorExtractor(const TensorEvaluator< - const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>& expr) - : ifExpr(expr.cond_impl()), - thenExpr(expr.then_impl()), - elseExpr(expr.else_impl()) {} +struct FunctorExtractor< TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > { + FunctorExtractor<TensorEvaluator<IfExpr, Dev> > ifExpr; + FunctorExtractor<TensorEvaluator<ThenExpr, Dev> > thenExpr; + FunctorExtractor<TensorEvaluator<ElseExpr, Dev> > elseExpr; + FunctorExtractor(const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>& expr) + : ifExpr(expr.cond_impl()), thenExpr(expr.then_impl()), elseExpr(expr.else_impl()) {} }; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorCwiseSelectOp +/// TensorCwiseSelectOp. This is an specialisation without OP so it has to be separated template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { - FunctorExtractor<IfExpr> ifExpr; - FunctorExtractor<ThenExpr> thenExpr; - FunctorExtractor<ElseExpr> elseExpr; - FunctorExtractor( - const TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>& - expr) - : ifExpr(expr.cond_impl()), - thenExpr(expr.then_impl()), - elseExpr(expr.else_impl()) {} -}; +struct FunctorExtractor<TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > +:FunctorExtractor< TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > {}; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorAssignOp +/// const TensorAssignOp. This is an specialisation without OP so it has to be separated. template <typename LHSExpr, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr; - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - FunctorExtractor( - const TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr) - : lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {} +struct FunctorExtractor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> > { + FunctorExtractor<TensorEvaluator<LHSExpr, Dev> > lhsExpr; + FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr; + FunctorExtractor(const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr) + : lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {} }; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorAssignOp +/// TensorAssignOp. This is an specialisation without OP so it has to be separated. template <typename LHSExpr, typename RHSExpr, typename Dev> -struct FunctorExtractor< - TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr; - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - FunctorExtractor( - const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr) - : lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {} -}; +struct FunctorExtractor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> > +:FunctorExtractor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{}; + /// specialisation of the \ref FunctorExtractor struct when the node type is -/// TensorEvalToOp +/// const TensorEvalToOp, This is an specialisation without OP so it has to be separated. template <typename RHSExpr, typename Dev> -struct FunctorExtractor<TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - FunctorExtractor(const TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()) {} +struct FunctorExtractor<TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev> > { + FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr; + FunctorExtractor(const TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev>& expr) + : rhsExpr(expr.impl()) {} }; /// specialisation of the \ref FunctorExtractor struct when the node type is -/// const TensorEvalToOp +/// TensorEvalToOp. This is a specialisation without OP so it has to be separated. template <typename RHSExpr, typename Dev> -struct FunctorExtractor<TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev>> { - FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; - FunctorExtractor( - const TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev>& expr) - : rhsExpr(expr.impl()) {} -}; +struct FunctorExtractor<TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev> > +: FunctorExtractor<TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev> > {}; + /// template deduction function for FunctorExtractor template <typename Evaluator> -auto extractFunctors(const Evaluator& evaluator) - -> FunctorExtractor<Evaluator> { +auto inline extractFunctors(const Evaluator& evaluator)-> FunctorExtractor<Evaluator> { return FunctorExtractor<Evaluator>(evaluator); } } // namespace internal } // namespace TensorSycl } // namespace Eigen -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_FUNCTORS_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h index 77e0e15e1..8d520d2da 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h @@ -19,8 +19,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_LEAF_COUNT_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_LEAF_COUNT_HPP namespace Eigen { namespace TensorSycl { @@ -31,158 +31,81 @@ namespace internal { template <typename Expr> struct LeafCount; -/// specialisation of the \ref LeafCount struct when the node type is const -/// TensorMap -template <typename PlainObjectType, int Options_, - template <class> class MakePointer_> -struct LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_>> { - static const size_t Count = 1; -}; - -/// specialisation of the \ref LeafCount struct when the node type is TensorMap -template <typename PlainObjectType, int Options_, - template <class> class MakePointer_> -struct LeafCount<TensorMap<PlainObjectType, Options_, MakePointer_>> { - static const size_t Count = 1; -}; +template<typename... Args> struct CategoryCount; -/// specialisation of the \ref LeafCount struct when the node type is const -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr> -struct LeafCount<const TensorCwiseNullaryOp<OP, RHSExpr>> { - static const size_t Count = LeafCount<RHSExpr>::Count; +template<> struct CategoryCount<> +{ + static const size_t Count =0; }; -/// specialisation of the \ref LeafCount struct when the node type is -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr> -struct LeafCount<TensorCwiseNullaryOp<OP, RHSExpr>> { - static const size_t Count = LeafCount<RHSExpr>::Count; +template<typename Arg, typename... Args> +struct CategoryCount<Arg,Args...>{ + static const size_t Count = LeafCount<Arg>::Count + CategoryCount<Args...>::Count; }; /// specialisation of the \ref LeafCount struct when the node type is const -/// TensorBroadcastingOp -template <typename OP, typename RHSExpr> -struct LeafCount<const TensorBroadcastingOp<OP, RHSExpr>> { - static const size_t Count = LeafCount<RHSExpr>::Count; -}; - -/// specialisation of the \ref LeafCount struct when the node type is -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr> -struct LeafCount<TensorBroadcastingOp<OP, RHSExpr>> { - static const size_t Count = LeafCount<RHSExpr>::Count; -}; - -// TensorCwiseUnaryOp -template <typename OP, typename RHSExpr> -struct LeafCount<const TensorCwiseUnaryOp<OP, RHSExpr>> { - static const size_t Count = LeafCount<RHSExpr>::Count; -}; - -// TensorCwiseUnaryOp -template <typename OP, typename RHSExpr> -struct LeafCount<TensorCwiseUnaryOp<OP, RHSExpr>> { - static const size_t Count = LeafCount<RHSExpr>::Count; +/// TensorMap +template <typename PlainObjectType, int Options_, template <class> class MakePointer_> +struct LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_> > { + static const size_t Count =1; }; -/// specialisation of the \ref LeafCount struct when the node type is const -/// TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr> -struct LeafCount<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> { - static const size_t Count = - LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count; -}; +/// specialisation of the \ref LeafCount struct when the node type is TensorMap +template <typename PlainObjectType, int Options_, template <class> class MakePointer_> +struct LeafCount<TensorMap<PlainObjectType, Options_, MakePointer_> > :LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_> >{}; -/// specialisation of the \ref LeafCount struct when the node type is -/// TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr> -struct LeafCount<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> { - static const size_t Count = - LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count; -}; +// const TensorCwiseUnaryOp, const TensorCwiseNullaryOp, const TensorCwiseBinaryOp, const TensorCwiseTernaryOp, and Const TensorBroadcastingOp +template <template <class, class...> class CategoryExpr, typename OP, typename... RHSExpr> +struct LeafCount<const CategoryExpr<OP, RHSExpr...> >: CategoryCount<RHSExpr...> {}; +// TensorCwiseUnaryOp, TensorCwiseNullaryOp, TensorCwiseBinaryOp, TensorCwiseTernaryOp, and TensorBroadcastingOp +template <template <class, class...> class CategoryExpr, typename OP, typename... RHSExpr> +struct LeafCount<CategoryExpr<OP, RHSExpr...> > :LeafCount<const CategoryExpr<OP, RHSExpr...> >{}; /// specialisation of the \ref LeafCount struct when the node type is -/// TensorCwiseTernaryOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr> -struct LeafCount<TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>> { - static const size_t Count = LeafCount<Arg1Expr>::Count + - LeafCount<Arg2Expr>::Count + - LeafCount<Arg3Expr>::Count; -}; - -/// specialisation of the \ref LeafCount struct when the node type is const -/// TensorCwiseTernaryOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr> -struct LeafCount<const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>> { - static const size_t Count = LeafCount<Arg1Expr>::Count + - LeafCount<Arg2Expr>::Count + - LeafCount<Arg3Expr>::Count; -}; - +/// const TensorSelectOp is an exception +template <typename IfExpr, typename ThenExpr, typename ElseExpr> +struct LeafCount<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr> > : CategoryCount<IfExpr, ThenExpr, ElseExpr> {}; /// specialisation of the \ref LeafCount struct when the node type is -/// TensorCwiseSelectOp +/// TensorSelectOp template <typename IfExpr, typename ThenExpr, typename ElseExpr> -struct LeafCount<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> { - static const size_t Count = LeafCount<IfExpr>::Count + - LeafCount<ThenExpr>::Count + - LeafCount<ElseExpr>::Count; -}; +struct LeafCount<TensorSelectOp<IfExpr, ThenExpr, ElseExpr> >: LeafCount<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr> > {}; -/// specialisation of the \ref LeafCount struct when the node type is const -/// TensorCwiseSelectOp -template <typename IfExpr, typename ThenExpr, typename ElseExpr> -struct LeafCount<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> { - static const size_t Count = LeafCount<IfExpr>::Count + - LeafCount<ThenExpr>::Count + - LeafCount<ElseExpr>::Count; -}; -/// specialisation of the \ref LeafCount struct when the node type is +/// specialisation of the \ref LeafCount struct when the node type is const /// TensorAssignOp template <typename LHSExpr, typename RHSExpr> -struct LeafCount<TensorAssignOp<LHSExpr, RHSExpr>> { - static const size_t Count = - LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count; -}; +struct LeafCount<const TensorAssignOp<LHSExpr, RHSExpr> >: CategoryCount<LHSExpr,RHSExpr> {}; -/// specialisation of the \ref LeafCount struct when the node type is const -/// TensorAssignOp +/// specialisation of the \ref LeafCount struct when the node type is +/// TensorAssignOp is an exception. It is not the same as Unary template <typename LHSExpr, typename RHSExpr> -struct LeafCount<const TensorAssignOp<LHSExpr, RHSExpr>> { - static const size_t Count = - LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count; -}; +struct LeafCount<TensorAssignOp<LHSExpr, RHSExpr> > :LeafCount<const TensorAssignOp<LHSExpr, RHSExpr> >{}; /// specialisation of the \ref LeafCount struct when the node type is const /// TensorForcedEvalOp template <typename Expr> -struct LeafCount<const TensorForcedEvalOp<Expr>> { - static const size_t Count = 1; +struct LeafCount<const TensorForcedEvalOp<Expr> > { + static const size_t Count =1; }; /// specialisation of the \ref LeafCount struct when the node type is /// TensorForcedEvalOp template <typename Expr> -struct LeafCount<TensorForcedEvalOp<Expr>> { - static const size_t Count = 1; -}; +struct LeafCount<TensorForcedEvalOp<Expr> >: LeafCount<const TensorForcedEvalOp<Expr> > {}; /// specialisation of the \ref LeafCount struct when the node type is const /// TensorEvalToOp template <typename Expr> -struct LeafCount<const TensorEvalToOp<Expr>> { - static const size_t Count = 1 + LeafCount<Expr>::Count; +struct LeafCount<const TensorEvalToOp<Expr> > { + static const size_t Count = 1 + CategoryCount<Expr>::Count; }; /// specialisation of the \ref LeafCount struct when the node type is /// TensorEvalToOp template <typename Expr> -struct LeafCount<TensorEvalToOp<Expr>> { - static const size_t Count = 1 + LeafCount<Expr>::Count; -}; +struct LeafCount<TensorEvalToOp<Expr> >: LeafCount<const TensorEvalToOp<Expr> >{}; } } } // namespace Eigen -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_LEAF_COUNT_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h index 87995a25e..43a63c73d 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h @@ -19,8 +19,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_HPP namespace Eigen { namespace internal { @@ -31,91 +31,52 @@ namespace internal { template <typename Scalar, size_t N> struct PlaceHolder { static constexpr size_t I = N; - using Type = Scalar; + typedef Scalar Type; }; -template <typename PlainObjectType, int Options_, - template <class> class MakePointer_, size_t N> -struct PlaceHolder<const TensorMap<PlainObjectType, Options_, MakePointer_>, - N> { - static constexpr size_t I = N; - - using Type = const TensorMap<PlainObjectType, Options_, MakePointer_>; - - typedef typename Type::Self Self; - typedef typename Type::Base Base; - typedef typename Type::Nested Nested; - typedef typename Type::StorageKind StorageKind; - typedef typename Type::Index Index; - typedef typename Type::Scalar Scalar; - typedef typename Type::RealScalar RealScalar; - typedef typename Type::CoeffReturnType CoeffReturnType; +/// \brief specialisation of the PlaceHolder node for const TensorMap +#define TENSORMAPPLACEHOLDER(CVQual)\ +template <typename PlainObjectType, int Options_, template <class> class MakePointer_, size_t N>\ +struct PlaceHolder<CVQual TensorMap<PlainObjectType, Options_, MakePointer_>, N> {\ + static const size_t I = N;\ + typedef CVQual TensorMap<PlainObjectType, Options_, MakePointer_> Type;\ + typedef typename Type::Self Self;\ + typedef typename Type::Base Base;\ + typedef typename Type::Nested Nested;\ + typedef typename Type::StorageKind StorageKind;\ + typedef typename Type::Index Index;\ + typedef typename Type::Scalar Scalar;\ + typedef typename Type::RealScalar RealScalar;\ + typedef typename Type::CoeffReturnType CoeffReturnType;\ }; -/// \brief specialisation of the PlaceHolder node for TensorForcedEvalOp. The -/// TensorForcedEvalOp act as a leaf node for its parent node. -template <typename Expression, size_t N> -struct PlaceHolder<const TensorForcedEvalOp<Expression>, N> { - static constexpr size_t I = N; - - using Type = const TensorForcedEvalOp<Expression>; - - typedef typename Type::Nested Nested; - typedef typename Type::StorageKind StorageKind; - typedef typename Type::Index Index; - - typedef typename Type::Scalar Scalar; - typedef typename Type::Packet Packet; +TENSORMAPPLACEHOLDER(const) +TENSORMAPPLACEHOLDER() +#undef TENSORMAPPLACEHOLDER - typedef typename Type::RealScalar RealScalar; - typedef typename Type::CoeffReturnType CoeffReturnType; - typedef typename Type::PacketReturnType PacketReturnType; +/// \brief specialisation of the PlaceHolder node for TensorForcedEvalOp. The +/// TensorForcedEvalOp acts as a leaf node for its parent node. +#define TENSORFORCEDEVALPLACEHOLDER(CVQual)\ +template <typename Expression, size_t N>\ +struct PlaceHolder<CVQual TensorForcedEvalOp<Expression>, N> {\ + static const size_t I = N;\ + typedef CVQual TensorForcedEvalOp<Expression> Type;\ + typedef typename Type::Nested Nested;\ + typedef typename Type::StorageKind StorageKind;\ + typedef typename Type::Index Index;\ + typedef typename Type::Scalar Scalar;\ + typedef typename Type::Packet Packet;\ + typedef typename Type::RealScalar RealScalar;\ + typedef typename Type::CoeffReturnType CoeffReturnType;\ + typedef typename Type::PacketReturnType PacketReturnType;\ }; -template <typename Expression, size_t N> -struct PlaceHolder<TensorForcedEvalOp<Expression>, N> { - static constexpr size_t I = N; - - using Type = TensorForcedEvalOp<Expression>; - - typedef typename Type::Nested Nested; - typedef typename Type::StorageKind StorageKind; - typedef typename Type::Index Index; +TENSORFORCEDEVALPLACEHOLDER(const) +TENSORFORCEDEVALPLACEHOLDER() +#undef TENSORFORCEDEVALPLACEHOLDER - typedef typename Type::Scalar Scalar; - typedef typename Type::Packet Packet; - - typedef typename Type::RealScalar RealScalar; - typedef typename Type::CoeffReturnType CoeffReturnType; - typedef typename Type::PacketReturnType PacketReturnType; -}; - -/// \brief specialisation of the PlaceHolder node for const TensorMap -template <typename PlainObjectType, int Options_, - template <class> class Makepointer_, size_t N> -struct PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N> { - static constexpr size_t I = N; - - using Type = TensorMap<PlainObjectType, Options_, Makepointer_>; - - typedef typename Type::Self Self; - typedef typename Type::Base Base; - typedef typename Type::Nested Nested; - typedef typename Type::StorageKind StorageKind; - typedef typename Type::Index Index; - typedef typename Type::Scalar Scalar; - typedef typename Type::Packet Packet; - typedef typename Type::RealScalar RealScalar; - typedef typename Type::CoeffReturnType CoeffReturnType; - typedef typename Base::PacketReturnType PacketReturnType; -}; - -/// specialisation of the traits struct for PlaceHolder -template <typename PlainObjectType, int Options_, - template <class> class Makepointer_, size_t N> -struct traits< - PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N>> - : public traits<PlainObjectType> { +template <typename PlainObjectType, int Options_, template <class> class Makepointer_, size_t N> +struct traits<PlaceHolder<const TensorMap<PlainObjectType, Options_, Makepointer_>, N> >: public traits<PlainObjectType> { typedef traits<PlainObjectType> BaseTraits; typedef typename BaseTraits::Scalar Scalar; typedef typename BaseTraits::StorageKind StorageKind; @@ -128,24 +89,11 @@ struct traits< }; }; -template <typename PlainObjectType, int Options_, - template <class> class Makepointer_, size_t N> -struct traits< - PlaceHolder<const TensorMap<PlainObjectType, Options_, Makepointer_>, N>> - : public traits<PlainObjectType> { - typedef traits<PlainObjectType> BaseTraits; - typedef typename BaseTraits::Scalar Scalar; - typedef typename BaseTraits::StorageKind StorageKind; - typedef typename BaseTraits::Index Index; - static const int NumDimensions = BaseTraits::NumDimensions; - static const int Layout = BaseTraits::Layout; - enum { - Options = Options_, - Flags = BaseTraits::Flags, - }; -}; +template <typename PlainObjectType, int Options_, template <class> class Makepointer_, size_t N> +struct traits<PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N> > +: traits<PlaceHolder<const TensorMap<PlainObjectType, Options_, Makepointer_>, N> > {}; -} // end namespoace internal +} // end namespace internal } // end namespoace Eigen -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h index dbd7a8544..f456c35aa 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h @@ -19,8 +19,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_EXPR_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_EXPR_HPP namespace Eigen { namespace TensorSycl { @@ -32,262 +32,127 @@ namespace internal { template <typename Expr, size_t N> struct PlaceHolderExpression; -/// specialisation of the \ref PlaceHolderExpression when the node is TensorMap -template <typename Scalar_, int Options_, int Options2_, int NumIndices_, - typename IndexType_, template <class> class MakePointer_, size_t N> -struct PlaceHolderExpression< - Eigen::TensorMap<Eigen::Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakePointer_>, - N> { - using Type = Eigen::internal::PlaceHolder< - Eigen::TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakePointer_>, - N>; -}; +template<size_t N, typename... Args> +struct CalculateIndex; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorMap -template <typename Scalar_, int Options_, int Options2_, int NumIndices_, - typename IndexType_, template <class> class MakePointer_, size_t N> -struct PlaceHolderExpression< - const Eigen::TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakePointer_>, - N> { - using Type = const Eigen::internal::PlaceHolder< - const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, - Options2_, MakePointer_>, - N>; +template<size_t N, typename Arg> +struct CalculateIndex<N, Arg>{ + typedef typename PlaceHolderExpression<Arg, N>::Type ArgType; + typedef utility::tuple::Tuple<ArgType> ArgsTuple; }; -/// specialisation of the \ref PlaceHolderExpression when the node is -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr, size_t N> -struct PlaceHolderExpression<TensorCwiseNullaryOp<OP, RHSExpr>, N> { - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - using Type = TensorCwiseNullaryOp<OP, RHSPlaceHolderType>; +template<size_t N, typename Arg1, typename Arg2> +struct CalculateIndex<N, Arg1, Arg2>{ + static const size_t Arg2LeafCount = LeafCount<Arg2>::Count; + typedef typename PlaceHolderExpression<Arg1, N - Arg2LeafCount>::Type Arg1Type; + typedef typename PlaceHolderExpression<Arg2, N>::Type Arg2Type; + typedef utility::tuple::Tuple<Arg1Type, Arg2Type> ArgsTuple; }; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorCwiseNullaryOp -template <typename OP, typename RHSExpr, size_t N> -struct PlaceHolderExpression<const TensorCwiseNullaryOp<OP, RHSExpr>, N> { - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - using Type = const TensorCwiseNullaryOp<OP, RHSPlaceHolderType>; +template<size_t N, typename Arg1, typename Arg2, typename Arg3> +struct CalculateIndex<N, Arg1, Arg2, Arg3> { + static const size_t Arg3LeafCount = LeafCount<Arg3>::Count; + static const size_t Arg2LeafCount = LeafCount<Arg2>::Count; + typedef typename PlaceHolderExpression<Arg1, N - Arg3LeafCount - Arg2LeafCount>::Type Arg1Type; + typedef typename PlaceHolderExpression<Arg2, N - Arg3LeafCount>::Type Arg2Type; + typedef typename PlaceHolderExpression<Arg3, N>::Type Arg3Type; + typedef utility::tuple::Tuple<Arg1Type, Arg2Type, Arg3Type> ArgsTuple; }; -/// specialisation of the \ref PlaceHolderExpression when the node is -/// TensorBroadcastingOp -template <typename OP, typename RHSExpr, size_t N> -struct PlaceHolderExpression<TensorBroadcastingOp<OP, RHSExpr>, N> { - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - using Type = TensorBroadcastingOp<OP, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorBroadcastingOp -template <typename OP, typename RHSExpr, size_t N> -struct PlaceHolderExpression<const TensorBroadcastingOp<OP, RHSExpr>, N> { - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - using Type = const TensorBroadcastingOp<OP, RHSPlaceHolderType>; -}; +template<template<class...> class Category , class OP, class TPL> +struct CategoryHelper; -/// specialisation of the \ref PlaceHolderExpression when the node is -/// TensorCwiseUnaryOp -template <typename OP, typename RHSExpr, size_t N> -struct PlaceHolderExpression<TensorCwiseUnaryOp<OP, RHSExpr>, N> { - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - using Type = TensorCwiseUnaryOp<OP, RHSPlaceHolderType>; +template<template<class...> class Category , class OP, class ...T > +struct CategoryHelper<Category, OP, utility::tuple::Tuple<T...> > { + typedef Category<OP, T... > Type; }; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorCwiseUnaryOp -template <typename OP, typename RHSExpr, size_t N> -struct PlaceHolderExpression<const TensorCwiseUnaryOp<OP, RHSExpr>, N> { - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - using Type = const TensorCwiseUnaryOp<OP, RHSPlaceHolderType>; +template<template<class...> class Category , class ...T > +struct CategoryHelper<Category, NoOP, utility::tuple::Tuple<T...> > { + typedef Category<T... > Type; }; /// specialisation of the \ref PlaceHolderExpression when the node is -/// TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr, size_t N> -struct PlaceHolderExpression<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, N> { - static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count; - - using LHSPlaceHolderType = - typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type; - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - - using Type = TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>; +/// TensorCwiseNullaryOp, TensorCwiseUnaryOp, TensorBroadcastingOp, TensorCwiseBinaryOp, TensorCwiseTernaryOp +#define OPEXPRCATEGORY(CVQual)\ +template <template <class, class... > class Category, typename OP, typename... SubExpr, size_t N>\ +struct PlaceHolderExpression<CVQual Category<OP, SubExpr...>, N>{\ + typedef CVQual typename CategoryHelper<Category, OP, typename CalculateIndex<N, SubExpr...>::ArgsTuple>::Type Type;\ }; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorCwiseBinaryOp -template <typename OP, typename LHSExpr, typename RHSExpr, size_t N> -struct PlaceHolderExpression<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, - N> { - static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count; - - using LHSPlaceHolderType = - typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type; - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - - using Type = - const TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>; -}; - -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorCwiseSelectOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, - size_t N> -struct PlaceHolderExpression< - const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, N> { - static const size_t Arg3LeafCount = LeafCount<Arg3Expr>::Count; - static const size_t Arg2LeafCount = LeafCount<Arg2Expr>::Count; - - using Arg1PlaceHolderType = - typename PlaceHolderExpression<Arg1Expr, - N - Arg3LeafCount - Arg2LeafCount>::Type; - using Arg2PlaceHolderType = - typename PlaceHolderExpression<Arg2Expr, N - Arg3LeafCount>::Type; - - using Arg3PlaceHolderType = typename PlaceHolderExpression<Arg3Expr, N>::Type; - - using Type = - const TensorCwiseTernaryOp<OP, Arg1PlaceHolderType, Arg2PlaceHolderType, - Arg3PlaceHolderType>; -}; +OPEXPRCATEGORY(const) +OPEXPRCATEGORY() +#undef OPEXPRCATEGORY /// specialisation of the \ref PlaceHolderExpression when the node is /// TensorCwiseSelectOp -template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, - size_t N> -struct PlaceHolderExpression< - TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, N> { - static const size_t Arg3LeafCount = LeafCount<Arg3Expr>::Count; - static const size_t Arg2LeafCount = LeafCount<Arg2Expr>::Count; - - using Arg1PlaceHolderType = - typename PlaceHolderExpression<Arg1Expr, - N - Arg3LeafCount - Arg2LeafCount>::Type; - using Arg2PlaceHolderType = - typename PlaceHolderExpression<Arg2Expr, N - Arg3LeafCount>::Type; - - using Arg3PlaceHolderType = typename PlaceHolderExpression<Arg3Expr, N>::Type; - - using Type = TensorCwiseTernaryOp<OP, Arg1PlaceHolderType, - Arg2PlaceHolderType, Arg3PlaceHolderType>; -}; - -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorCwiseSelectOp -template <typename IfExpr, typename ThenExpr, typename ElseExpr, size_t N> -struct PlaceHolderExpression<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, - N> { - static const size_t ElseLeafCount = LeafCount<ElseExpr>::Count; - static const size_t ThenLeafCount = LeafCount<ThenExpr>::Count; - - using IfPlaceHolderType = - typename PlaceHolderExpression<IfExpr, - N - ElseLeafCount - ThenLeafCount>::Type; - using ThenPlaceHolderType = - typename PlaceHolderExpression<ThenExpr, N - ElseLeafCount>::Type; - - using ElsePlaceHolderType = typename PlaceHolderExpression<ElseExpr, N>::Type; - - using Type = const TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType, - ElsePlaceHolderType>; +#define SELECTEXPR(CVQual)\ +template <typename IfExpr, typename ThenExpr, typename ElseExpr, size_t N>\ +struct PlaceHolderExpression<CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, N> {\ + typedef CVQual typename CategoryHelper<TensorSelectOp, NoOP, typename CalculateIndex<N, IfExpr, ThenExpr, ElseExpr>::ArgsTuple>::Type Type;\ }; -/// specialisation of the \ref PlaceHolderExpression when the node is -/// TensorCwiseSelectOp -template <typename IfExpr, typename ThenExpr, typename ElseExpr, size_t N> -struct PlaceHolderExpression<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, N> { - static const size_t ElseLeafCount = LeafCount<ElseExpr>::Count; - static const size_t ThenLeafCount = LeafCount<ThenExpr>::Count; - - using IfPlaceHolderType = - typename PlaceHolderExpression<IfExpr, - N - ElseLeafCount - ThenLeafCount>::Type; - using ThenPlaceHolderType = - typename PlaceHolderExpression<ThenExpr, N - ElseLeafCount>::Type; - - using ElsePlaceHolderType = typename PlaceHolderExpression<ElseExpr, N>::Type; - - using Type = TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType, - ElsePlaceHolderType>; -}; +SELECTEXPR(const) +SELECTEXPR() +#undef SELECTEXPR /// specialisation of the \ref PlaceHolderExpression when the node is /// TensorAssignOp -template <typename LHSExpr, typename RHSExpr, size_t N> -struct PlaceHolderExpression<TensorAssignOp<LHSExpr, RHSExpr>, N> { - static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count; - - using LHSPlaceHolderType = - typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type; - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; - - using Type = TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>; +#define ASSIGNEXPR(CVQual)\ +template <typename LHSExpr, typename RHSExpr, size_t N>\ +struct PlaceHolderExpression<CVQual TensorAssignOp<LHSExpr, RHSExpr>, N> {\ + typedef CVQual typename CategoryHelper<TensorAssignOp, NoOP, typename CalculateIndex<N, LHSExpr, RHSExpr>::ArgsTuple>::Type Type;\ }; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorAssignOp -template <typename LHSExpr, typename RHSExpr, size_t N> -struct PlaceHolderExpression<const TensorAssignOp<LHSExpr, RHSExpr>, N> { - static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count; - - using LHSPlaceHolderType = - typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type; - using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type; +ASSIGNEXPR(const) +ASSIGNEXPR() +#undef ASSIGNEXPR - using Type = const TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>; +/// specialisation of the \ref PlaceHolderExpression when the node is +/// TensorMap +#define TENSORMAPEXPR(CVQual)\ +template <typename Scalar_, int Options_, int Options2_, int NumIndices_, typename IndexType_, template <class> class MakePointer_, size_t N>\ +struct PlaceHolderExpression< CVQual TensorMap< Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_>, N> {\ + typedef CVQual Eigen::internal::PlaceHolder<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_>, N> Type;\ }; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorForcedEvalOp -template <typename Expr, size_t N> -struct PlaceHolderExpression<const TensorForcedEvalOp<Expr>, N> { - using Type = - const Eigen::internal::PlaceHolder<const TensorForcedEvalOp<Expr>, N>; -}; +TENSORMAPEXPR(const) +TENSORMAPEXPR() +#undef TENSORMAPEXPR /// specialisation of the \ref PlaceHolderExpression when the node is /// TensorForcedEvalOp -template <typename Expr, size_t N> -struct PlaceHolderExpression<TensorForcedEvalOp<Expr>, N> { - using Type = Eigen::internal::PlaceHolder<TensorForcedEvalOp<Expr>, N>; +#define FORCEDEVAL(CVQual)\ +template <typename Expr, size_t N>\ +struct PlaceHolderExpression<CVQual TensorForcedEvalOp<Expr>, N> {\ + typedef CVQual Eigen::internal::PlaceHolder<CVQual TensorForcedEvalOp<Expr>, N> Type;\ }; -/// specialisation of the \ref PlaceHolderExpression when the node is const -/// TensorEvalToOp -template <typename Expr, size_t N> -struct PlaceHolderExpression<const TensorEvalToOp<Expr>, N> { - static const size_t RHSLeafCount = LeafCount<Expr>::Count; - - using RHSPlaceHolderType = typename PlaceHolderExpression<Expr, N>::Type; - - using Type = const TensorEvalToOp<RHSPlaceHolderType>; -}; +FORCEDEVAL(const) +FORCEDEVAL() +#undef FORCEDEVAL /// specialisation of the \ref PlaceHolderExpression when the node is /// TensorEvalToOp -template <typename Expr, size_t N> -struct PlaceHolderExpression<TensorEvalToOp<Expr>, N> { - static const size_t RHSLeafCount = LeafCount<Expr>::Count; - - using RHSPlaceHolderType = typename PlaceHolderExpression<Expr, N>::Type; - - using Type = TensorEvalToOp<RHSPlaceHolderType>; +#define EVALTO(CVQual)\ +template <typename Expr, size_t N>\ +struct PlaceHolderExpression<CVQual TensorEvalToOp<Expr>, N> {\ + typedef CVQual TensorEvalToOp<typename CalculateIndex <N, Expr>::ArgType> Type;\ }; +EVALTO(const) +EVALTO() +#undef EVALTO + /// template deduction for \ref PlaceHolderExpression struct template <typename Expr> struct createPlaceHolderExpression { static const size_t TotalLeaves = LeafCount<Expr>::Count; - using Type = typename PlaceHolderExpression<Expr, TotalLeaves - 1>::Type; + typedef typename PlaceHolderExpression<Expr, TotalLeaves - 1>::Type Type; }; + } } } // namespace Eigen -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_EXPR_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h index 3758d46a0..57f2dda26 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h @@ -20,8 +20,8 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP namespace Eigen { namespace TensorSycl { @@ -34,17 +34,14 @@ void run(Expr &expr, Dev &dev) { Eigen::TensorEvaluator<Expr, Dev> evaluator(expr, dev); const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL); if (needs_assign) { - using PlaceHolderExpr = - typename internal::createPlaceHolderExpression<Expr>::Type; + typedef typename internal::createPlaceHolderExpression<Expr>::Type PlaceHolderExpr; auto functors = internal::extractFunctors(evaluator); dev.m_queue.submit([&](cl::sycl::handler &cgh) { // create a tuple of accessors from Evaluator - auto tuple_of_accessors = - internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator); - const auto range = - utility::tuple::get<0>(tuple_of_accessors).get_range()[0]; + auto tuple_of_accessors = internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator); + const auto range = utility::tuple::get<0>(tuple_of_accessors).get_range()[0]; size_t outTileSize = range; if (range > 64) outTileSize = 64; @@ -53,26 +50,14 @@ void run(Expr &expr, Dev &dev) { if (yMode != 0) yRange += (outTileSize - yMode); // run the kernel - cgh.parallel_for<PlaceHolderExpr>( - cl::sycl::nd_range<1>(cl::sycl::range<1>(yRange), - cl::sycl::range<1>(outTileSize)), - [=](cl::sycl::nd_item<1> itemID) { - using DevExpr = - typename internal::ConvertToDeviceExpression<Expr>::Type; - - auto device_expr = - internal::createDeviceExpression<DevExpr, PlaceHolderExpr>( - functors, tuple_of_accessors); - auto device_evaluator = - Eigen::TensorEvaluator<decltype(device_expr.expr), - Eigen::DefaultDevice>( - device_expr.expr, Eigen::DefaultDevice()); - - if (itemID.get_global_linear_id() < range) { - device_evaluator.evalScalar( - static_cast<int>(itemID.get_global_linear_id())); - } - }); + cgh.parallel_for<PlaceHolderExpr>( cl::sycl::nd_range<1>(cl::sycl::range<1>(yRange), cl::sycl::range<1>(outTileSize)), [=](cl::sycl::nd_item<1> itemID) { + typedef typename internal::ConvertToDeviceExpression<Expr>::Type DevExpr; + auto device_expr =internal::createDeviceExpression<DevExpr, PlaceHolderExpr>(functors, tuple_of_accessors); + auto device_evaluator = Eigen::TensorEvaluator<decltype(device_expr.expr), Eigen::DefaultDevice>(device_expr.expr, Eigen::DefaultDevice()); + if (itemID.get_global_linear_id() < range) { + device_evaluator.evalScalar(static_cast<int>(itemID.get_global_linear_id())); + } + }); }); dev.m_queue.throw_asynchronous(); } @@ -81,4 +66,4 @@ void run(Expr &expr, Dev &dev) { } // namespace TensorSycl } // namespace Eigen -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h index 8b9fc52c4..b443ebe1d 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h @@ -18,24 +18,23 @@ * *****************************************************************/ -#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP -#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP +#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_TUPLE_HPP +#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_TUPLE_HPP namespace utility { namespace tuple { -/// \struct EnableIf -/// \brief The EnableIf struct is used to statically define type based on the +/// \struct StaticIf +/// \brief The StaticIf struct is used to statically choose the type based on the /// condition. -template <bool, typename T = void> -struct EnableIf {}; -/// \brief specialisation of the \ref EnableIf when the condition is true +template <bool, typename T = void> struct StaticIf; +/// \brief specialisation of the \ref StaticIf when the condition is true template <typename T> -struct EnableIf<true, T> { +struct StaticIf<true, T> { typedef T type; }; /// \struct Tuple /// \brief is a fixed-size collection of heterogeneous values -/// \ztparam Ts... - the types of the elements that the tuple stores. +/// \ztparam Ts... - the types of the elements that the tuple stores. /// Empty list is supported. template <class... Ts> struct Tuple {}; @@ -47,12 +46,11 @@ struct Tuple {}; template <class T, class... Ts> struct Tuple<T, Ts...> { Tuple(T t, Ts... ts) : head(t), tail(ts...) {} - T head; Tuple<Ts...> tail; }; -/// \struct ElemTypeHolder +///\ struct ElemTypeHolder /// \brief ElemTypeHolder class is used to specify the types of the /// elements inside the tuple /// \tparam size_t the number of elements inside the tuple @@ -60,78 +58,59 @@ struct Tuple<T, Ts...> { template <size_t, class> struct ElemTypeHolder; -/// \brief specialisation of the \ref ElemTypeHolder class when the number +/// \brief specialisation of the \ref ElemTypeHolder class when the number of /// elements inside the tuple is 1 template <class T, class... Ts> -struct ElemTypeHolder<0, Tuple<T, Ts...>> { +struct ElemTypeHolder<0, Tuple<T, Ts...> > { typedef T type; }; /// \brief specialisation of the \ref ElemTypeHolder class when the number of -/// elements inside the tuple is bigger than 1. It recursively call itself to +/// elements inside the tuple is bigger than 1. It recursively calls itself to /// detect the type of each element in the tuple /// \tparam T : the type of the first element in the tuple. /// \tparam Ts... the rest of the elements in the tuple. Ts... can be empty. /// \tparam K is the Kth element in the tuple template <size_t k, class T, class... Ts> -struct ElemTypeHolder<k, Tuple<T, Ts...>> { - typedef typename ElemTypeHolder<k - 1, Tuple<Ts...>>::type type; +struct ElemTypeHolder<k, Tuple<T, Ts...> > { + typedef typename ElemTypeHolder<k - 1, Tuple<Ts...> >::type type; }; /// get /// \brief Extracts the first element from the tuple. /// K=0 represents the first element of the tuple. The tuple cannot be empty. -/// \tparam Ts... are the elements type in the tuple. +/// \tparam Ts... are the type of the elements in the tuple. /// \param t is the tuple whose contents to extract -/// \return typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type -template <size_t k, class... Ts> -typename EnableIf<k == 0, - typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type -get(Tuple<Ts...> &t) { - return t.head; +/// \return typename ElemTypeHolder<0, Tuple<Ts...> >::type &>::type + +#define TERMINATE_CONDS_TUPLE_GET(CVQual) \ +template <size_t k, class... Ts> \ +typename StaticIf<k == 0, CVQual typename ElemTypeHolder<0, Tuple<Ts...> >::type &>::type \ +get(CVQual Tuple<Ts...> &t) { \ + static_assert(sizeof...(Ts)!=0, "The requseted value is bigger than the size of the tuple"); \ + return t.head; \ } + +TERMINATE_CONDS_TUPLE_GET(const) +TERMINATE_CONDS_TUPLE_GET() +#undef TERMINATE_CONDS_TUPLE_GET /// get /// \brief Extracts the Kth element from the tuple. -/// \tparam K is an integer value in [0,sizeof...(Types)). +///\tparam K is an integer value in [0,sizeof...(Types)). /// \tparam T is the (sizeof...(Types) -(K+1)) element in the tuple -/// \tparam Ts... are the elements type in the tuple. +/// \tparam Ts... are the type of the elements in the tuple. /// \param t is the tuple whose contents to extract -/// \return typename ElemTypeHolder<K, Tuple<Ts...>>::type &>::type -template <size_t k, class T, class... Ts> -typename EnableIf<k != 0, - typename ElemTypeHolder<k, Tuple<T, Ts...>>::type &>::type -get(Tuple<T, Ts...> &t) { - return get<k - 1>(t.tail); +/// \return typename ElemTypeHolder<K, Tuple<Ts...> >::type &>::type +#define RECURSIVE_TUPLE_GET(CVQual) \ +template <size_t k, class T, class... Ts> \ +typename StaticIf<k != 0, CVQual typename ElemTypeHolder<k, Tuple<T, Ts...> >::type &>::type \ +get(CVQual Tuple<T, Ts...> &t) { \ + return get<k - 1>(t.tail); \ } +RECURSIVE_TUPLE_GET(const) +RECURSIVE_TUPLE_GET() +#undef RECURSIVE_TUPLE_GET -/// get -/// \brief Extracts the first element from the tuple when the tuple and all the -/// elements inside are const. -/// K=0 represents the first element of the tuple. The tuple cannot be empty. -/// \tparam Ts... are the elements type in the tuple. -/// \param t is the const tuple whose contents to extract -/// \return const typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type -template <size_t k, class... Ts> -typename EnableIf<k == 0, - const typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type -get(const Tuple<Ts...> &t) { - return t.head; -} - -/// get -/// \brief Extracts the Kth element from the tuple when the tuple and all the -/// elements inside are const. -/// \tparam K is an integer value in [0,sizeof...(Types)). -/// \tparam T is the (sizeof...(Types) -(K+1)) element in the tuple -/// \tparam Ts... are the elements type in the tuple. -/// \param t is the const tuple whose contents to extract -/// \return const typename ElemTypeHolder<K, Tuple<Ts...>>::type &>::type -template <size_t k, class T, class... Ts> -typename EnableIf< - k != 0, const typename ElemTypeHolder<k, Tuple<T, Ts...>>::type &>::type -get(const Tuple<T, Ts...> &t) { - return get<k - 1>(t.tail); -} /// make_tuple /// \brief Creates a tuple object, deducing the target type from the types of /// arguments. @@ -153,17 +132,17 @@ static constexpr size_t size(Tuple<Args...> &) { return sizeof...(Args); } -/// \struct Index_list +/// \struct IndexList /// \brief Creates a list of index from the elements in the tuple /// \tparam Is... a list of index from [0 to sizeof...(tuple elements)) template <size_t... Is> -struct Index_list {}; +struct IndexList {}; /// \struct RangeBuilder /// \brief Collects internal details for generating index ranges [MIN, MAX) /// Declare primary template for index range builder /// \tparam MIN is the starting index in the tuple -/// \tparam N represents sizeof..(elements)- sizeof...(Is) +/// \tparam N represents sizeof..(elemens)- sizeof...(Is) /// \tparam Is... are the list of generated index so far template <size_t MIN, size_t N, size_t... Is> struct RangeBuilder; @@ -174,14 +153,14 @@ struct RangeBuilder; /// \tparam Is is [0 to sizeof...(tuple elements)) template <size_t MIN, size_t... Is> struct RangeBuilder<MIN, MIN, Is...> { - typedef Index_list<Is...> type; + typedef IndexList<Is...> type; }; /// Induction step: Specialisation of the RangeBuilder class when N!=MIN -/// in this case we are recursively subtracting the N by one and adding one +/// in this case we are recursively subtracting N by one and adding one /// index to Is... list until MIN==N /// \tparam MIN is the starting index in the tuple -/// \tparam N represents sizeof..(elements)- sizeof...(Is) +/// \tparam N represents sizeof..(elemens)- sizeof...(Is) /// \tparam Is... are the list of generated index so far template <size_t MIN, size_t N, size_t... Is> struct RangeBuilder : public RangeBuilder<MIN, N - 1, N - 1, Is...> {}; @@ -190,75 +169,66 @@ struct RangeBuilder : public RangeBuilder<MIN, N - 1, N - 1, Is...> {}; /// \tparam MIN is the starting index in the tuple /// \tparam MAX is the size of the tuple template <size_t MIN, size_t MAX> -using Index_range = typename RangeBuilder<MIN, MAX>::type; +struct IndexRange: RangeBuilder<MIN, MAX>::type {}; -/// append_impl +/// append_base /// \brief unpacking the elements of the input tuple t and creating a new tuple /// by adding element a at the end of it. -/// \tparam Args... the type of the elements inside the tuple t +///\tparam Args... the type of the elements inside the tuple t /// \tparam T the type of the new element going to be added at the end of tuple /// \tparam I... is the list of index from [0 to sizeof...(t)) /// \param t the tuple on which we want to append a. /// \param a the new elements going to be added to the tuple /// \return Tuple<Args..., T> template <typename... Args, typename T, size_t... I> -Tuple<Args..., T> append_impl(utility::tuple::Tuple<Args...> t, T a, - utility::tuple::Index_list<I...>) { - return utility::tuple::make_tuple(get<I>(t)..., a); +Tuple<Args..., T> append_base(Tuple<Args...> t, T a,IndexList<I...>) { + return make_tuple(get<I>(t)..., a); } /// append -/// \brief the deduction function for \ref append_impl that automatically -/// generate the \ref Index_range -/// \tparam Args... the type of the elements inside the tuple t +/// \brief the deduction function for \ref append_base that automatically +/// generate the \ref IndexRange +///\tparam Args... the type of the elements inside the tuple t /// \tparam T the type of the new element going to be added at the end of tuple /// \param t the tuple on which we want to append a. /// \param a the new elements going to be added to the tuple /// \return Tuple<Args..., T> template <typename... Args, typename T> Tuple<Args..., T> append(Tuple<Args...> t, T a) { - return utility::tuple::append_impl( - t, a, utility::tuple::Index_range<0, sizeof...(Args)>()); + return append_base(t, a, IndexRange<0, sizeof...(Args)>()); } -/// append_impl -/// \brief This is an specialised of \ref append_impl when we want to +/// append_base +/// \brief This is a specialisation of \ref append_base when we want to /// concatenate -/// tuple t2 at the end of the tuple t1. Here we unpack both tuples, generate -/// the -/// Index_range for each of them and create an output tuple T that contains both +/// tuple t2 at the end of the tuple t1. Here we unpack both tuples, generate the +/// IndexRange for each of them and create an output tuple T that contains both /// elements of t1 and t2. -/// \tparam Args1... the type of the elements inside the tuple t1 -/// \tparam Args2... the type of the elements inside the tuple t2 +///\tparam Args1... the type of the elements inside the tuple t1 +///\tparam Args2... the type of the elements inside the tuple t2 /// \tparam I1... is the list of index from [0 to sizeof...(t1)) /// \tparam I2... is the list of index from [0 to sizeof...(t2)) /// \param t1 is the tuple on which we want to append t2. /// \param t2 is the tuple that is going to be added on t1. /// \return Tuple<Args1..., Args2...> template <typename... Args1, typename... Args2, size_t... I1, size_t... I2> -Tuple<Args1..., Args2...> append_impl(utility::tuple::Tuple<Args1...> t1, - utility::tuple::Tuple<Args2...> t2, - utility::tuple::Index_list<I1...>, - utility::tuple::Index_list<I2...>) { - return utility::tuple::make_tuple(utility::tuple::get<I1>(t1)..., - utility::tuple::get<I2>(t2)...); +Tuple<Args1..., Args2...> append_base(Tuple<Args1...> t1, Tuple<Args2...> t2, IndexList<I1...>, IndexList<I2...>) { + return make_tuple(get<I1>(t1)...,get<I2>(t2)...); } + /// append -/// \brief deduction function for \ref append_impl when we are appending tuple -/// t1 by tuple t2. In this case the \ref Index_range for both tuple are +/// \brief deduction function for \ref append_base when we are appending tuple +/// t1 by tuple t2. In this case the \ref IndexRange for both tuple are /// automatically generated. -/// \tparam Args1... the type of the elements inside the tuple t1 -/// \tparam Args2... the type of the elements inside the tuple t2 +///\tparam Args1... the type of the elements inside the tuple t1 +///\tparam Args2... the type of the elements inside the tuple t2 /// \param t1 is the tuple on which we want to append t2. /// \param t2 is the tuple that is going to be added on t1. /// \return Tuple<Args1..., Args2...> template <typename... Args1, typename... Args2> -Tuple<Args1..., Args2...> append(utility::tuple::Tuple<Args1...> t1, - utility::tuple::Tuple<Args2...> t2) { - return utility::tuple::append_impl( - t1, t2, utility::tuple::Index_range<0, sizeof...(Args1)>(), - utility::tuple::Index_range<0, sizeof...(Args2)>()); +Tuple<Args1..., Args2...> append(Tuple<Args1...> t1,Tuple<Args2...> t2) { + return append_base(t1, t2, IndexRange<0, sizeof...(Args1)>(), IndexRange<0, sizeof...(Args2)>()); } } // tuple } // utility -#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP +#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_TUPLE_HPP diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index 17073dfa7..aeaea162a 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt @@ -140,6 +140,12 @@ endif() endif() if(EIGEN_TEST_CXX11) + if(EIGEN_TEST_SYCL) + ei_add_test_sycl(cxx11_tensor_sycl "-std=c++11") + ei_add_test_sycl(cxx11_tensor_forced_eval_sycl "-std=c++11") + ei_add_test_sycl(cxx11_tensor_broadcast_sycl "-std=c++11") + ei_add_test_sycl(cxx11_tensor_device_sycl "-std=c++11") + endif(EIGEN_TEST_SYCL) # It should be safe to always run these tests as there is some fallback code for # older compiler that don't support cxx11. set(CMAKE_CXX_STANDARD 11) @@ -248,11 +254,3 @@ if(CUDA_FOUND AND EIGEN_TEST_CUDA) unset(EIGEN_ADD_TEST_FILENAME_EXTENSION) endif() - - -if(EIGEN_TEST_SYCL) - ei_add_test_sycl(cxx11_tensor_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_forced_eval_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_broadcast_sycl "-std=c++11") - ei_add_test_sycl(cxx11_tensor_device_sycl "-std=c++11") -endif(EIGEN_TEST_SYCL)
\ No newline at end of file diff --git a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp index 0179ab243..ecebf7d68 100644 --- a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp +++ b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp @@ -28,7 +28,7 @@ using Eigen::TensorMap; // Types used in tests: using TestTensor = Tensor<float, 3>; using TestTensorMap = TensorMap<Tensor<float, 3>>; -static void test_sycl_broadcast(){ +static void test_broadcast_sycl(){ cl::sycl::gpu_selector s; cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { @@ -75,5 +75,5 @@ static void test_sycl_broadcast(){ } void test_cxx11_tensor_broadcast_sycl() { - CALL_SUBTEST(test_sycl_broadcast()); + CALL_SUBTEST(test_broadcast_sycl()); } diff --git a/unsupported/test/cxx11_tensor_device_sycl.cpp b/unsupported/test/cxx11_tensor_device_sycl.cpp index b9368cb84..f54fc8786 100644 --- a/unsupported/test/cxx11_tensor_device_sycl.cpp +++ b/unsupported/test/cxx11_tensor_device_sycl.cpp @@ -20,7 +20,7 @@ #include "main.h" #include <unsupported/Eigen/CXX11/Tensor> -void test_sycl_device() { +void test_device_sycl() { cl::sycl::gpu_selector s; cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { for (const auto& e : l) { @@ -31,9 +31,9 @@ void test_sycl_device() { } } }); - SyclDevice sycl_device(q); + Eigen::SyclDevice sycl_device(q); printf("Helo from ComputeCpp: Device Exists\n"); } void test_cxx11_tensor_device_sycl() { - CALL_SUBTEST(test_sycl_device()); + CALL_SUBTEST(test_device_sycl()); } diff --git a/unsupported/test/cxx11_tensor_forced_eval_sycl.cpp b/unsupported/test/cxx11_tensor_forced_eval_sycl.cpp index 7c4446192..182ec7fa8 100644 --- a/unsupported/test/cxx11_tensor_forced_eval_sycl.cpp +++ b/unsupported/test/cxx11_tensor_forced_eval_sycl.cpp @@ -22,7 +22,7 @@ using Eigen::Tensor; -void test_sycl_gpu() { +void test_forced_eval_sycl() { cl::sycl::gpu_selector s; cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { for (const auto& e : l) { @@ -65,6 +65,4 @@ void test_sycl_gpu() { printf("(a+b)*b Test Passed\n"); } -void test_cxx11_tensor_forced_eval_sycl() { - CALL_SUBTEST(test_sycl_gpu()); -} +void test_cxx11_tensor_forced_eval_sycl() { CALL_SUBTEST(test_forced_eval_sycl()); } diff --git a/unsupported/test/cxx11_tensor_sycl.cpp b/unsupported/test/cxx11_tensor_sycl.cpp index a4dc9f9fa..0f66cd8f0 100644 --- a/unsupported/test/cxx11_tensor_sycl.cpp +++ b/unsupported/test/cxx11_tensor_sycl.cpp @@ -6,11 +6,13 @@ // Ralph Potter Codeplay Software Ltd. // Luke Iwanski Codeplay Software Ltd. // Contact: <eigen@codeplay.com> +// Benoit Steiner <benoit.steiner.goog@gmail.com> // // This Source Code Form is subject to the terms of the Mozilla // Public License v. 2.0. If a copy of the MPL was not distributed // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. + #define EIGEN_TEST_NO_LONGDOUBLE #define EIGEN_TEST_NO_COMPLEX #define EIGEN_TEST_FUNC cxx11_tensor_sycl @@ -30,7 +32,7 @@ using TestTensor = Tensor<float, 3>; using TestTensorMap = TensorMap<Tensor<float, 3>>; void test_sycl_cpu() { - cl::sycl::gpu_selector s; + cl::sycl::gpu_selector s; cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { for (const auto& e : l) { try { @@ -40,7 +42,7 @@ void test_sycl_cpu() { } } }); - SyclDevice sycl_device(q); + SyclDevice sycl_device(q); int sizeDim1 = 100; int sizeDim2 = 100; @@ -53,14 +55,14 @@ void test_sycl_cpu() { in1 = in1.random(); in2 = in2.random(); in3 = in3.random(); - TestTensorMap gpu_in1(in1.data(), tensorRange); - TestTensorMap gpu_in2(in2.data(), tensorRange); - TestTensorMap gpu_in3(in3.data(), tensorRange); - TestTensorMap gpu_out(out.data(), tensorRange); - - /// a=1.2f - gpu_in1.device(sycl_device) = gpu_in1.constant(1.2f); - sycl_device.deallocate(in1.data()); + TestTensorMap gpu_in1(in1.data(), tensorRange); + TestTensorMap gpu_in2(in2.data(), tensorRange); + TestTensorMap gpu_in3(in3.data(), tensorRange); + TestTensorMap gpu_out(out.data(), tensorRange); + + /// a=1.2f + gpu_in1.device(sycl_device) = gpu_in1.constant(1.2f); + sycl_device.deallocate(in1.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -68,11 +70,11 @@ void test_sycl_cpu() { } } } - printf("a=1.2f Test passed\n"); + printf("a=1.2f Test passed\n"); - /// a=b*1.2f - gpu_out.device(sycl_device) = gpu_in1 * 1.2f; - sycl_device.deallocate(out.data()); + /// a=b*1.2f + gpu_out.device(sycl_device) = gpu_in1 * 1.2f; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -81,11 +83,11 @@ void test_sycl_cpu() { } } } - printf("a=b*1.2f Test Passed\n"); + printf("a=b*1.2f Test Passed\n"); - /// c=a*b - gpu_out.device(sycl_device) = gpu_in1 * gpu_in2; - sycl_device.deallocate(out.data()); + /// c=a*b + gpu_out.device(sycl_device) = gpu_in1 * gpu_in2; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -95,11 +97,11 @@ void test_sycl_cpu() { } } } - printf("c=a*b Test Passed\n"); + printf("c=a*b Test Passed\n"); - /// c=a+b - gpu_out.device(sycl_device) = gpu_in1 + gpu_in2; - sycl_device.deallocate(out.data()); + /// c=a+b + gpu_out.device(sycl_device) = gpu_in1 + gpu_in2; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -109,11 +111,11 @@ void test_sycl_cpu() { } } } - printf("c=a+b Test Passed\n"); + printf("c=a+b Test Passed\n"); - /// c=a*a - gpu_out.device(sycl_device) = gpu_in1 * gpu_in1; - sycl_device.deallocate(out.data()); + /// c=a*a + gpu_out.device(sycl_device) = gpu_in1 * gpu_in1; + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -124,11 +126,11 @@ void test_sycl_cpu() { } } - printf("c= a*a Test Passed\n"); + printf("c= a*a Test Passed\n"); - //a*3.14f + b*2.7f - gpu_out.device(sycl_device) = gpu_in1 * gpu_in1.constant(3.14f) + gpu_in2 * gpu_in2.constant(2.7f); - sycl_device.deallocate(out.data()); + //a*3.14f + b*2.7f + gpu_out.device(sycl_device) = gpu_in1 * gpu_in1.constant(3.14f) + gpu_in2 * gpu_in2.constant(2.7f); + sycl_device.deallocate(out.data()); for (int i = 0; i < sizeDim1; ++i) { for (int j = 0; j < sizeDim2; ++j) { for (int k = 0; k < sizeDim3; ++k) { @@ -138,21 +140,21 @@ void test_sycl_cpu() { } } } - printf("a*3.14f + b*2.7f Test Passed\n"); - - ///d= (a>0.5? b:c) - gpu_out.device(sycl_device) =(gpu_in1 > gpu_in1.constant(0.5f)).select(gpu_in2, gpu_in3); - sycl_device.deallocate(out.data()); - for (int i = 0; i < sizeDim1; ++i) { - for (int j = 0; j < sizeDim2; ++j) { - for (int k = 0; k < sizeDim3; ++k) { - VERIFY_IS_APPROX(out(i, j, k), (in1(i, j, k) > 0.5f) - ? in2(i, j, k) - : in3(i, j, k)); - } - } - } - printf("d= (a>0.5? b:c) Test Passed\n"); + printf("a*3.14f + b*2.7f Test Passed\n"); + + ///d= (a>0.5? b:c) + gpu_out.device(sycl_device) =(gpu_in1 > gpu_in1.constant(0.5f)).select(gpu_in2, gpu_in3); + sycl_device.deallocate(out.data()); + for (int i = 0; i < sizeDim1; ++i) { + for (int j = 0; j < sizeDim2; ++j) { + for (int k = 0; k < sizeDim3; ++k) { + VERIFY_IS_APPROX(out(i, j, k), (in1(i, j, k) > 0.5f) + ? in2(i, j, k) + : in3(i, j, k)); + } + } + } + printf("d= (a>0.5? b:c) Test Passed\n"); } void test_cxx11_tensor_sycl() { |