GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-04-22 17:49:36 +08:00
Code Issues Packages Projects Releases Wiki Activity

Pulled latest update from trunk

Browse Source
This commit is contained in:
Benoit Steiner 2016-12-21 16:43:27 -08:00
commit 660da83e18
1 changed files with 21 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

42
unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h
View File

@ -26,25 +26,25 @@
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
template<typename Expr, typename FunctorExpr, typename TupleType > struct ExecExprFunctorKernel{
typedef typename internal::createPlaceHolderExpression<Expr>::Type PlaceHolderExpr;
template<typename Expr, typename FunctorExpr, typename TupleType > struct ExecExprFunctorKernel{ typedef typename Expr::Index Index;
typedef typename internal::createPlaceHolderExpression<Expr>::Type PlaceHolderExpr; 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::DefaultDevice>(device_expr.expr, Eigen::DefaultDevice());
typename DevExpr::Index gId = static_cast<typename DevExpr::Index>(itemID.get_global_linear_id());
if (gId < range)
device_evaluator.evalScalar(gId);
}
};
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::DefaultDevice>(device_expr.expr, Eigen::DefaultDevice());
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 /// The run function in tensor sycl convert the expression tree to a buffer
/// based expression tree; /// based expression tree;
/// creates the expression tree for the device with accessor to buffers; /// creates the expression tree for the device with accessor to buffers;
@ -54,12 +54,12 @@ void run(Expr &expr, Dev &dev) {
Eigen::TensorEvaluator<Expr, Dev> evaluator(expr, dev); Eigen::TensorEvaluator<Expr, Dev> evaluator(expr, dev);
const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL); const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);
if (needs_assign) { if (needs_assign) {
auto functors = internal::extractFunctors(evaluator); typedef decltype(internal::extractFunctors(evaluator)) FunctorExpr;
typedef decltype(functors) FunctorExpr; FunctorExpr functors = internal::extractFunctors(evaluator);
dev.sycl_queue().submit([&](cl::sycl::handler &cgh) { dev.sycl_queue().submit([&](cl::sycl::handler &cgh) {
// create a tuple of accessors from Evaluator // create a tuple of accessors from Evaluator
auto tuple_of_accessors = internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator); typedef decltype(internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator)) TupleType;
typedef decltype(tuple_of_accessors) TupleType; TupleType tuple_of_accessors = internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator);
typename Expr::Index range, GRange, tileSize; typename Expr::Index range, GRange, tileSize;
dev.parallel_for_setup(static_cast<typename Expr::Index>(evaluator.dimensions().TotalSize()), tileSize, range, GRange); dev.parallel_for_setup(static_cast<typename Expr::Index>(evaluator.dimensions().TotalSize()), tileSize, range, GRange);