Skip to content

Commit 08e79fd

Browse files
Refactor Thrust iterators 2/n (NVIDIA#3840)
* Refactor discard_iterator * Refactor permutation_iterator
1 parent 2f1f133 commit 08e79fd

File tree

4 files changed

+229
-340
lines changed

4 files changed

+229
-340
lines changed

thrust/thrust/iterator/detail/discard_iterator_base.h

-65
This file was deleted.

thrust/thrust/iterator/detail/permutation_iterator_base.h

-59
This file was deleted.

thrust/thrust/iterator/discard_iterator.h

+110-98
Original file line numberDiff line numberDiff line change
@@ -29,104 +29,123 @@
2929
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
3030
# pragma system_header
3131
#endif // no system header
32-
#include <thrust/iterator/detail/discard_iterator_base.h>
33-
#include <thrust/iterator/iterator_facade.h>
3432

35-
_CCCL_DIAG_PUSH
36-
_CCCL_DIAG_SUPPRESS_MSVC(4244 4267) // possible loss of data
33+
#include <thrust/iterator/counting_iterator.h>
34+
#include <thrust/iterator/detail/any_assign.h>
35+
#include <thrust/iterator/iterator_adaptor.h>
36+
37+
#include <cuda/std/cstddef>
3738

3839
THRUST_NAMESPACE_BEGIN
3940

40-
/*! \addtogroup iterators
41-
* \{
42-
*/
41+
template <typename>
42+
class discard_iterator;
4343

44-
/*! \addtogroup fancyiterator Fancy Iterators
45-
* \ingroup iterators
46-
* \{
47-
*/
44+
namespace detail
45+
{
46+
template <typename System>
47+
struct make_discard_iterator_base
48+
{
49+
// XXX value_type should actually be void but this interferes with zip_iterator<discard_iterator>
50+
using value_type = any_assign;
51+
using reference = any_assign&;
52+
using incrementable = ::cuda::std::ptrdiff_t;
53+
54+
using base_iterator = counting_iterator<incrementable, System, random_access_traversal_tag>;
55+
56+
using type =
57+
iterator_adaptor<discard_iterator<System>,
58+
base_iterator,
59+
value_type,
60+
iterator_system_t<base_iterator>,
61+
iterator_traversal_t<base_iterator>,
62+
reference>;
63+
};
64+
} // namespace detail
4865

49-
/*! \p discard_iterator is an iterator which represents a special kind of pointer that
50-
* ignores values written to it upon dereference. This iterator is useful for ignoring
51-
* the output of certain algorithms without wasting memory capacity or bandwidth.
52-
* \p discard_iterator may also be used to count the size of an algorithm's output which
53-
* may not be known a priori.
54-
*
55-
* The following code snippet demonstrates how to use \p discard_iterator to
56-
* ignore one of the output ranges of reduce_by_key
57-
*
58-
* \code
59-
* #include <thrust/iterator/discard_iterator.h>
60-
* #include <thrust/reduce.h>
61-
* #include <thrust/device_vector.h>
62-
*
63-
* int main()
64-
* {
65-
* thrust::device_vector<int> keys(7), values(7);
66-
*
67-
* keys[0] = 1;
68-
* keys[1] = 3;
69-
* keys[2] = 3;
70-
* keys[3] = 3;
71-
* keys[4] = 2;
72-
* keys[5] = 2;
73-
* keys[6] = 1;
74-
*
75-
* values[0] = 9;
76-
* values[1] = 8;
77-
* values[2] = 7;
78-
* values[3] = 6;
79-
* values[4] = 5;
80-
* values[5] = 4;
81-
* values[6] = 3;
82-
*
83-
* thrust::device_vector<int> result(4);
84-
*
85-
* // we are only interested in the reduced values
86-
* // use discard_iterator to ignore the output keys
87-
* thrust::reduce_by_key(keys.begin(), keys.end(),
88-
* values.begin(),
89-
* thrust::make_discard_iterator(),
90-
* result.begin());
91-
*
92-
* // result is now [9, 21, 9, 3]
93-
*
94-
* return 0;
95-
* }
96-
* \endcode
97-
*
98-
* \see make_discard_iterator
99-
*/
66+
_CCCL_DIAG_PUSH
67+
_CCCL_DIAG_SUPPRESS_MSVC(4244 4267) // possible loss of data
68+
69+
//! \addtogroup iterators
70+
//! \{
71+
72+
//! \addtogroup fancyiterator Fancy Iterators
73+
//! \ingroup iterators
74+
//! \{
75+
76+
//! \p discard_iterator is an iterator which represents a special kind of pointer that ignores values written to it upon
77+
//! dereference. This iterator is useful for ignoring the output of certain algorithms without wasting memory capacity
78+
//! or bandwidth. \p discard_iterator may also be used to count the size of an algorithm's output which may not be known
79+
//! a priori.
80+
//!
81+
//! The following code snippet demonstrates how to use \p discard_iterator to ignore one of the output ranges of
82+
//! reduce_by_key
83+
//!
84+
//! \code
85+
//! #include <thrust/iterator/discard_iterator.h>
86+
//! #include <thrust/reduce.h>
87+
//! #include <thrust/device_vector.h>
88+
//!
89+
//! int main()
90+
//! {
91+
//! thrust::device_vector<int> keys(7), values(7);
92+
//!
93+
//! keys[0] = 1;
94+
//! keys[1] = 3;
95+
//! keys[2] = 3;
96+
//! keys[3] = 3;
97+
//! keys[4] = 2;
98+
//! keys[5] = 2;
99+
//! keys[6] = 1;
100+
//!
101+
//! values[0] = 9;
102+
//! values[1] = 8;
103+
//! values[2] = 7;
104+
//! values[3] = 6;
105+
//! values[4] = 5;
106+
//! values[5] = 4;
107+
//! values[6] = 3;
108+
//!
109+
//! thrust::device_vector<int> result(4);
110+
//!
111+
//! // we are only interested in the reduced values
112+
//! // use discard_iterator to ignore the output keys
113+
//! thrust::reduce_by_key(keys.begin(), keys.end(),
114+
//! values.begin(),
115+
//! thrust::make_discard_iterator(),
116+
//! result.begin());
117+
//!
118+
//! // result is now [9, 21, 9, 3]
119+
//!
120+
//! return 0;
121+
//! }
122+
//! \endcode
123+
//!
124+
//! \see make_discard_iterator
100125
template <typename System = use_default>
101-
class discard_iterator : public detail::discard_iterator_base<System>::type
126+
class discard_iterator : public detail::make_discard_iterator_base<System>::type
102127
{
103-
/*! \cond
104-
*/
128+
//! \cond
105129
friend class iterator_core_access;
106-
using super_t = typename detail::discard_iterator_base<System>::type;
107-
using incrementable = typename detail::discard_iterator_base<System>::incrementable;
108-
using base_iterator = typename detail::discard_iterator_base<System>::base_iterator;
130+
using super_t = typename detail::make_discard_iterator_base<System>::type;
131+
using incrementable = typename detail::make_discard_iterator_base<System>::incrementable;
132+
using base_iterator = typename detail::make_discard_iterator_base<System>::base_iterator;
109133

110134
public:
111135
using reference = typename super_t::reference;
112136
using value_type = typename super_t::value_type;
113137

114-
/*! \endcond
115-
*/
138+
//! \endcond
116139

117-
/*! This constructor receives an optional index specifying the position of this
118-
* \p discard_iterator in a range.
119-
*
120-
* \p i The index of this \p discard_iterator in a range. Defaults to the
121-
* value returned by \c Incrementable's null constructor. For example,
122-
* when <tt>Incrementable == int</tt>, \c 0.
123-
*/
140+
//! This constructor receives an optional index specifying the position of this \p discard_iterator in a range.
141+
//!
142+
//! \p i The index of this \p discard_iterator in a range. Defaults to the value returned by \c Incrementable's null
143+
//! constructor. For example, when <tt>Incrementable == int</tt>, \c 0.
124144
_CCCL_HOST_DEVICE discard_iterator(incrementable const& i = incrementable())
125145
: super_t(base_iterator(i))
126146
{}
127147

128-
/*! \cond
129-
*/
148+
//! \cond
130149

131150
private: // Core iterator interface
132151
_CCCL_HOST_DEVICE reference dereference() const
@@ -136,31 +155,24 @@ class discard_iterator : public detail::discard_iterator_base<System>::type
136155

137156
mutable value_type m_element;
138157

139-
/*! \endcond
140-
*/
141-
}; // end constant_iterator
158+
//! \endcond
159+
};
142160

143-
/*! \p make_discard_iterator creates a \p discard_iterator from an optional index parameter.
144-
*
145-
* \param i The index of the returned \p discard_iterator within a range.
146-
* In the default case, the value of this parameter is \c 0.
147-
*
148-
* \return A new \p discard_iterator with index as given by \p i.
149-
*
150-
* \see constant_iterator
151-
*/
161+
//! \p make_discard_iterator creates a \p discard_iterator from an optional index parameter.
162+
//!
163+
//! \param i The index of the returned \p discard_iterator within a range. In the default case, the value of this
164+
//! parameter is \c 0.
165+
//! \return A new \p discard_iterator with index as given by \p i.
166+
//! \see constant_iterator
152167
inline _CCCL_HOST_DEVICE discard_iterator<>
153168
make_discard_iterator(discard_iterator<>::difference_type i = discard_iterator<>::difference_type(0))
154169
{
155170
return discard_iterator<>(i);
156-
} // end make_discard_iterator()
171+
}
157172

158-
/*! \} // end fancyiterators
159-
*/
173+
//! \} // end fancyiterators
174+
//! \} // end iterators
160175

161-
/*! \} // end iterators
162-
*/
176+
_CCCL_DIAG_POP
163177

164178
THRUST_NAMESPACE_END
165-
166-
_CCCL_DIAG_POP

0 commit comments

Comments
 (0)