// This file is part of Eigen, a lightweight C++ template library // for linear algebra. // // Mehdi Goli Codeplay Software Ltd. // Ralph Potter Codeplay Software Ltd. // Luke Iwanski Codeplay Software Ltd. // Contact: // // 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/. /***************************************************************** * TensorSyclExtractAccessor.h * * \brief: * ExtractAccessor takes Expression placeHolder expression and the tuple of sycl * buffers as an input. Using pre-order tree traversal, ExtractAccessor * recursively calls itself for its children in the expression tree. The * leaf node in the PlaceHolder expression is nothing but a container preserving * the order of the actual data in the tuple of sycl buffer. By invoking the * extract accessor for the PlaceHolder, an accessor is created for the Nth * buffer in the tuple of buffers. This accessor is then added as an Nth * element in the tuple of accessors. In this case we preserve the order of data * in the expression tree. * * This is the specialisation of extract accessor method for different operation * type in the PlaceHolder expression. * *****************************************************************/ #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP namespace Eigen { namespace TensorSycl { namespace internal { /// \struct ExtractAccessor: Extract Accessor Class is used to extract the /// accessor from a buffer. /// Depending on the type of the leaf node we can get a read accessor or a /// read_write accessor template struct ExtractAccessor; struct AccessorConstructor{ template static inline auto getTuple(cl::sycl::handler& cgh, Arg eval) -> decltype(ExtractAccessor::getTuple(cgh, eval)) { return ExtractAccessor::getTuple(cgh, eval); } template static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1, Arg2 eval2) -> decltype(utility::tuple::append(ExtractAccessor::getTuple(cgh, eval1), ExtractAccessor::getTuple(cgh, eval2))) { return utility::tuple::append(ExtractAccessor::getTuple(cgh, eval1), ExtractAccessor::getTuple(cgh, eval2)); } template static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1 , Arg2 eval2 , Arg3 eval3) -> decltype(utility::tuple::append(ExtractAccessor::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor::getTuple(cgh, eval2), ExtractAccessor::getTuple(cgh, eval3)))) { return utility::tuple::append(ExtractAccessor::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor::getTuple(cgh, eval2), ExtractAccessor::getTuple(cgh, eval3))); } 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::type>(eval.dimensions().TotalSize(), cgh,eval.data()))){ return utility::tuple::make_tuple(eval.device().template get_sycl_accessor::type>(eval.dimensions().TotalSize(), cgh,eval.data())); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorCwiseNullaryOp, const TensorCwiseUnaryOp and const TensorBroadcastingOp template class UnaryCategory, typename OP, typename RHSExpr, typename Dev> struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, 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 class UnaryCategory, typename OP, typename RHSExpr, typename Dev> struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> > {}; /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorCwiseBinaryOp template class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, 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 class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorCwiseTernaryOp template class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, 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 class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is /// const TensorCwiseSelectOp. This is a special case where there is no OP template struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, 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 struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorAssignOp template struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, 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 struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorMap #define TENSORMAPEXPR(CVQual, ACCType)\ template \ struct ExtractAccessor, Dev> > {\ static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator, Dev> eval)\ -> decltype(AccessorConstructor::template getAccessor(cgh, eval)){\ return AccessorConstructor::template getAccessor(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 struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, Dev> eval) -> decltype(AccessorConstructor::template getAccessor(cgh, eval)){ return AccessorConstructor::template getAccessor(cgh, eval); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is TensorForcedEvalOp template struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorEvalToOp template struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator, Dev> eval) -> decltype(utility::tuple::append(AccessorConstructor::template getAccessor(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl()))){ return utility::tuple::append(AccessorConstructor::template getAccessor(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl())); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is TensorEvalToOp template struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorReductionOp template struct ExtractAccessor, Dev> > { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator, Dev> eval) -> decltype(AccessorConstructor::template getAccessor(cgh, eval)){ return AccessorConstructor::template getAccessor(cgh, eval); } }; /// specialisation of the \ref ExtractAccessor struct when the node type is TensorReductionOp template struct ExtractAccessor, Dev> > : ExtractAccessor, Dev> >{}; /// template deduction for \ref ExtractAccessor template auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) -> decltype(ExtractAccessor::getTuple(cgh, expr)) { return ExtractAccessor::getTuple(cgh, expr); } } /// namespace TensorSycl } /// namespace internal } /// namespace Eigen #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP