11#ifndef EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
12#define EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
13#include "./InternalHeaderCheck.h"
24template<
typename ReverseDimensions,
typename XprType>
25struct traits<TensorReverseOp<ReverseDimensions,
26 XprType> > :
public traits<XprType>
28 typedef typename XprType::Scalar Scalar;
29 typedef traits<XprType> XprTraits;
30 typedef typename XprTraits::StorageKind StorageKind;
31 typedef typename XprTraits::Index
Index;
32 typedef typename XprType::Nested Nested;
33 typedef typename remove_reference<Nested>::type _Nested;
34 static const int NumDimensions = XprTraits::NumDimensions;
35 static const int Layout = XprTraits::Layout;
36 typedef typename XprTraits::PointerType PointerType;
39template<
typename ReverseDimensions,
typename XprType>
40struct eval<TensorReverseOp<ReverseDimensions, XprType>,
Eigen::Dense>
42 typedef const TensorReverseOp<ReverseDimensions, XprType>& type;
45template<
typename ReverseDimensions,
typename XprType>
46struct nested<TensorReverseOp<ReverseDimensions, XprType>, 1,
47 typename eval<TensorReverseOp<ReverseDimensions, XprType> >::type>
49 typedef TensorReverseOp<ReverseDimensions, XprType> type;
54template<
typename ReverseDimensions,
typename XprType>
55class TensorReverseOp :
public TensorBase<TensorReverseOp<ReverseDimensions,
56 XprType>, WriteAccessors>
59 typedef TensorBase<TensorReverseOp<ReverseDimensions, XprType>,
WriteAccessors>Base;
60 typedef typename Eigen::internal::traits<TensorReverseOp>::Scalar Scalar;
62 typedef typename XprType::CoeffReturnType CoeffReturnType;
63 typedef typename Eigen::internal::nested<TensorReverseOp>::type Nested;
64 typedef typename Eigen::internal::traits<TensorReverseOp>::StorageKind
66 typedef typename Eigen::internal::traits<TensorReverseOp>::Index
Index;
68 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorReverseOp(
69 const XprType& expr,
const ReverseDimensions& reverse_dims)
70 : m_xpr(expr), m_reverse_dims(reverse_dims) { }
73 const ReverseDimensions& reverse()
const {
return m_reverse_dims; }
76 const typename internal::remove_all<typename XprType::Nested>::type&
77 expression()
const {
return m_xpr; }
79 EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorReverseOp)
83 typename XprType::Nested m_xpr;
84 const ReverseDimensions m_reverse_dims;
88template<
typename ReverseDimensions,
typename ArgType,
typename Device>
89struct TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device>
91 typedef TensorReverseOp<ReverseDimensions, ArgType> XprType;
92 typedef typename XprType::Index Index;
93 static const int NumDims = internal::array_size<ReverseDimensions>::value;
94 typedef DSizes<Index, NumDims> Dimensions;
95 typedef typename XprType::Scalar Scalar;
96 typedef typename XprType::CoeffReturnType CoeffReturnType;
97 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
98 static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
99 typedef StorageMemory<CoeffReturnType, Device> Storage;
100 typedef typename Storage::Type EvaluatorPointerType;
104 PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
105 BlockAccess = NumDims > 0,
106 PreferBlockAccess =
true,
107 Layout = TensorEvaluator<ArgType, Device>::Layout,
112 typedef internal::TensorIntDivisor<Index> IndexDivisor;
115 typedef internal::TensorBlockDescriptor<NumDims, Index> TensorBlockDesc;
116 typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;
118 typedef typename TensorEvaluator<const ArgType, Device>::TensorBlock
121 typedef typename internal::TensorMaterializedBlock<CoeffReturnType, NumDims,
126 EIGEN_STRONG_INLINE TensorEvaluator(
const XprType& op,
const Device& device)
127 : m_impl(op.expression(), device),
128 m_reverse(op.reverse()),
132 EIGEN_STATIC_ASSERT((NumDims > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);
135 m_dimensions = m_impl.dimensions();
136 if (
static_cast<int>(Layout) ==
static_cast<int>(
ColMajor)) {
138 for (
int i = 1; i < NumDims; ++i) {
139 m_strides[i] = m_strides[i-1] * m_dimensions[i-1];
140 if (m_strides[i] > 0) m_fastStrides[i] = IndexDivisor(m_strides[i]);
143 m_strides[NumDims-1] = 1;
144 for (
int i = NumDims - 2; i >= 0; --i) {
145 m_strides[i] = m_strides[i+1] * m_dimensions[i+1];
146 if (m_strides[i] > 0) m_fastStrides[i] = IndexDivisor(m_strides[i]);
151 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
152 const Dimensions& dimensions()
const {
return m_dimensions; }
154 EIGEN_STRONG_INLINE
bool evalSubExprsIfNeeded(EvaluatorPointerType) {
155 m_impl.evalSubExprsIfNeeded(NULL);
159#ifdef EIGEN_USE_THREADS
160 template <
typename EvalSubExprsCallback>
161 EIGEN_STRONG_INLINE
void evalSubExprsIfNeededAsync(
162 EvaluatorPointerType, EvalSubExprsCallback done) {
163 m_impl.evalSubExprsIfNeededAsync(
nullptr, [done](
bool) { done(
true); });
167 EIGEN_STRONG_INLINE
void cleanup() {
171 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
Index reverseIndex(
173 eigen_assert(index < dimensions().TotalSize());
174 Index inputIndex = 0;
175 if (
static_cast<int>(Layout) ==
static_cast<int>(
ColMajor)) {
177 for (
int i = NumDims - 1; i > 0; --i) {
178 Index idx = index / m_fastStrides[i];
179 index -= idx * m_strides[i];
181 idx = m_dimensions[i] - idx - 1;
183 inputIndex += idx * m_strides[i] ;
186 inputIndex += (m_dimensions[0] - index - 1);
192 for (
int i = 0; i < NumDims - 1; ++i) {
193 Index idx = index / m_fastStrides[i];
194 index -= idx * m_strides[i];
196 idx = m_dimensions[i] - idx - 1;
198 inputIndex += idx * m_strides[i] ;
200 if (m_reverse[NumDims-1]) {
201 inputIndex += (m_dimensions[NumDims-1] - index - 1);
209 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(
211 return m_impl.coeff(reverseIndex(index));
214 template<
int LoadMode>
215 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
216 PacketReturnType packet(Index index)
const
218 EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
219 eigen_assert(index+PacketSize-1 < dimensions().TotalSize());
223 EIGEN_ALIGN_MAX
typename internal::remove_const<CoeffReturnType>::type
226 for (
int i = 0; i < PacketSize; ++i) {
227 values[i] = coeff(index+i);
229 PacketReturnType rslt = internal::pload<PacketReturnType>(values);
233 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
234 internal::TensorBlockResourceRequirements getResourceRequirements()
const {
235 const size_t target_size = m_device.lastLevelCacheSize();
238 return internal::TensorBlockResourceRequirements::skewed<Scalar>(
240 .addCostPerCoeff({0, 0, 24});
243 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorBlock
244 block(TensorBlockDesc& desc, TensorBlockScratch& scratch,
245 bool =
false)
const {
253 static const bool isColMajor =
254 static_cast<int>(Layout) ==
static_cast<int>(
ColMajor);
256 static const Index inner_dim_idx = isColMajor ? 0 : NumDims - 1;
257 const bool inner_dim_reversed = m_reverse[inner_dim_idx];
260 Index block_offset = 0;
263 Index input_offset = reverseIndex(desc.offset());
267 array<BlockIteratorState, NumDims> it;
268 for (
int i = 0; i < NumDims; ++i) {
269 const int dim = isColMajor ? i : NumDims - 1 - i;
270 it[i].size = desc.dimension(dim);
272 it[i].reverse = m_reverse[dim];
275 i == 0 ? 1 : (it[i - 1].size * it[i - 1].block_stride);
276 it[i].block_span = it[i].block_stride * (it[i].size - 1);
278 it[i].input_stride = m_strides[dim];
279 it[i].input_span = it[i].input_stride * (it[i].size - 1);
282 it[i].input_stride = -1 * it[i].input_stride;
283 it[i].input_span = -1 * it[i].input_span;
289 int effective_inner_dim = 0;
290 for (
int i = 1; i < NumDims; ++i) {
291 if (it[i].reverse != it[effective_inner_dim].reverse)
break;
292 if (it[i].block_stride != it[effective_inner_dim].size)
break;
293 if (it[i].block_stride != numext::abs(it[i].input_stride))
break;
295 it[i].size = it[effective_inner_dim].size * it[i].size;
297 it[i].block_stride = 1;
298 it[i].input_stride = (inner_dim_reversed ? -1 : 1);
300 it[i].block_span = it[i].block_stride * (it[i].size - 1);
301 it[i].input_span = it[i].input_stride * (it[i].size - 1);
303 effective_inner_dim = i;
306 eigen_assert(it[effective_inner_dim].block_stride == 1);
307 eigen_assert(it[effective_inner_dim].input_stride ==
308 (inner_dim_reversed ? -1 : 1));
310 const Index inner_dim_size = it[effective_inner_dim].size;
313 const typename TensorBlock::Storage block_storage =
314 TensorBlock::prepareStorage(desc, scratch);
315 CoeffReturnType* block_buffer = block_storage.data();
317 while (it[NumDims - 1].count < it[NumDims - 1].size) {
319 Index dst = block_offset;
320 Index src = input_offset;
324 if (inner_dim_reversed) {
325 for (Index i = 0; i < inner_dim_size; ++i) {
326 block_buffer[dst] = m_impl.coeff(src);
331 for (Index i = 0; i < inner_dim_size; ++i) {
332 block_buffer[dst] = m_impl.coeff(src);
339 if ((NumDims - effective_inner_dim) == 1)
break;
342 for (Index i = effective_inner_dim + 1; i < NumDims; ++i) {
343 if (++it[i].count < it[i].size) {
344 block_offset += it[i].block_stride;
345 input_offset += it[i].input_stride;
348 if (i != NumDims - 1) it[i].count = 0;
349 block_offset -= it[i].block_span;
350 input_offset -= it[i].input_span;
354 return block_storage.AsTensorMaterializedBlock();
357 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(
bool vectorized)
const {
358 double compute_cost = NumDims * (2 * TensorOpCost::AddCost<Index>() +
359 2 * TensorOpCost::MulCost<Index>() +
360 TensorOpCost::DivCost<Index>());
361 for (
int i = 0; i < NumDims; ++i) {
363 compute_cost += 2 * TensorOpCost::AddCost<Index>();
366 return m_impl.costPerCoeff(vectorized) +
367 TensorOpCost(0, 0, compute_cost,
false , PacketSize);
370 EIGEN_DEVICE_FUNC
typename Storage::Type data()
const {
return NULL; }
374 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
380 Dimensions m_dimensions;
381 array<Index, NumDims> m_strides;
382 array<IndexDivisor, NumDims> m_fastStrides;
383 TensorEvaluator<ArgType, Device> m_impl;
384 ReverseDimensions m_reverse;
385 const Device EIGEN_DEVICE_REF m_device;
388 struct BlockIteratorState {
410template <
typename ReverseDimensions,
typename ArgType,
typename Device>
411struct TensorEvaluator<TensorReverseOp<ReverseDimensions, ArgType>, Device>
412 :
public TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>,
414 typedef TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>,
416 typedef TensorReverseOp<ReverseDimensions, ArgType> XprType;
417 typedef typename XprType::Index
Index;
418 static const int NumDims = internal::array_size<ReverseDimensions>::value;
419 typedef DSizes<Index, NumDims> Dimensions;
423 PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
425 PreferBlockAccess =
false,
426 Layout = TensorEvaluator<ArgType, Device>::Layout,
430 EIGEN_STRONG_INLINE TensorEvaluator(
const XprType& op,
const Device& device)
431 : Base(op, device) {}
433 typedef typename XprType::Scalar Scalar;
434 typedef typename XprType::CoeffReturnType CoeffReturnType;
435 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
436 static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
439 typedef internal::TensorBlockNotImplemented TensorBlock;
442 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
443 const Dimensions& dimensions()
const {
return this->m_dimensions; }
445 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index index) {
446 return this->m_impl.coeffRef(this->reverseIndex(index));
449 template <
int StoreMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
450 void writePacket(Index index,
const PacketReturnType& x) {
451 EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
452 eigen_assert(index+PacketSize-1 < dimensions().TotalSize());
455 EIGEN_ALIGN_MAX CoeffReturnType values[PacketSize];
456 internal::pstore<CoeffReturnType, PacketReturnType>(values, x);
458 for (
int i = 0; i < PacketSize; ++i) {
459 this->coeffRef(index+i) = values[i];
Namespace containing all symbols from the Eigen library.
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index