GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-05-01 16:24:28 +08:00
Code Issues Packages Projects Releases Wiki Activity

Merged ComputeCpp into default.

Browse Source
This commit is contained in:
Luke Iwanski 2016-10-14 13:36:51 +01:00
commit e742da8b28
17 changed files with 755 additions and 1911 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

65
Eigen/Core
View File

@ -11,12 +11,20 @@
#ifndef EIGEN_CORE_H #ifndef EIGEN_CORE_H
#define EIGEN_CORE_H #define EIGEN_CORE_H
// first thing Eigen does: stop the compiler from committing suicide
#include "src/Core/util/DisableStupidWarnings.h"
// Handle NVCC/CUDA /// This will no longer be needed after the next release of the computecppCE
#ifdef __CUDACC__ #ifdef EIGEN_USE_SYCL
// Do not try asserts on CUDA! #undef min
#undef max
#undef isnan
#undef isinf
#undef isfinite
#include <SYCL/sycl.hpp>
#endif
// Handle NVCC/CUDA/SYCL
#if defined(__CUDACC__) || defined(__SYCL_DEVICE_ONLY__)
// Do not try asserts on CUDA and SYCL!
#ifndef EIGEN_NO_DEBUG #ifndef EIGEN_NO_DEBUG
#define EIGEN_NO_DEBUG #define EIGEN_NO_DEBUG
#endif #endif
@ -25,7 +33,7 @@
#undef EIGEN_INTERNAL_DEBUGGING #undef EIGEN_INTERNAL_DEBUGGING
#endif #endif
// Do not try to vectorize on CUDA! // Do not try to vectorize on CUDA and SYCL!
#ifndef EIGEN_DONT_VECTORIZE #ifndef EIGEN_DONT_VECTORIZE
#define EIGEN_DONT_VECTORIZE #define EIGEN_DONT_VECTORIZE
#endif #endif
@ -35,8 +43,12 @@
#endif #endif
// All functions callable from CUDA code must be qualified with __device__ // All functions callable from CUDA code must be qualified with __device__
#define EIGEN_DEVICE_FUNC __host__ __device__ #ifdef __CUDACC__
#define EIGEN_DEVICE_FUNC __host__ __device__
// For Sycl we dont need that
#else
#define EIGEN_DEVICE_FUNC
#endif
#else #else
#define EIGEN_DEVICE_FUNC #define EIGEN_DEVICE_FUNC
@ -51,43 +63,6 @@
#define EIGEN_USING_STD_MATH(FUNC) using std::FUNC; #define EIGEN_USING_STD_MATH(FUNC) using std::FUNC;
#endif #endif
#ifdef EIGEN_USE_SYCL
#undef min
#undef max
#undef isnan
#undef isinf
#undef isfinite
#include <SYCL/sycl.hpp>
#endif
// We need these predefines to determine if asserts need to be disabled for the device compiler
#if defined(__SYCL_DEVICE_ONLY__)
// Do not try asserts on SYCL!
#ifndef EIGEN_NO_DEBUG
#define EIGEN_NO_DEBUG
#endif
#ifdef EIGEN_INTERNAL_DEBUGGING
#undef EIGEN_INTERNAL_DEBUGGING
#endif
// Do not try to vectorize on SYCL!
#ifndef EIGEN_DONT_VECTORIZE
#define EIGEN_DONT_VECTORIZE
#endif
#ifdef EIGEN_EXCEPTIONS
#undef EIGEN_EXCEPTIONS
#endif
#define EIGEN_DEVICE_FUNC
#endif
#if (defined(_CPPUNWIND) || defined(__EXCEPTIONS)) && !defined(__CUDA_ARCH__) && !defined(EIGEN_EXCEPTIONS) && !defined(EIGEN_USE_SYCL)
#define EIGEN_EXCEPTIONS
#endif
#ifdef EIGEN_EXCEPTIONS #ifdef EIGEN_EXCEPTIONS
#include <new> #include <new>
#endif #endif

50
unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h
View File

@ -423,15 +423,15 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device>
// Precompute output strides. // Precompute output strides.
if (NumOutputDims > 0) { if (NumOutputDims > 0) {
if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) { if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
m_outputStrides[0] = 1; m_outputStrides[0] = 1;
for (int i = 1; i < NumOutputDims; ++i) { for (int i = 1; i < NumOutputDims; ++i) {
m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1]; m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1];
} }
} else { } else {
m_outputStrides.back() = 1; m_outputStrides.back() = 1;
for (int i = NumOutputDims - 2; i >= 0; --i) { for (int i = NumOutputDims - 2; i >= 0; --i) {
m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1]; m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1];
} }
} }
} }
@ -439,27 +439,27 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device>
if (NumInputDims > 0) { if (NumInputDims > 0) {
array<Index, NumInputDims> input_strides; array<Index, NumInputDims> input_strides;
if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) { if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
input_strides[0] = 1; input_strides[0] = 1;
for (int i = 1; i < NumInputDims; ++i) { for (int i = 1; i < NumInputDims; ++i) {
input_strides[i] = input_strides[i-1] * input_dims[i-1]; input_strides[i] = input_strides[i-1] * input_dims[i-1];
} }
} else { } else {
input_strides.back() = 1; input_strides.back() = 1;
for (int i = NumInputDims - 2; i >= 0; --i) { for (int i = NumInputDims - 2; i >= 0; --i) {
input_strides[i] = input_strides[i + 1] * input_dims[i + 1]; input_strides[i] = input_strides[i + 1] * input_dims[i + 1];
} }
} }
int outputIndex = 0; int outputIndex = 0;
int reduceIndex = 0; int reduceIndex = 0;
for (int i = 0; i < NumInputDims; ++i) { for (int i = 0; i < NumInputDims; ++i) {
if (m_reduced[i]) { if (m_reduced[i]) {
m_reducedStrides[reduceIndex] = input_strides[i]; m_reducedStrides[reduceIndex] = input_strides[i];
++reduceIndex; ++reduceIndex;
} else { } else {
m_preservedStrides[outputIndex] = input_strides[i]; m_preservedStrides[outputIndex] = input_strides[i];
++outputIndex; ++outputIndex;
} }
} }
} }
@ -578,7 +578,7 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device>
Op reducer(m_reducer); Op reducer(m_reducer);
if (ReducingInnerMostDims || RunningFullReduction) { if (ReducingInnerMostDims || RunningFullReduction) {
const Index num_values_to_reduce = const Index num_values_to_reduce =
(static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1]; (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1];
return internal::InnerMostDimReducer<Self, Op>::reduce(*this, firstInput(index), return internal::InnerMostDimReducer<Self, Op>::reduce(*this, firstInput(index),
num_values_to_reduce, reducer); num_values_to_reduce, reducer);
} else { } else {
@ -602,7 +602,7 @@ struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device>
EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[PacketSize]; EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[PacketSize];
if (ReducingInnerMostDims) { if (ReducingInnerMostDims) {
const Index num_values_to_reduce = const Index num_values_to_reduce =
(static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1]; (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1];
const Index firstIndex = firstInput(index); const Index firstIndex = firstInput(index);
for (Index i = 0; i < PacketSize; ++i) { for (Index i = 0; i < PacketSize; ++i) {
Op reducer(m_reducer); Op reducer(m_reducer);

27
unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h
View File

@ -11,20 +11,35 @@
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
// General include header of SYCL target for Tensor Module // General include header of SYCL target for Tensor Module
#ifndef TENSORSYCL_H #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H
#define TENSORSYCL_H #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H
#ifdef EIGEN_USE_SYCL #ifdef EIGEN_USE_SYCL
// trait class to extract different attribute contents
template <typename T>
struct Trait;
// global pointer to set different attribute state for a class // global pointer to set different attribute state for a class
template <class T> template <class T>
struct MakeGlobalPointer { struct MakeGlobalPointer {
typedef typename cl::sycl::global_ptr<T>::pointer_t Type; typedef typename cl::sycl::global_ptr<T>::pointer_t Type;
}; };
namespace Eigen {
namespace TensorSycl {
namespace internal {
/// This struct is used for special expression nodes with no operations (for example assign and selectOP).
struct NoOP;
template<bool IsConst, typename T> struct GetType{
typedef const T Type;
};
template<typename T> struct GetType<false, T>{
typedef T Type;
};
}
}
}
// tuple construction // tuple construction
#include "TensorSyclTuple.h" #include "TensorSyclTuple.h"
@ -59,4 +74,4 @@ struct MakeGlobalPointer {
#include "TensorSyclRun.h" #include "TensorSyclRun.h"
#endif // end of EIGEN_USE_SYCL #endif // end of EIGEN_USE_SYCL
#endif // TENSORSYCL_H #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_H

241
unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h
View File

@ -19,12 +19,13 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
namespace internal { namespace internal {
/// \struct ConvertToDeviceExpression /// \struct ConvertToDeviceExpression
/// \brief This struct is used to convert the MakePointer in the host expression /// \brief This struct is used to convert the MakePointer in the host expression
/// to the MakeGlobalPointer for the device expression. For the leafNodes /// to the MakeGlobalPointer for the device expression. For the leafNodes
@ -33,204 +34,74 @@ namespace internal {
template <typename Expr> template <typename Expr>
struct ConvertToDeviceExpression; struct ConvertToDeviceExpression;
template<template<class...> class NonOpCategory, bool IsConst, typename... Args>
struct NonOpConversion{
typedef typename GetType<IsConst, NonOpCategory<typename ConvertToDeviceExpression<Args>::Type...> >::Type Type;
};
template<template<class, template <class> class > class NonOpCategory, bool IsConst, typename Args>
struct DeviceConvertor{
typedef typename GetType<IsConst, NonOpCategory<typename ConvertToDeviceExpression<Args>::Type, MakeGlobalPointer> >::Type Type;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node /// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorMap /// type is TensorMap
template <typename Scalar_, int Options_, int Options2_, int NumIndices_, #define TENSORMAPCONVERT(CVQual)\
typename IndexType_, template <class> class MakePointer_> template <typename Scalar_, int Options_, int Options2_, int NumIndices_, typename IndexType_, template <class> class MakePointer_>\
struct ConvertToDeviceExpression< struct ConvertToDeviceExpression<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_> > {\
TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, typedef CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer> Type;\
MakePointer_>> {
using Type = TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakeGlobalPointer>;
}; };
/// specialisation of the \ref ConvertToDeviceExpression struct when the node TENSORMAPCONVERT(const)
/// type is const TensorMap TENSORMAPCONVERT()
template <typename Scalar_, int Options_, int Options2_, int NumIndices_, #undef TENSORMAPCONVERT
typename IndexType_, template <class> class MakePointer_>
struct ConvertToDeviceExpression<
const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakePointer_>> {
using Type =
const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakeGlobalPointer>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node /// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorCwiseNullaryOp /// type is TensorCwiseNullaryOp, TensorCwiseUnaryOp, TensorCwiseBinaryOp, TensorCwiseTernaryOp, TensorBroadcastingOp
template <typename OP, typename RHSExpr> #define CATEGORYCONVERT(CVQual)\
struct ConvertToDeviceExpression<const TensorCwiseNullaryOp<OP, RHSExpr>> { template <template<class, class...> class Category, typename OP, typename... subExprs>\
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; struct ConvertToDeviceExpression<CVQual Category<OP, subExprs...> > {\
using Type = const TensorCwiseNullaryOp<OP, RHSPlaceHolderType>; typedef CVQual Category<OP, typename ConvertToDeviceExpression<subExprs>::Type... > Type;\
}; };
CATEGORYCONVERT(const)
CATEGORYCONVERT()
#undef CATEGORYCONVERT
/// specialisation of the \ref ConvertToDeviceExpression struct when the node /// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorCwiseNullaryOp /// type is TensorCwiseSelectOp
template <typename OP, typename RHSExpr> #define SELECTOPCONVERT(CVQual, Res)\
struct ConvertToDeviceExpression<TensorCwiseNullaryOp<OP, RHSExpr>> { template <typename IfExpr, typename ThenExpr, typename ElseExpr>\
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; struct ConvertToDeviceExpression<CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr> >\
using Type = TensorCwiseNullaryOp<OP, RHSPlaceHolderType>; : NonOpConversion<TensorSelectOp, Res, IfExpr, ThenExpr, ElseExpr> {};
}; SELECTOPCONVERT(const, true)
SELECTOPCONVERT(, false)
/// specialisation of the \ref ConvertToDeviceExpression struct when the node #undef SELECTOPCONVERT
/// type is const TensorBroadcastingOp
template <typename OP, typename RHSExpr>
struct ConvertToDeviceExpression<const TensorBroadcastingOp<OP, RHSExpr>> {
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
using Type = const TensorBroadcastingOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorBroadcastingOp
template <typename OP, typename RHSExpr>
struct ConvertToDeviceExpression<TensorBroadcastingOp<OP, RHSExpr>> {
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
using Type = TensorBroadcastingOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorCwiseUnaryOp
template <typename OP, typename RHSExpr>
struct ConvertToDeviceExpression<const TensorCwiseUnaryOp<OP, RHSExpr>> {
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
using Type = const TensorCwiseUnaryOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorCwiseUnaryOp
template <typename OP, typename RHSExpr>
struct ConvertToDeviceExpression<TensorCwiseUnaryOp<OP, RHSExpr>> {
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
using Type = TensorCwiseUnaryOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr>
struct ConvertToDeviceExpression<
const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> {
using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type;
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
using Type =
const TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr>
struct ConvertToDeviceExpression<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> {
using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type;
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
using Type = TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorCwiseTernaryOp
template <typename OP, typename Arg1Impl, typename Arg2Impl, typename Arg3Impl>
struct ConvertToDeviceExpression<
const TensorCwiseTernaryOp<OP, Arg1Impl, Arg2Impl, Arg3Impl>> {
using Arg1PlaceHolderType =
typename ConvertToDeviceExpression<Arg1Impl>::Type;
using Arg2PlaceHolderType =
typename ConvertToDeviceExpression<Arg2Impl>::Type;
using Arg3PlaceHolderType =
typename ConvertToDeviceExpression<Arg3Impl>::Type;
using Type =
const TensorCwiseTernaryOp<OP, Arg1PlaceHolderType, Arg2PlaceHolderType,
Arg3PlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorCwiseTernaryOp
template <typename OP, typename Arg1Impl, typename Arg2Impl, typename Arg3Impl>
struct ConvertToDeviceExpression<
TensorCwiseTernaryOp<OP, Arg1Impl, Arg2Impl, Arg3Impl>> {
using Arg1PlaceHolderType =
typename ConvertToDeviceExpression<Arg1Impl>::Type;
using Arg2PlaceHolderType =
typename ConvertToDeviceExpression<Arg2Impl>::Type;
using Arg3PlaceHolderType =
typename ConvertToDeviceExpression<Arg3Impl>::Type;
using Type = TensorCwiseTernaryOp<OP, Arg1PlaceHolderType,
Arg2PlaceHolderType, Arg3PlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr>
struct ConvertToDeviceExpression<
const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> {
using IfPlaceHolderType = typename ConvertToDeviceExpression<IfExpr>::Type;
using ThenPlaceHolderType =
typename ConvertToDeviceExpression<ThenExpr>::Type;
using ElsePlaceHolderType =
typename ConvertToDeviceExpression<ElseExpr>::Type;
using Type = const TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType,
ElsePlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr>
struct ConvertToDeviceExpression<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> {
using IfPlaceHolderType = typename ConvertToDeviceExpression<IfExpr>::Type;
using ThenPlaceHolderType =
typename ConvertToDeviceExpression<ThenExpr>::Type;
using ElsePlaceHolderType =
typename ConvertToDeviceExpression<ElseExpr>::Type;
using Type = TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType,
ElsePlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node /// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const AssingOP /// type is const AssingOP
template <typename LHSExpr, typename RHSExpr> #define ASSIGNCONVERT(CVQual, Res)\
struct ConvertToDeviceExpression<const TensorAssignOp<LHSExpr, RHSExpr>> { template <typename LHSExpr, typename RHSExpr>\
using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type; struct ConvertToDeviceExpression<CVQual TensorAssignOp<LHSExpr, RHSExpr> >\
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; : NonOpConversion<TensorAssignOp, Res, LHSExpr, RHSExpr>{};
using Type = const TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>;
}; ASSIGNCONVERT(const, true)
ASSIGNCONVERT(, false)
#undef ASSIGNCONVERT
/// specialisation of the \ref ConvertToDeviceExpression struct when the node /// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is AssingOP /// type is either TensorForcedEvalOp or TensorEvalToOp
template <typename LHSExpr, typename RHSExpr> #define KERNELBROKERCONVERT(CVQual, Res, ExprNode)\
struct ConvertToDeviceExpression<TensorAssignOp<LHSExpr, RHSExpr>> { template <typename Expr>\
using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type; struct ConvertToDeviceExpression<CVQual ExprNode<Expr> > \
using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type; : DeviceConvertor<ExprNode, Res, Expr>{};
using Type = TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node KERNELBROKERCONVERT(const, true, TensorForcedEvalOp)
/// type is const TensorForcedEvalOp KERNELBROKERCONVERT(, false, TensorForcedEvalOp)
template <typename Expr> KERNELBROKERCONVERT(const, true, TensorEvalToOp)
struct ConvertToDeviceExpression<const TensorForcedEvalOp<Expr>> { KERNELBROKERCONVERT(, false, TensorEvalToOp)
using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type; #undef KERNELBROKERCONVERT
using Type = const TensorForcedEvalOp<PlaceHolderType, MakeGlobalPointer>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorForcedEvalOp
template <typename Expr>
struct ConvertToDeviceExpression<TensorForcedEvalOp<Expr>> {
using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type;
using Type = TensorForcedEvalOp<PlaceHolderType, MakeGlobalPointer>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is const TensorEvalToOp
template <typename Expr>
struct ConvertToDeviceExpression<const TensorEvalToOp<Expr>> {
using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type;
using Type = const TensorEvalToOp<PlaceHolderType, MakeGlobalPointer>;
};
/// specialisation of the \ref ConvertToDeviceExpression struct when the node
/// type is TensorEvalToOp
template <typename Expr>
struct ConvertToDeviceExpression<TensorEvalToOp<Expr>> {
using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type;
using Type = TensorEvalToOp<PlaceHolderType, MakeGlobalPointer>;
};
} // namespace internal } // namespace internal
} // namespace TensorSycl } // namespace TensorSycl
} // namespace Eigen } // namespace Eigen

550
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h
View File

@ -19,8 +19,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXPR_CONSTRUCTOR_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXPR_CONSTRUCTOR_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
@ -30,8 +30,7 @@ namespace internal {
template <typename PtrType, size_t N, typename... Params> template <typename PtrType, size_t N, typename... Params>
struct EvalToLHSConstructor { struct EvalToLHSConstructor {
PtrType expr; PtrType expr;
EvalToLHSConstructor(const utility::tuple::Tuple<Params...> &t) EvalToLHSConstructor(const utility::tuple::Tuple<Params...> &t): expr((&(*(utility::tuple::get<N>(t).get_pointer())))) {}
: expr((&(*(utility::tuple::get<N>(t).get_pointer())))) {}
}; };
/// \struct ExprConstructor is used to reconstruct the expression on the device /// \struct ExprConstructor is used to reconstruct the expression on the device
@ -44,447 +43,166 @@ struct EvalToLHSConstructor {
template <typename OrigExpr, typename IndexExpr, typename... Params> template <typename OrigExpr, typename IndexExpr, typename... Params>
struct ExprConstructor; struct ExprConstructor;
/// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorMap
template <typename Scalar_, int Options_, int Options2_, int Options3_,
int NumIndices_, typename IndexType_,
template <class> class MakePointer_, size_t N, typename... Params>
struct ExprConstructor<
const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakeGlobalPointer>,
const Eigen::internal::PlaceHolder<
const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options3_, MakePointer_>,
N>,
Params...> {
using Type =
const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakeGlobalPointer>;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)
: expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))),
fd.dimensions())) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorMap /// TensorMap
template <typename Scalar_, int Options_, int Options2_, int Options3_, #define TENSORMAP(CVQual)\
int NumIndices_, typename IndexType_, template <typename Scalar_, int Options_, int Options2_, int Options3_, int NumIndices_, typename IndexType_,\
template <class> class MakePointer_, size_t N, typename... Params> template <class> class MakePointer_, size_t N, typename... Params>\
struct ExprConstructor< struct ExprConstructor< CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer>,\
TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, CVQual Eigen::internal::PlaceHolder<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options3_, MakePointer_>, N>, Params...>{\
MakeGlobalPointer>, typedef CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakeGlobalPointer> Type;\
Eigen::internal::PlaceHolder< Type expr;\
TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options3_, template <typename FuncDetector>\
MakePointer_>, ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)\
N>, : expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), fd.dimensions())) {}\
Params...> {
using Type = TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakeGlobalPointer>;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)
: expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))),
fd.dimensions())) {}
}; };
TENSORMAP(const)
TENSORMAP()
#undef TENSORMAP
#define UNARYCATEGORY(CVQual)\
template <template<class, class> class UnaryCategory, typename OP, typename OrigRHSExpr, typename RHSExpr, typename... Params>\
struct ExprConstructor<CVQual UnaryCategory<OP, OrigRHSExpr>, CVQual UnaryCategory<OP, RHSExpr>, Params...> {\
typedef ExprConstructor<OrigRHSExpr, RHSExpr, Params...> my_type;\
my_type rhsExpr;\
typedef CVQual UnaryCategory<OP, typename my_type::Type> Type;\
Type expr;\
template <typename FuncDetector>\
ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}\
};
UNARYCATEGORY(const)
UNARYCATEGORY()
#undef UNARYCATEGORY
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorCwiseNullaryOp /// TensorBinaryOp
template <typename OP, typename OrigRHSExpr, typename RHSExpr, #define BINARYCATEGORY(CVQual)\
typename... Params> template <template<class, class, class> class BinaryCategory, typename OP, typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr,\
struct ExprConstructor<TensorCwiseNullaryOp<OP, OrigRHSExpr>, typename RHSExpr, typename... Params>\
TensorCwiseNullaryOp<OP, RHSExpr>, Params...> { struct ExprConstructor<CVQual BinaryCategory<OP, OrigLHSExpr, OrigRHSExpr>, CVQual BinaryCategory<OP, LHSExpr, RHSExpr>, Params...> {\
using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; typedef ExprConstructor<OrigLHSExpr, LHSExpr, Params...> my_left_type;\
my_type rhsExpr; typedef ExprConstructor<OrigRHSExpr, RHSExpr, Params...> my_right_type;\
using Type = TensorCwiseNullaryOp<OP, typename my_type::Type>; typedef CVQual BinaryCategory<OP, typename my_left_type::Type, typename my_right_type::Type> Type;\
Type expr; my_left_type lhsExpr;\
my_right_type rhsExpr;\
template <typename FuncDetector> Type expr;\
ExprConstructor(FuncDetector &funcD, template <typename FuncDetector>\
const utility::tuple::Tuple<Params...> &t) ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {} : lhsExpr(funcD.lhsExpr, t),rhsExpr(funcD.rhsExpr, t), expr(lhsExpr.expr, rhsExpr.expr, funcD.func) {}\
}; };
/// specialisation of the \ref ExprConstructor struct when the node type is BINARYCATEGORY(const)
/// const TensorCwiseNullaryOp BINARYCATEGORY()
template <typename OP, typename OrigRHSExpr, typename RHSExpr, #undef BINARYCATEGORY
typename... Params>
struct ExprConstructor<const TensorCwiseNullaryOp<OP, OrigRHSExpr>,
const TensorCwiseNullaryOp<OP, RHSExpr>, Params...> {
using my_type = const ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
my_type rhsExpr;
using Type = const TensorCwiseNullaryOp<OP, typename my_type::Type>;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorBroadcastingOp
template <typename OP, typename OrigRHSExpr, typename RHSExpr,
typename... Params>
struct ExprConstructor<TensorBroadcastingOp<OP, OrigRHSExpr>,
TensorBroadcastingOp<OP, RHSExpr>, Params...> {
using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
my_type rhsExpr;
using Type = TensorBroadcastingOp<OP, typename my_type::Type>;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorBroadcastingOp
template <typename OP, typename OrigRHSExpr, typename RHSExpr,
typename... Params>
struct ExprConstructor<const TensorBroadcastingOp<OP, OrigRHSExpr>,
const TensorBroadcastingOp<OP, RHSExpr>, Params...> {
using my_type = const ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
my_type rhsExpr;
using Type = const TensorBroadcastingOp<OP, typename my_type::Type>;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorCwiseUnaryOp
template <typename OP, typename OrigRHSExpr, typename RHSExpr,
typename... Params>
struct ExprConstructor<TensorCwiseUnaryOp<OP, OrigRHSExpr>,
TensorCwiseUnaryOp<OP, RHSExpr>, Params...> {
using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
using Type = TensorCwiseUnaryOp<OP, typename my_type::Type>;
my_type rhsExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD, utility::tuple::Tuple<Params...> &t)
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorCwiseUnaryOp
template <typename OP, typename OrigRHSExpr, typename RHSExpr,
typename... Params>
struct ExprConstructor<const TensorCwiseUnaryOp<OP, OrigRHSExpr>,
const TensorCwiseUnaryOp<OP, RHSExpr>, Params...> {
using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
using Type = const TensorCwiseUnaryOp<OP, typename my_type::Type>;
my_type rhsExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: rhsExpr(funcD.rhsExpr, t), expr(rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorCwiseBinaryOp
template <typename OP, typename OrigLHSExpr, typename OrigRHSExpr,
typename LHSExpr, typename RHSExpr, typename... Params>
struct ExprConstructor<TensorCwiseBinaryOp<OP, OrigLHSExpr, OrigRHSExpr>,
TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Params...> {
using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>;
using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
using Type = TensorCwiseBinaryOp<OP, typename my_left_type::Type,
typename my_right_type::Type>;
my_left_type lhsExpr;
my_right_type rhsExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: lhsExpr(funcD.lhsExpr, t),
rhsExpr(funcD.rhsExpr, t),
expr(lhsExpr.expr, rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorCwiseBinaryOp
template <typename OP, typename OrigLHSExpr, typename OrigRHSExpr,
typename LHSExpr, typename RHSExpr, typename... Params>
struct ExprConstructor<const TensorCwiseBinaryOp<OP, OrigLHSExpr, OrigRHSExpr>,
const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>,
Params...> {
using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>;
using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
using Type = const TensorCwiseBinaryOp<OP, typename my_left_type::Type,
typename my_right_type::Type>;
my_left_type lhsExpr;
my_right_type rhsExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: lhsExpr(funcD.lhsExpr, t),
rhsExpr(funcD.rhsExpr, t),
expr(lhsExpr.expr, rhsExpr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorCwiseTernaryOp
template <typename OP, typename OrigArg1Expr, typename OrigArg2Expr,
typename OrigArg3Expr, typename Arg1Expr, typename Arg2Expr,
typename Arg3Expr, typename... Params>
struct ExprConstructor<
const TensorCwiseTernaryOp<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>,
const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Params...> {
using my_arg1_type = ExprConstructor<OrigArg1Expr, Arg1Expr, Params...>;
using my_arg2_type = ExprConstructor<OrigArg2Expr, Arg2Expr, Params...>;
using my_arg3_type = ExprConstructor<OrigArg3Expr, Arg3Expr, Params...>;
using Type = const TensorCwiseTernaryOp<OP, typename my_arg1_type::Type,
typename my_arg2_type::Type,
typename my_arg3_type::Type>;
my_arg1_type arg1Expr;
my_arg2_type arg2Expr;
my_arg3_type arg3Expr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: arg1Expr(funcD.arg1Expr, t),
arg2Expr(funcD.arg2Expr, t),
arg3Expr(funcD.arg3Expr, t),
expr(arg1Expr.expr, arg2Expr.expr, arg3Expr.expr, funcD.func) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorCwiseTernaryOp /// TensorCwiseTernaryOp
template <typename OP, typename OrigArg1Expr, typename OrigArg2Expr, #define TERNARYCATEGORY(CVQual)\
typename OrigArg3Expr, typename Arg1Expr, typename Arg2Expr, template <template <class, class, class, class> class TernaryCategory, typename OP, typename OrigArg1Expr, typename OrigArg2Expr,typename OrigArg3Expr,\
typename Arg3Expr, typename... Params> typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename... Params>\
struct ExprConstructor< struct ExprConstructor<CVQual TernaryCategory<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>, CVQual TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Params...> {\
TensorCwiseTernaryOp<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>, typedef ExprConstructor<OrigArg1Expr, Arg1Expr, Params...> my_arg1_type;\
TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Params...> { typedef ExprConstructor<OrigArg2Expr, Arg2Expr, Params...> my_arg2_type;\
using my_arg1_type = ExprConstructor<OrigArg1Expr, Arg1Expr, Params...>; typedef ExprConstructor<OrigArg3Expr, Arg3Expr, Params...> my_arg3_type;\
using my_arg2_type = ExprConstructor<OrigArg2Expr, Arg2Expr, Params...>; typedef CVQual TernaryCategory<OP, typename my_arg1_type::Type, typename my_arg2_type::Type, typename my_arg3_type::Type> Type;\
using my_arg3_type = ExprConstructor<OrigArg3Expr, Arg3Expr, Params...>; my_arg1_type arg1Expr;\
using Type = TensorCwiseTernaryOp<OP, typename my_arg1_type::Type, my_arg2_type arg2Expr;\
typename my_arg2_type::Type, my_arg3_type arg3Expr;\
typename my_arg3_type::Type>; Type expr;\
template <typename FuncDetector>\
my_arg1_type arg1Expr; ExprConstructor(FuncDetector &funcD,const utility::tuple::Tuple<Params...> &t)\
my_arg2_type arg2Expr; : arg1Expr(funcD.arg1Expr, t), arg2Expr(funcD.arg2Expr, t), arg3Expr(funcD.arg3Expr, t), expr(arg1Expr.expr, arg2Expr.expr, arg3Expr.expr, funcD.func) {}\
my_arg3_type arg3Expr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: arg1Expr(funcD.arg1Expr, t),
arg2Expr(funcD.arg2Expr, t),
arg3Expr(funcD.arg3Expr, t),
expr(arg1Expr.expr, arg2Expr.expr, arg3Expr.expr, funcD.func) {}
}; };
/// specialisation of the \ref ExprConstructor struct when the node type is TERNARYCATEGORY(const)
/// const TensorCwiseSelectOp TERNARYCATEGORY()
template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr, #undef TERNARYCATEGORY
typename IfExpr, typename ThenExpr, typename ElseExpr,
typename... Params>
struct ExprConstructor<
const TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>,
const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Params...> {
using my_if_type = ExprConstructor<OrigIfExpr, IfExpr, Params...>;
using my_then_type = ExprConstructor<OrigThenExpr, ThenExpr, Params...>;
using my_else_type = ExprConstructor<OrigElseExpr, ElseExpr, Params...>;
using Type = const TensorSelectOp<typename my_if_type::Type,
typename my_then_type::Type,
typename my_else_type::Type>;
my_if_type ifExpr;
my_then_type thenExpr;
my_else_type elseExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: ifExpr(funcD.ifExpr, t),
thenExpr(funcD.thenExpr, t),
elseExpr(funcD.elseExpr, t),
expr(ifExpr.expr, thenExpr.expr, elseExpr.expr) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorCwiseSelectOp /// TensorCwiseSelectOp
template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr, #define SELECTOP(CVQual)\
typename IfExpr, typename ThenExpr, typename ElseExpr, template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr, typename IfExpr, typename ThenExpr, typename ElseExpr, typename... Params>\
typename... Params> struct ExprConstructor< CVQual TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>, CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Params...> {\
struct ExprConstructor<TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>, typedef ExprConstructor<OrigIfExpr, IfExpr, Params...> my_if_type;\
TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Params...> { typedef ExprConstructor<OrigThenExpr, ThenExpr, Params...> my_then_type;\
using my_if_type = ExprConstructor<OrigIfExpr, IfExpr, Params...>; typedef ExprConstructor<OrigElseExpr, ElseExpr, Params...> my_else_type;\
using my_then_type = ExprConstructor<OrigThenExpr, ThenExpr, Params...>; typedef CVQual TensorSelectOp<typename my_if_type::Type, typename my_then_type::Type, typename my_else_type::Type> Type;\
using my_else_type = ExprConstructor<OrigElseExpr, ElseExpr, Params...>; my_if_type ifExpr;\
using Type = my_then_type thenExpr;\
TensorSelectOp<typename my_if_type::Type, typename my_then_type::Type, my_else_type elseExpr;\
typename my_else_type::Type>; Type expr;\
template <typename FuncDetector>\
my_if_type ifExpr; ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
my_then_type thenExpr; : ifExpr(funcD.ifExpr, t), thenExpr(funcD.thenExpr, t), elseExpr(funcD.elseExpr, t), expr(ifExpr.expr, thenExpr.expr, elseExpr.expr) {}\
my_else_type elseExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: ifExpr(funcD.ifExpr, t),
thenExpr(funcD.thenExpr, t),
elseExpr(funcD.elseExpr, t),
expr(ifExpr.expr, thenExpr.expr, elseExpr.expr) {}
}; };
/// specialisation of the \ref ExprConstructor struct when the node type is SELECTOP(const)
/// TensorAssignOp SELECTOP()
template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr, #undef SELECTOP
typename RHSExpr, typename... Params>
struct ExprConstructor<TensorAssignOp<OrigLHSExpr, OrigRHSExpr>,
TensorAssignOp<LHSExpr, RHSExpr>, Params...> {
using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>;
using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
using Type =
TensorAssignOp<typename my_left_type::Type, typename my_right_type::Type>;
my_left_type lhsExpr;
my_right_type rhsExpr;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: lhsExpr(funcD.lhsExpr, t),
rhsExpr(funcD.rhsExpr, t),
expr(lhsExpr.expr, rhsExpr.expr) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorAssignOp /// const TensorAssignOp
template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr, #define ASSIGN(CVQual)\
typename RHSExpr, typename... Params> template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr, typename RHSExpr, typename... Params>\
struct ExprConstructor<const TensorAssignOp<OrigLHSExpr, OrigRHSExpr>, struct ExprConstructor<CVQual TensorAssignOp<OrigLHSExpr, OrigRHSExpr>, CVQual TensorAssignOp<LHSExpr, RHSExpr>, Params...> {\
const TensorAssignOp<LHSExpr, RHSExpr>, Params...> { typedef ExprConstructor<OrigLHSExpr, LHSExpr, Params...> my_left_type;\
using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>; typedef ExprConstructor<OrigRHSExpr, RHSExpr, Params...> my_right_type;\
using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>; typedef CVQual TensorAssignOp<typename my_left_type::Type, typename my_right_type::Type> Type;\
using Type = const TensorAssignOp<typename my_left_type::Type, my_left_type lhsExpr;\
typename my_right_type::Type>; my_right_type rhsExpr;\
Type expr;\
my_left_type lhsExpr; template <typename FuncDetector>\
my_right_type rhsExpr; ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
Type expr; : lhsExpr(funcD.lhsExpr, t), rhsExpr(funcD.rhsExpr, t), expr(lhsExpr.expr, rhsExpr.expr) {}\
template <typename FuncDetector> };
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: lhsExpr(funcD.lhsExpr, t),
rhsExpr(funcD.rhsExpr, t),
expr(lhsExpr.expr, rhsExpr.expr) {}
};
ASSIGN(const)
ASSIGN()
#undef ASSIGN
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorEvalToOp /// TensorEvalToOp
template <typename OrigExpr, typename Expr, typename... Params> #define EVALTO(CVQual)\
struct ExprConstructor<const TensorEvalToOp<OrigExpr, MakeGlobalPointer>, template <typename OrigExpr, typename Expr, typename... Params>\
const TensorEvalToOp<Expr>, Params...> { struct ExprConstructor<CVQual TensorEvalToOp<OrigExpr, MakeGlobalPointer>, CVQual TensorEvalToOp<Expr>, Params...> {\
using my_expr_type = ExprConstructor<OrigExpr, Expr, Params...>; typedef ExprConstructor<OrigExpr, Expr, Params...> my_expr_type;\
using my_buffer_type = typedef typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType my_buffer_type;\
typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType; typedef CVQual TensorEvalToOp<typename my_expr_type::Type, MakeGlobalPointer> Type;\
using Type = my_expr_type nestedExpression;\
const TensorEvalToOp<typename my_expr_type::Type, MakeGlobalPointer>; EvalToLHSConstructor<my_buffer_type, 0, Params...> buffer;\
my_expr_type nestedExpression; Type expr;\
EvalToLHSConstructor<my_buffer_type, 0, Params...> buffer; template <typename FuncDetector>\
Type expr; ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\
: nestedExpression(funcD.rhsExpr, t), buffer(t), expr(buffer.expr, nestedExpression.expr) {}\
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: nestedExpression(funcD.rhsExpr, t),
buffer(t),
expr(buffer.expr, nestedExpression.expr) {}
}; };
/// specialisation of the \ref ExprConstructor struct when the node type is EVALTO(const)
/// TensorEvalToOp EVALTO()
template <typename OrigExpr, typename Expr, typename... Params> #undef EVALTO
struct ExprConstructor<TensorEvalToOp<OrigExpr, MakeGlobalPointer>,
TensorEvalToOp<Expr>, Params...> {
using my_expr_type = ExprConstructor<OrigExpr, Expr, Params...>;
using my_buffer_type =
typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType;
using Type = TensorEvalToOp<typename my_expr_type::Type>;
my_expr_type nestedExpression;
EvalToLHSConstructor<my_buffer_type, 0, Params...> buffer;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &funcD,
const utility::tuple::Tuple<Params...> &t)
: nestedExpression(funcD.rhsExpr, t),
buffer(t),
expr(buffer.expr, nestedExpression.expr) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is
/// const TensorForcedEvalOp
template <typename OrigExpr, typename DevExpr, size_t N, typename... Params>
struct ExprConstructor<
const TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>,
const Eigen::internal::PlaceHolder<const TensorForcedEvalOp<DevExpr>, N>,
Params...> {
using Type = const TensorMap<
Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar,
TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::NumDimensions, 0,
typename TensorForcedEvalOp<DevExpr>::Index>,
0, MakeGlobalPointer>;
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)
: expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))),
fd.dimensions())) {}
};
/// specialisation of the \ref ExprConstructor struct when the node type is /// specialisation of the \ref ExprConstructor struct when the node type is
/// TensorForcedEvalOp /// TensorForcedEvalOp
template <typename OrigExpr, typename DevExpr, size_t N, typename... Params> #define FORCEDEVAL(CVQual)\
struct ExprConstructor< template <typename OrigExpr, typename DevExpr, size_t N, typename... Params>\
const TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>, struct ExprConstructor<CVQual TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>,\
const Eigen::internal::PlaceHolder<TensorForcedEvalOp<DevExpr>, N>, CVQual Eigen::internal::PlaceHolder<CVQual TensorForcedEvalOp<DevExpr>, N>, Params...> {\
Params...> { typedef CVQual TensorMap<Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar,\
using Type = TensorMap< TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::NumDimensions, 0, typename TensorForcedEvalOp<DevExpr>::Index>, 0, MakeGlobalPointer> Type;\
Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar, 1, Type expr;\
0, typename TensorForcedEvalOp<DevExpr>::Index>, template <typename FuncDetector>\
0, MakeGlobalPointer>; ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)\
: expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))), fd.dimensions())) {}\
Type expr;
template <typename FuncDetector>
ExprConstructor(FuncDetector &fd, const utility::tuple::Tuple<Params...> &t)
: expr(Type((&(*(utility::tuple::get<N>(t).get_pointer()))),
fd.dimensions())) {}
}; };
FORCEDEVAL(const)
FORCEDEVAL()
#undef FORCEDEVAL
/// template deduction for \ref ExprConstructor struct /// template deduction for \ref ExprConstructor struct
template <typename OrigExpr, typename IndexExpr, typename FuncD, template <typename OrigExpr, typename IndexExpr, typename FuncD, typename... Params>
typename... Params> auto createDeviceExpression(FuncD &funcD, const utility::tuple::Tuple<Params...> &t)
auto createDeviceExpression(FuncD &funcD,
const utility::tuple::Tuple<Params...> &t)
-> decltype(ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t)) { -> decltype(ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t)) {
return ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t); return ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t);
} }
@ -492,4 +210,4 @@ auto createDeviceExpression(FuncD &funcD,
} }
} // namespace Eigen } // namespace Eigen
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXPR_CONSTRUCTOR_HPP

447
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h
View File

@ -29,8 +29,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
@ -42,425 +42,160 @@ namespace internal {
template <typename Evaluator> template <typename Evaluator>
struct ExtractAccessor; struct ExtractAccessor;
/// specialisation of the \ref ExtractAccessor struct when the node type is struct AccessorConstructor{
/// const TensorMap template<typename Arg> static inline auto getTuple(cl::sycl::handler& cgh, Arg eval)
template <typename PlainObjectType, int Options_, typename Dev> -> decltype(ExtractAccessor<Arg>::getTuple(cgh, eval)) {
struct ExtractAccessor< return ExtractAccessor<Arg>::getTuple(cgh, eval);
TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>> { }
using actual_type = typename Eigen::internal::remove_all<
typename Eigen::internal::traits<PlainObjectType>::Scalar>::type; template<typename Arg1, typename Arg2> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1, Arg2 eval2)
static inline auto getTuple( -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2))) {
cl::sycl::handler& cgh, return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1), ExtractAccessor<Arg2>::getTuple(cgh, eval2));
const TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev> }
eval) template<typename Arg1, typename Arg2, typename Arg3> static inline auto getTuple(cl::sycl::handler& cgh, Arg1 eval1 , Arg2 eval2 , Arg3 eval3)
-> decltype(utility::tuple::make_tuple( -> decltype(utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3)))) {
(eval.device() return utility::tuple::append(ExtractAccessor<Arg1>::getTuple(cgh, eval1),utility::tuple::append(ExtractAccessor<Arg2>::getTuple(cgh, eval2), ExtractAccessor<Arg3>::getTuple(cgh, eval3)));
.template get_sycl_accessor<cl::sycl::access::mode::read, true, }
actual_type>( template< cl::sycl::access::mode AcM, typename Arg> static inline auto getAccessor(cl::sycl::handler& cgh, Arg eval)
eval.dimensions().TotalSize(), cgh, -> decltype(utility::tuple::make_tuple( eval.device().template get_sycl_accessor<AcM, true,
eval.derived().data())))) { typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data()))){
return utility::tuple::make_tuple( return utility::tuple::make_tuple(eval.device().template get_sycl_accessor<AcM, true, typename Eigen::internal::remove_all<typename Arg::CoeffReturnType>::type>(eval.dimensions().TotalSize(), cgh,eval.data()));
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::read, true,
actual_type>(
eval.dimensions().TotalSize(), cgh, eval.derived().data())));
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorMap /// const TensorCwiseNullaryOp, const TensorCwiseUnaryOp and const TensorBroadcastingOp
template <typename PlainObjectType, int Options_, typename Dev> template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor< struct ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > {
TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>> { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> eval)
using actual_type = typename Eigen::internal::remove_all< -> decltype(AccessorConstructor::getTuple(cgh, eval.impl())){
typename Eigen::internal::traits<PlainObjectType>::Scalar>::type; return AccessorConstructor::getTuple(cgh, eval.impl());
static inline auto getTuple(
cl::sycl::handler& cgh,
TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev> eval)
-> decltype(utility::tuple::make_tuple(
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::read_write,
true, actual_type>(
eval.dimensions().TotalSize(), cgh,
eval.derived().data())))) {
return utility::tuple::make_tuple(
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::read_write,
true, actual_type>(
eval.dimensions().TotalSize(), cgh, eval.derived().data())));
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorCwiseNullaryOp /// TensorCwiseNullaryOp, TensorCwiseUnaryOp and TensorBroadcastingOp
template <typename OP, typename RHSExpr, typename Dev> template <template<class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor< struct ExtractAccessor<TensorEvaluator<UnaryCategory<OP, RHSExpr>, Dev> >
TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> { : ExtractAccessor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > {};
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev> eval)
-> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl())) {
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl());
return RHSTuple;
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorCwiseNullaryOp
template <typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor<
TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> {
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev> eval)
-> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl())) {
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl());
return RHSTuple;
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorBroadcastingOp
template <typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor<
TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev>> {
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev> eval)
-> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl())) {
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl());
return RHSTuple;
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorBroadcastingOp
template <typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor<
TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev>> {
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev> eval)
-> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl())) {
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl());
return RHSTuple;
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TenosorCwiseUnary
template <typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor<
TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> {
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev> eval)
-> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl())) {
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl());
return RHSTuple;
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is
/// TenosorCwiseUnary
template <typename OP, typename RHSExpr, typename Dev>
struct ExtractAccessor<TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> {
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev> eval)
-> decltype(ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl())) {
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.impl());
return RHSTuple;
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorCwiseBinaryOp /// const TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
struct ExtractAccessor< struct ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > {
TensorEvaluator<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> eval)
static auto getTuple(cl::sycl::handler& cgh, -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){
const TensorEvaluator< return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl());
const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>
eval)
-> decltype(utility::tuple::append(
ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
cgh, eval.left_impl()),
ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.right_impl()))) {
auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
cgh, eval.left_impl());
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.right_impl());
return utility::tuple::append(LHSTuple, RHSTuple);
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorCwiseBinaryOp /// TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
struct ExtractAccessor< struct ExtractAccessor<TensorEvaluator<BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >
TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { : ExtractAccessor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >{};
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>
eval)
-> decltype(utility::tuple::append(
ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
cgh, eval.left_impl()),
ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.right_impl()))) {
auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
cgh, eval.left_impl());
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.right_impl());
return utility::tuple::append(LHSTuple, RHSTuple);
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorCwiseTernaryOp /// const TensorCwiseTernaryOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev>
typename Dev> struct ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > {
struct ExtractAccessor<TensorEvaluator< static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> eval)
const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> { -> decltype(AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl())){
static auto getTuple( return AccessorConstructor::getTuple(cgh, eval.arg1Impl(), eval.arg2Impl(), eval.arg3Impl());
cl::sycl::handler& cgh,
const TensorEvaluator<
const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>
eval)
-> decltype(utility::tuple::append(
ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple(
cgh, eval.arg1Impl()),
utility::tuple::append(
ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple(
cgh, eval.arg2Impl()),
ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple(
cgh, eval.arg3Impl())))) {
auto Arg1Tuple = ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple(
cgh, eval.arg1Impl());
auto Arg2Tuple = ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple(
cgh, eval.arg2Impl());
auto Arg3Tuple = ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple(
cgh, eval.arg3Impl());
return utility::tuple::append(Arg1Tuple,
utility::tuple::append(Arg2Tuple, Arg3Tuple));
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorCwiseTernaryOp /// TensorCwiseTernaryOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, template <template<class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev>
typename Dev> struct ExtractAccessor<TensorEvaluator<TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >
struct ExtractAccessor<TensorEvaluator< : ExtractAccessor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >{};
TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> {
static auto getTuple( /// specialisation of the \ref ExtractAccessor struct when the node type is
cl::sycl::handler& cgh, /// const TensorCwiseSelectOp. This is a special case where there is no OP
const TensorEvaluator< template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> struct ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > {
eval) static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> eval)
-> decltype(utility::tuple::append( -> decltype(AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl())){
ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple( return AccessorConstructor::getTuple(cgh, eval.cond_impl(), eval.then_impl(), eval.else_impl());
cgh, eval.arg1Impl()),
utility::tuple::append(
ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple(
cgh, eval.arg2Impl()),
ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple(
cgh, eval.arg3Impl())))) {
auto Arg1Tuple = ExtractAccessor<TensorEvaluator<Arg1Expr, Dev>>::getTuple(
cgh, eval.arg1Impl());
auto Arg2Tuple = ExtractAccessor<TensorEvaluator<Arg2Expr, Dev>>::getTuple(
cgh, eval.arg2Impl());
auto Arg3Tuple = ExtractAccessor<TensorEvaluator<Arg3Expr, Dev>>::getTuple(
cgh, eval.arg3Impl());
return utility::tuple::append(Arg1Tuple,
utility::tuple::append(Arg2Tuple, Arg3Tuple));
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorCwiseSelectOp /// TensorCwiseSelectOp. This is a special case where there is no OP
template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
struct ExtractAccessor< struct ExtractAccessor<TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> >
TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { : ExtractAccessor<TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> >{};
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>,
Dev>
eval)
-> decltype(utility::tuple::append(
ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple(
cgh, eval.cond_impl()),
utility::tuple::append(
ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple(
cgh, eval.then_impl()),
ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple(
cgh, eval.else_impl())))) {
auto IfTuple = ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple(
cgh, eval.cond_impl());
auto ThenTuple = ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple(
cgh, eval.then_impl());
auto ElseTuple = ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple(
cgh, eval.else_impl());
return utility::tuple::append(IfTuple,
utility::tuple::append(ThenTuple, ElseTuple));
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
struct ExtractAccessor<
TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> {
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>
eval)
-> decltype(utility::tuple::append(
ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple(
cgh, eval.cond_impl()),
utility::tuple::append(
ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple(
cgh, eval.then_impl()),
ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple(
cgh, eval.else_impl())))) {
auto IfTuple = ExtractAccessor<TensorEvaluator<IfExpr, Dev>>::getTuple(
cgh, eval.cond_impl());
auto ThenTuple = ExtractAccessor<TensorEvaluator<ThenExpr, Dev>>::getTuple(
cgh, eval.then_impl());
auto ElseTuple = ExtractAccessor<TensorEvaluator<ElseExpr, Dev>>::getTuple(
cgh, eval.else_impl());
return utility::tuple::append(IfTuple,
utility::tuple::append(ThenTuple, ElseTuple));
}
};
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorAssignOp /// const TensorAssignOp
template <typename LHSExpr, typename RHSExpr, typename Dev> template <typename LHSExpr, typename RHSExpr, typename Dev>
struct ExtractAccessor< struct ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> > {
TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval)
static auto getTuple( -> decltype(AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl())){
cl::sycl::handler& cgh, return AccessorConstructor::getTuple(cgh, eval.left_impl(), eval.right_impl());
const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) }
-> decltype(utility::tuple::append(
ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
cgh, eval.left_impl()),
ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.right_impl()))) {
auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
cgh, eval.left_impl());
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
cgh, eval.right_impl());
return utility::tuple::append(LHSTuple, RHSTuple);
}
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorAssignOp /// TensorAssignOp
template <typename LHSExpr, typename RHSExpr, typename Dev> template <typename LHSExpr, typename RHSExpr, typename Dev>
struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> >
static auto getTuple( : ExtractAccessor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{};
cl::sycl::handler& cgh,
const TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval) /// specialisation of the \ref ExtractAccessor struct when the node type is
-> decltype(utility::tuple::append( /// const TensorMap
ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( #define TENSORMAPEXPR(CVQual, ACCType)\
eval.left_impl()), template <typename PlainObjectType, int Options_, typename Dev>\
ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( struct ExtractAccessor<TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> > {\
eval.right_impl()))) { static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<CVQual TensorMap<PlainObjectType, Options_>, Dev> eval)\
auto LHSTuple = ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple( -> decltype(AccessorConstructor::template getAccessor<ACCType>(cgh, eval)){\
cgh, eval.left_impl()); return AccessorConstructor::template getAccessor<ACCType>(cgh, eval);\
auto RHSTuple = ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple( }\
cgh, eval.right_impl());
return utility::tuple::append(LHSTuple, RHSTuple);
}
}; };
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 /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorForcedEvalOp /// const TensorForcedEvalOp
template <typename Expr, typename Dev> template <typename Expr, typename Dev>
struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> { struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> > {
using actual_type = static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> eval)
typename Eigen::internal::remove_all<typename TensorEvaluator< -> decltype(AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval)){
const TensorForcedEvalOp<Expr>, Dev>::CoeffReturnType>::type; return AccessorConstructor::template getAccessor<cl::sycl::access::mode::read>(cgh, eval);
static auto getTuple(
cl::sycl::handler& cgh,
const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> eval)
-> decltype(utility::tuple::make_tuple(
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::read, false,
actual_type>(
eval.dimensions().TotalSize(), cgh, eval.data())))) {
return utility::tuple::make_tuple(
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::read, false,
actual_type>(
eval.dimensions().TotalSize(), cgh, eval.data())));
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorForcedEvalOp /// TensorForcedEvalOp
template <typename Expr, typename Dev> template <typename Expr, typename Dev>
struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>> struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev> >
: ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> {}; : ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev> >{};
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// const TensorEvalToOp /// const TensorEvalToOp
template <typename Expr, typename Dev> template <typename Expr, typename Dev>
struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev>> { struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> > {
using actual_type = static inline auto getTuple(cl::sycl::handler& cgh,const TensorEvaluator<const TensorEvalToOp<Expr>, Dev> eval)
typename Eigen::internal::remove_all<typename TensorEvaluator< -> decltype(utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl()))){
const TensorEvalToOp<Expr>, Dev>::CoeffReturnType>::type; return utility::tuple::append(AccessorConstructor::template getAccessor<cl::sycl::access::mode::write>(cgh, eval), AccessorConstructor::getTuple(cgh, eval.impl()));
static auto getTuple(cl::sycl::handler& cgh,
TensorEvaluator<const TensorEvalToOp<Expr>, Dev> eval)
-> decltype(utility::tuple::append(
utility::tuple::make_tuple(
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::write,
false, actual_type>(
eval.dimensions().TotalSize(), cgh, eval.data()))),
ExtractAccessor<TensorEvaluator<Expr, Dev>>::getTuple(cgh,
eval.impl()))) {
auto LHSTuple = utility::tuple::make_tuple(
(eval.device()
.template get_sycl_accessor<cl::sycl::access::mode::write, false,
actual_type>(
eval.dimensions().TotalSize(), cgh, eval.data())));
auto RHSTuple =
ExtractAccessor<TensorEvaluator<Expr, Dev>>::getTuple(cgh, eval.impl());
return utility::tuple::append(LHSTuple, RHSTuple);
} }
}; };
/// specialisation of the \ref ExtractAccessor struct when the node type is /// specialisation of the \ref ExtractAccessor struct when the node type is
/// TensorEvalToOp /// TensorEvalToOp
template <typename Expr, typename Dev> template <typename Expr, typename Dev>
struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev>> struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev> >
: ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev>> {}; : ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev> >{};
/// template deduction for \ref ExtractAccessor /// template deduction for \ref ExtractAccessor
template <typename Evaluator> template <typename Evaluator>
auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr) auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr)
-> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) { -> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) {
return ExtractAccessor<Evaluator>::getTuple(cgh, expr); return ExtractAccessor<Evaluator>::getTuple(cgh, expr);
} }
} }
} }
} }
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_ACCESSOR_HPP

305
unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h
View File

@ -19,8 +19,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_FUNCTORS_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_FUNCTORS_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
@ -31,283 +31,124 @@ namespace internal {
/// expression on the device. /// expression on the device.
/// We have to do that as in Eigen the functors are not stateless so we cannot /// We have to do that as in Eigen the functors are not stateless so we cannot
/// re-instantiate them on the device. /// re-instantiate them on the device.
/// We have to pass whatever instantiated to the device. /// We have to pass instantiated functors to the device.
template <typename Evaluator> // This struct is used for leafNode (TensorMap) and nodes behaving like leafNode (TensorForcedEval).
struct FunctorExtractor; template <typename Evaluator> struct FunctorExtractor{
typedef typename Evaluator::Dimensions Dimensions;
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorMap:
template <typename PlainObjectType, int Options_, typename Dev>
struct FunctorExtractor<
TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>> {
using Dimensions = typename PlainObjectType::Dimensions;
const Dimensions m_dimensions; const Dimensions m_dimensions;
const Dimensions& dimensions() const { return m_dimensions; } const Dimensions& dimensions() const { return m_dimensions; }
FunctorExtractor( FunctorExtractor(const Evaluator& expr)
const TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>& expr) : m_dimensions(expr.dimensions()) {}
: m_dimensions(expr.dimensions()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorMap /// const TensorCwiseNullaryOp, const TensorCwiseUnaryOp, and const TensorBroadcastingOp
template <typename PlainObjectType, int Options_, typename Dev> template <template <class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> > {
TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>> { FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr;
using Dimensions = typename PlainObjectType::Dimensions;
const Dimensions m_dimensions;
const Dimensions& dimensions() const { return m_dimensions; }
FunctorExtractor(
const TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>&
expr)
: m_dimensions(expr.dimensions()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorForcedEvalOp
template <typename Expr, typename Dev>
struct FunctorExtractor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>> {
using Dimensions = typename Expr::Dimensions;
const Dimensions m_dimensions;
const Dimensions& dimensions() const { return m_dimensions; }
FunctorExtractor(const TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>& expr)
: m_dimensions(expr.dimensions()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorForcedEvalOp
template <typename Expr, typename Dev>
struct FunctorExtractor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> {
using Dimensions =
typename TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>::Dimensions;
const Dimensions m_dimensions;
const Dimensions& dimensions() const { return m_dimensions; }
FunctorExtractor(
const TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>& expr)
: m_dimensions(expr.dimensions()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorCwiseNullaryOp
template <typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor<
TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func; OP func;
FunctorExtractor( FunctorExtractor(const TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev>& expr)
TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>& expr) : rhsExpr(expr.impl()), func(expr.functor()) {}
: rhsExpr(expr.impl()), func(expr.functor()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorCwiseNullaryOp, TensorCwiseUnaryOp, and TensorBroadcastingOp
template <template <class, class> class UnaryCategory, typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor<TensorEvaluator<UnaryCategory<OP, RHSExpr>, Dev> >
: FunctorExtractor<TensorEvaluator<const UnaryCategory<OP, RHSExpr>, Dev> >{};
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorCwiseNullaryOp /// const TensorCwiseBinaryOp
template <typename OP, typename RHSExpr, typename Dev> template <template<class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> > {
TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> { FunctorExtractor<TensorEvaluator<LHSExpr, Dev> > lhsExpr;
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr;
OP func; OP func;
FunctorExtractor( FunctorExtractor(const TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev>& expr)
const TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>& expr) : lhsExpr(expr.left_impl()),rhsExpr(expr.right_impl()),func(expr.functor()) {}
: rhsExpr(expr.impl()), func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorBroadcastingOp
template <typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor<
TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func;
FunctorExtractor(
const TensorEvaluator<TensorBroadcastingOp<OP, RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()), func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorBroadcastingOp
template <typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor<
TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func;
FunctorExtractor(
const TensorEvaluator<const TensorBroadcastingOp<OP, RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()), func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorCwiseUnaryOp
template <typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor<TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func;
FunctorExtractor(
const TensorEvaluator<TensorCwiseUnaryOp<OP, RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()), func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorCwiseUnaryOp
template <typename OP, typename RHSExpr, typename Dev>
struct FunctorExtractor<
TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func;
FunctorExtractor(
const TensorEvaluator<const TensorCwiseUnaryOp<OP, RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()), func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
struct FunctorExtractor<
TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr;
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func;
FunctorExtractor(
const TensorEvaluator<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>&
expr)
: lhsExpr(expr.left_impl()),
rhsExpr(expr.right_impl()),
func(expr.functor()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorCwiseBinaryOp /// const TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev> template <template <class, class, class> class BinaryCategory, typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor<TensorEvaluator<BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >
TensorEvaluator<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> { : FunctorExtractor<TensorEvaluator<const BinaryCategory<OP, LHSExpr, RHSExpr>, Dev> >{};
FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr;
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
OP func;
FunctorExtractor(const TensorEvaluator<
const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>& expr)
: lhsExpr(expr.left_impl()),
rhsExpr(expr.right_impl()),
func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorCwiseTernaryOp /// const TensorCwiseTernaryOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, template <template <class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,typename Dev>
typename Dev> struct FunctorExtractor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> > {
struct FunctorExtractor<TensorEvaluator< FunctorExtractor<TensorEvaluator<Arg1Expr, Dev> > arg1Expr;
const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> { FunctorExtractor<TensorEvaluator<Arg2Expr, Dev> > arg2Expr;
FunctorExtractor<TensorEvaluator<Arg1Expr, Dev>> arg1Expr; FunctorExtractor<TensorEvaluator<Arg3Expr, Dev> > arg3Expr;
FunctorExtractor<TensorEvaluator<Arg2Expr, Dev>> arg2Expr;
FunctorExtractor<TensorEvaluator<Arg3Expr, Dev>> arg3Expr;
OP func; OP func;
FunctorExtractor(const TensorEvaluator< FunctorExtractor(const TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>& expr)
const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, : arg1Expr(expr.arg1Impl()), arg2Expr(expr.arg2Impl()), arg3Expr(expr.arg3Impl()), func(expr.functor()) {}
Dev>& expr)
: arg1Expr(expr.arg1Impl()),
arg2Expr(expr.arg2Impl()),
arg3Expr(expr.arg3Impl()),
func(expr.functor()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorCwiseTernaryOp /// TensorCwiseTernaryOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, template <template <class, class, class, class> class TernaryCategory, typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr, typename Dev>
typename Dev> struct FunctorExtractor<TensorEvaluator< TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >
struct FunctorExtractor<TensorEvaluator< :FunctorExtractor<TensorEvaluator<const TernaryCategory<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev> >{};
TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> {
FunctorExtractor<TensorEvaluator<Arg1Expr, Dev>> arg1Expr;
FunctorExtractor<TensorEvaluator<Arg2Expr, Dev>> arg2Expr;
FunctorExtractor<TensorEvaluator<Arg3Expr, Dev>> arg3Expr;
OP func;
FunctorExtractor(
const TensorEvaluator<
TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>& expr)
: arg1Expr(expr.arg1Impl()),
arg2Expr(expr.arg2Impl()),
arg3Expr(expr.arg3Impl()),
func(expr.functor()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorCwiseSelectOp /// const TensorCwiseSelectOp. This is an specialisation without OP so it has to be separated.
template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor< TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > {
TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { FunctorExtractor<TensorEvaluator<IfExpr, Dev> > ifExpr;
FunctorExtractor<TensorEvaluator<IfExpr, Dev>> ifExpr; FunctorExtractor<TensorEvaluator<ThenExpr, Dev> > thenExpr;
FunctorExtractor<TensorEvaluator<ThenExpr, Dev>> thenExpr; FunctorExtractor<TensorEvaluator<ElseExpr, Dev> > elseExpr;
FunctorExtractor<TensorEvaluator<ElseExpr, Dev>> elseExpr; FunctorExtractor(const TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>& expr)
FunctorExtractor(const TensorEvaluator< : ifExpr(expr.cond_impl()), thenExpr(expr.then_impl()), elseExpr(expr.else_impl()) {}
const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>& expr)
: ifExpr(expr.cond_impl()),
thenExpr(expr.then_impl()),
elseExpr(expr.else_impl()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorCwiseSelectOp /// TensorCwiseSelectOp. This is an specialisation without OP so it has to be separated
template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev> template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor<TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> >
TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> { :FunctorExtractor< TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev> > {};
FunctorExtractor<IfExpr> ifExpr;
FunctorExtractor<ThenExpr> thenExpr;
FunctorExtractor<ElseExpr> elseExpr;
FunctorExtractor(
const TensorEvaluator<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>&
expr)
: ifExpr(expr.cond_impl()),
thenExpr(expr.then_impl()),
elseExpr(expr.else_impl()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorAssignOp /// const TensorAssignOp. This is an specialisation without OP so it has to be separated.
template <typename LHSExpr, typename RHSExpr, typename Dev> template <typename LHSExpr, typename RHSExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> > {
TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { FunctorExtractor<TensorEvaluator<LHSExpr, Dev> > lhsExpr;
FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr; FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr;
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; FunctorExtractor(const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr)
FunctorExtractor( : lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {}
const TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr)
: lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorAssignOp /// TensorAssignOp. This is an specialisation without OP so it has to be separated.
template <typename LHSExpr, typename RHSExpr, typename Dev> template <typename LHSExpr, typename RHSExpr, typename Dev>
struct FunctorExtractor< struct FunctorExtractor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> >
TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>> { :FunctorExtractor<TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev> >{};
FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr;
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
FunctorExtractor( /// specialisation of the \ref FunctorExtractor struct when the node type is
const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr) /// const TensorEvalToOp, This is an specialisation without OP so it has to be separated.
: lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {} template <typename RHSExpr, typename Dev>
struct FunctorExtractor<TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev> > {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev> > rhsExpr;
FunctorExtractor(const TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()) {}
}; };
/// specialisation of the \ref FunctorExtractor struct when the node type is /// specialisation of the \ref FunctorExtractor struct when the node type is
/// TensorEvalToOp /// TensorEvalToOp. This is a specialisation without OP so it has to be separated.
template <typename RHSExpr, typename Dev> template <typename RHSExpr, typename Dev>
struct FunctorExtractor<TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev>> { struct FunctorExtractor<TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev> >
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr; : FunctorExtractor<TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev> > {};
FunctorExtractor(const TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()) {}
};
/// specialisation of the \ref FunctorExtractor struct when the node type is
/// const TensorEvalToOp
template <typename RHSExpr, typename Dev>
struct FunctorExtractor<TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev>> {
FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
FunctorExtractor(
const TensorEvaluator<const TensorEvalToOp<RHSExpr>, Dev>& expr)
: rhsExpr(expr.impl()) {}
};
/// template deduction function for FunctorExtractor /// template deduction function for FunctorExtractor
template <typename Evaluator> template <typename Evaluator>
auto extractFunctors(const Evaluator& evaluator) auto inline extractFunctors(const Evaluator& evaluator)-> FunctorExtractor<Evaluator> {
-> FunctorExtractor<Evaluator> {
return FunctorExtractor<Evaluator>(evaluator); return FunctorExtractor<Evaluator>(evaluator);
} }
} // namespace internal } // namespace internal
} // namespace TensorSycl } // namespace TensorSycl
} // namespace Eigen } // namespace Eigen
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_EXTRACT_FUNCTORS_HPP

165
unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h
View File

@ -19,8 +19,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_LEAF_COUNT_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_LEAF_COUNT_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
@ -31,158 +31,81 @@ namespace internal {
template <typename Expr> template <typename Expr>
struct LeafCount; struct LeafCount;
template<typename... Args> struct CategoryCount;
template<> struct CategoryCount<>
{
static const size_t Count =0;
};
template<typename Arg, typename... Args>
struct CategoryCount<Arg,Args...>{
static const size_t Count = LeafCount<Arg>::Count + CategoryCount<Args...>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is const /// specialisation of the \ref LeafCount struct when the node type is const
/// TensorMap /// TensorMap
template <typename PlainObjectType, int Options_, template <typename PlainObjectType, int Options_, template <class> class MakePointer_>
template <class> class MakePointer_> struct LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_> > {
struct LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_>> { static const size_t Count =1;
static const size_t Count = 1;
}; };
/// specialisation of the \ref LeafCount struct when the node type is TensorMap /// specialisation of the \ref LeafCount struct when the node type is TensorMap
template <typename PlainObjectType, int Options_, template <typename PlainObjectType, int Options_, template <class> class MakePointer_>
template <class> class MakePointer_> struct LeafCount<TensorMap<PlainObjectType, Options_, MakePointer_> > :LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_> >{};
struct LeafCount<TensorMap<PlainObjectType, Options_, MakePointer_>> {
static const size_t Count = 1;
};
/// specialisation of the \ref LeafCount struct when the node type is const // const TensorCwiseUnaryOp, const TensorCwiseNullaryOp, const TensorCwiseBinaryOp, const TensorCwiseTernaryOp, and Const TensorBroadcastingOp
/// TensorCwiseNullaryOp template <template <class, class...> class CategoryExpr, typename OP, typename... RHSExpr>
template <typename OP, typename RHSExpr> struct LeafCount<const CategoryExpr<OP, RHSExpr...> >: CategoryCount<RHSExpr...> {};
struct LeafCount<const TensorCwiseNullaryOp<OP, RHSExpr>> { // TensorCwiseUnaryOp, TensorCwiseNullaryOp, TensorCwiseBinaryOp, TensorCwiseTernaryOp, and TensorBroadcastingOp
static const size_t Count = LeafCount<RHSExpr>::Count; template <template <class, class...> class CategoryExpr, typename OP, typename... RHSExpr>
}; struct LeafCount<CategoryExpr<OP, RHSExpr...> > :LeafCount<const CategoryExpr<OP, RHSExpr...> >{};
/// specialisation of the \ref LeafCount struct when the node type is /// specialisation of the \ref LeafCount struct when the node type is
/// TensorCwiseNullaryOp /// const TensorSelectOp is an exception
template <typename OP, typename RHSExpr>
struct LeafCount<TensorCwiseNullaryOp<OP, RHSExpr>> {
static const size_t Count = LeafCount<RHSExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is const
/// TensorBroadcastingOp
template <typename OP, typename RHSExpr>
struct LeafCount<const TensorBroadcastingOp<OP, RHSExpr>> {
static const size_t Count = LeafCount<RHSExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is
/// TensorCwiseNullaryOp
template <typename OP, typename RHSExpr>
struct LeafCount<TensorBroadcastingOp<OP, RHSExpr>> {
static const size_t Count = LeafCount<RHSExpr>::Count;
};
// TensorCwiseUnaryOp
template <typename OP, typename RHSExpr>
struct LeafCount<const TensorCwiseUnaryOp<OP, RHSExpr>> {
static const size_t Count = LeafCount<RHSExpr>::Count;
};
// TensorCwiseUnaryOp
template <typename OP, typename RHSExpr>
struct LeafCount<TensorCwiseUnaryOp<OP, RHSExpr>> {
static const size_t Count = LeafCount<RHSExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is const
/// TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr>
struct LeafCount<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> {
static const size_t Count =
LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is
/// TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr>
struct LeafCount<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>> {
static const size_t Count =
LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is
/// TensorCwiseTernaryOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr>
struct LeafCount<TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>> {
static const size_t Count = LeafCount<Arg1Expr>::Count +
LeafCount<Arg2Expr>::Count +
LeafCount<Arg3Expr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is const
/// TensorCwiseTernaryOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr>
struct LeafCount<const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>> {
static const size_t Count = LeafCount<Arg1Expr>::Count +
LeafCount<Arg2Expr>::Count +
LeafCount<Arg3Expr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is
/// TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr> template <typename IfExpr, typename ThenExpr, typename ElseExpr>
struct LeafCount<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> { struct LeafCount<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr> > : CategoryCount<IfExpr, ThenExpr, ElseExpr> {};
static const size_t Count = LeafCount<IfExpr>::Count +
LeafCount<ThenExpr>::Count +
LeafCount<ElseExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is const
/// TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr>
struct LeafCount<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>> {
static const size_t Count = LeafCount<IfExpr>::Count +
LeafCount<ThenExpr>::Count +
LeafCount<ElseExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is /// specialisation of the \ref LeafCount struct when the node type is
/// TensorAssignOp /// TensorSelectOp
template <typename LHSExpr, typename RHSExpr> template <typename IfExpr, typename ThenExpr, typename ElseExpr>
struct LeafCount<TensorAssignOp<LHSExpr, RHSExpr>> { struct LeafCount<TensorSelectOp<IfExpr, ThenExpr, ElseExpr> >: LeafCount<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr> > {};
static const size_t Count =
LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count;
};
/// specialisation of the \ref LeafCount struct when the node type is const /// specialisation of the \ref LeafCount struct when the node type is const
/// TensorAssignOp /// TensorAssignOp
template <typename LHSExpr, typename RHSExpr> template <typename LHSExpr, typename RHSExpr>
struct LeafCount<const TensorAssignOp<LHSExpr, RHSExpr>> { struct LeafCount<const TensorAssignOp<LHSExpr, RHSExpr> >: CategoryCount<LHSExpr,RHSExpr> {};
static const size_t Count =
LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count; /// specialisation of the \ref LeafCount struct when the node type is
}; /// TensorAssignOp is an exception. It is not the same as Unary
template <typename LHSExpr, typename RHSExpr>
struct LeafCount<TensorAssignOp<LHSExpr, RHSExpr> > :LeafCount<const TensorAssignOp<LHSExpr, RHSExpr> >{};
/// specialisation of the \ref LeafCount struct when the node type is const /// specialisation of the \ref LeafCount struct when the node type is const
/// TensorForcedEvalOp /// TensorForcedEvalOp
template <typename Expr> template <typename Expr>
struct LeafCount<const TensorForcedEvalOp<Expr>> { struct LeafCount<const TensorForcedEvalOp<Expr> > {
static const size_t Count = 1; static const size_t Count =1;
}; };
/// specialisation of the \ref LeafCount struct when the node type is /// specialisation of the \ref LeafCount struct when the node type is
/// TensorForcedEvalOp /// TensorForcedEvalOp
template <typename Expr> template <typename Expr>
struct LeafCount<TensorForcedEvalOp<Expr>> { struct LeafCount<TensorForcedEvalOp<Expr> >: LeafCount<const TensorForcedEvalOp<Expr> > {};
static const size_t Count = 1;
};
/// specialisation of the \ref LeafCount struct when the node type is const /// specialisation of the \ref LeafCount struct when the node type is const
/// TensorEvalToOp /// TensorEvalToOp
template <typename Expr> template <typename Expr>
struct LeafCount<const TensorEvalToOp<Expr>> { struct LeafCount<const TensorEvalToOp<Expr> > {
static const size_t Count = 1 + LeafCount<Expr>::Count; static const size_t Count = 1 + CategoryCount<Expr>::Count;
}; };
/// specialisation of the \ref LeafCount struct when the node type is /// specialisation of the \ref LeafCount struct when the node type is
/// TensorEvalToOp /// TensorEvalToOp
template <typename Expr> template <typename Expr>
struct LeafCount<TensorEvalToOp<Expr>> { struct LeafCount<TensorEvalToOp<Expr> >: LeafCount<const TensorEvalToOp<Expr> >{};
static const size_t Count = 1 + LeafCount<Expr>::Count;
};
} }
} }
} // namespace Eigen } // namespace Eigen
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_LEAF_COUNT_HPP

148
unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h
View File

@ -19,8 +19,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_HPP
namespace Eigen { namespace Eigen {
namespace internal { namespace internal {
@ -31,91 +31,52 @@ namespace internal {
template <typename Scalar, size_t N> template <typename Scalar, size_t N>
struct PlaceHolder { struct PlaceHolder {
static constexpr size_t I = N; static constexpr size_t I = N;
using Type = Scalar; typedef Scalar Type;
};
template <typename PlainObjectType, int Options_,
template <class> class MakePointer_, size_t N>
struct PlaceHolder<const TensorMap<PlainObjectType, Options_, MakePointer_>,
N> {
static constexpr size_t I = N;
using Type = const TensorMap<PlainObjectType, Options_, MakePointer_>;
typedef typename Type::Self Self;
typedef typename Type::Base Base;
typedef typename Type::Nested Nested;
typedef typename Type::StorageKind StorageKind;
typedef typename Type::Index Index;
typedef typename Type::Scalar Scalar;
typedef typename Type::RealScalar RealScalar;
typedef typename Type::CoeffReturnType CoeffReturnType;
};
/// \brief specialisation of the PlaceHolder node for TensorForcedEvalOp. The
/// TensorForcedEvalOp act as a leaf node for its parent node.
template <typename Expression, size_t N>
struct PlaceHolder<const TensorForcedEvalOp<Expression>, N> {
static constexpr size_t I = N;
using Type = const TensorForcedEvalOp<Expression>;
typedef typename Type::Nested Nested;
typedef typename Type::StorageKind StorageKind;
typedef typename Type::Index Index;
typedef typename Type::Scalar Scalar;
typedef typename Type::Packet Packet;
typedef typename Type::RealScalar RealScalar;
typedef typename Type::CoeffReturnType CoeffReturnType;
typedef typename Type::PacketReturnType PacketReturnType;
};
template <typename Expression, size_t N>
struct PlaceHolder<TensorForcedEvalOp<Expression>, N> {
static constexpr size_t I = N;
using Type = TensorForcedEvalOp<Expression>;
typedef typename Type::Nested Nested;
typedef typename Type::StorageKind StorageKind;
typedef typename Type::Index Index;
typedef typename Type::Scalar Scalar;
typedef typename Type::Packet Packet;
typedef typename Type::RealScalar RealScalar;
typedef typename Type::CoeffReturnType CoeffReturnType;
typedef typename Type::PacketReturnType PacketReturnType;
}; };
/// \brief specialisation of the PlaceHolder node for const TensorMap /// \brief specialisation of the PlaceHolder node for const TensorMap
template <typename PlainObjectType, int Options_, #define TENSORMAPPLACEHOLDER(CVQual)\
template <class> class Makepointer_, size_t N> template <typename PlainObjectType, int Options_, template <class> class MakePointer_, size_t N>\
struct PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N> { struct PlaceHolder<CVQual TensorMap<PlainObjectType, Options_, MakePointer_>, N> {\
static constexpr size_t I = N; static const size_t I = N;\
typedef CVQual TensorMap<PlainObjectType, Options_, MakePointer_> Type;\
using Type = TensorMap<PlainObjectType, Options_, Makepointer_>; typedef typename Type::Self Self;\
typedef typename Type::Base Base;\
typedef typename Type::Self Self; typedef typename Type::Nested Nested;\
typedef typename Type::Base Base; typedef typename Type::StorageKind StorageKind;\
typedef typename Type::Nested Nested; typedef typename Type::Index Index;\
typedef typename Type::StorageKind StorageKind; typedef typename Type::Scalar Scalar;\
typedef typename Type::Index Index; typedef typename Type::RealScalar RealScalar;\
typedef typename Type::Scalar Scalar; typedef typename Type::CoeffReturnType CoeffReturnType;\
typedef typename Type::Packet Packet;
typedef typename Type::RealScalar RealScalar;
typedef typename Type::CoeffReturnType CoeffReturnType;
typedef typename Base::PacketReturnType PacketReturnType;
}; };
/// specialisation of the traits struct for PlaceHolder TENSORMAPPLACEHOLDER(const)
template <typename PlainObjectType, int Options_, TENSORMAPPLACEHOLDER()
template <class> class Makepointer_, size_t N> #undef TENSORMAPPLACEHOLDER
struct traits<
PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N>> /// \brief specialisation of the PlaceHolder node for TensorForcedEvalOp. The
: public traits<PlainObjectType> { /// TensorForcedEvalOp acts as a leaf node for its parent node.
#define TENSORFORCEDEVALPLACEHOLDER(CVQual)\
template <typename Expression, size_t N>\
struct PlaceHolder<CVQual TensorForcedEvalOp<Expression>, N> {\
static const size_t I = N;\
typedef CVQual TensorForcedEvalOp<Expression> Type;\
typedef typename Type::Nested Nested;\
typedef typename Type::StorageKind StorageKind;\
typedef typename Type::Index Index;\
typedef typename Type::Scalar Scalar;\
typedef typename Type::Packet Packet;\
typedef typename Type::RealScalar RealScalar;\
typedef typename Type::CoeffReturnType CoeffReturnType;\
typedef typename Type::PacketReturnType PacketReturnType;\
};
TENSORFORCEDEVALPLACEHOLDER(const)
TENSORFORCEDEVALPLACEHOLDER()
#undef TENSORFORCEDEVALPLACEHOLDER
template <typename PlainObjectType, int Options_, template <class> class Makepointer_, size_t N>
struct traits<PlaceHolder<const TensorMap<PlainObjectType, Options_, Makepointer_>, N> >: public traits<PlainObjectType> {
typedef traits<PlainObjectType> BaseTraits; typedef traits<PlainObjectType> BaseTraits;
typedef typename BaseTraits::Scalar Scalar; typedef typename BaseTraits::Scalar Scalar;
typedef typename BaseTraits::StorageKind StorageKind; typedef typename BaseTraits::StorageKind StorageKind;
@ -128,24 +89,11 @@ struct traits<
}; };
}; };
template <typename PlainObjectType, int Options_, template <typename PlainObjectType, int Options_, template <class> class Makepointer_, size_t N>
template <class> class Makepointer_, size_t N> struct traits<PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N> >
struct traits< : traits<PlaceHolder<const TensorMap<PlainObjectType, Options_, Makepointer_>, N> > {};
PlaceHolder<const TensorMap<PlainObjectType, Options_, Makepointer_>, N>>
: public traits<PlainObjectType> {
typedef traits<PlainObjectType> BaseTraits;
typedef typename BaseTraits::Scalar Scalar;
typedef typename BaseTraits::StorageKind StorageKind;
typedef typename BaseTraits::Index Index;
static const int NumDimensions = BaseTraits::NumDimensions;
static const int Layout = BaseTraits::Layout;
enum {
Options = Options_,
Flags = BaseTraits::Flags,
};
};
} // end namespoace internal } // end namespace internal
} // end namespoace Eigen } // end namespoace Eigen
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_HPP

339
unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h
View File

@ -19,8 +19,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_EXPR_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_EXPR_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
@ -32,262 +32,127 @@ namespace internal {
template <typename Expr, size_t N> template <typename Expr, size_t N>
struct PlaceHolderExpression; struct PlaceHolderExpression;
/// specialisation of the \ref PlaceHolderExpression when the node is TensorMap template<size_t N, typename... Args>
template <typename Scalar_, int Options_, int Options2_, int NumIndices_, struct CalculateIndex;
typename IndexType_, template <class> class MakePointer_, size_t N>
struct PlaceHolderExpression< template<size_t N, typename Arg>
Eigen::TensorMap<Eigen::Tensor<Scalar_, NumIndices_, Options_, IndexType_>, struct CalculateIndex<N, Arg>{
Options2_, MakePointer_>, typedef typename PlaceHolderExpression<Arg, N>::Type ArgType;
N> { typedef utility::tuple::Tuple<ArgType> ArgsTuple;
using Type = Eigen::internal::PlaceHolder<
Eigen::TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakePointer_>,
N>;
}; };
/// specialisation of the \ref PlaceHolderExpression when the node is const template<size_t N, typename Arg1, typename Arg2>
struct CalculateIndex<N, Arg1, Arg2>{
static const size_t Arg2LeafCount = LeafCount<Arg2>::Count;
typedef typename PlaceHolderExpression<Arg1, N - Arg2LeafCount>::Type Arg1Type;
typedef typename PlaceHolderExpression<Arg2, N>::Type Arg2Type;
typedef utility::tuple::Tuple<Arg1Type, Arg2Type> ArgsTuple;
};
template<size_t N, typename Arg1, typename Arg2, typename Arg3>
struct CalculateIndex<N, Arg1, Arg2, Arg3> {
static const size_t Arg3LeafCount = LeafCount<Arg3>::Count;
static const size_t Arg2LeafCount = LeafCount<Arg2>::Count;
typedef typename PlaceHolderExpression<Arg1, N - Arg3LeafCount - Arg2LeafCount>::Type Arg1Type;
typedef typename PlaceHolderExpression<Arg2, N - Arg3LeafCount>::Type Arg2Type;
typedef typename PlaceHolderExpression<Arg3, N>::Type Arg3Type;
typedef utility::tuple::Tuple<Arg1Type, Arg2Type, Arg3Type> ArgsTuple;
};
template<template<class...> class Category , class OP, class TPL>
struct CategoryHelper;
template<template<class...> class Category , class OP, class ...T >
struct CategoryHelper<Category, OP, utility::tuple::Tuple<T...> > {
typedef Category<OP, T... > Type;
};
template<template<class...> class Category , class ...T >
struct CategoryHelper<Category, NoOP, utility::tuple::Tuple<T...> > {
typedef Category<T... > Type;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorCwiseNullaryOp, TensorCwiseUnaryOp, TensorBroadcastingOp, TensorCwiseBinaryOp, TensorCwiseTernaryOp
#define OPEXPRCATEGORY(CVQual)\
template <template <class, class... > class Category, typename OP, typename... SubExpr, size_t N>\
struct PlaceHolderExpression<CVQual Category<OP, SubExpr...>, N>{\
typedef CVQual typename CategoryHelper<Category, OP, typename CalculateIndex<N, SubExpr...>::ArgsTuple>::Type Type;\
};
OPEXPRCATEGORY(const)
OPEXPRCATEGORY()
#undef OPEXPRCATEGORY
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorCwiseSelectOp
#define SELECTEXPR(CVQual)\
template <typename IfExpr, typename ThenExpr, typename ElseExpr, size_t N>\
struct PlaceHolderExpression<CVQual TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, N> {\
typedef CVQual typename CategoryHelper<TensorSelectOp, NoOP, typename CalculateIndex<N, IfExpr, ThenExpr, ElseExpr>::ArgsTuple>::Type Type;\
};
SELECTEXPR(const)
SELECTEXPR()
#undef SELECTEXPR
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorAssignOp
#define ASSIGNEXPR(CVQual)\
template <typename LHSExpr, typename RHSExpr, size_t N>\
struct PlaceHolderExpression<CVQual TensorAssignOp<LHSExpr, RHSExpr>, N> {\
typedef CVQual typename CategoryHelper<TensorAssignOp, NoOP, typename CalculateIndex<N, LHSExpr, RHSExpr>::ArgsTuple>::Type Type;\
};
ASSIGNEXPR(const)
ASSIGNEXPR()
#undef ASSIGNEXPR
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorMap /// TensorMap
template <typename Scalar_, int Options_, int Options2_, int NumIndices_, #define TENSORMAPEXPR(CVQual)\
typename IndexType_, template <class> class MakePointer_, size_t N> template <typename Scalar_, int Options_, int Options2_, int NumIndices_, typename IndexType_, template <class> class MakePointer_, size_t N>\
struct PlaceHolderExpression< struct PlaceHolderExpression< CVQual TensorMap< Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_>, N> {\
const Eigen::TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, typedef CVQual Eigen::internal::PlaceHolder<CVQual TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_, MakePointer_>, N> Type;\
Options2_, MakePointer_>,
N> {
using Type = const Eigen::internal::PlaceHolder<
const TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
Options2_, MakePointer_>,
N>;
}; };
/// specialisation of the \ref PlaceHolderExpression when the node is TENSORMAPEXPR(const)
/// TensorCwiseNullaryOp TENSORMAPEXPR()
template <typename OP, typename RHSExpr, size_t N> #undef TENSORMAPEXPR
struct PlaceHolderExpression<TensorCwiseNullaryOp<OP, RHSExpr>, N> {
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = TensorCwiseNullaryOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorCwiseNullaryOp
template <typename OP, typename RHSExpr, size_t N>
struct PlaceHolderExpression<const TensorCwiseNullaryOp<OP, RHSExpr>, N> {
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = const TensorCwiseNullaryOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorBroadcastingOp
template <typename OP, typename RHSExpr, size_t N>
struct PlaceHolderExpression<TensorBroadcastingOp<OP, RHSExpr>, N> {
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = TensorBroadcastingOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorBroadcastingOp
template <typename OP, typename RHSExpr, size_t N>
struct PlaceHolderExpression<const TensorBroadcastingOp<OP, RHSExpr>, N> {
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = const TensorBroadcastingOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorCwiseUnaryOp
template <typename OP, typename RHSExpr, size_t N>
struct PlaceHolderExpression<TensorCwiseUnaryOp<OP, RHSExpr>, N> {
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = TensorCwiseUnaryOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorCwiseUnaryOp
template <typename OP, typename RHSExpr, size_t N>
struct PlaceHolderExpression<const TensorCwiseUnaryOp<OP, RHSExpr>, N> {
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = const TensorCwiseUnaryOp<OP, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr, size_t N>
struct PlaceHolderExpression<TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, N> {
static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count;
using LHSPlaceHolderType =
typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type;
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorCwiseBinaryOp
template <typename OP, typename LHSExpr, typename RHSExpr, size_t N>
struct PlaceHolderExpression<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>,
N> {
static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count;
using LHSPlaceHolderType =
typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type;
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type =
const TensorCwiseBinaryOp<OP, LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorCwiseSelectOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,
size_t N>
struct PlaceHolderExpression<
const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, N> {
static const size_t Arg3LeafCount = LeafCount<Arg3Expr>::Count;
static const size_t Arg2LeafCount = LeafCount<Arg2Expr>::Count;
using Arg1PlaceHolderType =
typename PlaceHolderExpression<Arg1Expr,
N - Arg3LeafCount - Arg2LeafCount>::Type;
using Arg2PlaceHolderType =
typename PlaceHolderExpression<Arg2Expr, N - Arg3LeafCount>::Type;
using Arg3PlaceHolderType = typename PlaceHolderExpression<Arg3Expr, N>::Type;
using Type =
const TensorCwiseTernaryOp<OP, Arg1PlaceHolderType, Arg2PlaceHolderType,
Arg3PlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorCwiseSelectOp
template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,
size_t N>
struct PlaceHolderExpression<
TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, N> {
static const size_t Arg3LeafCount = LeafCount<Arg3Expr>::Count;
static const size_t Arg2LeafCount = LeafCount<Arg2Expr>::Count;
using Arg1PlaceHolderType =
typename PlaceHolderExpression<Arg1Expr,
N - Arg3LeafCount - Arg2LeafCount>::Type;
using Arg2PlaceHolderType =
typename PlaceHolderExpression<Arg2Expr, N - Arg3LeafCount>::Type;
using Arg3PlaceHolderType = typename PlaceHolderExpression<Arg3Expr, N>::Type;
using Type = TensorCwiseTernaryOp<OP, Arg1PlaceHolderType,
Arg2PlaceHolderType, Arg3PlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr, size_t N>
struct PlaceHolderExpression<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>,
N> {
static const size_t ElseLeafCount = LeafCount<ElseExpr>::Count;
static const size_t ThenLeafCount = LeafCount<ThenExpr>::Count;
using IfPlaceHolderType =
typename PlaceHolderExpression<IfExpr,
N - ElseLeafCount - ThenLeafCount>::Type;
using ThenPlaceHolderType =
typename PlaceHolderExpression<ThenExpr, N - ElseLeafCount>::Type;
using ElsePlaceHolderType = typename PlaceHolderExpression<ElseExpr, N>::Type;
using Type = const TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType,
ElsePlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorCwiseSelectOp
template <typename IfExpr, typename ThenExpr, typename ElseExpr, size_t N>
struct PlaceHolderExpression<TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, N> {
static const size_t ElseLeafCount = LeafCount<ElseExpr>::Count;
static const size_t ThenLeafCount = LeafCount<ThenExpr>::Count;
using IfPlaceHolderType =
typename PlaceHolderExpression<IfExpr,
N - ElseLeafCount - ThenLeafCount>::Type;
using ThenPlaceHolderType =
typename PlaceHolderExpression<ThenExpr, N - ElseLeafCount>::Type;
using ElsePlaceHolderType = typename PlaceHolderExpression<ElseExpr, N>::Type;
using Type = TensorSelectOp<IfPlaceHolderType, ThenPlaceHolderType,
ElsePlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorAssignOp
template <typename LHSExpr, typename RHSExpr, size_t N>
struct PlaceHolderExpression<TensorAssignOp<LHSExpr, RHSExpr>, N> {
static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count;
using LHSPlaceHolderType =
typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type;
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorAssignOp
template <typename LHSExpr, typename RHSExpr, size_t N>
struct PlaceHolderExpression<const TensorAssignOp<LHSExpr, RHSExpr>, N> {
static const size_t RHSLeafCount = LeafCount<RHSExpr>::Count;
using LHSPlaceHolderType =
typename PlaceHolderExpression<LHSExpr, N - RHSLeafCount>::Type;
using RHSPlaceHolderType = typename PlaceHolderExpression<RHSExpr, N>::Type;
using Type = const TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is const
/// TensorForcedEvalOp
template <typename Expr, size_t N>
struct PlaceHolderExpression<const TensorForcedEvalOp<Expr>, N> {
using Type =
const Eigen::internal::PlaceHolder<const TensorForcedEvalOp<Expr>, N>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is /// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorForcedEvalOp /// TensorForcedEvalOp
template <typename Expr, size_t N> #define FORCEDEVAL(CVQual)\
struct PlaceHolderExpression<TensorForcedEvalOp<Expr>, N> { template <typename Expr, size_t N>\
using Type = Eigen::internal::PlaceHolder<TensorForcedEvalOp<Expr>, N>; struct PlaceHolderExpression<CVQual TensorForcedEvalOp<Expr>, N> {\
typedef CVQual Eigen::internal::PlaceHolder<CVQual TensorForcedEvalOp<Expr>, N> Type;\
}; };
/// specialisation of the \ref PlaceHolderExpression when the node is const FORCEDEVAL(const)
/// TensorEvalToOp FORCEDEVAL()
template <typename Expr, size_t N> #undef FORCEDEVAL
struct PlaceHolderExpression<const TensorEvalToOp<Expr>, N> {
static const size_t RHSLeafCount = LeafCount<Expr>::Count;
using RHSPlaceHolderType = typename PlaceHolderExpression<Expr, N>::Type;
using Type = const TensorEvalToOp<RHSPlaceHolderType>;
};
/// specialisation of the \ref PlaceHolderExpression when the node is /// specialisation of the \ref PlaceHolderExpression when the node is
/// TensorEvalToOp /// TensorEvalToOp
template <typename Expr, size_t N> #define EVALTO(CVQual)\
struct PlaceHolderExpression<TensorEvalToOp<Expr>, N> { template <typename Expr, size_t N>\
static const size_t RHSLeafCount = LeafCount<Expr>::Count; struct PlaceHolderExpression<CVQual TensorEvalToOp<Expr>, N> {\
typedef CVQual TensorEvalToOp<typename CalculateIndex <N, Expr>::ArgType> Type;\
using RHSPlaceHolderType = typename PlaceHolderExpression<Expr, N>::Type;
using Type = TensorEvalToOp<RHSPlaceHolderType>;
}; };
EVALTO(const)
EVALTO()
#undef EVALTO
/// template deduction for \ref PlaceHolderExpression struct /// template deduction for \ref PlaceHolderExpression struct
template <typename Expr> template <typename Expr>
struct createPlaceHolderExpression { struct createPlaceHolderExpression {
static const size_t TotalLeaves = LeafCount<Expr>::Count; static const size_t TotalLeaves = LeafCount<Expr>::Count;
using Type = typename PlaceHolderExpression<Expr, TotalLeaves - 1>::Type; typedef typename PlaceHolderExpression<Expr, TotalLeaves - 1>::Type Type;
}; };
} }
} }
} // namespace Eigen } // namespace Eigen
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_PLACEHOLDER_EXPR_HPP

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

@ -20,8 +20,8 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP
namespace Eigen { namespace Eigen {
namespace TensorSycl { namespace TensorSycl {
@ -34,17 +34,14 @@ 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) {
using PlaceHolderExpr = typedef typename internal::createPlaceHolderExpression<Expr>::Type PlaceHolderExpr;
typename internal::createPlaceHolderExpression<Expr>::Type;
auto functors = internal::extractFunctors(evaluator); auto functors = internal::extractFunctors(evaluator);
dev.m_queue.submit([&](cl::sycl::handler &cgh) { dev.m_queue.submit([&](cl::sycl::handler &cgh) {
// create a tuple of accessors from Evaluator // create a tuple of accessors from Evaluator
auto tuple_of_accessors = auto tuple_of_accessors = internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator);
internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator); const auto range = utility::tuple::get<0>(tuple_of_accessors).get_range()[0];
const auto range =
utility::tuple::get<0>(tuple_of_accessors).get_range()[0];
size_t outTileSize = range; size_t outTileSize = range;
if (range > 64) outTileSize = 64; if (range > 64) outTileSize = 64;
@ -53,26 +50,14 @@ void run(Expr &expr, Dev &dev) {
if (yMode != 0) yRange += (outTileSize - yMode); if (yMode != 0) yRange += (outTileSize - yMode);
// run the kernel // run the kernel
cgh.parallel_for<PlaceHolderExpr>( cgh.parallel_for<PlaceHolderExpr>( cl::sycl::nd_range<1>(cl::sycl::range<1>(yRange), cl::sycl::range<1>(outTileSize)), [=](cl::sycl::nd_item<1> itemID) {
cl::sycl::nd_range<1>(cl::sycl::range<1>(yRange), typedef typename internal::ConvertToDeviceExpression<Expr>::Type DevExpr;
cl::sycl::range<1>(outTileSize)), auto device_expr =internal::createDeviceExpression<DevExpr, PlaceHolderExpr>(functors, tuple_of_accessors);
[=](cl::sycl::nd_item<1> itemID) { auto device_evaluator = Eigen::TensorEvaluator<decltype(device_expr.expr), Eigen::DefaultDevice>(device_expr.expr, Eigen::DefaultDevice());
using DevExpr = if (itemID.get_global_linear_id() < range) {
typename internal::ConvertToDeviceExpression<Expr>::Type; device_evaluator.evalScalar(static_cast<int>(itemID.get_global_linear_id()));
}
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());
if (itemID.get_global_linear_id() < range) {
device_evaluator.evalScalar(
static_cast<int>(itemID.get_global_linear_id()));
}
});
}); });
dev.m_queue.throw_asynchronous(); dev.m_queue.throw_asynchronous();
} }
@ -81,4 +66,4 @@ void run(Expr &expr, Dev &dev) {
} // namespace TensorSycl } // namespace TensorSycl
} // namespace Eigen } // namespace Eigen
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_SYCLRUN_HPP

168
unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h
View File

@ -18,24 +18,23 @@
* *
*****************************************************************/ *****************************************************************/
#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_TUPLE_HPP
#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_TUPLE_HPP
namespace utility { namespace utility {
namespace tuple { namespace tuple {
/// \struct EnableIf /// \struct StaticIf
/// \brief The EnableIf struct is used to statically define type based on the /// \brief The StaticIf struct is used to statically choose the type based on the
/// condition. /// condition.
template <bool, typename T = void> template <bool, typename T = void> struct StaticIf;
struct EnableIf {}; /// \brief specialisation of the \ref StaticIf when the condition is true
/// \brief specialisation of the \ref EnableIf when the condition is true
template <typename T> template <typename T>
struct EnableIf<true, T> { struct StaticIf<true, T> {
typedef T type; typedef T type;
}; };
/// \struct Tuple /// \struct Tuple
/// \brief is a fixed-size collection of heterogeneous values /// \brief is a fixed-size collection of heterogeneous values
/// \ztparam Ts... - the types of the elements that the tuple stores. /// \ztparam Ts... - the types of the elements that the tuple stores.
/// Empty list is supported. /// Empty list is supported.
template <class... Ts> template <class... Ts>
struct Tuple {}; struct Tuple {};
@ -47,12 +46,11 @@ struct Tuple {};
template <class T, class... Ts> template <class T, class... Ts>
struct Tuple<T, Ts...> { struct Tuple<T, Ts...> {
Tuple(T t, Ts... ts) : head(t), tail(ts...) {} Tuple(T t, Ts... ts) : head(t), tail(ts...) {}
T head; T head;
Tuple<Ts...> tail; Tuple<Ts...> tail;
}; };
/// \struct ElemTypeHolder ///\ struct ElemTypeHolder
/// \brief ElemTypeHolder class is used to specify the types of the /// \brief ElemTypeHolder class is used to specify the types of the
/// elements inside the tuple /// elements inside the tuple
/// \tparam size_t the number of elements inside the tuple /// \tparam size_t the number of elements inside the tuple
@ -60,78 +58,59 @@ struct Tuple<T, Ts...> {
template <size_t, class> template <size_t, class>
struct ElemTypeHolder; struct ElemTypeHolder;
/// \brief specialisation of the \ref ElemTypeHolder class when the number /// \brief specialisation of the \ref ElemTypeHolder class when the number of
/// elements inside the tuple is 1 /// elements inside the tuple is 1
template <class T, class... Ts> template <class T, class... Ts>
struct ElemTypeHolder<0, Tuple<T, Ts...>> { struct ElemTypeHolder<0, Tuple<T, Ts...> > {
typedef T type; typedef T type;
}; };
/// \brief specialisation of the \ref ElemTypeHolder class when the number of /// \brief specialisation of the \ref ElemTypeHolder class when the number of
/// elements inside the tuple is bigger than 1. It recursively call itself to /// elements inside the tuple is bigger than 1. It recursively calls itself to
/// detect the type of each element in the tuple /// detect the type of each element in the tuple
/// \tparam T : the type of the first element in the tuple. /// \tparam T : the type of the first element in the tuple.
/// \tparam Ts... the rest of the elements in the tuple. Ts... can be empty. /// \tparam Ts... the rest of the elements in the tuple. Ts... can be empty.
/// \tparam K is the Kth element in the tuple /// \tparam K is the Kth element in the tuple
template <size_t k, class T, class... Ts> template <size_t k, class T, class... Ts>
struct ElemTypeHolder<k, Tuple<T, Ts...>> { struct ElemTypeHolder<k, Tuple<T, Ts...> > {
typedef typename ElemTypeHolder<k - 1, Tuple<Ts...>>::type type; typedef typename ElemTypeHolder<k - 1, Tuple<Ts...> >::type type;
}; };
/// get /// get
/// \brief Extracts the first element from the tuple. /// \brief Extracts the first element from the tuple.
/// K=0 represents the first element of the tuple. The tuple cannot be empty. /// K=0 represents the first element of the tuple. The tuple cannot be empty.
/// \tparam Ts... are the elements type in the tuple. /// \tparam Ts... are the type of the elements in the tuple.
/// \param t is the tuple whose contents to extract /// \param t is the tuple whose contents to extract
/// \return typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type /// \return typename ElemTypeHolder<0, Tuple<Ts...> >::type &>::type
template <size_t k, class... Ts>
typename EnableIf<k == 0, #define TERMINATE_CONDS_TUPLE_GET(CVQual) \
typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type template <size_t k, class... Ts> \
get(Tuple<Ts...> &t) { typename StaticIf<k == 0, CVQual typename ElemTypeHolder<0, Tuple<Ts...> >::type &>::type \
return t.head; get(CVQual Tuple<Ts...> &t) { \
static_assert(sizeof...(Ts)!=0, "The requseted value is bigger than the size of the tuple"); \
return t.head; \
} }
TERMINATE_CONDS_TUPLE_GET(const)
TERMINATE_CONDS_TUPLE_GET()
#undef TERMINATE_CONDS_TUPLE_GET
/// get /// get
/// \brief Extracts the Kth element from the tuple. /// \brief Extracts the Kth element from the tuple.
/// \tparam K is an integer value in [0,sizeof...(Types)). ///\tparam K is an integer value in [0,sizeof...(Types)).
/// \tparam T is the (sizeof...(Types) -(K+1)) element in the tuple /// \tparam T is the (sizeof...(Types) -(K+1)) element in the tuple
/// \tparam Ts... are the elements type in the tuple. /// \tparam Ts... are the type of the elements in the tuple.
/// \param t is the tuple whose contents to extract /// \param t is the tuple whose contents to extract
/// \return typename ElemTypeHolder<K, Tuple<Ts...>>::type &>::type /// \return typename ElemTypeHolder<K, Tuple<Ts...> >::type &>::type
template <size_t k, class T, class... Ts> #define RECURSIVE_TUPLE_GET(CVQual) \
typename EnableIf<k != 0, template <size_t k, class T, class... Ts> \
typename ElemTypeHolder<k, Tuple<T, Ts...>>::type &>::type typename StaticIf<k != 0, CVQual typename ElemTypeHolder<k, Tuple<T, Ts...> >::type &>::type \
get(Tuple<T, Ts...> &t) { get(CVQual Tuple<T, Ts...> &t) { \
return get<k - 1>(t.tail); return get<k - 1>(t.tail); \
} }
RECURSIVE_TUPLE_GET(const)
RECURSIVE_TUPLE_GET()
#undef RECURSIVE_TUPLE_GET
/// get
/// \brief Extracts the first element from the tuple when the tuple and all the
/// elements inside are const.
/// K=0 represents the first element of the tuple. The tuple cannot be empty.
/// \tparam Ts... are the elements type in the tuple.
/// \param t is the const tuple whose contents to extract
/// \return const typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type
template <size_t k, class... Ts>
typename EnableIf<k == 0,
const typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type
get(const Tuple<Ts...> &t) {
return t.head;
}
/// get
/// \brief Extracts the Kth element from the tuple when the tuple and all the
/// elements inside are const.
/// \tparam K is an integer value in [0,sizeof...(Types)).
/// \tparam T is the (sizeof...(Types) -(K+1)) element in the tuple
/// \tparam Ts... are the elements type in the tuple.
/// \param t is the const tuple whose contents to extract
/// \return const typename ElemTypeHolder<K, Tuple<Ts...>>::type &>::type
template <size_t k, class T, class... Ts>
typename EnableIf<
k != 0, const typename ElemTypeHolder<k, Tuple<T, Ts...>>::type &>::type
get(const Tuple<T, Ts...> &t) {
return get<k - 1>(t.tail);
}
/// make_tuple /// make_tuple
/// \brief Creates a tuple object, deducing the target type from the types of /// \brief Creates a tuple object, deducing the target type from the types of
/// arguments. /// arguments.
@ -153,17 +132,17 @@ static constexpr size_t size(Tuple<Args...> &) {
return sizeof...(Args); return sizeof...(Args);
} }
/// \struct Index_list /// \struct IndexList
/// \brief Creates a list of index from the elements in the tuple /// \brief Creates a list of index from the elements in the tuple
/// \tparam Is... a list of index from [0 to sizeof...(tuple elements)) /// \tparam Is... a list of index from [0 to sizeof...(tuple elements))
template <size_t... Is> template <size_t... Is>
struct Index_list {}; struct IndexList {};
/// \struct RangeBuilder /// \struct RangeBuilder
/// \brief Collects internal details for generating index ranges [MIN, MAX) /// \brief Collects internal details for generating index ranges [MIN, MAX)
/// Declare primary template for index range builder /// Declare primary template for index range builder
/// \tparam MIN is the starting index in the tuple /// \tparam MIN is the starting index in the tuple
/// \tparam N represents sizeof..(elements)- sizeof...(Is) /// \tparam N represents sizeof..(elemens)- sizeof...(Is)
/// \tparam Is... are the list of generated index so far /// \tparam Is... are the list of generated index so far
template <size_t MIN, size_t N, size_t... Is> template <size_t MIN, size_t N, size_t... Is>
struct RangeBuilder; struct RangeBuilder;
@ -174,14 +153,14 @@ struct RangeBuilder;
/// \tparam Is is [0 to sizeof...(tuple elements)) /// \tparam Is is [0 to sizeof...(tuple elements))
template <size_t MIN, size_t... Is> template <size_t MIN, size_t... Is>
struct RangeBuilder<MIN, MIN, Is...> { struct RangeBuilder<MIN, MIN, Is...> {
typedef Index_list<Is...> type; typedef IndexList<Is...> type;
}; };
/// Induction step: Specialisation of the RangeBuilder class when N!=MIN /// Induction step: Specialisation of the RangeBuilder class when N!=MIN
/// in this case we are recursively subtracting the N by one and adding one /// in this case we are recursively subtracting N by one and adding one
/// index to Is... list until MIN==N /// index to Is... list until MIN==N
/// \tparam MIN is the starting index in the tuple /// \tparam MIN is the starting index in the tuple
/// \tparam N represents sizeof..(elements)- sizeof...(Is) /// \tparam N represents sizeof..(elemens)- sizeof...(Is)
/// \tparam Is... are the list of generated index so far /// \tparam Is... are the list of generated index so far
template <size_t MIN, size_t N, size_t... Is> template <size_t MIN, size_t N, size_t... Is>
struct RangeBuilder : public RangeBuilder<MIN, N - 1, N - 1, Is...> {}; struct RangeBuilder : public RangeBuilder<MIN, N - 1, N - 1, Is...> {};
@ -190,75 +169,66 @@ struct RangeBuilder : public RangeBuilder<MIN, N - 1, N - 1, Is...> {};
/// \tparam MIN is the starting index in the tuple /// \tparam MIN is the starting index in the tuple
/// \tparam MAX is the size of the tuple /// \tparam MAX is the size of the tuple
template <size_t MIN, size_t MAX> template <size_t MIN, size_t MAX>
using Index_range = typename RangeBuilder<MIN, MAX>::type; struct IndexRange: RangeBuilder<MIN, MAX>::type {};
/// append_impl /// append_base
/// \brief unpacking the elements of the input tuple t and creating a new tuple /// \brief unpacking the elements of the input tuple t and creating a new tuple
/// by adding element a at the end of it. /// by adding element a at the end of it.
/// \tparam Args... the type of the elements inside the tuple t ///\tparam Args... the type of the elements inside the tuple t
/// \tparam T the type of the new element going to be added at the end of tuple /// \tparam T the type of the new element going to be added at the end of tuple
/// \tparam I... is the list of index from [0 to sizeof...(t)) /// \tparam I... is the list of index from [0 to sizeof...(t))
/// \param t the tuple on which we want to append a. /// \param t the tuple on which we want to append a.
/// \param a the new elements going to be added to the tuple /// \param a the new elements going to be added to the tuple
/// \return Tuple<Args..., T> /// \return Tuple<Args..., T>
template <typename... Args, typename T, size_t... I> template <typename... Args, typename T, size_t... I>
Tuple<Args..., T> append_impl(utility::tuple::Tuple<Args...> t, T a, Tuple<Args..., T> append_base(Tuple<Args...> t, T a,IndexList<I...>) {
utility::tuple::Index_list<I...>) { return make_tuple(get<I>(t)..., a);
return utility::tuple::make_tuple(get<I>(t)..., a);
} }
/// append /// append
/// \brief the deduction function for \ref append_impl that automatically /// \brief the deduction function for \ref append_base that automatically
/// generate the \ref Index_range /// generate the \ref IndexRange
/// \tparam Args... the type of the elements inside the tuple t ///\tparam Args... the type of the elements inside the tuple t
/// \tparam T the type of the new element going to be added at the end of tuple /// \tparam T the type of the new element going to be added at the end of tuple
/// \param t the tuple on which we want to append a. /// \param t the tuple on which we want to append a.
/// \param a the new elements going to be added to the tuple /// \param a the new elements going to be added to the tuple
/// \return Tuple<Args..., T> /// \return Tuple<Args..., T>
template <typename... Args, typename T> template <typename... Args, typename T>
Tuple<Args..., T> append(Tuple<Args...> t, T a) { Tuple<Args..., T> append(Tuple<Args...> t, T a) {
return utility::tuple::append_impl( return append_base(t, a, IndexRange<0, sizeof...(Args)>());
t, a, utility::tuple::Index_range<0, sizeof...(Args)>());
} }
/// append_impl /// append_base
/// \brief This is an specialised of \ref append_impl when we want to /// \brief This is a specialisation of \ref append_base when we want to
/// concatenate /// concatenate
/// tuple t2 at the end of the tuple t1. Here we unpack both tuples, generate /// tuple t2 at the end of the tuple t1. Here we unpack both tuples, generate the
/// the /// IndexRange for each of them and create an output tuple T that contains both
/// Index_range for each of them and create an output tuple T that contains both
/// elements of t1 and t2. /// elements of t1 and t2.
/// \tparam Args1... the type of the elements inside the tuple t1 ///\tparam Args1... the type of the elements inside the tuple t1
/// \tparam Args2... the type of the elements inside the tuple t2 ///\tparam Args2... the type of the elements inside the tuple t2
/// \tparam I1... is the list of index from [0 to sizeof...(t1)) /// \tparam I1... is the list of index from [0 to sizeof...(t1))
/// \tparam I2... is the list of index from [0 to sizeof...(t2)) /// \tparam I2... is the list of index from [0 to sizeof...(t2))
/// \param t1 is the tuple on which we want to append t2. /// \param t1 is the tuple on which we want to append t2.
/// \param t2 is the tuple that is going to be added on t1. /// \param t2 is the tuple that is going to be added on t1.
/// \return Tuple<Args1..., Args2...> /// \return Tuple<Args1..., Args2...>
template <typename... Args1, typename... Args2, size_t... I1, size_t... I2> template <typename... Args1, typename... Args2, size_t... I1, size_t... I2>
Tuple<Args1..., Args2...> append_impl(utility::tuple::Tuple<Args1...> t1, Tuple<Args1..., Args2...> append_base(Tuple<Args1...> t1, Tuple<Args2...> t2, IndexList<I1...>, IndexList<I2...>) {
utility::tuple::Tuple<Args2...> t2, return make_tuple(get<I1>(t1)...,get<I2>(t2)...);
utility::tuple::Index_list<I1...>,
utility::tuple::Index_list<I2...>) {
return utility::tuple::make_tuple(utility::tuple::get<I1>(t1)...,
utility::tuple::get<I2>(t2)...);
} }
/// append /// append
/// \brief deduction function for \ref append_impl when we are appending tuple /// \brief deduction function for \ref append_base when we are appending tuple
/// t1 by tuple t2. In this case the \ref Index_range for both tuple are /// t1 by tuple t2. In this case the \ref IndexRange for both tuple are
/// automatically generated. /// automatically generated.
/// \tparam Args1... the type of the elements inside the tuple t1 ///\tparam Args1... the type of the elements inside the tuple t1
/// \tparam Args2... the type of the elements inside the tuple t2 ///\tparam Args2... the type of the elements inside the tuple t2
/// \param t1 is the tuple on which we want to append t2. /// \param t1 is the tuple on which we want to append t2.
/// \param t2 is the tuple that is going to be added on t1. /// \param t2 is the tuple that is going to be added on t1.
/// \return Tuple<Args1..., Args2...> /// \return Tuple<Args1..., Args2...>
template <typename... Args1, typename... Args2> template <typename... Args1, typename... Args2>
Tuple<Args1..., Args2...> append(utility::tuple::Tuple<Args1...> t1, Tuple<Args1..., Args2...> append(Tuple<Args1...> t1,Tuple<Args2...> t2) {
utility::tuple::Tuple<Args2...> t2) { return append_base(t1, t2, IndexRange<0, sizeof...(Args1)>(), IndexRange<0, sizeof...(Args2)>());
return utility::tuple::append_impl(
t1, t2, utility::tuple::Index_range<0, sizeof...(Args1)>(),
utility::tuple::Index_range<0, sizeof...(Args2)>());
} }
} // tuple } // tuple
} // utility } // utility
#endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSORSYCL_TUPLE_HPP

14
unsupported/test/CMakeLists.txt
View File

@ -140,6 +140,12 @@ endif()
endif() endif()
if(EIGEN_TEST_CXX11) if(EIGEN_TEST_CXX11)
if(EIGEN_TEST_SYCL)
ei_add_test_sycl(cxx11_tensor_sycl "-std=c++11")
ei_add_test_sycl(cxx11_tensor_forced_eval_sycl "-std=c++11")
ei_add_test_sycl(cxx11_tensor_broadcast_sycl "-std=c++11")
ei_add_test_sycl(cxx11_tensor_device_sycl "-std=c++11")
endif(EIGEN_TEST_SYCL)
# It should be safe to always run these tests as there is some fallback code for # It should be safe to always run these tests as there is some fallback code for
# older compiler that don't support cxx11. # older compiler that don't support cxx11.
set(CMAKE_CXX_STANDARD 11) set(CMAKE_CXX_STANDARD 11)
@ -248,11 +254,3 @@ if(CUDA_FOUND AND EIGEN_TEST_CUDA)
unset(EIGEN_ADD_TEST_FILENAME_EXTENSION) unset(EIGEN_ADD_TEST_FILENAME_EXTENSION)
endif() endif()
if(EIGEN_TEST_SYCL)
ei_add_test_sycl(cxx11_tensor_sycl "-std=c++11")
ei_add_test_sycl(cxx11_tensor_forced_eval_sycl "-std=c++11")
ei_add_test_sycl(cxx11_tensor_broadcast_sycl "-std=c++11")
ei_add_test_sycl(cxx11_tensor_device_sycl "-std=c++11")
endif(EIGEN_TEST_SYCL)

4
unsupported/test/cxx11_tensor_broadcast_sycl.cpp
View File

@ -28,7 +28,7 @@ using Eigen::TensorMap;
// Types used in tests: // Types used in tests:
using TestTensor = Tensor<float, 3>; using TestTensor = Tensor<float, 3>;
using TestTensorMap = TensorMap<Tensor<float, 3>>; using TestTensorMap = TensorMap<Tensor<float, 3>>;
static void test_sycl_broadcast(){ static void test_broadcast_sycl(){
cl::sycl::gpu_selector s; cl::sycl::gpu_selector s;
cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { cl::sycl::queue q(s, [=](cl::sycl::exception_list l) {
@ -75,5 +75,5 @@ static void test_sycl_broadcast(){
} }
void test_cxx11_tensor_broadcast_sycl() { void test_cxx11_tensor_broadcast_sycl() {
CALL_SUBTEST(test_sycl_broadcast()); CALL_SUBTEST(test_broadcast_sycl());
} }

6
unsupported/test/cxx11_tensor_device_sycl.cpp
View File

@ -20,7 +20,7 @@
#include "main.h" #include "main.h"
#include <unsupported/Eigen/CXX11/Tensor> #include <unsupported/Eigen/CXX11/Tensor>
void test_sycl_device() { void test_device_sycl() {
cl::sycl::gpu_selector s; cl::sycl::gpu_selector s;
cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { cl::sycl::queue q(s, [=](cl::sycl::exception_list l) {
for (const auto& e : l) { for (const auto& e : l) {
@ -31,9 +31,9 @@ void test_sycl_device() {
} }
} }
}); });
SyclDevice sycl_device(q); Eigen::SyclDevice sycl_device(q);
printf("Helo from ComputeCpp: Device Exists\n"); printf("Helo from ComputeCpp: Device Exists\n");
} }
void test_cxx11_tensor_device_sycl() { void test_cxx11_tensor_device_sycl() {
CALL_SUBTEST(test_sycl_device()); CALL_SUBTEST(test_device_sycl());
} }

6
unsupported/test/cxx11_tensor_forced_eval_sycl.cpp
View File

@ -22,7 +22,7 @@
using Eigen::Tensor; using Eigen::Tensor;
void test_sycl_gpu() { void test_forced_eval_sycl() {
cl::sycl::gpu_selector s; cl::sycl::gpu_selector s;
cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { cl::sycl::queue q(s, [=](cl::sycl::exception_list l) {
for (const auto& e : l) { for (const auto& e : l) {
@ -65,6 +65,4 @@ void test_sycl_gpu() {
printf("(a+b)*b Test Passed\n"); printf("(a+b)*b Test Passed\n");
} }
void test_cxx11_tensor_forced_eval_sycl() { void test_cxx11_tensor_forced_eval_sycl() { CALL_SUBTEST(test_forced_eval_sycl()); }
CALL_SUBTEST(test_sycl_gpu());
}

88
unsupported/test/cxx11_tensor_sycl.cpp
View File

@ -6,11 +6,13 @@
// Ralph Potter Codeplay Software Ltd. // Ralph Potter Codeplay Software Ltd.
// Luke Iwanski Codeplay Software Ltd. // Luke Iwanski Codeplay Software Ltd.
// Contact: <eigen@codeplay.com> // Contact: <eigen@codeplay.com>
// Benoit Steiner <benoit.steiner.goog@gmail.com>
// //
// This Source Code Form is subject to the terms of the Mozilla // 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 // 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/. // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
#define EIGEN_TEST_NO_LONGDOUBLE #define EIGEN_TEST_NO_LONGDOUBLE
#define EIGEN_TEST_NO_COMPLEX #define EIGEN_TEST_NO_COMPLEX
#define EIGEN_TEST_FUNC cxx11_tensor_sycl #define EIGEN_TEST_FUNC cxx11_tensor_sycl
@ -30,7 +32,7 @@ using TestTensor = Tensor<float, 3>;
using TestTensorMap = TensorMap<Tensor<float, 3>>; using TestTensorMap = TensorMap<Tensor<float, 3>>;
void test_sycl_cpu() { void test_sycl_cpu() {
cl::sycl::gpu_selector s; cl::sycl::gpu_selector s;
cl::sycl::queue q(s, [=](cl::sycl::exception_list l) { cl::sycl::queue q(s, [=](cl::sycl::exception_list l) {
for (const auto& e : l) { for (const auto& e : l) {
try { try {
@ -40,7 +42,7 @@ void test_sycl_cpu() {
} }
} }
}); });
SyclDevice sycl_device(q); SyclDevice sycl_device(q);
int sizeDim1 = 100; int sizeDim1 = 100;
int sizeDim2 = 100; int sizeDim2 = 100;
@ -53,14 +55,14 @@ void test_sycl_cpu() {
in1 = in1.random(); in1 = in1.random();
in2 = in2.random(); in2 = in2.random();
in3 = in3.random(); in3 = in3.random();
TestTensorMap gpu_in1(in1.data(), tensorRange); TestTensorMap gpu_in1(in1.data(), tensorRange);
TestTensorMap gpu_in2(in2.data(), tensorRange); TestTensorMap gpu_in2(in2.data(), tensorRange);
TestTensorMap gpu_in3(in3.data(), tensorRange); TestTensorMap gpu_in3(in3.data(), tensorRange);
TestTensorMap gpu_out(out.data(), tensorRange); TestTensorMap gpu_out(out.data(), tensorRange);
/// a=1.2f /// a=1.2f
gpu_in1.device(sycl_device) = gpu_in1.constant(1.2f); gpu_in1.device(sycl_device) = gpu_in1.constant(1.2f);
sycl_device.deallocate(in1.data()); sycl_device.deallocate(in1.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
@ -68,11 +70,11 @@ void test_sycl_cpu() {
} }
} }
} }
printf("a=1.2f Test passed\n"); printf("a=1.2f Test passed\n");
/// a=b*1.2f /// a=b*1.2f
gpu_out.device(sycl_device) = gpu_in1 * 1.2f; gpu_out.device(sycl_device) = gpu_in1 * 1.2f;
sycl_device.deallocate(out.data()); sycl_device.deallocate(out.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
@ -81,11 +83,11 @@ void test_sycl_cpu() {
} }
} }
} }
printf("a=b*1.2f Test Passed\n"); printf("a=b*1.2f Test Passed\n");
/// c=a*b /// c=a*b
gpu_out.device(sycl_device) = gpu_in1 * gpu_in2; gpu_out.device(sycl_device) = gpu_in1 * gpu_in2;
sycl_device.deallocate(out.data()); sycl_device.deallocate(out.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
@ -95,11 +97,11 @@ void test_sycl_cpu() {
} }
} }
} }
printf("c=a*b Test Passed\n"); printf("c=a*b Test Passed\n");
/// c=a+b /// c=a+b
gpu_out.device(sycl_device) = gpu_in1 + gpu_in2; gpu_out.device(sycl_device) = gpu_in1 + gpu_in2;
sycl_device.deallocate(out.data()); sycl_device.deallocate(out.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
@ -109,11 +111,11 @@ void test_sycl_cpu() {
} }
} }
} }
printf("c=a+b Test Passed\n"); printf("c=a+b Test Passed\n");
/// c=a*a /// c=a*a
gpu_out.device(sycl_device) = gpu_in1 * gpu_in1; gpu_out.device(sycl_device) = gpu_in1 * gpu_in1;
sycl_device.deallocate(out.data()); sycl_device.deallocate(out.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
@ -124,11 +126,11 @@ void test_sycl_cpu() {
} }
} }
printf("c= a*a Test Passed\n"); printf("c= a*a Test Passed\n");
//a*3.14f + b*2.7f //a*3.14f + b*2.7f
gpu_out.device(sycl_device) = gpu_in1 * gpu_in1.constant(3.14f) + gpu_in2 * gpu_in2.constant(2.7f); gpu_out.device(sycl_device) = gpu_in1 * gpu_in1.constant(3.14f) + gpu_in2 * gpu_in2.constant(2.7f);
sycl_device.deallocate(out.data()); sycl_device.deallocate(out.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
@ -138,21 +140,21 @@ void test_sycl_cpu() {
} }
} }
} }
printf("a*3.14f + b*2.7f Test Passed\n"); printf("a*3.14f + b*2.7f Test Passed\n");
///d= (a>0.5? b:c) ///d= (a>0.5? b:c)
gpu_out.device(sycl_device) =(gpu_in1 > gpu_in1.constant(0.5f)).select(gpu_in2, gpu_in3); gpu_out.device(sycl_device) =(gpu_in1 > gpu_in1.constant(0.5f)).select(gpu_in2, gpu_in3);
sycl_device.deallocate(out.data()); sycl_device.deallocate(out.data());
for (int i = 0; i < sizeDim1; ++i) { for (int i = 0; i < sizeDim1; ++i) {
for (int j = 0; j < sizeDim2; ++j) { for (int j = 0; j < sizeDim2; ++j) {
for (int k = 0; k < sizeDim3; ++k) { for (int k = 0; k < sizeDim3; ++k) {
VERIFY_IS_APPROX(out(i, j, k), (in1(i, j, k) > 0.5f) VERIFY_IS_APPROX(out(i, j, k), (in1(i, j, k) > 0.5f)
? in2(i, j, k) ? in2(i, j, k)
: in3(i, j, k)); : in3(i, j, k));
} }
} }
} }
printf("d= (a>0.5? b:c) Test Passed\n"); printf("d= (a>0.5? b:c) Test Passed\n");
} }
void test_cxx11_tensor_sycl() { void test_cxx11_tensor_sycl() {