/* * 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 #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 #include #include 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 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 friend class iterator_facade; // iterator comparisons are our friends template inline _CCCL_HOST_DEVICE friend bool operator==(iterator_facade const& lhs, iterator_facade const& rhs); template inline _CCCL_HOST_DEVICE friend bool operator!=(iterator_facade const& lhs, iterator_facade const& rhs); template inline _CCCL_HOST_DEVICE friend bool operator<(iterator_facade const& lhs, iterator_facade const& rhs); template inline _CCCL_HOST_DEVICE friend bool operator>(iterator_facade const& lhs, iterator_facade const& rhs); template inline _CCCL_HOST_DEVICE friend bool operator<=(iterator_facade const& lhs, iterator_facade const& rhs); template inline _CCCL_HOST_DEVICE friend bool operator>=(iterator_facade const& lhs, iterator_facade const& rhs); // iterator difference is our friend template inline _CCCL_HOST_DEVICE friend typename thrust::detail::distance_from_result< iterator_facade, iterator_facade>::type operator-(iterator_facade const& lhs, iterator_facade const& rhs); template _CCCL_HOST_DEVICE static typename Facade::reference dereference(Facade const& f) { return f.dereference(); } template _CCCL_HOST_DEVICE static void increment(Facade& f) { f.increment(); } template _CCCL_HOST_DEVICE static void decrement(Facade& f) { f.decrement(); } template _CCCL_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 //__host__ __device__ // static bool equal(Facade1 const& f1, Facade2 const& f2, mpl::true_) //{ // return f1.equal(f2); //} // template //__host__ __device__ // static bool equal(Facade1 const& f1, Facade2 const& f2, mpl::false_) //{ // return f2.equal(f1); // } template _CCCL_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 _CCCL_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 _CCCL_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 _CCCL_HOST_DEVICE static typename thrust::detail::distance_from_result::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 ::cuda::std::is_convertible::type()); } // // Curiously Recurring Template interface. // template _CCCL_HOST_DEVICE static Derived& derived(iterator_facade& facade) { return *static_cast(&facade); } template _CCCL_HOST_DEVICE static Derived const& derived(iterator_facade const& facade) { return *static_cast(&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 boost::iterator_facade'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 class iterator_facade { private: /*! \cond */ // // Curiously Recurring Template interface. // _CCCL_HOST_DEVICE Derived& derived() { return *static_cast(this); } _CCCL_HOST_DEVICE Derived const& derived() const { return *static_cast(this); } /*! \endcond */ public: /*! The type of element pointed to by \p iterator_facade. */ typedef ::cuda::std::__remove_const_t value_type; /*! The return type of \p iterator_facade::operator*(). */ typedef 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 iter->member. \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. */ typedef void pointer; /*! The type of expressions of the form x - y where x and y * are of type \p iterator_facade. */ typedef Difference difference_type; /*! The type of iterator category of \p iterator_facade. */ typedef typename thrust::detail::iterator_facade_category::type iterator_category; /*! \p operator*() dereferences this \p iterator_facade. * \return A reference to the element pointed to by this \p iterator_facade. */ _CCCL_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. */ _CCCL_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 *this */ _CCCL_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 *this before increment. */ _CCCL_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 *this */ _CCCL_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 *this before decrement. */ _CCCL_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 *this */ _CCCL_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 * postition. \param n The quantity to decrement. \return *this */ _CCCL_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. */ _CCCL_HOST_DEVICE Derived operator-(difference_type n) const { Derived result(this->derived()); return result -= n; } }; // end iterator_facade /*! \cond */ // Comparison operators template inline _CCCL_HOST_DEVICE // XXX it might be nice to implement this at some point // typename enable_if_interoperable::type // exposition bool operator==(iterator_facade const& lhs, iterator_facade const& rhs) { return iterator_core_access ::equal(*static_cast(&lhs), *static_cast(&rhs)); } template inline _CCCL_HOST_DEVICE // XXX it might be nice to implement this at some point // typename enable_if_interoperable::type // exposition bool operator!=(iterator_facade const& lhs, iterator_facade const& rhs) { return !iterator_core_access ::equal(*static_cast(&lhs), *static_cast(&rhs)); } template inline _CCCL_HOST_DEVICE // XXX it might be nice to implement this at some point // typename enable_if_interoperable::type // exposition bool operator<(iterator_facade const& lhs, iterator_facade const& rhs) { return 0 > iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); } template inline _CCCL_HOST_DEVICE // XXX it might be nice to implement this at some point // typename enable_if_interoperable::type // exposition bool operator>(iterator_facade const& lhs, iterator_facade const& rhs) { return 0 < iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); } template inline _CCCL_HOST_DEVICE // XXX it might be nice to implement this at some point // typename enable_if_interoperable::type // exposition bool operator<=(iterator_facade const& lhs, iterator_facade const& rhs) { return 0 >= iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); } template inline _CCCL_HOST_DEVICE // XXX it might be nice to implement this at some point // typename enable_if_interoperable::type // exposition bool operator>=(iterator_facade const& lhs, iterator_facade const& rhs) { return 0 <= iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); } // Iterator difference template inline _CCCL_HOST_DEVICE // divine the type this operator returns typename thrust::detail::distance_from_result< iterator_facade, iterator_facade>::type operator-(iterator_facade const& lhs, iterator_facade const& rhs) { return iterator_core_access ::distance_from(static_cast(lhs), static_cast(rhs)); } // Iterator addition template inline _CCCL_HOST_DEVICE Derived operator+(iterator_facade const& i, typename Derived::difference_type n) { Derived tmp(static_cast(i)); return tmp += n; } template inline _CCCL_HOST_DEVICE Derived operator+(typename Derived::difference_type n, iterator_facade const& i) { Derived tmp(static_cast(i)); return tmp += n; } /*! \endcond */ /*! \} // end fancyiterators */ /*! \} // end iterators */ THRUST_NAMESPACE_END