1 // This file is part of Eigen, a lightweight C++ template library 2 // for linear algebra. 3 // 4 // Mehdi Goli Codeplay Software Ltd. 5 // Ralph Potter Codeplay Software Ltd. 6 // Luke Iwanski Codeplay Software Ltd. 7 // Contact: <eigen (at) codeplay.com> 8 // 9 // This Source Code Form is subject to the terms of the Mozilla 10 // Public License v. 2.0. If a copy of the MPL was not distributed 11 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 12 13 /***************************************************************** 14 * TensorSyclExtractAccessor.h 15 * 16 * \brief: 17 * ExtractAccessor takes Expression placeHolder expression and the tuple of sycl 18 * buffers as an input. Using pre-order tree traversal, ExtractAccessor 19 * recursively calls itself for its children in the expression tree. The 20 * leaf node in the PlaceHolder expression is nothing but a container preserving 21 * the order of the actual data in the tuple of sycl buffer. By invoking the 22 * extract accessor for the PlaceHolder<N>, an accessor is created for the Nth 23 * buffer in the tuple of buffers. This accessor is then added as an Nth 24 * element in the tuple of accessors. In this case we preserve the order of data 25 * in the expression tree. 26 * 27 * This is the specialisation of extract accessor method for different operation 28 * type in the PlaceHolder expression. 29 * 30 *****************************************************************/ 31 32 #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP 33 #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP 34 35 namespace Eigen { 36 namespace TensorSycl { 37 namespace internal { 38 /// \struct ExtractAccessor: Extract Accessor Class is used to extract the 39 /// accessor from a buffer. 40 /// Depending on the type of the leaf node we can get a read accessor or a 41 /// read_write accessor 42 template <typename Evaluator> 43 struct ExtractAccessor; 44 45 struct AccessorConstructor{ 46 template<typename Arg> static inline auto getTuple(cl::sycl::handler& cgh, Arg eval) 47 -> decltype(ExtractAccessor<Arg>::getTuple(cgh, eval)) { 48 return ExtractAccessor<Arg>::getTuple(cgh, eval); 49 } 50 51 template<typename Arg1, typename Arg2> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1, Arg2 eval2) 52 -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2))) { 53 return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2)); 54 } 55 template<typename Arg1, typename Arg2, typename Arg3> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1 , Arg2 eval2 , Arg3 eval3) 56 -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3)))) { 57 return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3))); 58 } 59 template< cl::sycl::access::mode AcM, typename Arg> static inline auto getAccessor(cl::sycl::handler& cgh, Arg eval) 60 -> decltype(utility::tuple::make_tuple( eval.device().template get_sycl_accessor<AcM, 61 typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data()))){ 62 return utility::tuple::make_tuple(eval.device().template get_sycl_accessor<AcM, typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data())); 63 } 64 }; 65 66 /// specialisation of the \ref ExtractAccessor struct when the node type is 67 /// const TensorCwiseNullaryOp, const TensorCwiseUnaryOp and const TensorBroadcastingOp 68 template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> 69 struct ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > { 70 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> eval) 71 -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){ 72 return AccessorConstructor::getTuple(cgh, eval.impl()); 73 } 74 }; 75 76 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorCwiseNullaryOp, TensorCwiseUnaryOp and TensorBroadcastingOp 77 template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> 78 struct ExtractAccessor<TensorEvaluator<UnaryCategory<OP, RHSExpr>, Dev> > 79 : ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > {}; 80 81 /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorCwiseBinaryOp 82 template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> 83 struct ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > { 84 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> eval) 85 -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){ 86 return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl()); 87 } 88 }; 89 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorCwiseBinaryOp 90 template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> 91 struct ExtractAccessor<TensorEvaluator<BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > 92 : ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >{}; 93 94 /// specialisation of the \ref ExtractAccessor struct when the node type is 95 /// const TensorCwiseTernaryOp 96 template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> 97 struct ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > { 98 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> eval) 99 -> decltype(AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl())){ 100 return AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl()); 101 } 102 }; 103 104 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorCwiseTernaryOp 105 template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> 106 struct ExtractAccessor<TensorEvaluator<TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > 107 : ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >{}; 108 109 /// specialisation of the \ref ExtractAccessor struct when the node type is 110 /// const TensorCwiseSelectOp. This is a special case where there is no OP 111 template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> 112 struct ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > { 113 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> eval) 114 -> decltype(AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl())){ 115 return AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl()); 116 } 117 }; 118 119 /// specialisation of the \ref ExtractAccessor struct when the node type is 120 /// TensorCwiseSelectOp. This is a special case where there is no OP 121 template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> 122 struct ExtractAccessor<TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > 123 : ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> >{}; 124 125 /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorAssignOp 126 template <typename LHSExpr, typename RHSExpr, typename Dev> 127 struct ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> > { 128 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) 129 -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){ 130 return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl()); 131 } 132 }; 133 134 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorAssignOp 135 template <typename LHSExpr, typename RHSExpr, typename Dev> 136 struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> > 137 : ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{}; 138 139 /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorMap 140 #define TENSORMAPEXPR(CVQual, ACCType)\ 141 template <typename PlainObjectType, int Options_, typename Dev>\ 142 struct ExtractAccessor<TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> > {\ 143 static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> eval)\ 144 -> decltype(AccessorConstructor::template getAccessor<ACCType>(cgh, eval)){\ 145 return AccessorConstructor::template getAccessor<ACCType>(cgh, eval);\ 146 }\ 147 }; 148 TENSORMAPEXPR(const, cl::sycl::access::mode::read) 149 TENSORMAPEXPR(, cl::sycl::access::mode::read_write) 150 #undef TENSORMAPEXPR 151 152 /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorForcedEvalOp 153 template <typename Expr, typename Dev> 154 struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> > { 155 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> eval) 156 -> decltype(AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval)){ 157 return AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval); 158 } 159 }; 160 161 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorForcedEvalOp 162 template <typename Expr, typename Dev> 163 struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev> > 164 : ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> >{}; 165 166 /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorEvalToOp 167 template <typename Expr, typename Dev> 168 struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> > { 169 static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<const TensorEvalToOp<Expr>, Dev> eval) 170 -> decltype(utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl()))){ 171 return utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl())); 172 } 173 }; 174 175 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorEvalToOp 176 template <typename Expr, typename Dev> 177 struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev> > 178 : ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> >{}; 179 180 /// specialisation of the \ref ExtractAccessor struct when the node type is const TensorReductionOp 181 template <typename OP, typename Dim, typename Expr, typename Dev> 182 struct ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> > { 183 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> eval) 184 -> decltype(AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval)){ 185 return AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval); 186 } 187 }; 188 189 /// specialisation of the \ref ExtractAccessor struct when the node type is TensorReductionOp 190 template <typename OP, typename Dim, typename Expr, typename Dev> 191 struct ExtractAccessor<TensorEvaluator<TensorReductionOp<OP, Dim, Expr>, Dev> > 192 : ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> >{}; 193 194 /// template deduction for \ref ExtractAccessor 195 template <typename Evaluator> 196 auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) 197 -> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) { 198 return ExtractAccessor<Evaluator>::getTuple(cgh, expr); 199 } 200 201 } /// namespace TensorSycl 202 } /// namespace internal 203 } /// namespace Eigen 204 #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP 205
