aboutsummaryrefslogtreecommitdiffhomepage
path: root/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h
diff options
context:
space:
mode:
Diffstat (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h')
-rw-r--r--unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h96
1 files changed, 0 insertions, 96 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h
deleted file mode 100644
index 29c78184d..000000000
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h
+++ /dev/null
@@ -1,96 +0,0 @@
-// 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.
-// Cummins Chris PhD student at The University of Edinburgh.
-// Contact: <eigen@codeplay.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/.
-
-/*****************************************************************
- * TensorSyclRun.h
- *
- * \brief:
- * Schedule_kernel invoke an specialised version of kernel struct. The
- * specialisation is based on the data dimension in sycl buffer
- *
-*****************************************************************/
-
-#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP
-#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP
-
-namespace Eigen {
-namespace TensorSycl {
-template<typename Expr, typename FunctorExpr, typename TupleType > struct ExecExprFunctorKernel{
- typedef typename internal::createPlaceHolderExpression<Expr>::Type PlaceHolderExpr;
-
- typedef typename Expr::Index Index;
- FunctorExpr functors;
- TupleType tuple_of_accessors;
- Index range;
- ExecExprFunctorKernel(Index range_, FunctorExpr functors_, TupleType tuple_of_accessors_)
- : functors(functors_), tuple_of_accessors(tuple_of_accessors_), range(range_){}
- void operator()(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::SyclKernelDevice>(device_expr.expr, Eigen::SyclKernelDevice());
- typename DevExpr::Index gId = static_cast<typename DevExpr::Index>(itemID.get_global_linear_id());
- if (gId < range)
- device_evaluator.evalScalar(gId);
- }
-};
-
-/// The run function in tensor sycl convert the expression tree to a buffer
-/// based expression tree;
-/// creates the expression tree for the device with accessor to buffers;
-/// construct the kernel and submit it to the sycl queue.
-/// std::array does not have TotalSize. So I have to get the size through template specialisation.
-template<typename , typename Dimensions> struct DimensionSize{
- static auto getDimSize(const Dimensions& dim)->decltype(dim.TotalSize()){
- return dim.TotalSize();
- }
-};
-#define DIMSIZEMACRO(CVQual)\
-template<typename Index, size_t NumDims> struct DimensionSize<Index, CVQual std::array<Index, NumDims>>{\
- static inline Index getDimSize(const std::array<Index, NumDims>& dim){\
- return (NumDims == 0) ? 1 : ::Eigen::internal::array_prod(dim);\
- }\
-};
-
-DIMSIZEMACRO(const)
-DIMSIZEMACRO()
-#undef DIMSIZEMACRO
-
-
-template <typename Expr, typename Dev>
-void run(Expr &expr, Dev &dev) {
- Eigen::TensorEvaluator<Expr, Dev> evaluator(expr, dev);
- const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);
- if (needs_assign) {
- typedef Eigen::TensorSycl::internal::FunctorExtractor<Eigen::TensorEvaluator<Expr, Dev> > FunctorExpr;
- FunctorExpr functors = internal::extractFunctors(evaluator);
- dev.sycl_queue().submit([&](cl::sycl::handler &cgh) {
- // create a tuple of accessors from Evaluator
- typedef decltype(internal::createTupleOfAccessors<Eigen::TensorEvaluator<Expr, Dev> >(cgh, evaluator)) TupleType;
- TupleType tuple_of_accessors = internal::createTupleOfAccessors<Eigen::TensorEvaluator<Expr, Dev> >(cgh, evaluator);
- typename Expr::Index range, GRange, tileSize;
- typename Expr::Index total_size = static_cast<typename Expr::Index>(DimensionSize<typename Expr::Index, typename Eigen::TensorEvaluator<Expr, Dev>::Dimensions>::getDimSize(evaluator.dimensions()));
- dev.parallel_for_setup(total_size, tileSize, range, GRange);
-
- cgh.parallel_for(cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)),
- ExecExprFunctorKernel<Expr,FunctorExpr,TupleType>(range
- , functors, tuple_of_accessors
- ));
- });
- dev.asynchronousExec();
- }
- evaluator.cleanup();
-}
-} // namespace TensorSycl
-} // namespace Eigen
-
-#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP