diff options
author | Mehdi Goli <mehdi.goli@codeplay.com> | 2017-02-13 17:25:12 +0000 |
---|---|---|
committer | Mehdi Goli <mehdi.goli@codeplay.com> | 2017-02-13 17:25:12 +0000 |
commit | 0d153ded29022021c4f7ac24b73a0adb1e423013 (patch) | |
tree | 54089cf3904762064383e5c5dff5e035d24a9813 /unsupported/Eigen/CXX11/src | |
parent | fad776492ff337b3bc0884715c5f80c980ed63a7 (diff) |
Adding TensorChippingOP for sycl backend; fixing the index value in the verification operation for cxx11_tensorChipping.cpp test
Diffstat (limited to 'unsupported/Eigen/CXX11/src')
7 files changed, 105 insertions, 9 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorChipping.h b/unsupported/Eigen/CXX11/src/Tensor/TensorChipping.h index 1ba7ef170..f335edf7d 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorChipping.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorChipping.h @@ -150,7 +150,7 @@ struct TensorEvaluator<const TensorChippingOp<DimId, ArgType>, Device> }; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device) - : m_impl(op.expression(), device), m_dim(op.dim()), m_device(device) + : m_impl(op.expression(), device), m_dim(op.dim()), m_device(device), m_offset(op.offset()) { EIGEN_STATIC_ASSERT((NumInputDims >= 1), YOU_MADE_A_PROGRAMMING_MISTAKE); eigen_assert(NumInputDims > m_dim.actualDim()); @@ -206,7 +206,7 @@ struct TensorEvaluator<const TensorChippingOp<DimId, ArgType>, Device> eigen_assert(index+PacketSize-1 < dimensions().TotalSize()); if ((static_cast<int>(Layout) == static_cast<int>(ColMajor) && m_dim.actualDim() == 0) || - (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == NumInputDims-1)) { + (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == NumInputDims-1)) { // m_stride is equal to 1, so let's avoid the integer division. eigen_assert(m_stride == 1); Index inputIndex = index * m_inputStride + m_inputOffset; @@ -218,7 +218,7 @@ struct TensorEvaluator<const TensorChippingOp<DimId, ArgType>, Device> PacketReturnType rslt = internal::pload<PacketReturnType>(values); return rslt; } else if ((static_cast<int>(Layout) == static_cast<int>(ColMajor) && m_dim.actualDim() == NumInputDims - 1) || - (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == 0)) { + (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == 0)) { // m_stride is aways greater than index, so let's avoid the integer division. eigen_assert(m_stride > index); return m_impl.template packet<LoadMode>(index + m_inputOffset); @@ -274,17 +274,29 @@ struct TensorEvaluator<const TensorChippingOp<DimId, ArgType>, Device> } } + /// used by sycl + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE DenseIndex dimId() const { + return m_dim.actualDim(); + } + + /// used by sycl + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const DenseIndex& offset() const { + return m_offset; + } + /// required by sycl in order to extract the accessor + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; } + protected: EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index srcCoeff(Index index) const { Index inputIndex; if ((static_cast<int>(Layout) == static_cast<int>(ColMajor) && m_dim.actualDim() == 0) || - (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == NumInputDims-1)) { + (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == NumInputDims-1)) { // m_stride is equal to 1, so let's avoid the integer division. eigen_assert(m_stride == 1); inputIndex = index * m_inputStride + m_inputOffset; } else if ((static_cast<int>(Layout) == static_cast<int>(ColMajor) && m_dim.actualDim() == NumInputDims-1) || - (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == 0)) { + (static_cast<int>(Layout) == static_cast<int>(RowMajor) && m_dim.actualDim() == 0)) { // m_stride is aways greater than index, so let's avoid the integer division. eigen_assert(m_stride > index); inputIndex = index + m_inputOffset; @@ -304,6 +316,9 @@ struct TensorEvaluator<const TensorChippingOp<DimId, ArgType>, Device> TensorEvaluator<ArgType, Device> m_impl; const internal::DimensionId<DimId> m_dim; const Device& m_device; +// required by sycl + const DenseIndex m_offset; + }; @@ -344,7 +359,7 @@ struct TensorEvaluator<TensorChippingOp<DimId, ArgType>, Device> EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE) if ((static_cast<int>(this->Layout) == static_cast<int>(ColMajor) && this->m_dim.actualDim() == 0) || - (static_cast<int>(this->Layout) == static_cast<int>(RowMajor) && this->m_dim.actualDim() == NumInputDims-1)) { + (static_cast<int>(this->Layout) == static_cast<int>(RowMajor) && this->m_dim.actualDim() == NumInputDims-1)) { // m_stride is equal to 1, so let's avoid the integer division. eigen_assert(this->m_stride == 1); EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[PacketSize]; @@ -355,7 +370,7 @@ struct TensorEvaluator<TensorChippingOp<DimId, ArgType>, Device> inputIndex += this->m_inputStride; } } else if ((static_cast<int>(this->Layout) == static_cast<int>(ColMajor) && this->m_dim.actualDim() == NumInputDims-1) || - (static_cast<int>(this->Layout) == static_cast<int>(RowMajor) && this->m_dim.actualDim() == 0)) { + (static_cast<int>(this->Layout) == static_cast<int>(RowMajor) && this->m_dim.actualDim() == 0)) { // m_stride is aways greater than index, so let's avoid the integer division. eigen_assert(this->m_stride > index); this->m_impl.template writePacket<StoreMode>(index + this->m_inputOffset, x); diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h index c0bcf26cd..ee8f3c9c2 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h @@ -146,6 +146,18 @@ KERNELBROKERCONVERTERSLICESTRIDEOP() #undef KERNELBROKERCONVERTERSLICESTRIDEOP +/// specialisation of the \ref ConvertToDeviceExpression struct when the node type is TensorChippingOp +#define KERNELBROKERCONVERTCHIPPINGOP(CVQual)\ +template <DenseIndex DimId, typename Expr>\ +struct ConvertToDeviceExpression<CVQual TensorChippingOp<DimId, Expr> > {\ + typedef CVQual TensorChippingOp<DimId, typename ConvertToDeviceExpression<Expr>::Type> Type;\ +}; +KERNELBROKERCONVERTCHIPPINGOP(const) +KERNELBROKERCONVERTCHIPPINGOP() +#undef KERNELBROKERCONVERTCHIPPINGOP + + + } // namespace internal } // namespace TensorSycl } // namespace Eigen diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h index 56ba82805..3b83b1d2c 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h @@ -368,6 +368,23 @@ SYCLPADDINGOPEXPRCONST(TensorPaddingOp, ) #undef SYCLPADDINGOPEXPRCONST +// TensorChippingOp +#define SYCLTENSORCHIPPINGOPEXPR(CVQual)\ +template<DenseIndex DimId, typename OrigXprType, typename XprType, typename... Params>\ +struct ExprConstructor<CVQual TensorChippingOp <DimId, OrigXprType> , CVQual TensorChippingOp<DimId, XprType>, Params... >{\ + typedef ExprConstructor<OrigXprType, XprType, Params...> my_xpr_type;\ + typedef CVQual TensorChippingOp<DimId, typename my_xpr_type::Type> Type;\ + my_xpr_type xprExpr;\ + Type expr;\ + template <typename FuncDetector>\ + ExprConstructor(FuncDetector &funcD, const utility::tuple::Tuple<Params...> &t)\ + : xprExpr(funcD.xprExpr, t), expr(xprExpr.expr, funcD.offset(), funcD.dimId()) {}\ +}; + +SYCLTENSORCHIPPINGOPEXPR(const) +SYCLTENSORCHIPPINGOPEXPR() +#undef SYCLTENSORCHIPPINGOPEXPR + /// template deduction for \ref ExprConstructor struct template <typename OrigExpr, typename IndexExpr, typename FuncD, typename... Params> diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h index 3fd607941..b512d43f6 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h @@ -188,7 +188,7 @@ SYCLCONTRACTIONCONVOLUTIONEXTACC(,TensorConvolutionOp) /// specialisation of the \ref ExtractAccessor struct when the node type is -/// const TensorSlicingOp. +/// const TensorSlicingOp. #define SYCLSLICEOPEXTACC(CVQual)\ template <typename StartIndices, typename Sizes, typename XprType, typename Dev>\ struct ExtractAccessor<TensorEvaluator<CVQual TensorSlicingOp<StartIndices, Sizes, XprType>, Dev> > {\ @@ -200,7 +200,7 @@ SYCLSLICEOPEXTACC(const) SYCLSLICEOPEXTACC() #undef SYCLSLICEOPEXTACC // specialisation of the \ref ExtractAccessor struct when the node type is -/// const TensorStridingSlicingOp. +/// TensorStridingSlicingOp. #define SYCLSLICESTRIDEOPEXTACC(CVQual)\ template<typename StartIndices, typename StopIndices, typename Strides, typename XprType, typename Dev>\ struct ExtractAccessor<TensorEvaluator<CVQual TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType>, Dev> >{\ @@ -212,6 +212,19 @@ SYCLSLICESTRIDEOPEXTACC(const) SYCLSLICESTRIDEOPEXTACC() #undef SYCLSLICESTRIDEOPEXTACC +// specialisation of the \ref ExtractAccessor struct when the node type is +/// TensorChippingOp. +#define SYCLTENSORCHIPPINGOPEXTACC(CVQual)\ +template<DenseIndex DimId, typename XprType, typename Dev>\ +struct ExtractAccessor<TensorEvaluator<CVQual TensorChippingOp<DimId, XprType>, Dev> >{\ + static inline auto getTuple(cl::sycl::handler& cgh, const TensorEvaluator<CVQual TensorChippingOp<DimId, XprType>, Dev>& eval)\ + RETURN_CPP11(AccessorConstructor::getTuple(cgh, eval.impl()))\ +}; + +SYCLTENSORCHIPPINGOPEXTACC(const) +SYCLTENSORCHIPPINGOPEXTACC() +#undef SYCLTENSORCHIPPINGOPEXTACC + /// template deduction for \ref ExtractAccessor template <typename Evaluator> diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h index e26cbdf6d..ee020184b 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h @@ -290,6 +290,22 @@ SYCLEXTRFUNCCONTRACTCONCAT(TensorConcatenationOp, axis(), const) SYCLEXTRFUNCCONTRACTCONCAT(TensorConcatenationOp, axis(),) #undef SYCLEXTRFUNCCONTRACTCONCAT +//TensorChippingOp +#define SYCLEXTRFUNCCHIPPINGOP(CVQual)\ +template<DenseIndex DimId, typename XprType, typename Device>\ +struct FunctorExtractor<TensorEvaluator<CVQual TensorChippingOp<DimId, XprType>, Device>>{\ + FunctorExtractor<Eigen::TensorEvaluator<XprType, Device> > xprExpr;\ + const DenseIndex m_dim;\ + const DenseIndex m_offset;\ + EIGEN_STRONG_INLINE const DenseIndex& dimId() const { return m_dim; }\ + EIGEN_STRONG_INLINE const DenseIndex& offset() const { return m_offset; }\ + FunctorExtractor(const TensorEvaluator<CVQual TensorChippingOp<DimId, XprType>, Device>& expr)\ + : xprExpr(expr.impl()), m_dim(expr.dimId()), m_offset(expr.offset()) {}\ +}; + +SYCLEXTRFUNCCHIPPINGOP(const) +SYCLEXTRFUNCCHIPPINGOP() +#undef SYCLEXTRFUNCCHIPPINGOP /// template deduction function for FunctorExtractor template <typename Evaluator> diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h index 0ac51e7bf..a1c112f4d 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h @@ -139,6 +139,17 @@ SLICEOPLEAFCOUNT(const) SLICEOPLEAFCOUNT() #undef SLICEOPLEAFCOUNT + +/// specialisation of the \ref LeafCount struct when the node type is TensorChippingOp +#define CHIPPINGOPLEAFCOUNT(CVQual)\ +template <DenseIndex DimId, typename XprType>\ +struct LeafCount<CVQual TensorChippingOp<DimId, XprType> >:CategoryCount<XprType>{}; + +CHIPPINGOPLEAFCOUNT(const) +CHIPPINGOPLEAFCOUNT() +#undef CHIPPINGOPLEAFCOUNT + + #define SLICESTRIDEOPLEAFCOUNT(CVQual)\ template<typename StartIndices, typename StopIndices, typename Strides, typename XprType>\ struct LeafCount<CVQual TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType> >:CategoryCount<XprType>{}; diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h index f6e3b4766..74566dcee 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h @@ -157,6 +157,18 @@ EVALTO() /// specialisation of the \ref PlaceHolderExpression when the node is +/// TensorChippingOp +#define CHIPPINGOP(CVQual)\ +template <DenseIndex DimId, typename Expr, size_t N>\ +struct PlaceHolderExpression<CVQual TensorChippingOp<DimId, Expr>, N> {\ + typedef CVQual TensorChippingOp< DimId, typename CalculateIndex <N, Expr>::ArgType> Type;\ +}; + +CHIPPINGOP(const) +CHIPPINGOP() +#undef CHIPPINGOP + +/// specialisation of the \ref PlaceHolderExpression when the node is /// TensorReductionOp #define SYCLREDUCTION(CVQual)\ template <typename OP, typename Dims, typename Expr, size_t N>\ |