Adding TensorChippingOP for sycl backend; fixing the index value in the verification operation for cxx11_tensorChipping.cpp test

2026-04-10 11:34:33 +08:00 · 2017-02-13 17:25:12 +00:00
parent fad776492f
commit 0d153ded29
10 changed files with 732 additions and 13 deletions
--- 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);
--- 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
--- 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>
--- 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>
--- 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>
--- 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>{};
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h
@@ -156,6 +156,18 @@ EVALTO()
 #undef 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)\