From 92d9b54bd41d635ed145bbc692fe1ecb81ae3b5e Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 18 Feb 2025 13:00:48 +0100 Subject: [PATCH] Refactor Thrust iterators 3/4 (#3842) * Refactor reverse_iterator * Refactor tabulate_output_iterator.h --- .../iterator/detail/reverse_iterator.inl | 85 ----- .../iterator/detail/reverse_iterator_base.h | 48 --- .../detail/tabulate_output_iterator.inl | 69 ---- thrust/thrust/iterator/reverse_iterator.h | 311 +++++++++--------- .../iterator/tabulate_output_iterator.h | 177 ++++++---- 5 files changed, 265 insertions(+), 425 deletions(-) delete mode 100644 thrust/thrust/iterator/detail/reverse_iterator.inl delete mode 100644 thrust/thrust/iterator/detail/reverse_iterator_base.h delete mode 100644 thrust/thrust/iterator/detail/tabulate_output_iterator.inl diff --git a/thrust/thrust/iterator/detail/reverse_iterator.inl b/thrust/thrust/iterator/detail/reverse_iterator.inl deleted file mode 100644 index 4a69bb10abf..00000000000 --- a/thrust/thrust/iterator/detail/reverse_iterator.inl +++ /dev/null @@ -1,85 +0,0 @@ -/* - * Copyright 2008-2021 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. - */ - -#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 - -THRUST_NAMESPACE_BEGIN - -namespace detail -{ - -_CCCL_EXEC_CHECK_DISABLE -template -_CCCL_HOST_DEVICE Iterator prior(Iterator x) -{ - return --x; -} // end prior() - -} // namespace detail - -template -_CCCL_HOST_DEVICE typename reverse_iterator::super_t::reference -reverse_iterator::dereference() const -{ - return *thrust::detail::prior(this->base()); -} // end reverse_iterator::increment() - -template -_CCCL_HOST_DEVICE void reverse_iterator::increment() -{ - --this->base_reference(); -} // end reverse_iterator::increment() - -template -_CCCL_HOST_DEVICE void reverse_iterator::decrement() -{ - ++this->base_reference(); -} // end reverse_iterator::decrement() - -template -_CCCL_HOST_DEVICE void reverse_iterator::advance(typename super_t::difference_type n) -{ - this->base_reference() += -n; -} // end reverse_iterator::advance() - -template -template -_CCCL_HOST_DEVICE typename reverse_iterator::super_t::difference_type -reverse_iterator::distance_to(reverse_iterator const& y) const -{ - return this->base_reference() - y.base(); -} // end reverse_iterator::distance_to() - -template -_CCCL_HOST_DEVICE reverse_iterator make_reverse_iterator(BidirectionalIterator x) -{ - return reverse_iterator(x); -} // end make_reverse_iterator() - -THRUST_NAMESPACE_END diff --git a/thrust/thrust/iterator/detail/reverse_iterator_base.h b/thrust/thrust/iterator/detail/reverse_iterator_base.h deleted file mode 100644 index 45c3daa133c..00000000000 --- a/thrust/thrust/iterator/detail/reverse_iterator_base.h +++ /dev/null @@ -1,48 +0,0 @@ -/* - * 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. - */ - -#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 - -THRUST_NAMESPACE_BEGIN - -template -class reverse_iterator; - -namespace detail -{ - -template -struct reverse_iterator_base -{ - using type = thrust::iterator_adaptor, BidirectionalIterator>; -}; // end reverse_iterator_base - -} // namespace detail - -THRUST_NAMESPACE_END diff --git a/thrust/thrust/iterator/detail/tabulate_output_iterator.inl b/thrust/thrust/iterator/detail/tabulate_output_iterator.inl deleted file mode 100644 index 56c093ff56a..00000000000 --- a/thrust/thrust/iterator/detail/tabulate_output_iterator.inl +++ /dev/null @@ -1,69 +0,0 @@ -// SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#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 - -template -class tabulate_output_iterator; - -namespace detail -{ - -// Proxy reference that invokes a BinaryFunction with the index of the dereferenced iterator and the assigned value -template -class tabulate_output_iterator_proxy -{ -public: - _CCCL_HOST_DEVICE tabulate_output_iterator_proxy(BinaryFunction fun, DifferenceT index) - : fun(fun) - , index(index) - {} - - _CCCL_EXEC_CHECK_DISABLE - template - _CCCL_HOST_DEVICE tabulate_output_iterator_proxy operator=(const T& x) - { - fun(index, x); - return *this; - } - -private: - BinaryFunction fun; - DifferenceT index; -}; - -// Alias template for the iterator_adaptor instantiation to be used for tabulate_output_iterator -template -using tabulate_output_iterator_base = - thrust::iterator_adaptor, - counting_iterator, - void, - System, - thrust::use_default, - tabulate_output_iterator_proxy>; - -// Register tabulate_output_iterator_proxy with 'is_proxy_reference' from -// type_traits to enable its use with algorithms. -template -struct is_proxy_reference> - : public thrust::detail::true_type -{}; - -} // namespace detail -THRUST_NAMESPACE_END diff --git a/thrust/thrust/iterator/reverse_iterator.h b/thrust/thrust/iterator/reverse_iterator.h index eb94b6eeb31..077f32be509 100644 --- a/thrust/thrust/iterator/reverse_iterator.h +++ b/thrust/thrust/iterator/reverse_iterator.h @@ -14,9 +14,8 @@ * limitations under the License. */ -/*! \file thrust/iterator/reverse_iterator.h - * \brief An iterator adaptor which adapts another iterator to traverse backwards - */ +//! \file thrust/iterator/reverse_iterator.h +//! \brief An iterator adaptor which adapts another iterator to traverse backwards /* * (C) Copyright David Abrahams 2002. @@ -40,186 +39,196 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header + #include -#include -#include +#include +#include THRUST_NAMESPACE_BEGIN -/*! \addtogroup iterators - * \{ - */ - -/*! \addtogroup fancyiterator Fancy Iterators - * \ingroup iterators - * \{ - */ +template +class reverse_iterator; -/*! \p reverse_iterator is an iterator which represents a pointer into a - * reversed view of a given range. In this way, \p reverse_iterator allows - * backwards iteration through a bidirectional input range. - * - * It is important to note that although \p reverse_iterator is constructed - * from a given iterator, it points to the element preceding it. In this way, - * the past-the-end \p reverse_iterator of a given range points to the element - * preceding the first element of the input range. By the same token, the first - * \p reverse_iterator of a given range is constructed from a past-the-end iterator - * of the original range yet points to the last element of the input. - * - * The following code snippet demonstrates how to create a \p reverse_iterator - * which represents a reversed view of the contents of a \p device_vector. - * - * \code - * #include - * #include - * ... - * thrust::device_vector v(4); - * v[0] = 0.0f; - * v[1] = 1.0f; - * v[2] = 2.0f; - * v[3] = 3.0f; - * - * using Iterator = thrust::device_vector::iterator; - * - * // note that we point the iterator to the *end* of the device_vector - * thrust::reverse_iterator iter(values.end()); - * - * *iter; // returns 3.0f; - * iter[0]; // returns 3.0f; - * iter[1]; // returns 2.0f; - * iter[2]; // returns 1.0f; - * iter[3]; // returns 0.0f; - * - * // iter[4] is an out-of-bounds error - * \endcode - * - * Since reversing a range is a common operation, containers like \p device_vector - * have nested aliases for declaration shorthand and methods for constructing - * reverse_iterators. The following code snippet is equivalent to the previous: - * - * \code - * #include - * ... - * thrust::device_vector v(4); - * v[0] = 0.0f; - * v[1] = 1.0f; - * v[2] = 2.0f; - * v[3] = 3.0f; - * - * // we use the nested type reverse_iterator to refer to a reversed view of - * // a device_vector and the method rbegin() to create a reverse_iterator pointing - * // to the beginning of the reversed device_vector - * thrust::device_iterator::reverse_iterator iter = values.rbegin(); - * - * *iter; // returns 3.0f; - * iter[0]; // returns 3.0f; - * iter[1]; // returns 2.0f; - * iter[2]; // returns 1.0f; - * iter[3]; // returns 0.0f; - * - * // iter[4] is an out-of-bounds error - * - * // similarly, rend() points to the end of the reversed sequence: - * assert(values.rend() == (iter + 4)); - * \endcode - * - * Finally, the following code snippet demonstrates how to use reverse_iterator to - * perform a reversed prefix sum operation on the contents of a device_vector: - * - * \code - * #include - * #include - * ... - * thrust::device_vector v(5); - * v[0] = 0; - * v[1] = 1; - * v[2] = 2; - * v[3] = 3; - * v[4] = 4; - * - * thrust::device_vector result(5); - * - * // exclusive scan v into result in reverse - * thrust::exclusive_scan(v.rbegin(), v.rend(), result.begin()); - * - * // result is now {0, 4, 7, 9, 10} - * \endcode - * - * \see make_reverse_iterator - */ +namespace detail +{ +template +struct make_reverse_iterator_base +{ + using type = iterator_adaptor, BidirectionalIterator>; +}; +} // namespace detail + +//! \addtogroup iterators +//! \{ + +//! \addtogroup fancyiterator Fancy Iterators +//! \ingroup iterators +//! \{ + +//! \p reverse_iterator is an iterator which represents a pointer into a reversed view of a given range. In this way, \p +//! reverse_iterator allows backwards iteration through a bidirectional input range. +//! +//! It is important to note that although \p reverse_iterator is constructed from a given iterator, it points to the +//! element preceding it. In this way, the past-the-end \p reverse_iterator of a given range points to the element +//! preceding the first element of the input range. By the same token, the first \p reverse_iterator of a given range is +//! constructed from a past-the-end iterator of the original range yet points to the last element of the input. +//! +//! The following code snippet demonstrates how to create a \p reverse_iterator which represents a reversed view of the +//! contents of a \p device_vector. +//! +//! \code +//! #include +//! #include +//! ... +//! thrust::device_vector v(4); +//! v[0] = 0.0f; +//! v[1] = 1.0f; +//! v[2] = 2.0f; +//! v[3] = 3.0f; +//! +//! using Iterator = thrust::device_vector::iterator; +//! +//! // note that we point the iterator to the *end* of the device_vector +//! thrust::reverse_iterator iter(values.end()); +//! +//! *iter; // returns 3.0f; +//! iter[0]; // returns 3.0f; +//! iter[1]; // returns 2.0f; +//! iter[2]; // returns 1.0f; +//! iter[3]; // returns 0.0f; +//! +//! // iter[4] is an out-of-bounds error +//! \endcode +//! +//! Since reversing a range is a common operation, containers like \p device_vector have nested aliases for declaration +//! shorthand and methods for constructing reverse_iterators. The following code snippet is equivalent to the previous: +//! +//! \code +//! #include +//! ... +//! thrust::device_vector v(4); +//! v[0] = 0.0f; +//! v[1] = 1.0f; +//! v[2] = 2.0f; +//! v[3] = 3.0f; +//! +//! // we use the nested type reverse_iterator to refer to a reversed view of a device_vector and the method rbegin() to +//! // create a reverse_iterator pointing to the beginning of the reversed device_vector +//! thrust::device_iterator::reverse_iterator iter = values.rbegin(); +//! +//! *iter; // returns 3.0f; +//! iter[0]; // returns 3.0f; +//! iter[1]; // returns 2.0f; +//! iter[2]; // returns 1.0f; +//! iter[3]; // returns 0.0f; +//! +//! // iter[4] is an out-of-bounds error +//! +//! // similarly, rend() points to the end of the reversed sequence: +//! assert(values.rend() == (iter + 4)); +//! \endcode +//! +//! Finally, the following code snippet demonstrates how to use reverse_iterator to perform a reversed prefix sum +//! operation on the contents of a device_vector: +//! +//! \code +//! #include +//! #include +//! ... +//! thrust::device_vector v(5); +//! v[0] = 0; +//! v[1] = 1; +//! v[2] = 2; +//! v[3] = 3; +//! v[4] = 4; +//! +//! thrust::device_vector result(5); +//! +//! // exclusive scan v into result in reverse +//! thrust::exclusive_scan(v.rbegin(), v.rend(), result.begin()); +//! +//! // result is now {0, 4, 7, 9, 10} +//! \endcode +//! +//! \see make_reverse_iterator template -class reverse_iterator : public detail::reverse_iterator_base::type +class reverse_iterator : public detail::make_reverse_iterator_base::type { - /*! \cond - */ + //! \cond private: - using super_t = typename thrust::detail::reverse_iterator_base::type; + using super_t = typename detail::make_reverse_iterator_base::type; friend class iterator_core_access; - /*! \endcond - */ + //! \endcond public: - /*! Default constructor does nothing. - */ reverse_iterator() = default; - /*! \p Constructor accepts a \c BidirectionalIterator pointing to a range - * for this \p reverse_iterator to reverse. - * - * \param x A \c BidirectionalIterator pointing to a range to reverse. - */ + //! \p Constructor accepts a \c BidirectionalIterator pointing to a range + //! for this \p reverse_iterator to reverse. + //! + //! \param x A \c BidirectionalIterator pointing to a range to reverse. _CCCL_HOST_DEVICE explicit reverse_iterator(BidirectionalIterator x) : super_t(x) {} - /*! \p Copy constructor allows construction from a related compatible - * \p reverse_iterator. - * - * \param r A \p reverse_iterator to copy from. - */ + //! \p Copy constructor allows construction from a related compatible + //! \p reverse_iterator. + //! + //! \param r A \p reverse_iterator to copy from. template = 0> _CCCL_HOST_DEVICE reverse_iterator(reverse_iterator const& rhs) : super_t(rhs.base()) {} - /*! \cond - */ + //! \cond private: _CCCL_EXEC_CHECK_DISABLE - _CCCL_HOST_DEVICE typename super_t::reference dereference() const; - - _CCCL_HOST_DEVICE void increment(); - - _CCCL_HOST_DEVICE void decrement(); - - _CCCL_HOST_DEVICE void advance(typename super_t::difference_type n); + _CCCL_HOST_DEVICE typename super_t::reference dereference() const + { + auto b = this->base(); + return *--b; + } + + _CCCL_HOST_DEVICE void increment() + { + --this->base_reference(); + } + + _CCCL_HOST_DEVICE void decrement() + { + ++this->base_reference(); + } + + _CCCL_HOST_DEVICE void advance(typename super_t::difference_type n) + { + this->base_reference() += -n; + } template _CCCL_HOST_DEVICE typename super_t::difference_type - distance_to(reverse_iterator const& y) const; - /*! \endcond - */ -}; // end reverse_iterator - -/*! \p make_reverse_iterator creates a \p reverse_iterator - * from a \c BidirectionalIterator pointing to a range of elements to reverse. - * - * \param x A \c BidirectionalIterator pointing to a range to reverse. - * \return A new \p reverse_iterator which reverses the range \p x. - */ + distance_to(reverse_iterator const& y) const + { + return this->base_reference() - y.base(); + } + //! \endcond +}; + +//! \p make_reverse_iterator creates a \p reverse_iterator +//! from a \c BidirectionalIterator pointing to a range of elements to reverse. +//! +//! \param x A \c BidirectionalIterator pointing to a range to reverse. +//! \return A new \p reverse_iterator which reverses the range \p x. template -_CCCL_HOST_DEVICE reverse_iterator make_reverse_iterator(BidirectionalIterator x); - -/*! \} // end fancyiterators - */ +_CCCL_HOST_DEVICE reverse_iterator make_reverse_iterator(BidirectionalIterator x) +{ + return reverse_iterator(x); +} -/*! \} // end iterators - */ +//! \} // end fancyiterators +//! \} // end iterators THRUST_NAMESPACE_END - -#include diff --git a/thrust/thrust/iterator/tabulate_output_iterator.h b/thrust/thrust/iterator/tabulate_output_iterator.h index 5f430d7b6eb..c5beb6c710a 100644 --- a/thrust/thrust/iterator/tabulate_output_iterator.h +++ b/thrust/thrust/iterator/tabulate_output_iterator.h @@ -13,105 +13,138 @@ # pragma system_header #endif // no system header -#include +#include + +#include +#include +#include THRUST_NAMESPACE_BEGIN -/*! \addtogroup iterators - * \{ - */ - -/*! \addtogroup fancyiterator Fancy Iterators - * \ingroup iterators - * \{ - */ - -/*! \p tabulate_output_iterator is a special kind of output iterator which, whenever a value is assigned to a - * dereferenced iterator, calls the given callable with the index that corresponds to the offset of the dereferenced - * iterator and the assigned value. - * - * The following code snippet demonstrated how to create a \p tabulate_output_iterator which prints the index and the - * assigned value. - * - * \code - * #include - * - * // note: functor inherits form binary function - * struct print_op - * { - * __host__ __device__ - * void operator()(int index, float value) const - * { - * printf("%d: %f\n", index, value); - * } - * }; - * - * int main() - * { - * auto tabulate_it = thrust::make_tabulate_output_iterator(print_op{}); - * - * tabulate_it[0] = 1.0f; // prints: 0: 1.0 - * tabulate_it[1] = 3.0f; // prints: 1: 3.0 - * tabulate_it[9] = 5.0f; // prints: 9: 5.0 - * } - * \endcode - * - * \see make_tabulate_output_iterator - */ +template +class tabulate_output_iterator; -template -class tabulate_output_iterator : public detail::tabulate_output_iterator_base +namespace detail { - /*! \cond - */ +// Proxy reference that invokes a BinaryFunction with the index of the dereferenced iterator and the assigned value +template +class tabulate_output_iterator_proxy +{ +public: + _CCCL_HOST_DEVICE tabulate_output_iterator_proxy(BinaryFunction fun, DifferenceT index) + : fun(fun) + , index(index) + {} + _CCCL_EXEC_CHECK_DISABLE + template + _CCCL_HOST_DEVICE tabulate_output_iterator_proxy operator=(const T& x) + { + fun(index, x); + return *this; + } + +private: + BinaryFunction fun; + DifferenceT index; +}; + +// Alias template for the iterator_adaptor instantiation to be used for tabulate_output_iterator +template +using make_tabulate_output_iterator_base = + iterator_adaptor, + counting_iterator, + void, + System, + use_default, + tabulate_output_iterator_proxy>; + +// Register tabulate_output_iterator_proxy with 'is_proxy_reference' from +// type_traits to enable its use with algorithms. +template +struct is_proxy_reference> : true_type +{}; +} // namespace detail + +//! \addtogroup iterators +//! \{ + +//! \addtogroup fancyiterator Fancy Iterators +//! \ingroup iterators +//! \{ + +//! \p tabulate_output_iterator is a special kind of output iterator which, whenever a value is assigned to a +//! dereferenced iterator, calls the given callable with the index that corresponds to the offset of the dereferenced +//! iterator and the assigned value. +//! +//! The following code snippet demonstrated how to create a \p tabulate_output_iterator which prints the index and the +//! assigned value. +//! +//! \code +//! #include +//! +//! // note: functor inherits form binary function +//! struct print_op +//! { +//! __host__ __device__ +//! void operator()(int index, float value) const +//! { +//! printf("%d: %f\n", index, value); +//! } +//! }; +//! +//! int main() +//! { +//! auto tabulate_it = thrust::make_tabulate_output_iterator(print_op{}); +//! +//! tabulate_it[0] = 1.0f; // prints: 0: 1.0 +//! tabulate_it[1] = 3.0f; // prints: 1: 3.0 +//! tabulate_it[9] = 5.0f; // prints: 9: 5.0 +//! } +//! \endcode +//! +//! \see make_tabulate_output_iterator +template +class tabulate_output_iterator : public detail::make_tabulate_output_iterator_base +{ public: - using super_t = detail::tabulate_output_iterator_base; + //! \cond + using super_t = detail::make_tabulate_output_iterator_base; friend class iterator_core_access; - /*! \endcond - */ + //! \endcond tabulate_output_iterator() = default; - /*! This constructor takes as argument a \c BinaryFunction and copies it to a new \p tabulate_output_iterator - * - * \param fun A \c BinaryFunction called whenever a value is assigned to this \p tabulate_output_iterator. - */ + //! This constructor takes as argument a \c BinaryFunction and copies it to a new \p tabulate_output_iterator + //! + //! \param fun A \c BinaryFunction called whenever a value is assigned to this \p tabulate_output_iterator. _CCCL_HOST_DEVICE tabulate_output_iterator(BinaryFunction fun) : fun(fun) {} - /*! \cond - */ - private: + //! \cond _CCCL_HOST_DEVICE typename super_t::reference dereference() const { - return detail::tabulate_output_iterator_proxy(fun, *this->base()); + return detail::tabulate_output_iterator_proxy{fun, *this->base()}; } BinaryFunction fun; + //! \endcond +}; - /*! \endcond - */ -}; // end tabulate_output_iterator - -/*! \p make_tabulate_output_iterator creates a \p tabulate_output_iterator from a \c BinaryFunction. - * - * \param fun The \c BinaryFunction invoked whenever assigning to a dereferenced \p tabulate_output_iterator - * \see tabulate_output_iterator - */ +//! \p make_tabulate_output_iterator creates a \p tabulate_output_iterator from a \c BinaryFunction. +//! +//! \param fun The \c BinaryFunction invoked whenever assigning to a dereferenced \p tabulate_output_iterator +//! \see tabulate_output_iterator template tabulate_output_iterator _CCCL_HOST_DEVICE make_tabulate_output_iterator(BinaryFunction fun) { return tabulate_output_iterator(fun); -} // end make_tabulate_output_iterator - -/*! \} // end fancyiterators - */ +} -/*! \} // end iterators - */ +//! \} // end fancyiterators +//! \} // end iterators THRUST_NAMESPACE_END