// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Copyright (C) 2014 Navdeep Jaitly <ndjaitly@google.com>
//                    Benoit Steiner <benoit.steiner.goog@gmail.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/.

#ifndef EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
#define EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
namespace Eigen {

/** \class TensorReverse
  * \ingroup CXX11_Tensor_Module
  *
  * \brief Tensor reverse elements class.
  *
  */
namespace internal {
template<typename ReverseDimensions, typename XprType>
struct traits<TensorReverseOp<ReverseDimensions,
                              XprType> > : public traits<XprType>
{
  typedef typename XprType::Scalar Scalar;
  typedef traits<XprType> XprTraits;
  typedef typename XprTraits::StorageKind StorageKind;
  typedef typename XprTraits::Index Index;
  typedef typename XprType::Nested Nested;
  typedef typename remove_reference<Nested>::type _Nested;
  static const int NumDimensions = XprTraits::NumDimensions;
  static const int Layout = XprTraits::Layout;
  typedef typename XprTraits::PointerType PointerType;
};

template<typename ReverseDimensions, typename XprType>
struct eval<TensorReverseOp<ReverseDimensions, XprType>, Eigen::Dense>
{
  typedef const TensorReverseOp<ReverseDimensions, XprType>& type;
};

template<typename ReverseDimensions, typename XprType>
struct nested<TensorReverseOp<ReverseDimensions, XprType>, 1,
            typename eval<TensorReverseOp<ReverseDimensions, XprType> >::type>
{
  typedef TensorReverseOp<ReverseDimensions, XprType> type;
};

}  // end namespace internal

template<typename ReverseDimensions, typename XprType>
class TensorReverseOp : public TensorBase<TensorReverseOp<ReverseDimensions,
                                          XprType>, WriteAccessors>
{
  public:
    typedef TensorBase<TensorReverseOp<ReverseDimensions, XprType>, WriteAccessors>Base;
    typedef typename Eigen::internal::traits<TensorReverseOp>::Scalar Scalar;
    typedef typename Eigen::NumTraits<Scalar>::Real RealScalar;
    typedef typename XprType::CoeffReturnType CoeffReturnType;
    typedef typename Eigen::internal::nested<TensorReverseOp>::type Nested;
    typedef typename Eigen::internal::traits<TensorReverseOp>::StorageKind
                                                                      StorageKind;
    typedef typename Eigen::internal::traits<TensorReverseOp>::Index Index;

    EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorReverseOp(
      const XprType& expr, const ReverseDimensions& reverse_dims)
      : m_xpr(expr), m_reverse_dims(reverse_dims) { }

    EIGEN_DEVICE_FUNC
    const ReverseDimensions& reverse() const { return m_reverse_dims; }

    EIGEN_DEVICE_FUNC
    const typename internal::remove_all<typename XprType::Nested>::type&
    expression() const { return m_xpr; }

    EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorReverseOp)


  protected:
    typename XprType::Nested m_xpr;
    const ReverseDimensions m_reverse_dims;
};

// Eval as rvalue
template<typename ReverseDimensions, typename ArgType, typename Device>
struct TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device>
{
  typedef TensorReverseOp<ReverseDimensions, ArgType> XprType;
  typedef typename XprType::Index Index;
  static const int NumDims = internal::array_size<ReverseDimensions>::value;
  typedef DSizes<Index, NumDims> Dimensions;
  typedef typename XprType::Scalar Scalar;
  typedef typename XprType::CoeffReturnType CoeffReturnType;
  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
  static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
  typedef StorageMemory<CoeffReturnType, Device> Storage;
  typedef typename Storage::Type EvaluatorPointerType;

  enum {
    IsAligned         = false,
    PacketAccess      = TensorEvaluator<ArgType, Device>::PacketAccess,
    BlockAccess       = NumDims > 0,
    PreferBlockAccess = true,
    Layout            = TensorEvaluator<ArgType, Device>::Layout,
    CoordAccess       = false,  // to be implemented
    RawAccess         = false
  };

  typedef internal::TensorIntDivisor<Index> IndexDivisor;

  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
  typedef internal::TensorBlockDescriptor<NumDims, Index> TensorBlockDesc;
  typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;

  typedef typename TensorEvaluator<const ArgType, Device>::TensorBlock
      ArgTensorBlock;

  typedef typename internal::TensorMaterializedBlock<CoeffReturnType, NumDims,
                                                     Layout, Index>
      TensorBlock;
  //===--------------------------------------------------------------------===//

  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
      : m_impl(op.expression(), device),
        m_reverse(op.reverse()),
        m_device(device)
  {
    // Reversing a scalar isn't supported yet. It would be a no-op anyway.
    EIGEN_STATIC_ASSERT((NumDims > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);

    // Compute strides
    m_dimensions = m_impl.dimensions();
    if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
      m_strides[0] = 1;
      for (int i = 1; i < NumDims; ++i) {
        m_strides[i] = m_strides[i-1] * m_dimensions[i-1];
        if (m_strides[i] > 0) m_fastStrides[i] = IndexDivisor(m_strides[i]);
      }
    } else {
      m_strides[NumDims-1] = 1;
      for (int i = NumDims - 2; i >= 0; --i) {
        m_strides[i] = m_strides[i+1] * m_dimensions[i+1];
        if (m_strides[i] > 0) m_fastStrides[i] = IndexDivisor(m_strides[i]);
      }
    }
  }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
  const Dimensions& dimensions() const { return m_dimensions; }

  EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(EvaluatorPointerType) {
    m_impl.evalSubExprsIfNeeded(NULL);
    return true;
  }

#ifdef EIGEN_USE_THREADS
  template <typename EvalSubExprsCallback>
  EIGEN_STRONG_INLINE void evalSubExprsIfNeededAsync(
      EvaluatorPointerType, EvalSubExprsCallback done) {
    m_impl.evalSubExprsIfNeededAsync(nullptr, [done](bool) { done(true); });
  }
#endif  // EIGEN_USE_THREADS

  EIGEN_STRONG_INLINE void cleanup() {
    m_impl.cleanup();
  }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index reverseIndex(
      Index index) const {
    eigen_assert(index < dimensions().TotalSize());
    Index inputIndex = 0;
    if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
      EIGEN_UNROLL_LOOP
      for (int i = NumDims - 1; i > 0; --i) {
        Index idx = index / m_fastStrides[i];
        index -= idx * m_strides[i];
        if (m_reverse[i]) {
          idx = m_dimensions[i] - idx - 1;
        }
        inputIndex += idx * m_strides[i] ;
      }
      if (m_reverse[0]) {
        inputIndex += (m_dimensions[0] - index - 1);
      } else {
        inputIndex += index;
      }
    } else {
      EIGEN_UNROLL_LOOP
      for (int i = 0; i < NumDims - 1; ++i) {
        Index idx = index / m_fastStrides[i];
        index -= idx * m_strides[i];
        if (m_reverse[i]) {
          idx = m_dimensions[i] - idx - 1;
        }
        inputIndex += idx * m_strides[i] ;
      }
      if (m_reverse[NumDims-1]) {
        inputIndex += (m_dimensions[NumDims-1] - index - 1);
      } else {
        inputIndex += index;
      }
    }
    return inputIndex;
  }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(
      Index index) const  {
    return m_impl.coeff(reverseIndex(index));
  }

  template<int LoadMode>
  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
  PacketReturnType packet(Index index) const
  {
    EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
    eigen_assert(index+PacketSize-1 < dimensions().TotalSize());

    // TODO(ndjaitly): write a better packing routine that uses
    // local structure.
    EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type
                                                            values[PacketSize];
    EIGEN_UNROLL_LOOP
    for (int i = 0; i < PacketSize; ++i) {
      values[i] = coeff(index+i);
    }
    PacketReturnType rslt = internal::pload<PacketReturnType>(values);
    return rslt;
  }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
  internal::TensorBlockResourceRequirements getResourceRequirements() const {
    const size_t target_size = m_device.lastLevelCacheSize();
    // Block evaluation reads underlying memory in reverse order, and default
    // cost model does not properly catch this in bytes stored/loaded.
    return internal::TensorBlockResourceRequirements::skewed<Scalar>(
               target_size)
        .addCostPerCoeff({0, 0, 24});
  }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorBlock
  block(TensorBlockDesc& desc, TensorBlockScratch& scratch,
          bool /*root_of_expr_ast*/ = false) const {
    // TODO(ezhulenev): If underlying tensor expression supports and prefers
    // block evaluation we must use it. Currently we use coeff and packet
    // access into the underlying tensor expression.
    // static const bool useBlockAccessForArgType =
    //     TensorEvaluator<ArgType, Device>::BlockAccess &&
    //     TensorEvaluator<ArgType, Device>::PreferBlockAccess;

    static const bool isColMajor =
        static_cast<int>(Layout) == static_cast<int>(ColMajor);

    static const Index inner_dim_idx = isColMajor ? 0 : NumDims - 1;
    const bool inner_dim_reversed = m_reverse[inner_dim_idx];

    // Offset in the output block.
    Index block_offset = 0;

    // Offset in the input Tensor.
    Index input_offset = reverseIndex(desc.offset());

    // Initialize output block iterator state. Dimension in this array are
    // always in inner_most -> outer_most order (col major layout).
    array<BlockIteratorState, NumDims> it;
    for (int i = 0; i < NumDims; ++i) {
      const int dim = isColMajor ? i : NumDims - 1 - i;
      it[i].size = desc.dimension(dim);
      it[i].count = 0;
      it[i].reverse = m_reverse[dim];

      it[i].block_stride =
          i == 0 ? 1 : (it[i - 1].size * it[i - 1].block_stride);
      it[i].block_span = it[i].block_stride * (it[i].size - 1);

      it[i].input_stride = m_strides[dim];
      it[i].input_span = it[i].input_stride * (it[i].size - 1);

      if (it[i].reverse) {
        it[i].input_stride = -1 * it[i].input_stride;
        it[i].input_span = -1 * it[i].input_span;
      }
    }

    // If multiple inner dimensions have the same reverse flag, check if we can
    // merge them into a single virtual inner dimension.
    int effective_inner_dim = 0;
    for (int i = 1; i < NumDims; ++i) {
      if (it[i].reverse != it[effective_inner_dim].reverse) break;
      if (it[i].block_stride != it[effective_inner_dim].size) break;
      if (it[i].block_stride != numext::abs(it[i].input_stride)) break;

      it[i].size = it[effective_inner_dim].size * it[i].size;

      it[i].block_stride = 1;
      it[i].input_stride = (inner_dim_reversed ? -1 : 1);

      it[i].block_span = it[i].block_stride * (it[i].size - 1);
      it[i].input_span = it[i].input_stride * (it[i].size - 1);

      effective_inner_dim = i;
    }

    eigen_assert(it[effective_inner_dim].block_stride == 1);
    eigen_assert(it[effective_inner_dim].input_stride ==
                 (inner_dim_reversed ? -1 : 1));

    const Index inner_dim_size = it[effective_inner_dim].size;

    // Prepare storage for the materialized reverse result.
    const typename TensorBlock::Storage block_storage =
        TensorBlock::prepareStorage(desc, scratch);
    CoeffReturnType* block_buffer = block_storage.data();

    while (it[NumDims - 1].count < it[NumDims - 1].size) {
      // Copy inner-most dimension data from reversed location in input.
      Index dst = block_offset;
      Index src = input_offset;

      // NOTE(ezhulenev): Adding vectorized path with internal::preverse showed
      // worse results in benchmarks than a simple coefficient loop.
      if (inner_dim_reversed) {
        for (Index i = 0; i < inner_dim_size; ++i) {
          block_buffer[dst] = m_impl.coeff(src);
          ++dst;
          --src;
        }
      } else {
        for (Index i = 0; i < inner_dim_size; ++i) {
          block_buffer[dst] = m_impl.coeff(src);
          ++dst;
          ++src;
        }
      }

      // For the 1d tensor we need to generate only one inner-most dimension.
      if ((NumDims - effective_inner_dim) == 1) break;

      // Update offset.
      for (Index i = effective_inner_dim + 1; i < NumDims; ++i) {
        if (++it[i].count < it[i].size) {
          block_offset += it[i].block_stride;
          input_offset += it[i].input_stride;
          break;
        }
        if (i != NumDims - 1) it[i].count = 0;
        block_offset -= it[i].block_span;
        input_offset -= it[i].input_span;
      }
    }

    return block_storage.AsTensorMaterializedBlock();
  }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const {
    double compute_cost = NumDims * (2 * TensorOpCost::AddCost<Index>() +
                                     2 * TensorOpCost::MulCost<Index>() +
                                     TensorOpCost::DivCost<Index>());
    for (int i = 0; i < NumDims; ++i) {
      if (m_reverse[i]) {
        compute_cost += 2 * TensorOpCost::AddCost<Index>();
      }
    }
    return m_impl.costPerCoeff(vectorized) +
           TensorOpCost(0, 0, compute_cost, false /* vectorized */, PacketSize);
  }

  EIGEN_DEVICE_FUNC typename Storage::Type data() const { return NULL; }

#ifdef EIGEN_USE_SYCL
  // binding placeholder accessors to a command group handler for SYCL
  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void bind(cl::sycl::handler &cgh) const {
    m_impl.bind(cgh);
  }
#endif

 protected:
  Dimensions m_dimensions;
  array<Index, NumDims> m_strides;
  array<IndexDivisor, NumDims> m_fastStrides;
  TensorEvaluator<ArgType, Device> m_impl;
  ReverseDimensions m_reverse;
  const Device EIGEN_DEVICE_REF m_device;

 private:
  struct BlockIteratorState {
    BlockIteratorState()
        : size(0),
          count(0),
          reverse(false),
          block_stride(0),
          block_span(0),
          input_stride(0),
          input_span(0) {}

    Index size;
    Index count;
    bool reverse;
    Index block_stride;
    Index block_span;
    Index input_stride;
    Index input_span;
  };
};

// Eval as lvalue

template <typename ReverseDimensions, typename ArgType, typename Device>
struct TensorEvaluator<TensorReverseOp<ReverseDimensions, ArgType>, Device>
    : public TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>,
                             Device> {
  typedef TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>,
                          Device> Base;
  typedef TensorReverseOp<ReverseDimensions, ArgType> XprType;
  typedef typename XprType::Index Index;
  static const int NumDims = internal::array_size<ReverseDimensions>::value;
  typedef DSizes<Index, NumDims> Dimensions;

  enum {
    IsAligned = false,
    PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
    BlockAccess = false,
    PreferBlockAccess = false,
    Layout = TensorEvaluator<ArgType, Device>::Layout,
    CoordAccess = false,  // to be implemented
    RawAccess = false
  };
  EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
      : Base(op, device) {}

  typedef typename XprType::Scalar Scalar;
  typedef typename XprType::CoeffReturnType CoeffReturnType;
  typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
  static const int PacketSize = PacketType<CoeffReturnType, Device>::size;

  //===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
  typedef internal::TensorBlockNotImplemented TensorBlock;
  //===--------------------------------------------------------------------===//

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
  const Dimensions& dimensions() const { return this->m_dimensions; }

  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index index) {
    return this->m_impl.coeffRef(this->reverseIndex(index));
  }

  template <int StoreMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
  void writePacket(Index index, const PacketReturnType& x) {
    EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
    eigen_assert(index+PacketSize-1 < dimensions().TotalSize());

    // This code is pilfered from TensorMorphing.h
    EIGEN_ALIGN_MAX CoeffReturnType values[PacketSize];
    internal::pstore<CoeffReturnType, PacketReturnType>(values, x);
    EIGEN_UNROLL_LOOP
    for (int i = 0; i < PacketSize; ++i) {
      this->coeffRef(index+i) = values[i];
    }
  }
};


}  // end namespace Eigen

#endif // EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H