diff options
Diffstat (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h')
| -rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h | 235 |
1 files changed, 128 insertions, 107 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h index 06e2d5ae0..f5ef05e36 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h @@ -43,172 +43,193 @@ template <typename Evaluator> struct ExtractAccessor; struct AccessorConstructor{ - template<typename Arg> static inline auto getTuple(cl::sycl::handler& cgh, Arg eval) + template<typename Arg> static inline auto getTuple(cl::sycl::handler& cgh, const Arg& eval) -> decltype(ExtractAccessor<Arg>::getTuple(cgh, eval)) { return ExtractAccessor<Arg>::getTuple(cgh, eval); } - template<typename Arg1, typename Arg2> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1, Arg2 eval2) + template<typename Arg1, typename Arg2> static inline auto getTuple(cl::sycl::handler& cgh, const Arg1& eval1, const 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)); } - template<typename Arg1, typename Arg2, typename Arg3> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1 , Arg2 eval2 , Arg3 eval3) + template<typename Arg1, typename Arg2, typename Arg3> static inline auto getTuple(cl::sycl::handler& cgh, const Arg1& eval1 , const Arg2& eval2 , const 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))); } - template< cl::sycl::access::mode AcM, typename Arg> static inline auto getAccessor(cl::sycl::handler& cgh, Arg eval) + template< cl::sycl::access::mode AcM, typename Arg> static inline auto getAccessor(cl::sycl::handler& cgh, const Arg& eval) -> decltype(utility::tuple::make_tuple( eval.device().template get_sycl_accessor<AcM>(cgh,eval.data()))){ return utility::tuple::make_tuple(eval.device().template get_sycl_accessor<AcM>(cgh,eval.data())); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is -/// 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()); - } +/// TensorCwiseNullaryOp, TensorCwiseUnaryOp and TensorBroadcastingOp +#define SYCLUNARYCATEGORYEXTACC(CVQual)\ +template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual UnaryCategory<OP, RHSExpr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual 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 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 <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()); - } -}; +SYCLUNARYCATEGORYEXTACC(const) +SYCLUNARYCATEGORYEXTACC() +#undef SYCLUNARYCATEGORYEXTACC + + /// specialisation of the \ref ExtractAccessor struct when the node type is TensorCwiseBinaryOp -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> >{}; +#define SYCLBINARYCATEGORYEXTACC(CVQual)\ +template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual 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());\ + }\ +}; + +SYCLBINARYCATEGORYEXTACC(const) +SYCLBINARYCATEGORYEXTACC() +#undef SYCLBINARYCATEGORYEXTACC /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorCwiseTernaryOp -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()); - } +#define SYCLTERNARYCATEGORYEXTACC(CVQual)\ +template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual 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 <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> >{}; +SYCLTERNARYCATEGORYEXTACC(const) +SYCLTERNARYCATEGORYEXTACC() +#undef SYCLTERNARYCATEGORYEXTACC -/// specialisation of the \ref ExtractAccessor struct when the node type is -/// 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 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. 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> > -: 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 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()); - } +#define SYCLSELECTOPEXTACC(CVQual)\ +template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual 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());\ + }\ }; +SYCLSELECTOPEXTACC(const) +SYCLSELECTOPEXTACC() +#undef SYCLSELECTOPEXTACC + /// 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> > -: ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{}; +#define SYCLTENSORASSIGNOPEXTACC(CVQual)\ +template <typename LHSExpr, typename RHSExpr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorAssignOp<LHSExpr, RHSExpr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual 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());\ + }\ +}; + + SYCLTENSORASSIGNOPEXTACC(const) + SYCLTENSORASSIGNOPEXTACC() + #undef SYCLTENSORASSIGNOPEXTACC /// 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)\ + 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> > { - 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> >{}; - -/// 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> > { - 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())); - } +#define SYCLFORCEDEVALEXTACC(CVQual)\ +template <typename Expr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorForcedEvalOp<Expr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual 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);\ + }\ }; +SYCLFORCEDEVALEXTACC(const) +SYCLFORCEDEVALEXTACC() +#undef SYCLFORCEDEVALEXTACC + + /// 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> >{}; - -/// specialisation of the \ref ExtractAccessor struct when the node type is const TensorReductionOp -template <typename OP, typename Dim, typename Expr, typename Dev> -struct ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> > { - static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorReductionOp<OP, Dim, 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); - } +#define SYCLEVALTOEXTACC(CVQual)\ +template <typename Expr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorEvalToOp<Expr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<CVQual 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()));\ + }\ }; +SYCLEVALTOEXTACC(const) +SYCLEVALTOEXTACC() +#undef SYCLEVALTOEXTACC + /// specialisation of the \ref ExtractAccessor struct when the node type is TensorReductionOp -template <typename OP, typename Dim, typename Expr, typename Dev> -struct ExtractAccessor<TensorEvaluator<TensorReductionOp<OP, Dim, Expr>, Dev> > -: ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> >{}; +#define SYCLREDUCTIONEXTACC(CVQual)\ +template <typename OP, typename Dim, typename Expr, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorReductionOp<OP, Dim, Expr>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual TensorReductionOp<OP, Dim, 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);\ + }\ +}; +SYCLREDUCTIONEXTACC(const) +SYCLREDUCTIONEXTACC() +#undef SYCLREDUCTIONEXTACC /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorSlicingOp. This is a special case where there is no OP -template <typename StartIndices, typename Sizes, typename XprType, typename Dev> -struct ExtractAccessor<TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > { - static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> eval) - -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){ - return AccessorConstructor::getTuple(cgh, eval.impl()); - } +#define SYCLSLICEOPEXTACC(CVQual)\ +template <typename StartIndices, typename Sizes, typename XprType, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual TensorSlicingOp<StartIndices, Sizes, XprType>, Dev>& eval)\ + -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){\ + return AccessorConstructor::getTuple(cgh, eval.impl());\ + }\ +}; + +SYCLSLICEOPEXTACC(const) +SYCLSLICEOPEXTACC() +#undef SYCLSLICEOPEXTACC + +#define RESHAPEANDSHUFFOPEXTRACC(OPEXPR, CVQual)\ +template<typename Param, typename XprType, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual OPEXPR<Param, XprType>, Dev> > {\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual OPEXPR<Param, XprType>, Dev>& eval)\ + -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){\ + return AccessorConstructor::getTuple(cgh, eval.impl());\ + }\ }; +// tensor reshaping +RESHAPEANDSHUFFOPEXTRACC(TensorReshapingOp, const) +RESHAPEANDSHUFFOPEXTRACC(TensorReshapingOp, ) +/// Tensor shuffling +RESHAPEANDSHUFFOPEXTRACC(TensorShufflingOp, const) +RESHAPEANDSHUFFOPEXTRACC(TensorShufflingOp, ) +#undef RESHAPEANDSHUFFOPEXTRACC -template <typename StartIndices, typename Sizes, typename XprType, typename Dev> -struct ExtractAccessor<TensorEvaluator<TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > -:ExtractAccessor<TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> >{}; /// template deduction for \ref ExtractAccessor template <typename Evaluator> -auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) --> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) { - return ExtractAccessor<Evaluator>::getTuple(cgh, expr); +auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& eval) +-> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, eval)) { + return ExtractAccessor<Evaluator>::getTuple(cgh, eval); } } /// namespace TensorSycl |