Partial OpenCL support via SYCL compatible with ComputeCpp CE.

22 files changed, 2861 insertions, 62 deletions

diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorAssign.h b/unsupported/Eigen/CXX11/src/Tensor/TensorAssign.h
index cb615c75b..166be200c 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorAssign.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorAssign.h
@@ -163,6 +163,11 @@ struct TensorEvaluator<const TensorAssignOp<LeftArgType, RightArgType>, Device>
            TensorOpCost(0, sizeof(CoeffReturnType), 0, vectorized, PacketSize);
   }
 
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<RightArgType, Device>& right_impl() const { return m_rightImpl; }
+
   EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return m_leftImpl.data(); }
 
  private:
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorBase.h b/unsupported/Eigen/CXX11/src/Tensor/TensorBase.h
index 19d2b50b5..e3880d2e0 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorBase.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorBase.h
@@ -811,7 +811,7 @@ class TensorBase<Derived, ReadOnlyAccessors>
 
   protected:
     template <typename Scalar, int NumIndices, int Options, typename IndexType> friend class Tensor;
-    template <typename Scalar, typename Dimensions, int Option, typename IndexTypes> friend class TensorFixedSize;
+    template <typename Scalar, typename Dimensions, int Option, typename IndexTypes, template <class> class MakePointer_> friend class TensorFixedSize;
     template <typename OtherDerived, int AccessLevel> friend class TensorBase;
     EIGEN_DEVICE_FUNC
     EIGEN_STRONG_INLINE const Derived& derived() const { return *static_cast<const Derived*>(this); }
@@ -827,7 +827,7 @@ class TensorBase : public TensorBase<Derived, ReadOnlyAccessors> {
     static const int NumDimensions = DerivedTraits::NumDimensions;
 
     template <typename Scalar, int NumIndices, int Options, typename IndexType> friend class Tensor;
-    template <typename Scalar, typename Dimensions, int Option, typename IndexTypes> friend class TensorFixedSize;
+    template <typename Scalar, typename Dimensions, int Option, typename IndexTypes, template <class> class MakePointer_> friend class TensorFixedSize;
     template <typename OtherDerived, int OtherAccessLevel> friend class TensorBase;
 
     EIGEN_DEVICE_FUNC
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorBroadcasting.h b/unsupported/Eigen/CXX11/src/Tensor/TensorBroadcasting.h
index 5d67f69f3..4cfe300eb 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorBroadcasting.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorBroadcasting.h
@@ -113,7 +113,7 @@ struct TensorEvaluator<const TensorBroadcastingOp<Broadcast, ArgType>, Device>
   };
 
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
-    : m_impl(op.expression(), device)
+    : m_broadcast(op.broadcast()),m_impl(op.expression(), device)
   {
     // The broadcasting op doesn't change the rank of the tensor. One can't broadcast a scalar
     // and store the result in a scalar. Instead one should reshape the scalar into a a N-D
@@ -374,7 +374,12 @@ struct TensorEvaluator<const TensorBroadcastingOp<Broadcast, ArgType>, Device>
 
   EIGEN_DEVICE_FUNC Scalar* data() const { return NULL; }
 
+  const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
+
+  Broadcast functor() const { return m_broadcast; }
+
  protected:
+  const Broadcast m_broadcast;
   Dimensions m_dimensions;
   array<Index, NumDims> m_outputStrides;
   array<Index, NumDims> m_inputStrides;
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h
new file mode 100644
index 000000000..bfd36f5aa
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h
@@ -0,0 +1,122 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Copyright (C) 2016 Benoit Steiner <benoit.steiner.goog@gmail.com>
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Cummins Chris PhD student at The University of Edinburgh.
+// Contact: <eigen@codeplay.com>
+
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+#if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H)
+#define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H
+
+namespace Eigen {
+/// \struct BufferT is used to specialise add_sycl_buffer function for
+//  two types of buffer we have. When the MapAllocator is true, we create the
+//  sycl buffer with MapAllocator.
+/// We have to const_cast the input pointer in order to work around the fact
+/// that sycl does not accept map allocator for const pointer.
+template <typename T, bool MapAllocator>
+struct BufferT {
+  using Type = cl::sycl::buffer<T, 1, cl::sycl::map_allocator<T>>;
+  static inline void add_sycl_buffer(
+      const T *ptr, size_t num_bytes,
+      std::map<const void *, std::shared_ptr<void>> &buffer_map) {
+    buffer_map.insert(std::pair<const void *, std::shared_ptr<void>>(
+        ptr, std::shared_ptr<void>(std::make_shared<Type>(
+                 Type(const_cast<T *>(ptr), cl::sycl::range<1>(num_bytes))))));
+  }
+};
+
+/// specialisation of the \ref BufferT when the MapAllocator is false. In this
+/// case we only create the device-only buffer.
+template <typename T>
+struct BufferT<T, false> {
+  using Type = cl::sycl::buffer<T, 1>;
+  static inline void add_sycl_buffer(
+      const T *ptr, size_t num_bytes,
+      std::map<const void *, std::shared_ptr<void>> &buffer_map) {
+    buffer_map.insert(std::pair<const void *, std::shared_ptr<void>>(
+        ptr, std::shared_ptr<void>(
+                 std::make_shared<Type>(Type(cl::sycl::range<1>(num_bytes))))));
+  }
+};
+
+struct SyclDevice {
+  /// class members
+  /// sycl queue
+  cl::sycl::queue &m_queue;
+  /// std::map is the container used to make sure that we create only one buffer
+  /// per pointer. The lifespan of the buffer
+  /// now depends on the lifespan of SyclDevice. If a non-read-only pointer is
+  /// needed to be accessed on the host we should manually deallocate it.
+  mutable std::map<const void *, std::shared_ptr<void>> buffer_map;
+
+  SyclDevice(cl::sycl::queue &q) : m_queue(q) {}
+  // destructor
+  ~SyclDevice() { deallocate_all(); }
+
+  template <typename T>
+  void deallocate(const T *p) const {
+    auto it = buffer_map.find(p);
+    if (it != buffer_map.end()) {
+      buffer_map.erase(it);
+    }
+  }
+  void deallocate_all() const { buffer_map.clear(); }
+
+  /// creation of sycl accessor for a buffer. This function first tries to find
+  /// the buffer in the buffer_map.
+  /// If found it gets the accessor from it, if not, the function then adds an
+  /// entry by creating a sycl buffer
+  /// for that particular pointer.
+  template <cl::sycl::access::mode AcMd, bool MapAllocator, typename T>
+  inline cl::sycl::accessor<T, 1, AcMd, cl::sycl::access::target::global_buffer>
+  get_sycl_accessor(size_t num_bytes, cl::sycl::handler &cgh,
+                    const T *ptr) const {
+    auto it = buffer_map.find(ptr);
+    if (it == buffer_map.end()) {
+      BufferT<T, MapAllocator>::add_sycl_buffer(ptr, num_bytes, buffer_map);
+    }
+    return (
+        ((typename BufferT<T, MapAllocator>::Type *)(buffer_map.at(ptr).get()))
+            ->template get_access<AcMd>(cgh));
+  }
+
+  /// allocating memory on the cpu
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void *allocate(size_t num_bytes) const {
+    return internal::aligned_malloc(num_bytes);
+  }
+
+  // some runtime conditions that can be applied here
+  bool isDeviceSuitable() const { return true; }
+
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void deallocate(void *buffer) const {
+    internal::aligned_free(buffer);
+  }
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void memcpy(void *dst, const void *src,
+                                                    size_t n) const {
+    ::memcpy(dst, src, n);
+  }
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void memcpyHostToDevice(
+      void *dst, const void *src, size_t n) const {
+    memcpy(dst, src, n);
+  }
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void memcpyDeviceToHost(
+      void *dst, const void *src, size_t n) const {
+    memcpy(dst, src, n);
+  }
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void memset(void *buffer, int c,
+                                                    size_t n) const {
+    ::memset(buffer, c, n);
+  }
+};
+}  // end namespace Eigen
+
+#endif  // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorEvalTo.h b/unsupported/Eigen/CXX11/src/Tensor/TensorEvalTo.h
index a08dfa7c3..3dab6da99 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorEvalTo.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorEvalTo.h
@@ -20,8 +20,8 @@ namespace Eigen {
   *
   */
 namespace internal {
-template<typename XprType>
-struct traits<TensorEvalToOp<XprType> >
+template<typename XprType, template <class> class MakePointer_>
+struct traits<TensorEvalToOp<XprType, MakePointer_> >
 {
   // Type promotion to handle the case where the types of the lhs and the rhs are different.
   typedef typename XprType::Scalar Scalar;
@@ -36,16 +36,20 @@ struct traits<TensorEvalToOp<XprType> >
   enum {
     Flags = 0
   };
+  template <class T>
+  struct MakePointer {
+    typedef typename MakePointer_<T>::Type Type;
+  };
 };
 
-template<typename XprType>
-struct eval<TensorEvalToOp<XprType>, Eigen::Dense>
+template<typename XprType, template <class> class MakePointer_>
+struct eval<TensorEvalToOp<XprType, MakePointer_>, Eigen::Dense>
 {
   typedef const TensorEvalToOp<XprType>& type;
 };
 
-template<typename XprType>
-struct nested<TensorEvalToOp<XprType>, 1, typename eval<TensorEvalToOp<XprType> >::type>
+template<typename XprType, template <class> class MakePointer_>
+struct nested<TensorEvalToOp<XprType, MakePointer_>, 1, typename eval<TensorEvalToOp<XprType, MakePointer_> >::type>
 {
   typedef TensorEvalToOp<XprType> type;
 };
@@ -55,37 +59,38 @@ struct nested<TensorEvalToOp<XprType>, 1, typename eval<TensorEvalToOp<XprType>
 
 
 
-template<typename XprType>
-class TensorEvalToOp : public TensorBase<TensorEvalToOp<XprType>, ReadOnlyAccessors>
+template<typename XprType, template <class> class MakePointer_>
+class TensorEvalToOp : public TensorBase<TensorEvalToOp<XprType, MakePointer_>, ReadOnlyAccessors>
 {
   public:
   typedef typename Eigen::internal::traits<TensorEvalToOp>::Scalar Scalar;
   typedef typename Eigen::NumTraits<Scalar>::Real RealScalar;
   typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
+  typedef typename MakePointer_<CoeffReturnType>::Type PointerType;
   typedef typename Eigen::internal::nested<TensorEvalToOp>::type Nested;
   typedef typename Eigen::internal::traits<TensorEvalToOp>::StorageKind StorageKind;
   typedef typename Eigen::internal::traits<TensorEvalToOp>::Index Index;
 
-  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvalToOp(CoeffReturnType* buffer, const XprType& expr)
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvalToOp(PointerType buffer, const XprType& expr)
       : m_xpr(expr), m_buffer(buffer) {}
 
     EIGEN_DEVICE_FUNC
     const typename internal::remove_all<typename XprType::Nested>::type&
     expression() const { return m_xpr; }
 
-    EIGEN_DEVICE_FUNC CoeffReturnType* buffer() const { return m_buffer; }
+    EIGEN_DEVICE_FUNC PointerType buffer() const { return m_buffer; }
 
   protected:
     typename XprType::Nested m_xpr;
-    CoeffReturnType* m_buffer;
+    PointerType m_buffer;
 };
 
 
 
-template<typename ArgType, typename Device>
-struct TensorEvaluator<const TensorEvalToOp<ArgType>, Device>
+template<typename ArgType, typename Device, template <class> class MakePointer_>
+struct TensorEvaluator<const TensorEvalToOp<ArgType, MakePointer_>, Device>
 {
-  typedef TensorEvalToOp<ArgType> XprType;
+  typedef TensorEvalToOp<ArgType, MakePointer_> XprType;
   typedef typename ArgType::Scalar Scalar;
   typedef typename TensorEvaluator<ArgType, Device>::Dimensions Dimensions;
   typedef typename XprType::Index Index;
@@ -102,15 +107,22 @@ struct TensorEvaluator<const TensorEvalToOp<ArgType>, Device>
   };
 
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
-      : m_impl(op.expression(), device), m_device(device), m_buffer(op.buffer())
+      : m_impl(op.expression(), device), m_device(device),
+          m_buffer(op.buffer()), m_op(op), m_expression(op.expression())
   { }
 
+  // Used for accessor extraction in SYCL Managed TensorMap:
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const XprType& op() const {
+    return m_op;
+  }
+  
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ~TensorEvaluator() {
   }
 
+  typedef typename internal::traits<const TensorEvalToOp<ArgType, MakePointer_>>::template MakePointer<CoeffReturnType>::Type DevicePointer;
   EIGEN_DEVICE_FUNC const Dimensions& dimensions() const { return m_impl.dimensions(); }
 
-  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType* scalar) {
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(DevicePointer scalar) {
     EIGEN_UNUSED_VARIABLE(scalar);
     eigen_assert(scalar == NULL);
     return m_impl.evalSubExprsIfNeeded(m_buffer);
@@ -145,12 +157,20 @@ struct TensorEvaluator<const TensorEvalToOp<ArgType>, Device>
         TensorOpCost(0, sizeof(CoeffReturnType), 0, vectorized, PacketSize);
   }
 
-  EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return m_buffer; }
+  EIGEN_DEVICE_FUNC DevicePointer data() const { return m_buffer; }
+  ArgType expression() const { return m_expression; }
+
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
+  /// added for sycl in order to construct the buffer from the sycl device
+  const Device& device() const{return m_device;}
 
  private:
   TensorEvaluator<ArgType, Device> m_impl;
   const Device& m_device;
-  CoeffReturnType* m_buffer;
+  DevicePointer m_buffer;
+  const XprType& m_op;
+  const ArgType m_expression;
 };
 
 
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorEvaluator.h b/unsupported/Eigen/CXX11/src/Tensor/TensorEvaluator.h
index c2a327bf0..b2b4bcf62 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorEvaluator.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorEvaluator.h
@@ -46,9 +46,11 @@ struct TensorEvaluator
   };
 
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const Derived& m, const Device& device)
-      : m_data(const_cast<Scalar*>(m.data())), m_dims(m.dimensions()), m_device(device)
+      : m_data(const_cast<typename internal::traits<Derived>::template MakePointer<Scalar>::Type>(m.data())), m_dims(m.dimensions()), m_device(device), m_impl(m)
   { }
 
+  // Used for accessor extraction in SYCL Managed TensorMap:
+  const Derived& derived() const { return m_impl; }
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dims; }
 
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType* dest) {
@@ -106,12 +108,16 @@ struct TensorEvaluator
                         internal::unpacket_traits<PacketReturnType>::size);
   }
 
-  EIGEN_DEVICE_FUNC Scalar* data() const { return m_data; }
+  EIGEN_DEVICE_FUNC typename internal::traits<Derived>::template MakePointer<Scalar>::Type data() const { return m_data; }
+
+  /// required by sycl in order to construct sycl buffer from raw pointer
+  const Device& device() const{return m_device;}
 
  protected:
-  Scalar* m_data;
+  typename internal::traits<Derived>::template MakePointer<Scalar>::Type m_data;
   Dimensions m_dims;
   const Device& m_device;
+  const Derived& m_impl;
 };
 
 namespace {
@@ -159,8 +165,11 @@ struct TensorEvaluator<const Derived, Device>
     RawAccess = true
   };
 
+  // Used for accessor extraction in SYCL Managed TensorMap:
+  const Derived& derived() const { return m_impl; }
+
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const Derived& m, const Device& device)
-      : m_data(m.data()), m_dims(m.dimensions()), m_device(device)
+      : m_data(m.data()), m_dims(m.dimensions()), m_device(device), m_impl(m)
   { }
 
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dims; }
@@ -198,12 +207,16 @@ struct TensorEvaluator<const Derived, Device>
                         internal::unpacket_traits<PacketReturnType>::size);
   }
 
-  EIGEN_DEVICE_FUNC const Scalar* data() const { return m_data; }
+  EIGEN_DEVICE_FUNC typename internal::traits<Derived>::template MakePointer<const Scalar>::Type data() const { return m_data; }
+
+  /// added for sycl in order to construct the buffer from the sycl device
+  const Device& device() const{return m_device;}
 
  protected:
-  const Scalar* m_data;
+  typename internal::traits<Derived>::template MakePointer<const Scalar>::Type m_data;
   Dimensions m_dims;
   const Device& m_device;
+  const Derived& m_impl;
 };
 
 
@@ -260,6 +273,12 @@ struct TensorEvaluator<const TensorCwiseNullaryOp<NullaryOp, ArgType>, Device>
 
   EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; }
 
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<ArgType, Device>& impl() const { return m_argImpl; }
+  /// required by sycl in order to extract the accessor
+  NullaryOp functor() const { return m_functor; }
+
+
  private:
   const NullaryOp m_functor;
   TensorEvaluator<ArgType, Device> m_argImpl;
@@ -323,6 +342,12 @@ struct TensorEvaluator<const TensorCwiseUnaryOp<UnaryOp, ArgType>, Device>
 
   EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; }
 
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<ArgType, Device> & impl() const { return m_argImpl; }
+  /// added for sycl in order to construct the buffer from sycl device
+  UnaryOp functor() const { return m_functor; }
+
+
  private:
   const UnaryOp m_functor;
   TensorEvaluator<ArgType, Device> m_argImpl;
@@ -396,6 +421,12 @@ struct TensorEvaluator<const TensorCwiseBinaryOp<BinaryOp, LeftArgType, RightArg
   }
 
   EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<RightArgType, Device>& right_impl() const { return m_rightImpl; }
+  /// required by sycl in order to extract the accessor
+  BinaryOp functor() const { return m_functor; }
 
  private:
   const BinaryOp m_functor;
@@ -491,10 +522,17 @@ struct TensorEvaluator<const TensorCwiseTernaryOp<TernaryOp, Arg1Type, Arg2Type,
 
   EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; }
 
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<Arg1Type, Device> & arg1Impl() const { return m_arg1Impl; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<Arg2Type, Device>& arg2Impl() const { return m_arg2Impl; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<Arg3Type, Device>& arg3Impl() const { return m_arg3Impl; }
+
  private:
   const TernaryOp m_functor;
   TensorEvaluator<Arg1Type, Device> m_arg1Impl;
-  TensorEvaluator<Arg1Type, Device> m_arg2Impl;
+  TensorEvaluator<Arg2Type, Device> m_arg2Impl;
   TensorEvaluator<Arg3Type, Device> m_arg3Impl;
 };
 
@@ -575,6 +613,12 @@ struct TensorEvaluator<const TensorSelectOp<IfArgType, ThenArgType, ElseArgType>
   }
 
   EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<IfArgType, Device> & cond_impl() const { return m_condImpl; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<ThenArgType, Device>& then_impl() const { return m_thenImpl; }
+  /// required by sycl in order to extract the accessor
+  const TensorEvaluator<ElseArgType, Device>& else_impl() const { return m_elseImpl; }
 
  private:
   TensorEvaluator<IfArgType, Device> m_condImpl;
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h
index a116bf17f..9b99af641 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h
@@ -272,6 +272,20 @@ inline void TensorExecutor<Expression, GpuDevice, Vectorizable>::run(
 #endif  // __CUDACC__
 #endif  // EIGEN_USE_GPU
 
+// SYCL Executor policy
+#ifdef EIGEN_USE_SYCL
+
+template <typename Expression, bool Vectorizable>
+class TensorExecutor<Expression, SyclDevice, Vectorizable> {
+public:
+  static inline void run(const Expression &expr, const SyclDevice &device) {
+    // call TensorSYCL module
+    TensorSycl::run(expr, device);
+  }
+};
+
+#endif
+
 } // end namespace internal
 
 } // end namespace Eigen
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorFixedSize.h b/unsupported/Eigen/CXX11/src/Tensor/TensorFixedSize.h
index fcee5f60d..415e459b9 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorFixedSize.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorFixedSize.h
@@ -23,8 +23,8 @@ namespace Eigen {
   * Eigen::TensorFixedSize<float, Size<3,5,7>> t;
   */
 
-template<typename Scalar_, typename Dimensions_, int Options_, typename IndexType>
-class TensorFixedSize : public TensorBase<TensorFixedSize<Scalar_, Dimensions_, Options_, IndexType> >
+template<typename Scalar_, typename Dimensions_, int Options_, typename IndexType, template <class> class MakePointer_>
+class TensorFixedSize : public TensorBase<TensorFixedSize<Scalar_, Dimensions_, Options_, IndexType, MakePointer_> >
 {
   public:
     typedef TensorFixedSize<Scalar_, Dimensions_, Options_, IndexType> Self;
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h b/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h
index c23ecdbc4..9cf4a07e5 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorForcedEval.h
@@ -19,9 +19,15 @@ namespace Eigen {
   *
   *
   */
+/// template <class> class MakePointer_ is added to convert the host pointer to the device pointer.
+/// It is added due to the fact that for our device compiler T* is not allowed.
+/// If we wanted to use the same Evaluator functions we have to convert that type to our pointer T.
+/// This is done through our MakePointer_ class. By default the Type in the MakePointer_<T> is T* .
+/// Therefore, by adding the default value, we managed to convert the type and it does not break any
+/// existing code as its default value is T*.
 namespace internal {
-template<typename XprType>
-struct traits<TensorForcedEvalOp<XprType> >
+template<typename XprType, template <class> class MakePointer_>
+struct traits<TensorForcedEvalOp<XprType, MakePointer_> >
 {
   // Type promotion to handle the case where the types of the lhs and the rhs are different.
   typedef typename XprType::Scalar Scalar;
@@ -36,26 +42,30 @@ struct traits<TensorForcedEvalOp<XprType> >
   enum {
     Flags = 0
   };
+  template <class T>
+  struct MakePointer {
+    typedef typename MakePointer_<T>::Type Type;
+  };
 };
 
-template<typename XprType>
-struct eval<TensorForcedEvalOp<XprType>, Eigen::Dense>
+template<typename XprType, template <class> class MakePointer_>
+struct eval<TensorForcedEvalOp<XprType, MakePointer_>, Eigen::Dense>
 {
-  typedef const TensorForcedEvalOp<XprType>& type;
+  typedef const TensorForcedEvalOp<XprType, MakePointer_>& type;
 };
 
-template<typename XprType>
-struct nested<TensorForcedEvalOp<XprType>, 1, typename eval<TensorForcedEvalOp<XprType> >::type>
+template<typename XprType, template <class> class MakePointer_>
+struct nested<TensorForcedEvalOp<XprType, MakePointer_>, 1, typename eval<TensorForcedEvalOp<XprType, MakePointer_> >::type>
 {
-  typedef TensorForcedEvalOp<XprType> type;
+  typedef TensorForcedEvalOp<XprType, MakePointer_> type;
 };
 
 }  // end namespace internal
 
 
 
-template<typename XprType>
-class TensorForcedEvalOp : public TensorBase<TensorForcedEvalOp<XprType>, ReadOnlyAccessors>
+template<typename XprType, template <class> class MakePointer_>
+class TensorForcedEvalOp : public TensorBase<TensorForcedEvalOp<XprType, MakePointer_>, ReadOnlyAccessors>
 {
   public:
   typedef typename Eigen::internal::traits<TensorForcedEvalOp>::Scalar Scalar;
@@ -77,10 +87,10 @@ class TensorForcedEvalOp : public TensorBase<TensorForcedEvalOp<XprType>, ReadOn
 };
 
 
-template<typename ArgType, typename Device>
-struct TensorEvaluator<const TensorForcedEvalOp<ArgType>, Device>
+template<typename ArgType, typename Device, template <class> class MakePointer_>
+struct TensorEvaluator<const TensorForcedEvalOp<ArgType, MakePointer_>, Device>
 {
-  typedef TensorForcedEvalOp<ArgType> XprType;
+  typedef TensorForcedEvalOp<ArgType, MakePointer_> XprType;
   typedef typename ArgType::Scalar Scalar;
   typedef typename TensorEvaluator<ArgType, Device>::Dimensions Dimensions;
   typedef typename XprType::Index Index;
@@ -96,6 +106,7 @@ struct TensorEvaluator<const TensorForcedEvalOp<ArgType>, Device>
   };
 
   EIGEN_DEVICE_FUNC TensorEvaluator(const XprType& op, const Device& device)
+	/// op_ is used for sycl
       : m_impl(op.expression(), device), m_op(op.expression()), m_device(device), m_buffer(NULL)
   { }
 
@@ -110,10 +121,10 @@ struct TensorEvaluator<const TensorForcedEvalOp<ArgType>, Device>
         new(m_buffer+i) CoeffReturnType();
       }
     }
-    typedef TensorEvalToOp<const ArgType> EvalTo;
+    typedef TensorEvalToOp< const typename internal::remove_const<ArgType>::type > EvalTo;
     EvalTo evalToTmp(m_buffer, m_op);
     const bool PacketAccess = internal::IsVectorizable<Device, const ArgType>::value;
-    internal::TensorExecutor<const EvalTo, Device, PacketAccess>::run(evalToTmp, m_device);
+    internal::TensorExecutor<const EvalTo, typename internal::remove_const<Device>::type, PacketAccess>::run(evalToTmp, m_device);
     return true;
   }
   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() {
@@ -136,13 +147,17 @@ struct TensorEvaluator<const TensorForcedEvalOp<ArgType>, Device>
     return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
   }
 
-  EIGEN_DEVICE_FUNC Scalar* data() const { return m_buffer; }
+  EIGEN_DEVICE_FUNC typename MakePointer<Scalar>::Type data() const { return m_buffer; }
 
+  /// required by sycl in order to extract the sycl accessor
+  const TensorEvaluator<ArgType, Device>& impl() { return m_impl; }
+  /// used by sycl in order to build the sycl buffer
+  const Device& device() const{return m_device;}
  private:
   TensorEvaluator<ArgType, Device> m_impl;
   const ArgType m_op;
   const Device& m_device;
-  CoeffReturnType* m_buffer;
+  typename MakePointer<CoeffReturnType>::Type m_buffer;
 };
 
 
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h b/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h
index 490ddd8bd..83c690133 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorForwardDeclarations.h
@@ -12,9 +12,19 @@
 
 namespace Eigen {
 
+// MakePointer class is used as a container of the adress space of the pointer
+// on the host and on the device. From the host side it generates the T* pointer
+// and when EIGEN_USE_SYCL is used it construct a buffer with a map_allocator to
+// T* m_data on the host. It is always called on the device.
+// Specialisation of MakePointer class for creating the sycl buffer with
+// map_allocator.
+template<class T> struct MakePointer{
+	typedef T* Type;
+};
+
+template<typename PlainObjectType, int Options_ = Unaligned, template <class> class MakePointer_ = MakePointer> class TensorMap;
 template<typename Scalar_, int NumIndices_, int Options_ = 0, typename IndexType = DenseIndex> class Tensor;
-template<typename Scalar_, typename Dimensions, int Options_ = 0, typename IndexType = DenseIndex> class TensorFixedSize;
-template<typename PlainObjectType, int Options_ = Unaligned> class TensorMap;
+template<typename Scalar_, typename Dimensions, int Options_ = 0, typename IndexType = DenseIndex, template <class> class MakePointer_ = MakePointer> class TensorFixedSize;
 template<typename PlainObjectType> class TensorRef;
 template<typename Derived, int AccessLevel> class TensorBase;
 
@@ -52,8 +62,8 @@ template<typename Op, typename XprType> class TensorScanOp;
 template<typename CustomUnaryFunc, typename XprType> class TensorCustomUnaryOp;
 template<typename CustomBinaryFunc, typename LhsXprType, typename RhsXprType> class TensorCustomBinaryOp;
 
-template<typename XprType> class TensorEvalToOp;
-template<typename XprType> class TensorForcedEvalOp;
+template<typename XprType, template <class> class MakePointer_ = MakePointer> class TensorEvalToOp;
+template<typename XprType, template <class> class MakePointer_ = MakePointer> class TensorForcedEvalOp;
 
 template<typename ExpressionType, typename DeviceType> class TensorDevice;
 template<typename Derived, typename Device> struct TensorEvaluator;
@@ -61,6 +71,7 @@ template<typename Derived, typename Device> struct TensorEvaluator;
 struct DefaultDevice;
 struct ThreadPoolDevice;
 struct GpuDevice;
+struct SyclDevice;
 
 enum FFTResultType {
   RealPart = 0,
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorMap.h b/unsupported/Eigen/CXX11/src/Tensor/TensorMap.h
index 6fb4f4a31..298a49138 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorMap.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorMap.h
@@ -18,11 +18,16 @@ namespace Eigen {
   * \brief A tensor expression mapping an existing array of data.
   *
   */
-
-template<typename PlainObjectType, int Options_> class TensorMap : public TensorBase<TensorMap<PlainObjectType, Options_> >
+/// template <class> class MakePointer_ is added to convert the host pointer to the device pointer.
+/// It is added due to the fact that for our device compiler T* is not allowed.
+/// If we wanted to use the same Evaluator functions we have to convert that type to our pointer T.
+/// This is done through our MakePointer_ class. By default the Type in the MakePointer_<T> is T* .
+/// Therefore, by adding the default value, we managed to convert the type and it does not break any
+/// existing code as its default value is T*.
+template<typename PlainObjectType, int Options_, template <class> class MakePointer_> class TensorMap : public TensorBase<TensorMap<PlainObjectType, Options_, MakePointer_> >
 {
   public:
-    typedef TensorMap<PlainObjectType, Options_> Self;
+    typedef TensorMap<PlainObjectType, Options_, MakePointer_> Self;
     typedef typename PlainObjectType::Base Base;
     typedef typename Eigen::internal::nested<Self>::type Nested;
     typedef typename internal::traits<PlainObjectType>::StorageKind StorageKind;
@@ -36,7 +41,7 @@ template<typename PlainObjectType, int Options_> class TensorMap : public Tensor
                          Scalar *,
                          const Scalar *>::type
                      PointerType;*/
-    typedef Scalar* PointerType;
+    typedef typename MakePointer_<Scalar>::Type PointerType;
     typedef PointerType PointerArgType;
 
     static const int Options = Options_;
@@ -109,9 +114,9 @@ template<typename PlainObjectType, int Options_> class TensorMap : public Tensor
     EIGEN_DEVICE_FUNC
     EIGEN_STRONG_INLINE Index size() const { return m_dimensions.TotalSize(); }
     EIGEN_DEVICE_FUNC
-    EIGEN_STRONG_INLINE Scalar* data() { return m_data; }
+    EIGEN_STRONG_INLINE PointerType data() { return m_data; }
     EIGEN_DEVICE_FUNC
-    EIGEN_STRONG_INLINE const Scalar* data() const { return m_data; }
+    EIGEN_STRONG_INLINE const PointerType data() const { return m_data; }
 
     EIGEN_DEVICE_FUNC
     EIGEN_STRONG_INLINE const Scalar& operator()(const array<Index, NumIndices>& indices) const
@@ -307,8 +312,9 @@ template<typename PlainObjectType, int Options_> class TensorMap : public Tensor
     }
 
   private:
-    Scalar* m_data;
+    typename MakePointer_<Scalar>::Type m_data;
     Dimensions m_dimensions;
+    size_t is_coverted= size_t(0);
 };
 
 } // end namespace Eigen
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h
new file mode 100644
index 000000000..277dd739c
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSycl.h
@@ -0,0 +1,62 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: eigen@codeplay.com
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+// General include header of SYCL target for Tensor Module
+#ifndef TENSORSYCL_H
+#define TENSORSYCL_H
+
+#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
+template <class T>
+struct MakeGlobalPointer {
+  typedef typename cl::sycl::global_ptr<T>::pointer_t Type;
+};
+
+// tuple construction
+#include "TensorSyclTuple.h"
+
+// This file contains the PlaceHolder that replaces the actual data
+#include "TensorSyclPlaceHolder.h"
+
+#include "TensorSyclLeafCount.h"
+
+// The index PlaceHolder takes the actual expression and replaces the actual
+// data on it with the place holder. It uses the same pre-order expression tree
+// traverse as the leaf count in order to give the right access number to each
+// node in the expression
+#include "TensorSyclPlaceHolderExpr.h"
+
+// creation of an accessor tuple from a tuple of SYCL buffers
+#include "TensorSyclExtractAccessor.h"
+
+// actual data extraction using accessors
+//#include "GetDeviceData.h"
+
+// this is used to change the address space type in tensor map for GPU
+#include "TensorSyclConvertToDeviceExpression.h"
+
+// this is used to extract the functors
+#include "TensorSyclExtractFunctors.h"
+
+// this is used to create tensormap on the device
+// this is used to construct the expression on the device
+#include "TensorSyclExprConstructor.h"
+
+// kernel execution using fusion
+#include "TensorSyclRun.h"
+
+#endif  // end of EIGEN_USE_SYCL
+#endif  // TENSORSYCL_H
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h
new file mode 100644
index 000000000..b3748131b
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclConvertToDeviceExpression.h
@@ -0,0 +1,238 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclConvertToDeviceExpression.h
+ *
+ * \brief:
+ *  Conversion from host pointer to device pointer
+ *  inside leaf nodes of the expression.
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_TENSORSYCL_CONVERT_TO_DEVICE_EXPRESSION_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+namespace internal {
+/// \struct ConvertToDeviceExpression
+/// \brief This struct is used to convert the MakePointer in the host expression
+/// to the MakeGlobalPointer for the device expression. For the leafNodes
+/// containing the pointer. This is due to the fact that the address space of
+/// the pointer T* is different on the host and the device.
+template <typename Expr>
+struct ConvertToDeviceExpression;
+
+/// specialisation of the \ref ConvertToDeviceExpression struct when the node
+/// type is TensorMap
+template <typename Scalar_, int Options_, int Options2_, int NumIndices_,
+          typename IndexType_, template <class> class MakePointer_>
+struct ConvertToDeviceExpression<
+    TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_,
+              MakePointer_>> {
+  using Type = TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
+                         Options2_, MakeGlobalPointer>;
+};
+
+/// specialisation of the \ref ConvertToDeviceExpression struct when the node
+/// type is const TensorMap
+template <typename Scalar_, int Options_, int Options2_, int NumIndices_,
+          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
+/// type is const TensorCwiseNullaryOp
+template <typename OP, typename RHSExpr>
+struct ConvertToDeviceExpression<const TensorCwiseNullaryOp<OP, RHSExpr>> {
+  using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
+  using Type = const TensorCwiseNullaryOp<OP, RHSPlaceHolderType>;
+};
+
+/// specialisation of the \ref ConvertToDeviceExpression struct when the node
+/// type is TensorCwiseNullaryOp
+template <typename OP, typename RHSExpr>
+struct ConvertToDeviceExpression<TensorCwiseNullaryOp<OP, RHSExpr>> {
+  using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
+  using Type = TensorCwiseNullaryOp<OP, RHSPlaceHolderType>;
+};
+
+/// specialisation of the \ref ConvertToDeviceExpression struct when the node
+/// 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
+/// type is const AssingOP
+template <typename LHSExpr, typename RHSExpr>
+struct ConvertToDeviceExpression<const TensorAssignOp<LHSExpr, RHSExpr>> {
+  using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type;
+  using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
+  using Type = const TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>;
+};
+
+/// specialisation of the \ref ConvertToDeviceExpression struct when the node
+/// type is AssingOP
+template <typename LHSExpr, typename RHSExpr>
+struct ConvertToDeviceExpression<TensorAssignOp<LHSExpr, RHSExpr>> {
+  using LHSPlaceHolderType = typename ConvertToDeviceExpression<LHSExpr>::Type;
+  using RHSPlaceHolderType = typename ConvertToDeviceExpression<RHSExpr>::Type;
+  using Type = TensorAssignOp<LHSPlaceHolderType, RHSPlaceHolderType>;
+};
+
+/// specialisation of the \ref ConvertToDeviceExpression struct when the node
+/// type is const TensorForcedEvalOp
+template <typename Expr>
+struct ConvertToDeviceExpression<const TensorForcedEvalOp<Expr>> {
+  using PlaceHolderType = typename ConvertToDeviceExpression<Expr>::Type;
+  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 TensorSycl
+}  // namespace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX1
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h
new file mode 100644
index 000000000..fe3994175
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExprConstructor.h
@@ -0,0 +1,495 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclExprConstructor.h
+ *
+ * \brief:
+ *  This file re-create an expression on the SYCL device in order
+ *  to use the original tensor evaluator.
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+namespace internal {
+/// this class is used by EvalToOp in order to create an lhs expression which is
+/// a pointer from an accessor on device-only buffer
+template <typename PtrType, size_t N, typename... Params>
+struct EvalToLHSConstructor {
+  PtrType expr;
+  EvalToLHSConstructor(const utility::tuple::Tuple<Params...> &t)
+      : expr((&(*(utility::tuple::get<N>(t).get_pointer())))) {}
+};
+
+/// \struct ExprConstructor is used to reconstruct the expression on the device
+/// and
+/// recreate the expression with MakeGlobalPointer containing the device address
+/// space for the TensorMap pointers used in eval function.
+/// It receives the original expression type, the functor of the node, the tuple
+/// of accessors, and the device expression type to re-instantiate the
+/// expression tree for the device
+template <typename OrigExpr, typename IndexExpr, typename... Params>
+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
+/// 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<
+    TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options2_,
+              MakeGlobalPointer>,
+    Eigen::internal::PlaceHolder<
+        TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>, Options3_,
+                  MakePointer_>,
+        N>,
+    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())) {}
+};
+
+/// specialisation of the \ref ExprConstructor struct when the node type is
+/// TensorCwiseNullaryOp
+template <typename OP, typename OrigRHSExpr, typename RHSExpr,
+          typename... Params>
+struct ExprConstructor<TensorCwiseNullaryOp<OP, OrigRHSExpr>,
+                       TensorCwiseNullaryOp<OP, RHSExpr>, Params...> {
+  using my_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
+  my_type rhsExpr;
+  using Type = 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
+/// const TensorCwiseNullaryOp
+template <typename OP, typename OrigRHSExpr, typename RHSExpr,
+          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
+/// TensorCwiseTernaryOp
+template <typename OP, typename OrigArg1Expr, typename OrigArg2Expr,
+          typename OrigArg3Expr, typename Arg1Expr, typename Arg2Expr,
+          typename Arg3Expr, typename... Params>
+struct ExprConstructor<
+    TensorCwiseTernaryOp<OP, OrigArg1Expr, OrigArg2Expr, OrigArg3Expr>,
+    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 = 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
+/// const TensorCwiseSelectOp
+template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr,
+          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
+/// TensorCwiseSelectOp
+template <typename OrigIfExpr, typename OrigThenExpr, typename OrigElseExpr,
+          typename IfExpr, typename ThenExpr, typename ElseExpr,
+          typename... Params>
+struct ExprConstructor<TensorSelectOp<OrigIfExpr, OrigThenExpr, OrigElseExpr>,
+                       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 =
+      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
+/// TensorAssignOp
+template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr,
+          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
+/// const TensorAssignOp
+template <typename OrigLHSExpr, typename OrigRHSExpr, typename LHSExpr,
+          typename RHSExpr, typename... Params>
+struct ExprConstructor<const TensorAssignOp<OrigLHSExpr, OrigRHSExpr>,
+                       const TensorAssignOp<LHSExpr, RHSExpr>, Params...> {
+  using my_left_type = ExprConstructor<OrigLHSExpr, LHSExpr, Params...>;
+  using my_right_type = ExprConstructor<OrigRHSExpr, RHSExpr, Params...>;
+  using Type = const 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
+/// const TensorEvalToOp
+template <typename OrigExpr, typename Expr, typename... Params>
+struct ExprConstructor<const TensorEvalToOp<OrigExpr, MakeGlobalPointer>,
+                       const TensorEvalToOp<Expr>, Params...> {
+  using my_expr_type = ExprConstructor<OrigExpr, Expr, Params...>;
+  using my_buffer_type =
+      typename TensorEvalToOp<OrigExpr, MakeGlobalPointer>::PointerType;
+  using Type =
+      const TensorEvalToOp<typename my_expr_type::Type, MakeGlobalPointer>;
+  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
+/// TensorEvalToOp
+template <typename OrigExpr, typename Expr, typename... Params>
+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
+/// TensorForcedEvalOp
+template <typename OrigExpr, typename DevExpr, size_t N, typename... Params>
+struct ExprConstructor<
+    const TensorForcedEvalOp<OrigExpr, MakeGlobalPointer>,
+    const Eigen::internal::PlaceHolder<TensorForcedEvalOp<DevExpr>, N>,
+    Params...> {
+  using Type = TensorMap<
+      Tensor<typename TensorForcedEvalOp<DevExpr, MakeGlobalPointer>::Scalar, 1,
+             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())) {}
+};
+
+/// template deduction for \ref ExprConstructor struct
+template <typename OrigExpr, typename IndexExpr, typename FuncD,
+          typename... Params>
+auto createDeviceExpression(FuncD &funcD,
+                            const utility::tuple::Tuple<Params...> &t)
+    -> decltype(ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t)) {
+  return ExprConstructor<OrigExpr, IndexExpr, Params...>(funcD, t);
+}
+}
+}
+}  // namespace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXPR_CONSTRUCTOR_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h
new file mode 100644
index 000000000..cb0ac131d
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractAccessor.h
@@ -0,0 +1,466 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclExtractAccessor.h
+ *
+ * \brief:
+ * ExtractAccessor takes Expression placeHolder expression and the tuple of sycl
+ * buffers as an input. Using pre-order tree traversal, ExtractAccessor
+ * recursively calls itself for its children in the expression tree. The
+ * leaf node in the PlaceHolder expression is nothing but a container preserving
+ * the order of the actual data in the tuple of sycl buffer. By invoking the
+ * extract accessor for the PlaceHolder<N>, an accessor is created for the Nth
+ * buffer in the tuple of buffers. This accessor is then added as an Nth
+ * element in the tuple of accessors. In this case we preserve the order of data
+ * in the expression tree.
+ *
+ * This is the specialisation of extract accessor method for different operation
+ * type in the PlaceHolder expression.
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+namespace internal {
+/// \struct ExtractAccessor: Extract Accessor Class is used to extract the
+/// accessor from a buffer.
+/// Depending on the type of the leaf node we can get a read accessor or a
+/// read_write accessor
+template <typename Evaluator>
+struct ExtractAccessor;
+
+/// specialisation of the \ref ExtractAccessor struct when the node type is
+/// const TensorMap
+template <typename PlainObjectType, int Options_, typename Dev>
+struct ExtractAccessor<
+    TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>> {
+  using actual_type = typename Eigen::internal::remove_all<
+      typename Eigen::internal::traits<PlainObjectType>::Scalar>::type;
+  static inline auto getTuple(
+      cl::sycl::handler& cgh,
+      const TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>
+          eval)
+      -> decltype(utility::tuple::make_tuple(
+          (eval.device()
+               .template get_sycl_accessor<cl::sycl::access::mode::read, 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, true,
+                                         actual_type>(
+                 eval.dimensions().TotalSize(), cgh, eval.derived().data())));
+  }
+};
+
+/// specialisation of the \ref ExtractAccessor struct when the node type is
+/// TensorMap
+template <typename PlainObjectType, int Options_, typename Dev>
+struct ExtractAccessor<
+    TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>> {
+  using actual_type = typename Eigen::internal::remove_all<
+      typename Eigen::internal::traits<PlainObjectType>::Scalar>::type;
+
+  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
+/// const TensorCwiseNullaryOp
+template <typename OP, typename RHSExpr, typename Dev>
+struct ExtractAccessor<
+    TensorEvaluator<const TensorCwiseNullaryOp<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
+/// const TensorCwiseBinaryOp
+template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
+struct ExtractAccessor<
+    TensorEvaluator<const TensorCwiseBinaryOp<OP, LHSExpr, RHSExpr>, Dev>> {
+  static auto getTuple(cl::sycl::handler& cgh,
+                       const TensorEvaluator<
+                           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
+/// TensorCwiseBinaryOp
+template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
+struct ExtractAccessor<
+    TensorEvaluator<TensorCwiseBinaryOp<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
+/// const TensorCwiseTernaryOp
+template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,
+          typename Dev>
+struct ExtractAccessor<TensorEvaluator<
+    const TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> {
+  static auto getTuple(
+      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
+/// TensorCwiseTernaryOp
+template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,
+          typename Dev>
+struct ExtractAccessor<TensorEvaluator<
+    TensorCwiseTernaryOp<OP, Arg1Expr, Arg2Expr, Arg3Expr>, Dev>> {
+  static auto getTuple(
+      cl::sycl::handler& cgh,
+      const TensorEvaluator<
+          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
+/// const TensorCwiseSelectOp
+template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
+struct 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
+/// const TensorAssignOp
+template <typename LHSExpr, typename RHSExpr, typename Dev>
+struct ExtractAccessor<
+    TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>> {
+  static auto getTuple(
+      cl::sycl::handler& cgh,
+      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
+/// TensorAssignOp
+template <typename LHSExpr, typename RHSExpr, typename Dev>
+struct ExtractAccessor<TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>> {
+  static auto getTuple(
+      cl::sycl::handler& cgh,
+      const TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev> eval)
+      -> decltype(utility::tuple::append(
+          ExtractAccessor<TensorEvaluator<LHSExpr, Dev>>::getTuple(
+              eval.left_impl()),
+          ExtractAccessor<TensorEvaluator<RHSExpr, Dev>>::getTuple(
+              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
+/// const TensorForcedEvalOp
+template <typename Expr, typename Dev>
+struct ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> {
+  using actual_type =
+      typename Eigen::internal::remove_all<typename TensorEvaluator<
+          const TensorForcedEvalOp<Expr>, Dev>::CoeffReturnType>::type;
+  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
+/// TensorForcedEvalOp
+template <typename Expr, typename Dev>
+struct ExtractAccessor<TensorEvaluator<TensorForcedEvalOp<Expr>, Dev>>
+    : ExtractAccessor<TensorEvaluator<const TensorForcedEvalOp<Expr>, Dev>> {};
+
+/// specialisation of the \ref ExtractAccessor struct when the node type is
+/// const TensorEvalToOp
+template <typename Expr, typename Dev>
+struct ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev>> {
+  using actual_type =
+      typename Eigen::internal::remove_all<typename TensorEvaluator<
+          const TensorEvalToOp<Expr>, Dev>::CoeffReturnType>::type;
+
+  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
+/// TensorEvalToOp
+template <typename Expr, typename Dev>
+struct ExtractAccessor<TensorEvaluator<TensorEvalToOp<Expr>, Dev>>
+    : ExtractAccessor<TensorEvaluator<const TensorEvalToOp<Expr>, Dev>> {};
+
+/// template deduction for \ref ExtractAccessor
+template <typename Evaluator>
+auto createTupleOfAccessors(cl::sycl::handler& cgh, const Evaluator& expr)
+    -> decltype(ExtractAccessor<Evaluator>::getTuple(cgh, expr)) {
+  return ExtractAccessor<Evaluator>::getTuple(cgh, expr);
+}
+}
+}
+}
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_ACCESSOR_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h
new file mode 100644
index 000000000..f69c5afcb
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclExtractFunctors.h
@@ -0,0 +1,313 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclextractFunctors.h
+ *
+ * \brief:
+ *  Used to extract all the functors allocated to each node of the expression
+*tree.
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+namespace internal {
+/// \struct FunctorExtractor:  This struct is used to extract the functors
+/// constructed on
+/// the host-side, to pack them and reuse them in reconstruction of the
+/// expression on the device.
+/// We have to do that as in Eigen the functors are not stateless so we cannot
+/// re-instantiate them on the device.
+/// We have to pass whatever instantiated to the device.
+template <typename Evaluator>
+struct FunctorExtractor;
+
+/// 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& dimensions() const { return m_dimensions; }
+  FunctorExtractor(
+      const TensorEvaluator<TensorMap<PlainObjectType, Options_>, Dev>& expr)
+      : m_dimensions(expr.dimensions()) {}
+};
+
+/// specialisation of the \ref FunctorExtractor struct when the node type is
+/// const TensorMap
+template <typename PlainObjectType, int Options_, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<const TensorMap<PlainObjectType, Options_>, Dev>> {
+  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;
+  FunctorExtractor(
+      TensorEvaluator<TensorCwiseNullaryOp<OP, RHSExpr>, Dev>& expr)
+      : rhsExpr(expr.impl()), func(expr.functor()) {}
+};
+
+/// specialisation of the \ref FunctorExtractor struct when the node type is
+/// const TensorCwiseNullaryOp
+template <typename OP, typename RHSExpr, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>> {
+  FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
+  OP func;
+  FunctorExtractor(
+      const TensorEvaluator<const TensorCwiseNullaryOp<OP, RHSExpr>, Dev>& expr)
+      : 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
+/// const TensorCwiseBinaryOp
+template <typename OP, typename LHSExpr, typename RHSExpr, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<const TensorCwiseBinaryOp<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
+/// const TensorCwiseTernaryOp
+template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,
+          typename Dev>
+struct FunctorExtractor<TensorEvaluator<
+    const 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<
+                   const 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
+/// TensorCwiseTernaryOp
+template <typename OP, typename Arg1Expr, typename Arg2Expr, typename Arg3Expr,
+          typename Dev>
+struct FunctorExtractor<TensorEvaluator<
+    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
+/// const TensorCwiseSelectOp
+template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<const TensorSelectOp<IfExpr, ThenExpr, ElseExpr>, Dev>> {
+  FunctorExtractor<TensorEvaluator<IfExpr, Dev>> ifExpr;
+  FunctorExtractor<TensorEvaluator<ThenExpr, Dev>> thenExpr;
+  FunctorExtractor<TensorEvaluator<ElseExpr, Dev>> elseExpr;
+  FunctorExtractor(const TensorEvaluator<
+                   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
+/// TensorCwiseSelectOp
+template <typename IfExpr, typename ThenExpr, typename ElseExpr, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<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
+/// TensorAssignOp
+template <typename LHSExpr, typename RHSExpr, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<TensorAssignOp<LHSExpr, RHSExpr>, Dev>> {
+  FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr;
+  FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
+  FunctorExtractor(
+      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
+/// const TensorAssignOp
+template <typename LHSExpr, typename RHSExpr, typename Dev>
+struct FunctorExtractor<
+    TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>> {
+  FunctorExtractor<TensorEvaluator<LHSExpr, Dev>> lhsExpr;
+  FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
+  FunctorExtractor(
+      const TensorEvaluator<const TensorAssignOp<LHSExpr, RHSExpr>, Dev>& expr)
+      : lhsExpr(expr.left_impl()), rhsExpr(expr.right_impl()) {}
+};
+
+/// specialisation of the \ref FunctorExtractor struct when the node type is
+/// TensorEvalToOp
+template <typename RHSExpr, typename Dev>
+struct FunctorExtractor<TensorEvaluator<TensorEvalToOp<RHSExpr>, Dev>> {
+  FunctorExtractor<TensorEvaluator<RHSExpr, Dev>> rhsExpr;
+  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 <typename Evaluator>
+auto extractFunctors(const Evaluator& evaluator)
+    -> FunctorExtractor<Evaluator> {
+  return FunctorExtractor<Evaluator>(evaluator);
+}
+}  // namespace internal
+}  // namespace TensorSycl
+}  // namespace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_EXTRACT_FUNCTORS_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h
new file mode 100644
index 000000000..77e0e15e1
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclLeafCount.h
@@ -0,0 +1,188 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclLeafCount.h
+ *
+ * \brief:
+ *  The leaf count used the pre-order expression tree traverse in order to name
+ *  count the number of leaf nodes in the expression
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+namespace internal {
+/// \brief LeafCount used to counting terminal nodes. The total number of
+/// leaf nodes is used by MakePlaceHolderExprHelper to find the order
+/// of the leaf node in a expression tree at compile time.
+template <typename Expr>
+struct LeafCount;
+
+/// specialisation of the \ref LeafCount struct when the node type is const
+/// TensorMap
+template <typename PlainObjectType, int Options_,
+          template <class> class MakePointer_>
+struct LeafCount<const TensorMap<PlainObjectType, Options_, MakePointer_>> {
+  static const size_t Count = 1;
+};
+
+/// specialisation of the \ref LeafCount struct when the node type is TensorMap
+template <typename PlainObjectType, int Options_,
+          template <class> class 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
+/// TensorCwiseNullaryOp
+template <typename OP, typename RHSExpr>
+struct LeafCount<const TensorCwiseNullaryOp<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<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>
+struct LeafCount<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 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
+/// TensorAssignOp
+template <typename LHSExpr, typename RHSExpr>
+struct LeafCount<TensorAssignOp<LHSExpr, RHSExpr>> {
+  static const size_t Count =
+      LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count;
+};
+
+/// specialisation of the \ref LeafCount struct when the node type is const
+/// TensorAssignOp
+template <typename LHSExpr, typename RHSExpr>
+struct LeafCount<const TensorAssignOp<LHSExpr, RHSExpr>> {
+  static const size_t Count =
+      LeafCount<LHSExpr>::Count + LeafCount<RHSExpr>::Count;
+};
+
+/// specialisation of the \ref LeafCount struct when the node type is const
+/// TensorForcedEvalOp
+template <typename Expr>
+struct LeafCount<const TensorForcedEvalOp<Expr>> {
+  static const size_t Count = 1;
+};
+
+/// specialisation of the \ref LeafCount struct when the node type is
+/// TensorForcedEvalOp
+template <typename Expr>
+struct LeafCount<TensorForcedEvalOp<Expr>> {
+  static const size_t Count = 1;
+};
+
+/// specialisation of the \ref LeafCount struct when the node type is const
+/// TensorEvalToOp
+template <typename Expr>
+struct LeafCount<const TensorEvalToOp<Expr>> {
+  static const size_t Count = 1 + LeafCount<Expr>::Count;
+};
+
+/// specialisation of the \ref LeafCount struct when the node type is
+/// TensorEvalToOp
+template <typename Expr>
+struct LeafCount<TensorEvalToOp<Expr>> {
+  static const size_t Count = 1 + LeafCount<Expr>::Count;
+};
+}
+}
+}  // namespace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_LEAF_COUNT_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h
new file mode 100644
index 000000000..87995a25e
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolder.h
@@ -0,0 +1,151 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclPlaceHolder.h
+ *
+ * \brief:
+ *  The PlaceHolder expression are nothing but a container preserving
+ *  the order of actual data in the tuple of sycl buffer.
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP
+
+namespace Eigen {
+namespace internal {
+/// \struct PlaceHolder
+/// \brief PlaceHolder is used to replace the \ref TensorMap in the expression
+/// tree.
+/// PlaceHolder contains the order of the leaf node in the expression tree.
+template <typename Scalar, size_t N>
+struct PlaceHolder {
+  static constexpr size_t I = N;
+  using Type = Scalar;
+};
+
+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
+template <typename PlainObjectType, int Options_,
+          template <class> class Makepointer_, size_t N>
+struct PlaceHolder<TensorMap<PlainObjectType, Options_, Makepointer_>, N> {
+  static constexpr size_t I = N;
+
+  using Type = 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::Packet Packet;
+  typedef typename Type::RealScalar RealScalar;
+  typedef typename Type::CoeffReturnType CoeffReturnType;
+  typedef typename Base::PacketReturnType PacketReturnType;
+};
+
+/// specialisation of the traits struct for PlaceHolder
+template <typename PlainObjectType, int Options_,
+          template <class> class Makepointer_, size_t N>
+struct traits<
+    PlaceHolder<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,
+  };
+};
+
+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 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 namespoace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h
new file mode 100644
index 000000000..dbd7a8544
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclPlaceHolderExpr.h
@@ -0,0 +1,293 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclPlaceHolderExpr.h
+ *
+ * \brief:
+ *  This is the specialisation of the placeholder expression based on the
+ * operation type
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+namespace internal {
+/// \sttruct PlaceHolderExpression
+/// \brief it is used to create the PlaceHolder expression. The PlaceHolder
+/// expression is a copy of expression type in which the TensorMap of the has
+/// been replaced with PlaceHolder.
+template <typename Expr, size_t N>
+struct PlaceHolderExpression;
+
+/// specialisation of the \ref PlaceHolderExpression when the node is TensorMap
+template <typename Scalar_, int Options_, int Options2_, int NumIndices_,
+          typename IndexType_, template <class> class MakePointer_, size_t N>
+struct PlaceHolderExpression<
+    Eigen::TensorMap<Eigen::Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
+                     Options2_, MakePointer_>,
+    N> {
+  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
+/// TensorMap
+template <typename Scalar_, int Options_, int Options2_, int NumIndices_,
+          typename IndexType_, template <class> class MakePointer_, size_t N>
+struct PlaceHolderExpression<
+    const Eigen::TensorMap<Tensor<Scalar_, NumIndices_, Options_, IndexType_>,
+                           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
+/// TensorCwiseNullaryOp
+template <typename OP, typename RHSExpr, size_t N>
+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
+/// TensorForcedEvalOp
+template <typename Expr, size_t N>
+struct PlaceHolderExpression<TensorForcedEvalOp<Expr>, N> {
+  using Type = Eigen::internal::PlaceHolder<TensorForcedEvalOp<Expr>, N>;
+};
+
+/// specialisation of the \ref PlaceHolderExpression when the node is const
+/// TensorEvalToOp
+template <typename Expr, size_t N>
+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
+/// TensorEvalToOp
+template <typename Expr, size_t N>
+struct PlaceHolderExpression<TensorEvalToOp<Expr>, N> {
+  static const size_t RHSLeafCount = LeafCount<Expr>::Count;
+
+  using RHSPlaceHolderType = typename PlaceHolderExpression<Expr, N>::Type;
+
+  using Type = TensorEvalToOp<RHSPlaceHolderType>;
+};
+
+/// template deduction for \ref PlaceHolderExpression struct
+template <typename Expr>
+struct createPlaceHolderExpression {
+  static const size_t TotalLeaves = LeafCount<Expr>::Count;
+  using Type = typename PlaceHolderExpression<Expr, TotalLeaves - 1>::Type;
+};
+}
+}
+}  // namespace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORYSYCL_PLACEHOLDER_EXPR_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h
new file mode 100644
index 000000000..3758d46a0
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclRun.h
@@ -0,0 +1,84 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Cummins Chris PhD student at The University of Edinburgh.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensorSyclRun.h
+ *
+ * \brief:
+ * Schedule_kernel invoke an specialised version of kernel struct. The
+ * specialisation is based on the data dimension in sycl buffer
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP
+
+namespace Eigen {
+namespace TensorSycl {
+/// The run function in tensor sycl convert the expression tree to a buffer
+/// based expression tree;
+/// creates the expression tree for the device with accessor to buffers;
+/// construct the kernel and submit it to the sycl queue.
+template <typename Expr, typename Dev>
+void run(Expr &expr, Dev &dev) {
+  Eigen::TensorEvaluator<Expr, Dev> evaluator(expr, dev);
+  const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);
+  if (needs_assign) {
+    using PlaceHolderExpr =
+        typename internal::createPlaceHolderExpression<Expr>::Type;
+    auto functors = internal::extractFunctors(evaluator);
+
+    dev.m_queue.submit([&](cl::sycl::handler &cgh) {
+
+      // create a tuple of accessors from Evaluator
+      auto tuple_of_accessors =
+          internal::createTupleOfAccessors<decltype(evaluator)>(cgh, evaluator);
+      const auto range =
+          utility::tuple::get<0>(tuple_of_accessors).get_range()[0];
+
+      size_t outTileSize = range;
+      if (range > 64) outTileSize = 64;
+      size_t yMode = range % outTileSize;
+      int yRange = static_cast<int>(range);
+      if (yMode != 0) yRange += (outTileSize - yMode);
+
+      // run the kernel
+      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) {
+            using DevExpr =
+                typename internal::ConvertToDeviceExpression<Expr>::Type;
+
+            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();
+  }
+  evaluator.cleanup();
+}
+}  // namespace TensorSycl
+}  // namespace Eigen
+
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_SYCLRUN_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h
new file mode 100644
index 000000000..8b9fc52c4
--- /dev/null
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorSyclTuple.h
@@ -0,0 +1,264 @@
+// This file is part of Eigen, a lightweight C++ template library
+// for linear algebra.
+//
+// Mehdi Goli    Codeplay Software Ltd.
+// Ralph Potter  Codeplay Software Ltd.
+// Luke Iwanski  Codeplay Software Ltd.
+// Contact: <eigen@codeplay.com>
+//
+// This Source Code Form is subject to the terms of the Mozilla
+// Public License v. 2.0. If a copy of the MPL was not distributed
+// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
+
+/*****************************************************************
+ * TensroSyclTuple.h
+ *
+ * \brief:
+ *  Minimal implementation of std::tuple that can be used inside a SYCL kernel.
+ *
+*****************************************************************/
+
+#ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP
+#define UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP
+namespace utility {
+namespace tuple {
+/// \struct EnableIf
+/// \brief The EnableIf struct is used to statically define type based on the
+/// condition.
+template <bool, typename T = void>
+struct EnableIf {};
+/// \brief specialisation of the \ref EnableIf when the condition is true
+template <typename T>
+struct EnableIf<true, T> {
+  typedef T type;
+};
+
+/// \struct Tuple
+/// \brief is a fixed-size collection of heterogeneous values
+/// \ztparam Ts... - the types of the elements that the tuple stores.
+/// Empty list is supported.
+template <class... Ts>
+struct Tuple {};
+
+/// \brief specialisation of the \ref Tuple class when the tuple has at least
+/// one element.
+/// \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.
+template <class T, class... Ts>
+struct Tuple<T, Ts...> {
+  Tuple(T t, Ts... ts) : head(t), tail(ts...) {}
+
+  T head;
+  Tuple<Ts...> tail;
+};
+
+/// \struct ElemTypeHolder
+/// \brief ElemTypeHolder class is used to specify the types of the
+/// elements inside the tuple
+/// \tparam size_t the number of elements inside the tuple
+/// \tparam class the tuple class
+template <size_t, class>
+struct ElemTypeHolder;
+
+/// \brief specialisation of the \ref ElemTypeHolder class when the number
+/// elements inside the tuple is 1
+template <class T, class... Ts>
+struct ElemTypeHolder<0, Tuple<T, Ts...>> {
+  typedef T type;
+};
+
+/// \brief specialisation of the \ref ElemTypeHolder class when the number of
+/// elements inside the tuple is bigger than 1. It recursively call itself to
+/// detect the type of each 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 K is the Kth element in the tuple
+template <size_t k, class T, class... Ts>
+struct ElemTypeHolder<k, Tuple<T, Ts...>> {
+  typedef typename ElemTypeHolder<k - 1, Tuple<Ts...>>::type type;
+};
+
+/// get
+/// \brief Extracts the first element from the tuple.
+/// 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 tuple whose contents to extract
+/// \return  typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type
+template <size_t k, class... Ts>
+typename EnableIf<k == 0,
+                  typename ElemTypeHolder<0, Tuple<Ts...>>::type &>::type
+get(Tuple<Ts...> &t) {
+  return t.head;
+}
+/// get
+/// \brief Extracts the Kth element from the tuple.
+/// \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 tuple whose contents to extract
+/// \return  typename ElemTypeHolder<K, Tuple<Ts...>>::type &>::type
+template <size_t k, class T, class... Ts>
+typename EnableIf<k != 0,
+                  typename ElemTypeHolder<k, Tuple<T, Ts...>>::type &>::type
+get(Tuple<T, Ts...> &t) {
+  return get<k - 1>(t.tail);
+}
+
+/// 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
+/// \brief Creates a tuple object, deducing the target type from the types of
+/// arguments.
+/// \tparam Args the type of the arguments to construct the tuple from
+/// \param args zero or more arguments to construct the tuple from
+/// \return Tuple<Args...>
+template <typename... Args>
+Tuple<Args...> make_tuple(Args... args) {
+  return Tuple<Args...>(args...);
+}
+
+/// size
+/// \brief Provides access to the number of elements in a tuple as a
+/// compile-time constant expression.
+/// \tparam Args the type of the arguments to construct the tuple from
+/// \return size_t
+template <typename... Args>
+static constexpr size_t size(Tuple<Args...> &) {
+  return sizeof...(Args);
+}
+
+/// \struct Index_list
+/// \brief Creates a list of index from the elements in the tuple
+/// \tparam Is... a list of index from [0 to sizeof...(tuple elements))
+template <size_t... Is>
+struct Index_list {};
+
+/// \struct RangeBuilder
+/// \brief Collects internal details for generating index ranges [MIN, MAX)
+/// Declare primary template for index range builder
+/// \tparam MIN is the starting index in the tuple
+/// \tparam N represents sizeof..(elements)- sizeof...(Is)
+/// \tparam Is... are the list of generated index so far
+template <size_t MIN, size_t N, size_t... Is>
+struct RangeBuilder;
+
+/// \brief base Step: Specialisation of the \ref RangeBuilder when the
+/// MIN==MAX. In this case the Is... is [0 to sizeof...(tuple elements))
+/// \tparam MIN is the starting index of the tuple
+/// \tparam Is is [0 to sizeof...(tuple elements))
+template <size_t MIN, size_t... Is>
+struct RangeBuilder<MIN, MIN, Is...> {
+  typedef Index_list<Is...> type;
+};
+
+/// Induction step: Specialisation of the RangeBuilder class when N!=MIN
+/// in this case we are recursively subtracting the N by one and adding one
+/// index to Is... list until MIN==N
+/// \tparam MIN is the starting index in the tuple
+/// \tparam N represents sizeof..(elements)- sizeof...(Is)
+/// \tparam Is... are the list of generated index so far
+template <size_t MIN, size_t N, size_t... Is>
+struct RangeBuilder : public RangeBuilder<MIN, N - 1, N - 1, Is...> {};
+
+/// \brief IndexRange that returns a [MIN, MAX) index range
+/// \tparam MIN is the starting index in the tuple
+/// \tparam MAX is the size of the tuple
+template <size_t MIN, size_t MAX>
+using Index_range = typename RangeBuilder<MIN, MAX>::type;
+
+/// append_impl
+/// \brief unpacking the elements of the input tuple t and creating a new tuple
+/// by adding element a at the end of it.
+/// \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 I... is the list of index from [0 to sizeof...(t))
+/// \param t the tuple on which we want to append a.
+/// \param a the new elements going to be added to the tuple
+/// \return Tuple<Args..., T>
+template <typename... Args, typename T, size_t... I>
+Tuple<Args..., T> append_impl(utility::tuple::Tuple<Args...> t, T a,
+                              utility::tuple::Index_list<I...>) {
+  return utility::tuple::make_tuple(get<I>(t)..., a);
+}
+
+/// append
+/// \brief the deduction function for \ref append_impl that automatically
+/// generate the \ref Index_range
+/// \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
+/// \param t the tuple on which we want to append a.
+/// \param a the new elements going to be added to the tuple
+/// \return Tuple<Args..., T>
+template <typename... Args, typename T>
+Tuple<Args..., T> append(Tuple<Args...> t, T a) {
+  return utility::tuple::append_impl(
+      t, a, utility::tuple::Index_range<0, sizeof...(Args)>());
+}
+
+/// append_impl
+/// \brief This is an specialised of \ref append_impl when we want to
+/// concatenate
+/// tuple t2 at the end of the tuple t1. Here we unpack both tuples, generate
+/// the
+/// Index_range for each of them and create an output tuple T that contains both
+/// elements of t1 and t2.
+/// \tparam Args1... the type of the elements inside the tuple t1
+/// \tparam Args2... the type of the elements inside the tuple t2
+/// \tparam I1... is the list of index from [0 to sizeof...(t1))
+/// \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 t2 is the tuple that is going to be added on t1.
+/// \return Tuple<Args1..., Args2...>
+template <typename... Args1, typename... Args2, size_t... I1, size_t... I2>
+Tuple<Args1..., Args2...> append_impl(utility::tuple::Tuple<Args1...> t1,
+                                      utility::tuple::Tuple<Args2...> 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
+/// \brief deduction function for \ref append_impl when we are appending tuple
+/// t1 by tuple t2. In this case the \ref Index_range for both tuple are
+/// automatically generated.
+/// \tparam Args1... the type of the elements inside the tuple t1
+/// \tparam Args2... the type of the elements inside the tuple 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.
+/// \return Tuple<Args1..., Args2...>
+template <typename... Args1, typename... Args2>
+Tuple<Args1..., Args2...> append(utility::tuple::Tuple<Args1...> t1,
+                                 utility::tuple::Tuple<Args2...> t2) {
+  return utility::tuple::append_impl(
+      t1, t2, utility::tuple::Index_range<0, sizeof...(Args1)>(),
+      utility::tuple::Index_range<0, sizeof...(Args2)>());
+}
+}  // tuple
+}  // utility
+#endif  // UNSUPPORTED_EIGEN_CXX11_SRC_TENSORSYCL_TUPLE_HPP
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h b/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h
index b7597b3a5..62c5caf6c 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorTraits.h
@@ -56,11 +56,12 @@ struct traits<Tensor<Scalar_, NumIndices_, Options_, IndexType_> >
     Options = Options_,
     Flags = compute_tensor_flags<Scalar_, Options_>::ret | (is_const<Scalar_>::value ? 0 : LvalueBit)
   };
+  template <class T> using MakePointer = MakePointer<T>;
 };
 
 
-template<typename Scalar_, typename Dimensions, int Options_, typename IndexType_>
-struct traits<TensorFixedSize<Scalar_, Dimensions, Options_, IndexType_> >
+template<typename Scalar_, typename Dimensions, int Options_, typename IndexType_, template <class> class MakePointer_>
+struct traits<TensorFixedSize<Scalar_, Dimensions, Options_, IndexType_, MakePointer_> >
 {
   typedef Scalar_ Scalar;
   typedef Dense StorageKind;
@@ -71,11 +72,12 @@ struct traits<TensorFixedSize<Scalar_, Dimensions, Options_, IndexType_> >
     Options = Options_,
     Flags = compute_tensor_flags<Scalar_, Options_>::ret | (is_const<Scalar_>::value ? 0: LvalueBit)
   };
+  template <class T> using MakePointer = MakePointer_<T>;
 };
 
 
-template<typename PlainObjectType, int Options_>
-struct traits<TensorMap<PlainObjectType, Options_> >
+template<typename PlainObjectType, int Options_ , template <class> class MakePointer_>
+struct traits<TensorMap<PlainObjectType, Options_ , MakePointer_> >
   : public traits<PlainObjectType>
 {
   typedef traits<PlainObjectType> BaseTraits;
@@ -88,6 +90,7 @@ struct traits<TensorMap<PlainObjectType, Options_> >
     Options = Options_,
     Flags = BaseTraits::Flags
   };
+  template <class T> using MakePointer = MakePointer_<T>;
 };
 
 template<typename PlainObjectType>