TheRock/rocm-libraries/projects/rocthrust/thrust/iterator/iterator_facade.h

/*
 *  Copyright 2008-2013 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

/*! \file thrust/iterator/iterator_facade.h
 *  \brief A class which exposes a public interface for iterators
 */

/*
 * (C) Copyright David Abrahams 2002.
 * (C) Copyright Jeremy Siek    2002.
 * (C) Copyright Thomas Witt    2002.
 *
 * Distributed under the Boost Software License, Version 1.0.
 * (See accompanying NOTICE file for the complete license)
 *
 * For more information, see http://www.boost.org
 */

#pragma once

#include <thrust/detail/config.h>

#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
#  pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
#  pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
#  pragma system_header
#endif // no system header
#include <thrust/detail/type_traits.h>
#include <thrust/iterator/detail/distance_from_result.h>
#include <thrust/iterator/detail/iterator_facade_category.h>

#if !_THRUST_HAS_DEVICE_SYSTEM_STD
#  include <cstddef>
#  include <type_traits>
#endif

THRUST_NAMESPACE_BEGIN

/*! \addtogroup iterators
 *  \{
 */

/*! \addtogroup fancyiterator Fancy Iterators
 *  \ingroup iterators
 *  \{
 */

// This forward declaration is required for the friend declaration
// in iterator_core_access
template <typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference>
class iterator_facade;

/*! \p iterator_core_access is the class which user iterator types derived from \p thrust::iterator_adaptor
 *  or \p thrust::iterator_facade must befriend to allow it to access their private interface.
 */
class iterator_core_access
{
  /*! \cond
   */

  // declare our friends
  template <typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference>
  friend class iterator_facade;

  // iterator comparisons are our friends
  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend bool
  operator==(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend bool
  operator!=(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend bool
  operator<(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
            iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend bool
  operator>(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
            iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend bool
  operator<=(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend bool
  operator>=(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  // iterator difference is our friend
  template <typename Derived1,
            typename Value1,
            typename System1,
            typename Traversal1,
            typename Reference1,
            typename Difference1,
            typename Derived2,
            typename Value2,
            typename System2,
            typename Traversal2,
            typename Reference2,
            typename Difference2>
  inline THRUST_HOST_DEVICE friend typename thrust::detail::distance_from_result<
    iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1>,
    iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2>>::type
  operator-(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
            iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs);

  template <typename Facade>
  THRUST_HOST_DEVICE static typename Facade::reference dereference(Facade const& f)
  {
    return f.dereference();
  }

  template <typename Facade>
  THRUST_HOST_DEVICE static void increment(Facade& f)
  {
    f.increment();
  }

  template <typename Facade>
  THRUST_HOST_DEVICE static void decrement(Facade& f)
  {
    f.decrement();
  }

  template <class Facade1, class Facade2>
  THRUST_HOST_DEVICE static bool equal(Facade1 const& f1, Facade2 const& f2)
  {
    return f1.equal(f2);
  }

  // XXX TODO: Investigate whether we need both of these cases
  // template <class Facade1, class Facade2>
  //__host__ __device__
  // static bool equal(Facade1 const& f1, Facade2 const& f2, mpl::true_)
  //{
  //  return f1.equal(f2);
  //}

  // template <class Facade1, class Facade2>
  //__host__ __device__
  // static bool equal(Facade1 const& f1, Facade2 const& f2, mpl::false_)
  //{
  //   return f2.equal(f1);
  // }

  template <class Facade>
  THRUST_HOST_DEVICE static void advance(Facade& f, typename Facade::difference_type n)
  {
    f.advance(n);
  }

  // Facade2 is convertible to Facade1,
  // so return Facade1's difference_type
  template <class Facade1, class Facade2>
  THRUST_HOST_DEVICE static typename Facade1::difference_type
  distance_from(Facade1 const& f1, Facade2 const& f2, thrust::detail::true_type)
  {
    return -f1.distance_to(f2);
  }

  // Facade2 is not convertible to Facade1,
  // so return Facade2's difference_type
  template <class Facade1, class Facade2>
  THRUST_HOST_DEVICE static typename Facade2::difference_type
  distance_from(Facade1 const& f1, Facade2 const& f2, thrust::detail::false_type)
  {
    return f2.distance_to(f1);
  }

  template <class Facade1, class Facade2>
  THRUST_HOST_DEVICE static typename thrust::detail::distance_from_result<Facade1, Facade2>::type
  distance_from(Facade1 const& f1, Facade2 const& f2)
  {
    // dispatch the implementation of this method upon whether or not
    // Facade2 is convertible to Facade1
    return distance_from(f1, f2, typename _THRUST_STD::is_convertible<Facade2, Facade1>::type());
  }

  //
  // Curiously Recurring Template interface.
  //
  template <typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference>
  THRUST_HOST_DEVICE static Derived&
  derived(iterator_facade<Derived, Value, System, Traversal, Reference, Difference>& facade)
  {
    return *static_cast<Derived*>(&facade);
  }

  template <typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference>
  THRUST_HOST_DEVICE static Derived const&
  derived(iterator_facade<Derived, Value, System, Traversal, Reference, Difference> const& facade)
  {
    return *static_cast<Derived const*>(&facade);
  }

  /*! \endcond
   */
}; // end iterator_core_access

/*! \p iterator_facade is a template which allows the programmer to define a novel iterator with a standards-conforming
 * interface which Thrust can use to reason about algorithm acceleration opportunities.
 *
 *  Because most of a standard iterator's interface is defined in terms of a small set of core primitives, \p
 * iterator_facade defines the non-primitive portion mechanically. In principle a novel iterator could explicitly
 * provide the entire interface in an ad hoc fashion but doing so might be tedious and prone to subtle errors.
 *
 *  Often \p iterator_facade is too primitive a tool to use for defining novel iterators. In these cases, \p
 * iterator_adaptor or a specific fancy iterator should be used instead.
 *
 *  \p iterator_facade's functionality is derived from and generally equivalent to \p boost::iterator_facade.
 *  The exception is Thrust's addition of the template parameter \p System, which is necessary to allow Thrust
 *  to dispatch an algorithm to one of several parallel backend systems. An additional exception is Thrust's omission
 *  of the \c operator-> member function.
 *
 *  Interested users may refer to <tt>boost::iterator_facade</tt>'s documentation for usage examples.
 *
 *  \note \p iterator_facade's arithmetic operator free functions exist with the usual meanings but are omitted here for
 * brevity.
 */
template <typename Derived,
          typename Value,
          typename System,
          typename Traversal,
          typename Reference,
          typename Difference = _THRUST_STD::ptrdiff_t>
class iterator_facade
{
private:
  /*! \cond
   */

  //
  // Curiously Recurring Template interface.
  //
  THRUST_HOST_DEVICE Derived& derived()
  {
    return *static_cast<Derived*>(this);
  }

  THRUST_HOST_DEVICE Derived const& derived() const
  {
    return *static_cast<Derived const*>(this);
  }
  /*! \endcond
   */

public:
  /*! The type of element pointed to by \p iterator_facade.
   */
  using value_type = _THRUST_STD::remove_const_t<Value>;

  /*! The return type of \p iterator_facade::operator*().
   */
  using reference = Reference;

  /*! The return type of \p iterator_facade's non-existent \c operator->()
   *  member function. Unlike \c boost::iterator_facade, \p iterator_facade
   *  disallows access to the \p value_type's members through expressions of the
   *  form <tt>iter->member</tt>. \p pointer is defined to \c void to indicate
   *  that these expressions are not allowed. This limitation may be relaxed in a
   *  future version of Thrust.
   */
  using pointer = void;

  /*! The type of expressions of the form <tt>x - y</tt> where <tt>x</tt> and <tt>y</tt>
   *  are of type \p iterator_facade.
   */
  using difference_type = Difference;

  /*! The type of iterator category of \p iterator_facade.
   */
  using iterator_category =
    typename thrust::detail::iterator_facade_category<System, Traversal, Value, Reference>::type;

  /*! \p operator*() dereferences this \p iterator_facade.
   *  \return A reference to the element pointed to by this \p iterator_facade.
   */
  THRUST_HOST_DEVICE reference operator*() const
  {
    return iterator_core_access::dereference(this->derived());
  }

  // XXX unimplemented for now, consider implementing it later
  // pointer operator->() const
  //{
  //  return;
  //}

  // XXX investigate whether or not we need to go to the lengths
  //     boost does to determine the return type

  /*! \p operator[] performs indexed dereference.
   *  \return A reference to the element \p n distance away from this \p iterator_facade.
   */
  THRUST_HOST_DEVICE reference operator[](difference_type n) const
  {
    return *(this->derived() + n);
  }

  /*! \p operator++ pre-increments this \p iterator_facade to refer to the element in the next position.
   *  \return <tt>*this</tt>
   */
  THRUST_HOST_DEVICE Derived& operator++()
  {
    iterator_core_access::increment(this->derived());
    return this->derived();
  }

  /*! \p operator++ post-increments this \p iterator_facade and returns a new \p iterator_facade referring to the
   * element in the next position. \return A copy of <tt>*this</tt> before increment.
   */
  THRUST_HOST_DEVICE Derived operator++(int)
  {
    Derived tmp(this->derived());
    ++*this;
    return tmp;
  }

  /*! \p operator-- pre-decrements this \p iterator_facade to refer to the element in the previous position.
   *  \return <tt>*this</tt>
   */
  THRUST_HOST_DEVICE Derived& operator--()
  {
    iterator_core_access::decrement(this->derived());
    return this->derived();
  }

  /*! \p operator-- post-decrements this \p iterator_facade and returns a new \p iterator_facade referring to the
   * element in the previous position. \return A copy of <tt>*this</tt> before decrement.
   */
  THRUST_HOST_DEVICE Derived operator--(int)
  {
    Derived tmp(this->derived());
    --*this;
    return tmp;
  }

  /*! \p operator+= increments this \p iterator_facade to refer to an element a given distance after its current
   * position. \param n The quantity to increment. \return <tt>*this</tt>
   */
  THRUST_HOST_DEVICE Derived& operator+=(difference_type n)
  {
    iterator_core_access::advance(this->derived(), n);
    return this->derived();
  }

  /*! \p operator-= decrements this \p iterator_facade to refer to an element a given distance before its current
   * position. \param n The quantity to decrement. \return <tt>*this</tt>
   */
  THRUST_HOST_DEVICE Derived& operator-=(difference_type n)
  {
    iterator_core_access::advance(this->derived(), -n);
    return this->derived();
  }

  /*! \p operator- subtracts a given quantity from this \p iterator_facade and returns a new \p iterator_facade
   * referring to the element at the given position before this \p iterator_facade. \param n The quantity to decrement
   *  \return An \p iterator_facade pointing \p n elements before this \p iterator_facade.
   */
  THRUST_HOST_DEVICE Derived operator-(difference_type n) const
  {
    Derived result(this->derived());
    return result -= n;
  }
}; // end iterator_facade

/*! \cond
 */

// Comparison operators
template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE
  // XXX it might be nice to implement this at some point
  // typename enable_if_interoperable<Dr1,Dr2,bool>::type // exposition
  bool
  operator==(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return iterator_core_access ::equal(*static_cast<Derived1 const*>(&lhs), *static_cast<Derived2 const*>(&rhs));
}

template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE
  // XXX it might be nice to implement this at some point
  // typename enable_if_interoperable<Dr1,Dr2,bool>::type // exposition
  bool
  operator!=(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return !iterator_core_access ::equal(*static_cast<Derived1 const*>(&lhs), *static_cast<Derived2 const*>(&rhs));
}

template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE
  // XXX it might be nice to implement this at some point
  // typename enable_if_interoperable<Dr1,Dr2,bool>::type // exposition
  bool
  operator<(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
            iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return 0
       > iterator_core_access ::distance_from(*static_cast<Derived1 const*>(&lhs), *static_cast<Derived2 const*>(&rhs));
}

template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE
  // XXX it might be nice to implement this at some point
  // typename enable_if_interoperable<Dr1,Dr2,bool>::type // exposition
  bool
  operator>(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
            iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return 0
       < iterator_core_access ::distance_from(*static_cast<Derived1 const*>(&lhs), *static_cast<Derived2 const*>(&rhs));
}

template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE
  // XXX it might be nice to implement this at some point
  // typename enable_if_interoperable<Dr1,Dr2,bool>::type // exposition
  bool
  operator<=(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return 0
      >= iterator_core_access ::distance_from(*static_cast<Derived1 const*>(&lhs), *static_cast<Derived2 const*>(&rhs));
}

template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE
  // XXX it might be nice to implement this at some point
  // typename enable_if_interoperable<Dr1,Dr2,bool>::type // exposition
  bool
  operator>=(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
             iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return 0
      <= iterator_core_access ::distance_from(*static_cast<Derived1 const*>(&lhs), *static_cast<Derived2 const*>(&rhs));
}

// Iterator difference
template <typename Derived1,
          typename Value1,
          typename System1,
          typename Traversal1,
          typename Reference1,
          typename Difference1,
          typename Derived2,
          typename Value2,
          typename System2,
          typename Traversal2,
          typename Reference2,
          typename Difference2>
inline THRUST_HOST_DEVICE

  // divine the type this operator returns
  typename thrust::detail::distance_from_result<
    iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1>,
    iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2>>::type

  operator-(iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> const& lhs,
            iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> const& rhs)
{
  return iterator_core_access ::distance_from(static_cast<Derived1 const&>(lhs), static_cast<Derived2 const&>(rhs));
}

// Iterator addition
template <typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference>
inline THRUST_HOST_DEVICE Derived
operator+(iterator_facade<Derived, Value, System, Traversal, Reference, Difference> const& i,
          typename Derived::difference_type n)
{
  Derived tmp(static_cast<Derived const&>(i));
  return tmp += n;
}

template <typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference>
inline THRUST_HOST_DEVICE Derived
operator+(typename Derived::difference_type n,
          iterator_facade<Derived, Value, System, Traversal, Reference, Difference> const& i)
{
  Derived tmp(static_cast<Derived const&>(i));
  return tmp += n;
}

/*! \endcond
 */

/*! \} // end fancyiterators
 */

/*! \} // end iterators
 */

THRUST_NAMESPACE_END