Please, help us to better know about our user community by answering the following short survey: https://forms.gle/wpyrxWi18ox9Z5ae9
TensorMorphing.h
1 // This file is part of Eigen, a lightweight C++ template library
2 // for linear algebra.
3 //
4 // Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com>
5 //
6 // This Source Code Form is subject to the terms of the Mozilla
7 // Public License v. 2.0. If a copy of the MPL was not distributed
8 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
9 
10 #ifndef EIGEN_CXX11_TENSOR_TENSOR_MORPHING_H
11 #define EIGEN_CXX11_TENSOR_TENSOR_MORPHING_H
12 
13 namespace Eigen {
14 
22 namespace internal {
23 template<typename NewDimensions, typename XprType>
24 struct traits<TensorReshapingOp<NewDimensions, XprType> > : public traits<XprType>
25 {
26  typedef typename XprType::Scalar Scalar;
27  typedef traits<XprType> XprTraits;
28  typedef typename XprTraits::StorageKind StorageKind;
29  typedef typename XprTraits::Index Index;
30  typedef typename XprType::Nested Nested;
31  typedef typename remove_reference<Nested>::type _Nested;
32  static const int NumDimensions = array_size<NewDimensions>::value;
33  static const int Layout = XprTraits::Layout;
34  typedef typename XprTraits::PointerType PointerType;
35 };
36 
37 template<typename NewDimensions, typename XprType>
38 struct eval<TensorReshapingOp<NewDimensions, XprType>, Eigen::Dense>
39 {
40  typedef const TensorReshapingOp<NewDimensions, XprType>EIGEN_DEVICE_REF type;
41 };
42 
43 template<typename NewDimensions, typename XprType>
44 struct nested<TensorReshapingOp<NewDimensions, XprType>, 1, typename eval<TensorReshapingOp<NewDimensions, XprType> >::type>
45 {
46  typedef TensorReshapingOp<NewDimensions, XprType> type;
47 };
48 
49 } // end namespace internal
50 
51 
52 
53 template<typename NewDimensions, typename XprType>
54 class TensorReshapingOp : public TensorBase<TensorReshapingOp<NewDimensions, XprType>, WriteAccessors>
55 {
56  public:
57  typedef TensorBase<TensorReshapingOp<NewDimensions, XprType>, WriteAccessors> Base;
58  typedef typename Eigen::internal::traits<TensorReshapingOp>::Scalar Scalar;
59  typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
60  typedef typename Eigen::internal::nested<TensorReshapingOp>::type Nested;
61  typedef typename Eigen::internal::traits<TensorReshapingOp>::StorageKind StorageKind;
62  typedef typename Eigen::internal::traits<TensorReshapingOp>::Index Index;
63 
64  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorReshapingOp(const XprType& expr, const NewDimensions& dims)
65  : m_xpr(expr), m_dims(dims) {}
66 
67  EIGEN_DEVICE_FUNC
68  const NewDimensions& dimensions() const { return m_dims; }
69 
70  EIGEN_DEVICE_FUNC
71  const typename internal::remove_all<typename XprType::Nested>::type&
72  expression() const { return m_xpr; }
73 
74  EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorReshapingOp)
75 
76  protected:
77  typename XprType::Nested m_xpr;
78  const NewDimensions m_dims;
79 };
80 
81 
82 // Eval as rvalue
83 template<typename NewDimensions, typename ArgType, typename Device>
84 struct TensorEvaluator<const TensorReshapingOp<NewDimensions, ArgType>, Device>
85 {
86  typedef TensorReshapingOp<NewDimensions, ArgType> XprType;
87  typedef NewDimensions Dimensions;
88 
89  typedef typename XprType::Index Index;
90  typedef typename XprType::Scalar Scalar;
91  typedef typename XprType::CoeffReturnType CoeffReturnType;
92  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
93  typedef StorageMemory<CoeffReturnType, Device> Storage;
94  typedef typename Storage::Type EvaluatorPointerType;
95  typedef StorageMemory<typename internal::remove_const<CoeffReturnType>::type, Device> ConstCastStorage;
96 
97  static const int NumOutputDims = internal::array_size<Dimensions>::value;
98  static const int NumInputDims = internal::array_size<typename TensorEvaluator<ArgType, Device>::Dimensions>::value;
99 
100  enum ReshapingKind {
101  // We do not use layout information to determine reshaping kind.
102  // Depending on the layout `N` can be inner or outer dimension.
103  OneByN = 0, // expr.reshape(1, N)
104  NByOne = 1, // expr.reshape(N, 1)
105  Runtime = 2 // Reshape dimensions are dynamic (specified at runtime).
106  };
107 
108  // clang-format off
109  static const ReshapingKind kind =
110 #if defined(EIGEN_HAS_INDEX_LIST)
111  (NumOutputDims == 2 && internal::index_statically_eq<NewDimensions>(/*index=*/0, /*value=*/1)) ? OneByN
112  : (NumOutputDims == 2 && internal::index_statically_eq<NewDimensions>(/*index=*/1, /*value=*/1)) ? NByOne
113  : Runtime;
114 #else
115  Runtime;
116 #endif
117  // clang-format on
118 
119  enum {
120  IsAligned = TensorEvaluator<ArgType, Device>::IsAligned,
121  PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
122  // For trivial reshapes with raw access to underlying data we will provide
123  // zero overhead block access.
124  // TODO(ezhulenev): Consider adding block access without raw access?
125  BlockAccess = TensorEvaluator<ArgType, Device>::RawAccess &&
126  NumInputDims > 0 && NumOutputDims > 0,
127  PreferBlockAccess = false,
128  Layout = TensorEvaluator<ArgType, Device>::Layout,
129  CoordAccess = false, // to be implemented
130  RawAccess = TensorEvaluator<ArgType, Device>::RawAccess
131  };
132 
133  typedef typename internal::remove_const<Scalar>::type ScalarNoConst;
134 
135  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
136  typedef internal::TensorBlockDescriptor<NumOutputDims, Index> TensorBlockDesc;
137  typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;
138 
139  typedef
140  typename internal::TensorMaterializedBlock<ScalarNoConst, NumOutputDims,
141  Layout, Index>
142  TensorBlock;
143  //===--------------------------------------------------------------------===//
144 
145  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
146  : m_impl(op.expression(), device), m_dimensions(op.dimensions())
147  {
148  // The total size of the reshaped tensor must be equal to the total size
149  // of the input tensor.
150  eigen_assert(internal::array_prod(m_impl.dimensions()) == internal::array_prod(op.dimensions()));
151  }
152 
153  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dimensions; }
154 
155 #ifdef EIGEN_USE_THREADS
156  template <typename EvalSubExprsCallback>
157  EIGEN_STRONG_INLINE void evalSubExprsIfNeededAsync(
158  EvaluatorPointerType data, EvalSubExprsCallback done) {
159  m_impl.evalSubExprsIfNeededAsync(data, std::move(done));
160  }
161 #endif
162 
163  EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(EvaluatorPointerType data) {
164  return m_impl.evalSubExprsIfNeeded(data);
165  }
166  EIGEN_STRONG_INLINE void cleanup() {
167  m_impl.cleanup();
168  }
169 
170  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const
171  {
172  return m_impl.coeff(index);
173  }
174 
175  template<int LoadMode>
176  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const
177  {
178  return m_impl.template packet<LoadMode>(index);
179  }
180 
181  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const {
182  return m_impl.costPerCoeff(vectorized);
183  }
184 
185  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
186  internal::TensorBlockResourceRequirements getResourceRequirements() const {
187  return internal::TensorBlockResourceRequirements::any();
188  }
189 
190  // required in block(OutputTensorBlock* output_block) const
191  // For C++03 compatibility this must be defined outside the method
192  struct BlockIteratorState {
193  Index stride;
194  Index span;
195  Index size;
196  Index count;
197  };
198 
199  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorBlock
200  block(TensorBlockDesc& desc, TensorBlockScratch& scratch,
201  bool /*root_of_expr_ast*/ = false) const {
202  eigen_assert(m_impl.data() != NULL);
203  eigen_assert((kind == Runtime) ||
204  (kind == OneByN && desc.dimensions()[0] == 1) ||
205  (kind == NByOne && desc.dimensions()[1] == 1));
206 
207  if (kind == OneByN || kind == NByOne) {
208  // We can guarantee at compile time that block is just a contiguous slice
209  // of the underlying expression memory buffer.
210  return TensorBlock(internal::TensorBlockKind::kView,
211  m_impl.data() + desc.offset(), desc.dimensions());
212  } else {
213  // This will do additional runtime checks, and in the end it might be also
214  // a view, or it might be a block materialized in the temporary buffer.
215  return TensorBlock::materialize(m_impl.data(), m_dimensions, desc,
216  scratch);
217  }
218  }
219 
220  EIGEN_DEVICE_FUNC typename Storage::Type data() const {
221  return constCast(m_impl.data());
222  }
223 
224  EIGEN_DEVICE_FUNC const TensorEvaluator<ArgType, Device>& impl() const { return m_impl; }
225 
226  #ifdef EIGEN_USE_SYCL
227  // binding placeholder accessors to a command group handler for SYCL
228  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void bind(cl::sycl::handler &cgh) const {
229  m_impl.bind(cgh);
230  }
231  #endif
232  protected:
233  TensorEvaluator<ArgType, Device> m_impl;
234  NewDimensions m_dimensions;
235 };
236 
237 
238 // Eval as lvalue
239 template<typename NewDimensions, typename ArgType, typename Device>
240  struct TensorEvaluator<TensorReshapingOp<NewDimensions, ArgType>, Device>
241  : public TensorEvaluator<const TensorReshapingOp<NewDimensions, ArgType>, Device>
242 
243 {
244  typedef TensorEvaluator<const TensorReshapingOp<NewDimensions, ArgType>, Device> Base;
245  typedef TensorReshapingOp<NewDimensions, ArgType> XprType;
246  typedef NewDimensions Dimensions;
247 
248  enum {
249  IsAligned = TensorEvaluator<ArgType, Device>::IsAligned,
250  PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
251  BlockAccess = TensorEvaluator<ArgType, Device>::RawAccess,
252  PreferBlockAccess = false,
253  Layout = TensorEvaluator<ArgType, Device>::Layout,
254  CoordAccess = false, // to be implemented
255  RawAccess = TensorEvaluator<ArgType, Device>::RawAccess
256  };
257 
258  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
259  : Base(op, device)
260  { }
261 
262  typedef typename XprType::Index Index;
263  typedef typename XprType::Scalar Scalar;
264  typedef typename XprType::CoeffReturnType CoeffReturnType;
265  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
266 
267  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
268  typedef internal::TensorBlockDescriptor<TensorEvaluator::NumOutputDims, Index>
269  TensorBlockDesc;
270  //===--------------------------------------------------------------------===//
271 
272  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType& coeffRef(Index index)
273  {
274  return this->m_impl.coeffRef(index);
275  }
276 
277  template <int StoreMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
278  void writePacket(Index index, const PacketReturnType& x)
279  {
280  this->m_impl.template writePacket<StoreMode>(index, x);
281  }
282 
283  template <typename TensorBlock>
284  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void writeBlock(
285  const TensorBlockDesc& desc, const TensorBlock& block) {
286  assert(this->m_impl.data() != NULL);
287 
288  typedef typename TensorBlock::XprType TensorBlockExpr;
289  typedef internal::TensorBlockAssignment<
290  Scalar, TensorEvaluator::NumOutputDims, TensorBlockExpr, Index>
291  TensorBlockAssign;
292 
293  TensorBlockAssign::Run(
294  TensorBlockAssign::target(desc.dimensions(),
295  internal::strides<Layout>(this->dimensions()),
296  this->m_impl.data(), desc.offset()),
297  block.expr());
298  }
299 };
300 
301 
309 namespace internal {
310 template<typename StartIndices, typename Sizes, typename XprType>
311 struct traits<TensorSlicingOp<StartIndices, Sizes, XprType> > : public traits<XprType>
312 {
313  typedef typename XprType::Scalar Scalar;
314  typedef traits<XprType> XprTraits;
315  typedef typename XprTraits::StorageKind StorageKind;
316  typedef typename XprTraits::Index Index;
317  typedef typename XprType::Nested Nested;
318  typedef typename remove_reference<Nested>::type _Nested;
319  static const int NumDimensions = array_size<StartIndices>::value;
320  static const int Layout = XprTraits::Layout;
321  typedef typename XprTraits::PointerType PointerType;
322 };
323 
324 template<typename StartIndices, typename Sizes, typename XprType>
325 struct eval<TensorSlicingOp<StartIndices, Sizes, XprType>, Eigen::Dense>
326 {
327  typedef const TensorSlicingOp<StartIndices, Sizes, XprType>EIGEN_DEVICE_REF type;
328 };
329 
330 template<typename StartIndices, typename Sizes, typename XprType>
331 struct nested<TensorSlicingOp<StartIndices, Sizes, XprType>, 1, typename eval<TensorSlicingOp<StartIndices, Sizes, XprType> >::type>
332 {
333  typedef TensorSlicingOp<StartIndices, Sizes, XprType> type;
334 };
335 
336 } // end namespace internal
337 
338 
339 
340 template<typename StartIndices, typename Sizes, typename XprType>
341 class TensorSlicingOp : public TensorBase<TensorSlicingOp<StartIndices, Sizes, XprType> >
342 {
343  public:
344  typedef TensorBase<TensorSlicingOp<StartIndices, Sizes, XprType> > Base;
345  typedef typename Eigen::internal::traits<TensorSlicingOp>::Scalar Scalar;
346  typedef typename XprType::CoeffReturnType CoeffReturnType;
347  typedef typename Eigen::internal::nested<TensorSlicingOp>::type Nested;
348  typedef typename Eigen::internal::traits<TensorSlicingOp>::StorageKind StorageKind;
349  typedef typename Eigen::internal::traits<TensorSlicingOp>::Index Index;
350 
351  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorSlicingOp(const XprType& expr, const StartIndices& indices, const Sizes& sizes)
352  : m_xpr(expr), m_indices(indices), m_sizes(sizes) {}
353 
354  EIGEN_DEVICE_FUNC
355  const StartIndices& startIndices() const { return m_indices; }
356  EIGEN_DEVICE_FUNC
357  const Sizes& sizes() const { return m_sizes; }
358 
359  EIGEN_DEVICE_FUNC
360  const typename internal::remove_all<typename XprType::Nested>::type&
361  expression() const { return m_xpr; }
362 
363  EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorSlicingOp)
364 
365  protected:
366  typename XprType::Nested m_xpr;
367  const StartIndices m_indices;
368  const Sizes m_sizes;
369 };
370 
371 
372 // Fixme: figure out the exact threshold
373 namespace {
374 template <typename Index, typename Device, bool BlockAccess> struct MemcpyTriggerForSlicing {
375  EIGEN_DEVICE_FUNC MemcpyTriggerForSlicing(const Device& device) : threshold_(2 * device.numThreads()) { }
376  EIGEN_DEVICE_FUNC bool operator ()(Index total, Index contiguous) const {
377  const bool prefer_block_evaluation = BlockAccess && total > 32*1024;
378  return !prefer_block_evaluation && contiguous > threshold_;
379  }
380 
381  private:
382  Index threshold_;
383 };
384 
385 // It is very expensive to start the memcpy kernel on GPU: we therefore only
386 // use it for large copies.
387 #ifdef EIGEN_USE_GPU
388 template <typename Index, bool BlockAccess> struct MemcpyTriggerForSlicing<Index, GpuDevice, BlockAccess> {
389  EIGEN_DEVICE_FUNC MemcpyTriggerForSlicing(const GpuDevice&) { }
390  EIGEN_DEVICE_FUNC bool operator ()(Index, Index contiguous) const { return contiguous > 4*1024*1024; }
391 };
392 #endif
393 
394 // It is very expensive to start the memcpy kernel on GPU: we therefore only
395 // use it for large copies.
396 #ifdef EIGEN_USE_SYCL
397 template <typename Index, bool BlockAccess> struct MemcpyTriggerForSlicing<Index, Eigen::SyclDevice, BlockAccess> {
398  EIGEN_DEVICE_FUNC MemcpyTriggerForSlicing(const SyclDevice&) { }
399  EIGEN_DEVICE_FUNC bool operator ()(Index, Index contiguous) const { return contiguous > 4*1024*1024; }
400 };
401 #endif
402 
403 }
404 
405 // Eval as rvalue
406 template<typename StartIndices, typename Sizes, typename ArgType, typename Device>
407 struct TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, ArgType>, Device>
408 {
409  typedef TensorSlicingOp<StartIndices, Sizes, ArgType> XprType;
410  static const int NumDims = internal::array_size<Sizes>::value;
411 
412  typedef typename XprType::Index Index;
413  typedef typename XprType::Scalar Scalar;
414  typedef typename XprType::CoeffReturnType CoeffReturnType;
415  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
416  typedef Sizes Dimensions;
417  typedef StorageMemory<CoeffReturnType, Device> Storage;
418  typedef StorageMemory<typename internal::remove_const<CoeffReturnType>::type, Device> ConstCastStorage;
419  typedef typename Storage::Type EvaluatorPointerType;
420 
421  enum {
422  // Alignment can't be guaranteed at compile time since it depends on the
423  // slice offsets and sizes.
424  IsAligned = false,
425  PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
426  BlockAccess = TensorEvaluator<ArgType, Device>::BlockAccess &&
427  // FIXME: Temporary workaround for bug in slicing of bool tensors.
428  !internal::is_same<typename internal::remove_const<Scalar>::type, bool>::value,
429  PreferBlockAccess = true,
430  Layout = TensorEvaluator<ArgType, Device>::Layout,
431  CoordAccess = false,
432  RawAccess = false
433  };
434 
435  typedef typename internal::remove_const<Scalar>::type ScalarNoConst;
436 
437  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
438  typedef internal::TensorBlockDescriptor<NumDims, Index> TensorBlockDesc;
439  typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;
440 
441  // Tensor slicing does not change the block type.
442  typedef typename TensorEvaluator<const ArgType, Device>::TensorBlock
443  TensorBlock;
444  //===--------------------------------------------------------------------===//
445 
446  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
447  : m_impl(op.expression(), device), m_device(device), m_dimensions(op.sizes()), m_offsets(op.startIndices())
448  {
449  m_is_identity = true;
450  for (int i = 0; i < internal::array_size<Dimensions>::value; ++i) {
451  eigen_assert(m_impl.dimensions()[i] >=
452  op.sizes()[i] + op.startIndices()[i]);
453  if (m_impl.dimensions()[i] != op.sizes()[i] ||
454  op.startIndices()[i] != 0) {
455  m_is_identity = false;
456  }
457  }
458 
459  // No strides for scalars.
460  if (NumDims == 0) return;
461 
462  const typename TensorEvaluator<ArgType, Device>::Dimensions& input_dims = m_impl.dimensions();
463  const Sizes& output_dims = op.sizes();
464  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
465  m_inputStrides[0] = 1;
466  for (int i = 1; i < NumDims; ++i) {
467  m_inputStrides[i] = m_inputStrides[i-1] * input_dims[i-1];
468  }
469 
470  // Don't initialize m_fastOutputStrides[0] since it won't ever be accessed.
471  m_outputStrides[0] = 1;
472  for (int i = 1; i < NumDims; ++i) {
473  m_outputStrides[i] = m_outputStrides[i-1] * output_dims[i-1];
474  m_fastOutputStrides[i] = internal::TensorIntDivisor<Index>(m_outputStrides[i] > 0 ? m_outputStrides[i] : 1);
475  }
476  } else {
477  m_inputStrides[NumDims-1] = 1;
478  for (int i = NumDims - 2; i >= 0; --i) {
479  m_inputStrides[i] = m_inputStrides[i+1] * input_dims[i+1];
480  }
481 
482  // Don't initialize m_fastOutputStrides[NumDims-1] since it won't ever be accessed.
483  m_outputStrides[NumDims-1] = 1;
484  for (int i = NumDims - 2; i >= 0; --i) {
485  m_outputStrides[i] = m_outputStrides[i+1] * output_dims[i+1];
486  m_fastOutputStrides[i] = internal::TensorIntDivisor<Index>(m_outputStrides[i] > 0 ? m_outputStrides[i] : 1);
487  }
488  }
489  }
490 
491  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dimensions; }
492 
493  EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(EvaluatorPointerType data) {
494  m_impl.evalSubExprsIfNeeded(NULL);
495  if (!NumTraits<typename internal::remove_const<Scalar>::type>::RequireInitialization
496  && data && m_impl.data()) {
497  Index contiguous_values = 1;
498  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
499  for (int i = 0; i < NumDims; ++i) {
500  contiguous_values *= dimensions()[i];
501  if (dimensions()[i] != m_impl.dimensions()[i]) {
502  break;
503  }
504  }
505  } else {
506  for (int i = NumDims-1; i >= 0; --i) {
507  contiguous_values *= dimensions()[i];
508  if (dimensions()[i] != m_impl.dimensions()[i]) {
509  break;
510  }
511  }
512  }
513  // Use memcpy if it's going to be faster than using the regular evaluation.
514  const MemcpyTriggerForSlicing<Index, Device, BlockAccess> trigger(m_device);
515  if (trigger(internal::array_prod(dimensions()), contiguous_values)) {
516  EvaluatorPointerType src = (EvaluatorPointerType)m_impl.data();
517  for (Index i = 0; i < internal::array_prod(dimensions()); i += contiguous_values) {
518  Index offset = srcCoeff(i);
519  m_device.memcpy((void*)(m_device.get(data + i)), m_device.get(src+offset), contiguous_values * sizeof(Scalar));
520  }
521  return false;
522  }
523  }
524  return true;
525  }
526 
527 #ifdef EIGEN_USE_THREADS
528  template <typename EvalSubExprsCallback>
529  EIGEN_STRONG_INLINE void evalSubExprsIfNeededAsync(
530  EvaluatorPointerType /*data*/, EvalSubExprsCallback done) {
531  m_impl.evalSubExprsIfNeededAsync(nullptr, [done](bool) { done(true); });
532  }
533 #endif // EIGEN_USE_THREADS
534 
535  EIGEN_STRONG_INLINE void cleanup() {
536  m_impl.cleanup();
537  }
538 
539  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const
540  {
541  if (m_is_identity) {
542  return m_impl.coeff(index);
543  } else {
544  return m_impl.coeff(srcCoeff(index));
545  }
546  }
547 
548  template<int LoadMode>
549  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const
550  {
551  const int packetSize = PacketType<CoeffReturnType, Device>::size;
552  EIGEN_STATIC_ASSERT((packetSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
553  eigen_assert(index+packetSize-1 < internal::array_prod(dimensions()));
554 
555  if (m_is_identity) {
556  return m_impl.template packet<LoadMode>(index);
557  }
558 
559  Index inputIndices[] = {0, 0};
560  Index indices[] = {index, index + packetSize - 1};
561  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
562  EIGEN_UNROLL_LOOP
563  for (int i = NumDims - 1; i > 0; --i) {
564  const Index idx0 = indices[0] / m_fastOutputStrides[i];
565  const Index idx1 = indices[1] / m_fastOutputStrides[i];
566  inputIndices[0] += (idx0 + m_offsets[i]) * m_inputStrides[i];
567  inputIndices[1] += (idx1 + m_offsets[i]) * m_inputStrides[i];
568  indices[0] -= idx0 * m_outputStrides[i];
569  indices[1] -= idx1 * m_outputStrides[i];
570  }
571  inputIndices[0] += (indices[0] + m_offsets[0]);
572  inputIndices[1] += (indices[1] + m_offsets[0]);
573  } else {
574  EIGEN_UNROLL_LOOP
575  for (int i = 0; i < NumDims - 1; ++i) {
576  const Index idx0 = indices[0] / m_fastOutputStrides[i];
577  const Index idx1 = indices[1] / m_fastOutputStrides[i];
578  inputIndices[0] += (idx0 + m_offsets[i]) * m_inputStrides[i];
579  inputIndices[1] += (idx1 + m_offsets[i]) * m_inputStrides[i];
580  indices[0] -= idx0 * m_outputStrides[i];
581  indices[1] -= idx1 * m_outputStrides[i];
582  }
583  inputIndices[0] += (indices[0] + m_offsets[NumDims-1]);
584  inputIndices[1] += (indices[1] + m_offsets[NumDims-1]);
585  }
586  if (inputIndices[1] - inputIndices[0] == packetSize - 1) {
587  PacketReturnType rslt = m_impl.template packet<Unaligned>(inputIndices[0]);
588  return rslt;
589  }
590  else {
591  EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[packetSize];
592  values[0] = m_impl.coeff(inputIndices[0]);
593  values[packetSize-1] = m_impl.coeff(inputIndices[1]);
594  EIGEN_UNROLL_LOOP
595  for (int i = 1; i < packetSize-1; ++i) {
596  values[i] = coeff(index+i);
597  }
598  PacketReturnType rslt = internal::pload<PacketReturnType>(values);
599  return rslt;
600  }
601  }
602 
603  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const {
604  return m_impl.costPerCoeff(vectorized) + TensorOpCost(0, 0, m_is_identity ? 1 : NumDims);
605  }
606 
607  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
608  internal::TensorBlockResourceRequirements getResourceRequirements() const {
609  const size_t target_size = m_device.lastLevelCacheSize();
610  return internal::TensorBlockResourceRequirements::merge(
611  internal::TensorBlockResourceRequirements::skewed<Scalar>(target_size),
612  m_impl.getResourceRequirements());
613  }
614 
615  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorBlock
616  block(TensorBlockDesc& desc, TensorBlockScratch& scratch,
617  bool /*root_of_expr_ast*/ = false) const {
618  TensorBlockDesc arg_desc = desc.WithOffset(srcCoeff(desc.offset()));
619  TensorBlock block = m_impl.block(arg_desc, scratch);
620  if (!arg_desc.HasDestinationBuffer()) desc.DropDestinationBuffer();
621  return block;
622  }
623 
624  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Storage::Type data() const {
625  typename Storage::Type result = constCast(m_impl.data());
626  if (result) {
627  Index offset = 0;
628  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
629  for (int i = 0; i < NumDims; ++i) {
630  if (m_dimensions[i] != m_impl.dimensions()[i]) {
631  offset += m_offsets[i] * m_inputStrides[i];
632  for (int j = i+1; j < NumDims; ++j) {
633  if (m_dimensions[j] > 1) {
634  return NULL;
635  }
636  offset += m_offsets[j] * m_inputStrides[j];
637  }
638  break;
639  }
640  }
641  } else {
642  for (int i = NumDims - 1; i >= 0; --i) {
643  if (m_dimensions[i] != m_impl.dimensions()[i]) {
644  offset += m_offsets[i] * m_inputStrides[i];
645  for (int j = i-1; j >= 0; --j) {
646  if (m_dimensions[j] > 1) {
647  return NULL;
648  }
649  offset += m_offsets[j] * m_inputStrides[j];
650  }
651  break;
652  }
653  }
654  }
655  return result + offset;
656  }
657  return NULL;
658  }
659 #ifdef EIGEN_USE_SYCL
660  // binding placeholder accessors to a command group handler for SYCL
661  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void bind(cl::sycl::handler &cgh) const {
662  m_impl.bind(cgh);
663  }
664 #endif
665 
666  protected:
667  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index srcCoeff(Index index) const
668  {
669  Index inputIndex = 0;
670  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
671  EIGEN_UNROLL_LOOP
672  for (int i = NumDims - 1; i > 0; --i) {
673  const Index idx = index / m_fastOutputStrides[i];
674  inputIndex += (idx + m_offsets[i]) * m_inputStrides[i];
675  index -= idx * m_outputStrides[i];
676  }
677  inputIndex += (index + m_offsets[0]);
678  } else {
679  EIGEN_UNROLL_LOOP
680  for (int i = 0; i < NumDims - 1; ++i) {
681  const Index idx = index / m_fastOutputStrides[i];
682  inputIndex += (idx + m_offsets[i]) * m_inputStrides[i];
683  index -= idx * m_outputStrides[i];
684  }
685  inputIndex += (index + m_offsets[NumDims-1]);
686  }
687  return inputIndex;
688  }
689 
690  array<Index, NumDims> m_outputStrides;
691  array<internal::TensorIntDivisor<Index>, NumDims> m_fastOutputStrides;
692  array<Index, NumDims> m_inputStrides;
693  TensorEvaluator<ArgType, Device> m_impl;
694  const Device EIGEN_DEVICE_REF m_device;
695  Dimensions m_dimensions;
696  bool m_is_identity;
697  const StartIndices m_offsets;
698 };
699 
700 
701 // Eval as lvalue
702 template<typename StartIndices, typename Sizes, typename ArgType, typename Device>
703 struct TensorEvaluator<TensorSlicingOp<StartIndices, Sizes, ArgType>, Device>
704  : public TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, ArgType>, Device>
705 {
706  typedef TensorEvaluator<const TensorSlicingOp<StartIndices, Sizes, ArgType>, Device> Base;
707  typedef TensorSlicingOp<StartIndices, Sizes, ArgType> XprType;
708  static const int NumDims = internal::array_size<Sizes>::value;
709 
710  typedef typename XprType::Index Index;
711  typedef typename XprType::Scalar Scalar;
712  typedef typename XprType::CoeffReturnType CoeffReturnType;
713  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
714  typedef Sizes Dimensions;
715 
716  enum {
717  IsAligned = false,
718  PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
719  BlockAccess = TensorEvaluator<ArgType, Device>::BlockAccess,
720  PreferBlockAccess = true,
721  Layout = TensorEvaluator<ArgType, Device>::Layout,
722  CoordAccess = false,
723  RawAccess = (NumDims == 1) & TensorEvaluator<ArgType, Device>::RawAccess
724  };
725 
726  typedef typename internal::remove_const<Scalar>::type ScalarNoConst;
727 
728  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
729  typedef internal::TensorBlockDescriptor<NumDims, Index> TensorBlockDesc;
730  typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;
731  //===--------------------------------------------------------------------===//
732 
733  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
734  : Base(op, device)
735  { }
736 
737  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType& coeffRef(Index index)
738  {
739  if (this->m_is_identity) {
740  return this->m_impl.coeffRef(index);
741  } else {
742  return this->m_impl.coeffRef(this->srcCoeff(index));
743  }
744  }
745 
746  template <int StoreMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
747  void writePacket(Index index, const PacketReturnType& x)
748  {
749  if (this->m_is_identity) {
750  this->m_impl.template writePacket<StoreMode>(index, x);
751  return;
752  }
753 
754  const int packetSize = PacketType<CoeffReturnType, Device>::size;
755  Index inputIndices[] = {0, 0};
756  Index indices[] = {index, index + packetSize - 1};
757  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
758  EIGEN_UNROLL_LOOP
759  for (int i = NumDims - 1; i > 0; --i) {
760  const Index idx0 = indices[0] / this->m_fastOutputStrides[i];
761  const Index idx1 = indices[1] / this->m_fastOutputStrides[i];
762  inputIndices[0] += (idx0 + this->m_offsets[i]) * this->m_inputStrides[i];
763  inputIndices[1] += (idx1 + this->m_offsets[i]) * this->m_inputStrides[i];
764  indices[0] -= idx0 * this->m_outputStrides[i];
765  indices[1] -= idx1 * this->m_outputStrides[i];
766  }
767  inputIndices[0] += (indices[0] + this->m_offsets[0]);
768  inputIndices[1] += (indices[1] + this->m_offsets[0]);
769  } else {
770  EIGEN_UNROLL_LOOP
771  for (int i = 0; i < NumDims - 1; ++i) {
772  const Index idx0 = indices[0] / this->m_fastOutputStrides[i];
773  const Index idx1 = indices[1] / this->m_fastOutputStrides[i];
774  inputIndices[0] += (idx0 + this->m_offsets[i]) * this->m_inputStrides[i];
775  inputIndices[1] += (idx1 + this->m_offsets[i]) * this->m_inputStrides[i];
776  indices[0] -= idx0 * this->m_outputStrides[i];
777  indices[1] -= idx1 * this->m_outputStrides[i];
778  }
779  inputIndices[0] += (indices[0] + this->m_offsets[NumDims-1]);
780  inputIndices[1] += (indices[1] + this->m_offsets[NumDims-1]);
781  }
782  if (inputIndices[1] - inputIndices[0] == packetSize - 1) {
783  this->m_impl.template writePacket<StoreMode>(inputIndices[0], x);
784  }
785  else {
786  EIGEN_ALIGN_MAX CoeffReturnType values[packetSize];
787  internal::pstore<CoeffReturnType, PacketReturnType>(values, x);
788  this->m_impl.coeffRef(inputIndices[0]) = values[0];
789  this->m_impl.coeffRef(inputIndices[1]) = values[packetSize-1];
790  EIGEN_UNROLL_LOOP
791  for (int i = 1; i < packetSize-1; ++i) {
792  this->coeffRef(index+i) = values[i];
793  }
794  }
795  }
796 
797  template<typename TensorBlock>
798  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void writeBlock(
799  const TensorBlockDesc& desc, const TensorBlock& block) {
800  TensorBlockDesc arg_desc = desc.WithOffset(this->srcCoeff(desc.offset()));
801  this->m_impl.writeBlock(arg_desc, block);
802  }
803 };
804 
805 namespace internal {
806 template<typename StartIndices, typename StopIndices, typename Strides, typename XprType>
807 struct traits<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType> > : public traits<XprType>
808 {
809  typedef typename XprType::Scalar Scalar;
810  typedef traits<XprType> XprTraits;
811  typedef typename XprTraits::StorageKind StorageKind;
812  typedef typename XprTraits::Index Index;
813  typedef typename XprType::Nested Nested;
814  typedef typename remove_reference<Nested>::type _Nested;
815  static const int NumDimensions = array_size<StartIndices>::value;
816  static const int Layout = XprTraits::Layout;
817  typedef typename XprTraits::PointerType PointerType;
818 };
819 
820 template<typename StartIndices, typename StopIndices, typename Strides, typename XprType>
821 struct eval<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType>, Eigen::Dense>
822 {
823  typedef const TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType>EIGEN_DEVICE_REF type;
824 };
825 
826 template<typename StartIndices, typename StopIndices, typename Strides, typename XprType>
827 struct nested<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType>, 1, typename eval<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType> >::type>
828 {
829  typedef TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType> type;
830 };
831 
832 } // end namespace internal
833 
834 
835 template<typename StartIndices, typename StopIndices, typename Strides, typename XprType>
836 class TensorStridingSlicingOp : public TensorBase<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType> >
837 {
838  public:
839  typedef TensorBase<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, XprType> > Base;
840  typedef typename internal::traits<TensorStridingSlicingOp>::Scalar Scalar;
841  typedef typename XprType::CoeffReturnType CoeffReturnType;
842  typedef typename internal::nested<TensorStridingSlicingOp>::type Nested;
843  typedef typename internal::traits<TensorStridingSlicingOp>::StorageKind StorageKind;
844  typedef typename internal::traits<TensorStridingSlicingOp>::Index Index;
845 
846  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorStridingSlicingOp(
847  const XprType& expr, const StartIndices& startIndices,
848  const StopIndices& stopIndices, const Strides& strides)
849  : m_xpr(expr), m_startIndices(startIndices), m_stopIndices(stopIndices),
850  m_strides(strides) {}
851 
852  EIGEN_DEVICE_FUNC
853  const StartIndices& startIndices() const { return m_startIndices; }
854  EIGEN_DEVICE_FUNC
855  const StartIndices& stopIndices() const { return m_stopIndices; }
856  EIGEN_DEVICE_FUNC
857  const StartIndices& strides() const { return m_strides; }
858 
859  EIGEN_DEVICE_FUNC
860  const typename internal::remove_all<typename XprType::Nested>::type&
861  expression() const { return m_xpr; }
862 
863  EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorStridingSlicingOp)
864 
865  protected:
866  typename XprType::Nested m_xpr;
867  const StartIndices m_startIndices;
868  const StopIndices m_stopIndices;
869  const Strides m_strides;
870 };
871 
872 // Eval as rvalue
873 template<typename StartIndices, typename StopIndices, typename Strides, typename ArgType, typename Device>
874 struct TensorEvaluator<const TensorStridingSlicingOp<StartIndices, StopIndices, Strides, ArgType>, Device>
875 {
876  typedef TensorStridingSlicingOp<StartIndices, StopIndices, Strides, ArgType> XprType;
877  static const int NumDims = internal::array_size<Strides>::value;
878  typedef typename XprType::Index Index;
879  typedef typename XprType::Scalar Scalar;
880  typedef typename XprType::CoeffReturnType CoeffReturnType;
881  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
882  typedef StorageMemory<CoeffReturnType, Device> Storage;
883  typedef typename Storage::Type EvaluatorPointerType;
884  typedef Strides Dimensions;
885 
886  enum {
887  // Alignment can't be guaranteed at compile time since it depends on the
888  // slice offsets and sizes.
889  IsAligned = false,
890  PacketAccess = false,
891  BlockAccess = false,
892  PreferBlockAccess = TensorEvaluator<ArgType, Device>::PreferBlockAccess,
893  Layout = TensorEvaluator<ArgType, Device>::Layout,
894  RawAccess = false
895  };
896 
897  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
898  typedef internal::TensorBlockNotImplemented TensorBlock;
899  //===--------------------------------------------------------------------===//
900 
901  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
902  : m_impl(op.expression(), device),
903  m_device(device),
904  m_strides(op.strides())
905  {
906  // Handle degenerate intervals by gracefully clamping and allowing m_dimensions to be zero
907  DSizes<Index, NumDims> startIndicesClamped, stopIndicesClamped;
908  for (ptrdiff_t i = 0; i < internal::array_size<Dimensions>::value; ++i) {
909  eigen_assert(m_strides[i] != 0 && "0 stride is invalid");
910  if (m_strides[i] > 0) {
911  startIndicesClamped[i] =
912  clamp(op.startIndices()[i], 0, m_impl.dimensions()[i]);
913  stopIndicesClamped[i] =
914  clamp(op.stopIndices()[i], 0, m_impl.dimensions()[i]);
915  } else {
916  /* implies m_strides[i] < 0 by assert */
917  startIndicesClamped[i] =
918  clamp(op.startIndices()[i], -1, m_impl.dimensions()[i] - 1);
919  stopIndicesClamped[i] =
920  clamp(op.stopIndices()[i], -1, m_impl.dimensions()[i] - 1);
921  }
922  m_startIndices[i] = startIndicesClamped[i];
923  }
924 
925  typedef typename TensorEvaluator<ArgType, Device>::Dimensions InputDimensions;
926  const InputDimensions& input_dims = m_impl.dimensions();
927 
928  // compute output tensor shape
929  m_is_identity = true;
930  for (int i = 0; i < NumDims; i++) {
931  Index interval = stopIndicesClamped[i] - startIndicesClamped[i];
932  if (interval == 0 || ((interval < 0) != (m_strides[i] < 0))) {
933  m_dimensions[i] = 0;
934  } else {
935  m_dimensions[i] =
936  (interval / m_strides[i]) + (interval % m_strides[i] != 0 ? 1 : 0);
937  eigen_assert(m_dimensions[i] >= 0);
938  }
939  if (m_strides[i] != 1 || interval != m_impl.dimensions()[i]) {
940  m_is_identity = false;
941  }
942  }
943 
944  Strides output_dims = m_dimensions;
945 
946  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
947  m_inputStrides[0] = m_strides[0];
948  m_offsets[0] = startIndicesClamped[0];
949  Index previousDimProduct = 1;
950  for (int i = 1; i < NumDims; ++i) {
951  previousDimProduct *= input_dims[i-1];
952  m_inputStrides[i] = previousDimProduct * m_strides[i];
953  m_offsets[i] = startIndicesClamped[i] * previousDimProduct;
954  }
955 
956  // Don't initialize m_fastOutputStrides[0] since it won't ever be accessed.
957  m_outputStrides[0] = 1;
958  for (int i = 1; i < NumDims; ++i) {
959  m_outputStrides[i] = m_outputStrides[i-1] * output_dims[i-1];
960  m_fastOutputStrides[i] = internal::TensorIntDivisor<Index>(m_outputStrides[i] > 0 ? m_outputStrides[i] : 1);
961  }
962  } else {
963  m_inputStrides[NumDims-1] = m_strides[NumDims-1];
964  m_offsets[NumDims-1] = startIndicesClamped[NumDims-1];
965  Index previousDimProduct = 1;
966  for (int i = NumDims - 2; i >= 0; --i) {
967  previousDimProduct *= input_dims[i+1];
968  m_inputStrides[i] = previousDimProduct * m_strides[i];
969  m_offsets[i] = startIndicesClamped[i] * previousDimProduct;
970  }
971 
972  m_outputStrides[NumDims-1] = 1;
973  for (int i = NumDims - 2; i >= 0; --i) {
974  m_outputStrides[i] = m_outputStrides[i+1] * output_dims[i+1];
975  m_fastOutputStrides[i] = internal::TensorIntDivisor<Index>(m_outputStrides[i] > 0 ? m_outputStrides[i] : 1);
976  }
977  }
978  }
979 
980  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dimensions; }
981 
982 
983  EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(EvaluatorPointerType) {
984  m_impl.evalSubExprsIfNeeded(NULL);
985  return true;
986  }
987 
988  EIGEN_STRONG_INLINE void cleanup() {
989  m_impl.cleanup();
990  }
991 
992  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const
993  {
994  if (m_is_identity) {
995  return m_impl.coeff(index);
996  } else {
997  return m_impl.coeff(srcCoeff(index));
998  }
999  }
1000 
1001  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const {
1002  return m_impl.costPerCoeff(vectorized) + TensorOpCost(0, 0, m_is_identity ? 1 : NumDims);
1003  }
1004 
1005  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename Storage::Type data() const {
1006  return NULL;
1007  }
1008 #ifdef EIGEN_USE_SYCL
1009  // binding placeholder accessors to a command group handler for SYCL
1010  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void bind(cl::sycl::handler &cgh) const {
1011  m_impl.bind(cgh);
1012  }
1013 #endif
1014  protected:
1015  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index srcCoeff(Index index) const
1016  {
1017  Index inputIndex = 0;
1018  if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
1019  EIGEN_UNROLL_LOOP
1020  for (int i = NumDims - 1; i >= 0; --i) {
1021  const Index idx = index / m_fastOutputStrides[i];
1022  inputIndex += idx * m_inputStrides[i] + m_offsets[i];
1023  index -= idx * m_outputStrides[i];
1024  }
1025  } else {
1026  EIGEN_UNROLL_LOOP
1027  for (int i = 0; i < NumDims; ++i) {
1028  const Index idx = index / m_fastOutputStrides[i];
1029  inputIndex += idx * m_inputStrides[i] + m_offsets[i];
1030  index -= idx * m_outputStrides[i];
1031  }
1032  }
1033  return inputIndex;
1034  }
1035 
1036  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index clamp(Index value, Index min, Index max) {
1037 #ifndef SYCL_DEVICE_ONLY
1038  return numext::maxi(min, numext::mini(max,value));
1039 #else
1040  return cl::sycl::clamp(value, min, max);
1041 #endif
1042  }
1043 
1044  array<Index, NumDims> m_outputStrides;
1045  array<internal::TensorIntDivisor<Index>, NumDims> m_fastOutputStrides;
1046  array<Index, NumDims> m_inputStrides;
1047  bool m_is_identity;
1048  TensorEvaluator<ArgType, Device> m_impl;
1049  const Device EIGEN_DEVICE_REF m_device;
1050  DSizes<Index, NumDims> m_startIndices; // clamped startIndices
1051  DSizes<Index, NumDims> m_dimensions;
1052  DSizes<Index, NumDims> m_offsets; // offset in a flattened shape
1053  const Strides m_strides;
1054 };
1055 
1056 // Eval as lvalue
1057 template<typename StartIndices, typename StopIndices, typename Strides, typename ArgType, typename Device>
1058 struct TensorEvaluator<TensorStridingSlicingOp<StartIndices, StopIndices, Strides, ArgType>, Device>
1059  : public TensorEvaluator<const TensorStridingSlicingOp<StartIndices, StopIndices, Strides, ArgType>, Device>
1060 {
1061  typedef TensorEvaluator<const TensorStridingSlicingOp<StartIndices, StopIndices, Strides, ArgType>, Device> Base;
1062  typedef TensorStridingSlicingOp<StartIndices, StopIndices, Strides, ArgType> XprType;
1063  static const int NumDims = internal::array_size<Strides>::value;
1064 
1065  enum {
1066  IsAligned = false,
1067  PacketAccess = false,
1068  BlockAccess = false,
1069  PreferBlockAccess = TensorEvaluator<ArgType, Device>::PreferBlockAccess,
1070  Layout = TensorEvaluator<ArgType, Device>::Layout,
1071  CoordAccess = TensorEvaluator<ArgType, Device>::CoordAccess,
1072  RawAccess = false
1073  };
1074 
1075  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
1076  typedef internal::TensorBlockNotImplemented TensorBlock;
1077  //===--------------------------------------------------------------------===//
1078 
1079  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
1080  : Base(op, device)
1081  { }
1082 
1083  typedef typename XprType::Index Index;
1084  typedef typename XprType::Scalar Scalar;
1085  typedef typename XprType::CoeffReturnType CoeffReturnType;
1086  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
1087  typedef Strides Dimensions;
1088 
1089  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType& coeffRef(Index index)
1090  {
1091  if (this->m_is_identity) {
1092  return this->m_impl.coeffRef(index);
1093  } else {
1094  return this->m_impl.coeffRef(this->srcCoeff(index));
1095  }
1096  }
1097 };
1098 
1099 
1100 } // end namespace Eigen
1101 
1102 #endif // EIGEN_CXX11_TENSOR_TENSOR_MORPHING_H
WriteAccessors
Namespace containing all symbols from the Eigen library.
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index