![]() |
Eigen-unsupported
3.3.3
|
00001 // This file is part of Eigen, a lightweight C++ template library 00002 // for linear algebra. 00003 // 00004 // Mehdi Goli Codeplay Software Ltd. 00005 // Ralph Potter Codeplay Software Ltd. 00006 // Luke Iwanski Codeplay Software Ltd. 00007 // Contact: <eigen@codeplay.com> 00008 // 00009 // This Source Code Form is subject to the terms of the Mozilla 00010 // Public License v. 2.0. If a copy of the MPL was not distributed 00011 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 00012 00013 /***************************************************************** 00014 * TensorSyclExtractAccessor.h 00015 * 00016 * \brief: 00017 * ExtractAccessor takes Expression placeHolder expression and the tuple of sycl 00018 * buffers as an input. Using pre-order tree traversal, ExtractAccessor 00019 * recursively calls itself for its children in the expression tree. The 00020 * leaf node in the PlaceHolder expression is nothing but a container preserving 00021 * the order of the actual data in the tuple of sycl buffer. By invoking the 00022 * extract accessor for the PlaceHolder<N>, an accessor is created for the Nth 00023 * buffer in the tuple of buffers. This accessor is then added as an Nth 00024 * element in the tuple of accessors. In this case we preserve the order of data 00025 * in the expression tree. 00026 * 00027 * This is the specialisation of extract accessor method for different operation 00028 * type in the PlaceHolder expression. 00029 * 00030 *****************************************************************/ 00031 00032 #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP 00033 #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP 00034 00035 namespace Eigen { 00036 namespace TensorSycl { 00037 namespace internal { 00042 template <typename Evaluator> 00043 struct ExtractAccessor; 00044 00045 struct AccessorConstructor{ 00046 template<typename Arg> static inline auto getTuple(cl::sycl::handler& cgh, Arg eval) 00047 -> decltype(ExtractAccessor<Arg>::getTuple(cgh, eval)) { 00048 return ExtractAccessor<Arg>::getTuple(cgh, eval); 00049 } 00050 00051 template<typename Arg1, typename Arg2> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1, Arg2 eval2) 00052 -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2))) { 00053 return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2)); 00054 } 00055 template<typename Arg1, typename Arg2, typename Arg3> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1 , Arg2 eval2 , Arg3 eval3) 00056 -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3)))) { 00057 return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3))); 00058 } 00059 template< cl::sycl::access::mode AcM, typename Arg> static inline auto getAccessor(cl::sycl::handler& cgh, Arg eval) 00060 -> decltype(utility::tuple::make_tuple( eval.device().template get_sycl_accessor<AcM, 00061 typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data()))){ 00062 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())); 00063 } 00064 }; 00065 00068 template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> 00069 struct ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > { 00070 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> eval) 00071 -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){ 00072 return AccessorConstructor::getTuple(cgh, eval.impl()); 00073 } 00074 }; 00075 00077 template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev> 00078 struct ExtractAccessor<TensorEvaluator<UnaryCategory<OP, RHSExpr>, Dev> > 00079 : ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > {}; 00080 00082 template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> 00083 struct ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > { 00084 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> eval) 00085 -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){ 00086 return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl()); 00087 } 00088 }; 00090 template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev> 00091 struct ExtractAccessor<TensorEvaluator<BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > 00092 : ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >{}; 00093 00096 template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> 00097 struct ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > { 00098 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> eval) 00099 -> decltype(AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl())){ 00100 return AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl()); 00101 } 00102 }; 00103 00105 template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev> 00106 struct ExtractAccessor<TensorEvaluator<TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > 00107 : ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >{}; 00108 00111 template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> 00112 struct ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > { 00113 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> eval) 00114 -> decltype(AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl())){ 00115 return AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl()); 00116 } 00117 }; 00118 00121 template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> 00122 struct ExtractAccessor<TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > 00123 : ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> >{}; 00124 00126 template <typename LHSExpr, typename RHSExpr, typename Dev> 00127 struct ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> > { 00128 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) 00129 -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){ 00130 return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl()); 00131 } 00132 }; 00133 00135 template <typename LHSExpr, typename RHSExpr, typename Dev> 00136 struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> > 00137 : ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{}; 00138 00140 #define TENSORMAPEXPR(CVQual, ACCType)\ 00141 template <typename PlainObjectType, int Options_, typename Dev>\ 00142 struct ExtractAccessor<TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> > {\ 00143 static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> eval)\ 00144 -> decltype(AccessorConstructor::template getAccessor<ACCType>(cgh, eval)){\ 00145 return AccessorConstructor::template getAccessor<ACCType>(cgh, eval);\ 00146 }\ 00147 }; 00148 TENSORMAPEXPR(const, cl::sycl::access::mode::read) 00149 TENSORMAPEXPR(, cl::sycl::access::mode::read_write) 00150 #undef TENSORMAPEXPR 00151 00153 template <typename Expr, typename Dev> 00154 struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> > { 00155 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> eval) 00156 -> decltype(AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval)){ 00157 return AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval); 00158 } 00159 }; 00160 00162 template <typename Expr, typename Dev> 00163 struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev> > 00164 : ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> >{}; 00165 00167 template <typename Expr, typename Dev> 00168 struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> > { 00169 static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<const TensorEvalToOp<Expr>, Dev> eval) 00170 -> decltype(utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl()))){ 00171 return utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl())); 00172 } 00173 }; 00174 00176 template <typename Expr, typename Dev> 00177 struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev> > 00178 : ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> >{}; 00179 00181 template <typename OP, typename Dim, typename Expr, typename Dev> 00182 struct ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> > { 00183 static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> eval) 00184 -> decltype(AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval)){ 00185 return AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval); 00186 } 00187 }; 00188 00190 template <typename OP, typename Dim, typename Expr, typename Dev> 00191 struct ExtractAccessor<TensorEvaluator<TensorReductionOp<OP, Dim, Expr>, Dev> > 00192 : ExtractAccessor<TensorEvaluator<const TensorReductionOp<OP, Dim, Expr>, Dev> >{}; 00193 00195 template <typename Evaluator> 00196 auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) 00197 -> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) { 00198 return ExtractAccessor<Evaluator>::getTuple(cgh, expr); 00199 } 00200 00201 } 00202 } 00203 } 00204 #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP