diff --git a/thrust/testing/copy.cu b/thrust/testing/copy.cu index d4808142b1b..b8f967e52f8 100644 --- a/thrust/testing/copy.cu +++ b/thrust/testing/copy.cu @@ -717,15 +717,30 @@ template <> struct is_non_const_reference : thrust::true_type {}; } // end namespace detail +THRUST_NAMESPACE_END +namespace std +{ template <> struct iterator_traits { using value_type = long long; using reference = only_set_when_expected_it; using iterator_category = thrust::random_access_device_iterator_tag; + using difference_type = ::cuda::std::ptrdiff_t; }; -THRUST_NAMESPACE_END +} // namespace std + +_LIBCUDACXX_BEGIN_NAMESPACE_STD +template <> +struct iterator_traits +{ + using value_type = long long; + using reference = only_set_when_expected_it; + using iterator_category = thrust::random_access_device_iterator_tag; + using difference_type = ::cuda::std::ptrdiff_t; +}; +_LIBCUDACXX_END_NAMESPACE_STD void TestCopyWithBigIndexesHelper(int magnitude) { diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 2de233d7eaf..df4aee7ae7d 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -565,24 +565,27 @@ struct only_set_when_expected_it } }; -THRUST_NAMESPACE_BEGIN +namespace std +{ template <> struct iterator_traits { - using value_type = long long; - using reference = only_set_when_expected_it; + using value_type = long long; + using reference = only_set_when_expected_it; + using difference_type = ::cuda::std::ptrdiff_t; }; -THRUST_NAMESPACE_END +} // namespace std -namespace std -{ +_LIBCUDACXX_BEGIN_NAMESPACE_STD template <> struct iterator_traits { - using value_type = long long; - using reference = only_set_when_expected_it; + using value_type = long long; + using reference = only_set_when_expected_it; + using iterator_category = thrust::random_access_device_iterator_tag; + using difference_type = ::cuda::std::ptrdiff_t; }; -} // namespace std +_LIBCUDACXX_END_NAMESPACE_STD void TestInclusiveScanWithBigIndexesHelper(int magnitude) { diff --git a/thrust/thrust/iterator/constant_iterator.h b/thrust/thrust/iterator/constant_iterator.h index 10c0448a026..020e82679f6 100644 --- a/thrust/thrust/iterator/constant_iterator.h +++ b/thrust/thrust/iterator/constant_iterator.h @@ -14,10 +14,8 @@ * limitations under the License. */ -/*! \file thrust/iterator/constant_iterator.h - * \brief An iterator which returns a constant value when - * dereferenced - */ +//! \file thrust/iterator/constant_iterator.h +//! \brief An iterator which returns a constant value when dereferenced #pragma once @@ -30,152 +28,157 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include +#include +#include #include -THRUST_NAMESPACE_BEGIN +#include -/*! \addtogroup iterators - * \{ - */ +THRUST_NAMESPACE_BEGIN -/*! \addtogroup fancyiterator Fancy Iterators - * \ingroup iterators - * \{ - */ +template +class constant_iterator; -/*! \p constant_iterator is an iterator which represents a pointer into a range - * of constant values. This iterator is useful for creating a range filled with the same - * value without explicitly storing it in memory. Using \p constant_iterator saves both - * memory capacity and bandwidth. - * - * The following code snippet demonstrates how to create a \p constant_iterator whose - * \c value_type is \c int and whose value is \c 10. - * - * \code - * #include - * - * thrust::constant_iterator iter(10); - * - * *iter; // returns 10 - * iter[0]; // returns 10 - * iter[1]; // returns 10 - * iter[13]; // returns 10 - * - * // and so on... - * \endcode - * - * This next example demonstrates how to use a \p constant_iterator with the - * \p thrust::transform function to increment all elements of a sequence by the - * same value. We will create a temporary \p constant_iterator with the function - * \p make_constant_iterator function in order to avoid explicitly specifying - * its type: - * - * \code - * #include - * #include - * #include - * #include - * - * int main() - * { - * thrust::device_vector data(4); - * data[0] = 3; - * data[1] = 7; - * data[2] = 2; - * data[3] = 5; - * - * // add 10 to all values in data - * thrust::transform(data.begin(), data.end(), - * thrust::make_constant_iterator(10), - * data.begin(), - * thrust::plus()); - * - * // data is now [13, 17, 12, 15] - * - * return 0; - * } - * \endcode - * - * \see make_constant_iterator - */ +namespace detail +{ +template +struct make_constant_iterator_base +{ + using incrementable = typename ia_dflt_help>::type; + using base_iterator = counting_iterator; + using type = + iterator_adaptor, + base_iterator, + Value, + iterator_system_t, + iterator_traversal_t, + Value>; +}; +} // namespace detail + +//! \addtogroup iterators +//! \{ + +//! \addtogroup fancyiterator Fancy Iterators +//! \ingroup iterators +//! \{ + +//! \p constant_iterator is an iterator which represents a pointer into a range of constant values. This iterator is +//! useful for creating a range filled with the same value without explicitly storing it in memory. Using \p +//! constant_iterator saves both memory capacity and bandwidth. +//! +//! The following code snippet demonstrates how to create a \p constant_iterator whose \c value_type is \c int and whose +//! value is \c 10. +//! +//! \code +//! #include +//! +//! thrust::constant_iterator iter(10); +//! +//! *iter; // returns 10 +//! iter[0]; // returns 10 +//! iter[1]; // returns 10 +//! iter[13]; // returns 10 +//! +//! // and so on... +//! \endcode +//! +//! This next example demonstrates how to use a \p constant_iterator with the \p thrust::transform function to increment +//! all elements of a sequence by the same value. We will create a temporary \p constant_iterator with the function \p +//! make_constant_iterator function in order to avoid explicitly specifying its type: +//! +//! \code +//! #include +//! #include +//! #include +//! #include +//! +//! int main() +//! { +//! thrust::device_vector data(4); +//! data[0] = 3; +//! data[1] = 7; +//! data[2] = 2; +//! data[3] = 5; +//! +//! // add 10 to all values in data +//! thrust::transform(data.begin(), data.end(), +//! thrust::make_constant_iterator(10), +//! data.begin(), +//! thrust::plus()); +//! +//! // data is now [13, 17, 12, 15] +//! +//! return 0; +//! } +//! \endcode +//! +//! \see make_constant_iterator template -class constant_iterator : public detail::constant_iterator_base::type +class constant_iterator : public detail::make_constant_iterator_base::type { - /*! \cond - */ + //! \cond friend class iterator_core_access; - using super_t = typename detail::constant_iterator_base::type; - using incrementable = typename detail::constant_iterator_base::incrementable; - using base_iterator = typename detail::constant_iterator_base::base_iterator; + using super_t = typename detail::make_constant_iterator_base::type; + using incrementable = typename detail::make_constant_iterator_base::incrementable; + using base_iterator = typename detail::make_constant_iterator_base::base_iterator; public: using reference = typename super_t::reference; using value_type = typename super_t::value_type; - /*! \endcond - */ + //! \endcond - /*! Default constructor initializes this \p constant_iterator's constant using its default constructor - */ + //! Default constructor initializes this \p constant_iterator's constant using its default constructor _CCCL_HOST_DEVICE constant_iterator() : super_t() , m_value() {} - /*! Copy constructor copies the value of another \p constant_iterator with related - * System type. - * - * \param rhs The \p constant_iterator to copy. - */ - template >::type, - typename thrust::iterator_system::type, - int> = 0> + //! Copy constructor copies the value of another \p constant_iterator with related System type. + //! + //! \param rhs The \p constant_iterator to copy. + template < + class OtherSystem, + detail::enable_if_convertible_t>::type, + typename iterator_system::type, + int> = 0> _CCCL_HOST_DEVICE constant_iterator(constant_iterator const& rhs) : super_t(rhs.base()) , m_value(rhs.value()) {} - /*! This constructor receives a value to use as the constant value of this - * \p constant_iterator and an index specifying the location of this - * \p constant_iterator in a sequence. - * - * \p v The value of this \p constant_iterator's constant value. - * \p i The index of this \p constant_iterator in a sequence. Defaults to the - * value returned by \c Incrementable's null constructor. For example, - * when Incrementable == int, \c 0. - */ + //! This constructor receives a value to use as the constant value of this \p constant_iterator and an index + //! specifying the location of this \p constant_iterator in a sequence. + //! + //! \p v The value of this \p constant_iterator's constant value. + //! \p i The index of this \p constant_iterator in a sequence. Defaults to the value returned by \c Incrementable's + //! null constructor. For example, when Incrementable == int, \c 0. _CCCL_HOST_DEVICE constant_iterator(value_type const& v, incrementable const& i = incrementable()) : super_t(base_iterator(i)) , m_value(v) {} - /*! This constructor is templated to allow construction from a value type and - * incrementable type related this this \p constant_iterator's respective types. - * - * \p v The value of this \p constant_iterator's constant value. - * \p i The index of this \p constant_iterator in a sequence. Defaults to the - * value returned by \c Incrementable's null constructor. For example, - * when Incrementable == int, \c 0. - */ + //! This constructor is templated to allow construction from a value type and incrementable type related this this \p + //! constant_iterator's respective types. + //! + //! \p v The value of this \p constant_iterator's constant value. + //! \p i The index of this \p constant_iterator in a sequence. Defaults to the value returned by \c Incrementable's + //! null constructor. For example, when Incrementable == int, \c 0. template _CCCL_HOST_DEVICE constant_iterator(OtherValue const& v, OtherIncrementable const& i = incrementable()) : super_t(base_iterator(i)) , m_value(v) {} - /*! This method returns the value of this \p constant_iterator's constant value. - * \return A \c const reference to this \p constant_iterator's constant value. - */ + //! This method returns the value of this \p constant_iterator's constant value. \return A \c const reference to this + //! \p constant_iterator's constant value. _CCCL_HOST_DEVICE Value const& value() const { return m_value; } - /*! \cond - */ + //! \cond protected: _CCCL_HOST_DEVICE Value const& value_reference() const @@ -194,52 +197,38 @@ class constant_iterator : public detail::constant_iterator_base inline _CCCL_HOST_DEVICE constant_iterator make_constant_iterator(ValueT x, IndexT i = int()) { return constant_iterator(x, i); } // end make_constant_iterator() -/*! This version of \p make_constant_iterator creates a \p constant_iterator - * using only a parameter for the desired constant value. The value of the - * returned \p constant_iterator's index is set to \c 0. - * - * \param x The value of the returned \p constant_iterator's constant value. - * \return A new \p constant_iterator with constant value equal to \p x and - * index equal to \c 0. - * \see constant_iterator - */ +//! This version of \p make_constant_iterator creates a \p constant_iterator using only a parameter for the desired +//! constant value. The value of the returned \p constant_iterator's index is set to \c 0. +//! +//! \param x The value of the returned \p constant_iterator's constant value. +//! \return A new \p constant_iterator with constant value equal to \p x and index equal to \c 0. +//! \see constant_iterator template inline _CCCL_HOST_DEVICE constant_iterator make_constant_iterator(V x) { return constant_iterator(x, 0); } // end make_constant_iterator() -/*! \} // end fancyiterators - */ - -/*! \} // end iterators - */ +//! \} // end fancyiterators +//! \} // end iterators THRUST_NAMESPACE_END diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index d6a9e62762f..e2f712967f8 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -40,152 +40,276 @@ # pragma system_header #endif // no system header +#include +#include #include -#include -#include +#include -// #include the details first -#include +#include +#include +#include +#include THRUST_NAMESPACE_BEGIN -/*! \addtogroup iterators - * \{ - */ +// forward declaration of counting_iterator +template +class counting_iterator; -/*! \addtogroup fancyiterator Fancy Iterators - * \ingroup iterators - * \{ - */ +namespace detail +{ +template +struct num_digits + : eval_if<::cuda::std::numeric_limits::is_specialized, + integral_constant::digits>, + integral_constant::digits + - (::cuda::std::numeric_limits::is_signed ? 1 : 0)>>::type +{}; // end num_digits -/*! \p counting_iterator is an iterator which represents a pointer into a range - * of sequentially changing values. This iterator is useful for creating a range - * filled with a sequence without explicitly storing it in memory. Using - * \p counting_iterator saves memory capacity and bandwidth. - * - * The following code snippet demonstrates how to create a \p counting_iterator whose - * \c value_type is \c int and which sequentially increments by \c 1. - * - * \code - * #include - * ... - * // create iterators - * thrust::counting_iterator first(10); - * thrust::counting_iterator last = first + 3; - * - * first[0] // returns 10 - * first[1] // returns 11 - * first[100] // returns 110 - * - * // sum of [first, last) - * thrust::reduce(first, last); // returns 33 (i.e. 10 + 11 + 12) - * - * // initialize vector to [0,1,2,..] - * thrust::counting_iterator iter(0); - * thrust::device_vector vec(500); - * thrust::copy(iter, iter + vec.size(), vec.begin()); - * \endcode - * - * This next example demonstrates how to use a \p counting_iterator with the - * \p thrust::copy_if function to compute the indices of the non-zero elements - * of a \p device_vector. In this example, we use the \p make_counting_iterator - * function to avoid specifying the type of the \p counting_iterator. - * - * \code - * #include - * #include - * #include - * #include - * - * int main() - * { - * // this example computes indices for all the nonzero values in a sequence - * - * // sequence of zero and nonzero values - * thrust::device_vector stencil(8); - * stencil[0] = 0; - * stencil[1] = 1; - * stencil[2] = 1; - * stencil[3] = 0; - * stencil[4] = 0; - * stencil[5] = 1; - * stencil[6] = 0; - * stencil[7] = 1; - * - * // storage for the nonzero indices - * thrust::device_vector indices(8); - * - * // compute indices of nonzero elements - * using IndexIterator = thrust::device_vector::iterator; - * - * // use make_counting_iterator to define the sequence [0, 8) - * IndexIterator indices_end = thrust::copy_if(thrust::make_counting_iterator(0), - * thrust::make_counting_iterator(8), - * stencil.begin(), - * indices.begin(), - * ::cuda::std::identity{}); - * // indices now contains [1,2,5,7] - * - * return 0; - * } - * \endcode - * - * \see make_counting_iterator - */ +template +struct integer_difference +//: eval_if< +// sizeof(Integer) >= sizeof(intmax_t), +// eval_if< +// is_signed::value, +// identity_, +// identity_ +// >, +// eval_if< +// sizeof(Integer) < sizeof(std::ptrdiff_t), +// identity_, +// identity_ +// > +// > +{ +private: + +public: + using type = + typename eval_if<::cuda::std::numeric_limits::is_signed + && (!::cuda::std::numeric_limits::is_bounded + || (int(::cuda::std::numeric_limits::digits) + 1 >= num_digits::value)), + identity_, + eval_if::digits) + 1 < num_digits::value, + identity_, + eval_if::digits) + 1 < num_digits::value, + identity_, + identity_>>>::type; +}; // end integer_difference + +template +struct numeric_difference + : eval_if<::cuda::std::is_integral::value, integer_difference, identity_> +{}; // end numeric_difference + +template +_CCCL_HOST_DEVICE typename numeric_difference::type numeric_distance(Number x, Number y) +{ + using difference_type = typename numeric_difference::type; + return difference_type(y) - difference_type(x); +} // end numeric_distance + +template +struct make_counting_iterator_base +{ + using system = + typename eval_if<::cuda::std::is_same::value, identity_, identity_>::type; + + using traversal = + typename ia_dflt_help::value, + identity_, + iterator_traversal>>::type; + + // unlike Boost, we explicitly use std::ptrdiff_t as the difference type + // for floating point counting_iterators + using difference = + typename ia_dflt_help::value, + eval_if<::cuda::std::is_integral::value, + numeric_difference, + identity_<::cuda::std::ptrdiff_t>>, + iterator_difference>>::type; + + // our implementation departs from Boost's in that counting_iterator::dereference + // returns a copy of its counter, rather than a reference to it. returning a reference + // to the internal state of an iterator causes subtle bugs (consider the temporary + // iterator created in the expression *(iter + i)) and has no compelling use case + using type = + iterator_adaptor, + Incrementable, + Incrementable, + system, + traversal, + Incrementable, + difference>; +}; // end counting_iterator_base + +template +struct iterator_distance +{ + _CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y) + { + return y - x; + } +}; + +template +struct number_distance +{ + _CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y) + { + return static_cast(numeric_distance(x, y)); + } +}; + +template +struct counting_iterator_equal +{ + _CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y) + { + return x == y; + } +}; + +// specialization for floating point equality +template +struct counting_iterator_equal::value + || ::cuda::std::is_floating_point::value>> +{ + _CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y) + { + using d = number_distance; + return d::distance(x, y) == 0; + } +}; + +} // namespace detail + +//! \addtogroup iterators +//! \{ + +//! \addtogroup fancyiterator Fancy Iterators +//! \ingroup iterators +//! \{ + +//! \p counting_iterator is an iterator which represents a pointer into a range of sequentially changing values. This +//! iterator is useful for creating a range filled with a sequence without explicitly storing it in memory. Using \p +//! counting_iterator saves memory capacity and bandwidth. +//! +//! The following code snippet demonstrates how to create a \p counting_iterator whose \c value_type is \c int and which +//! sequentially increments by \c 1. +//! +//! \code +//! #include +//! ... +//! // create iterators +//! thrust::counting_iterator first(10); +//! thrust::counting_iterator last = first + 3; +//! +//! first[0] // returns 10 +//! first[1] // returns 11 +//! first[100] // returns 110 +//! +//! // sum of [first, last) +//! thrust::reduce(first, last); // returns 33 (i.e. 10 + 11 + 12) +//! +//! // initialize vector to [0,1,2,..] +//! thrust::counting_iterator iter(0); +//! thrust::device_vector vec(500); +//! thrust::copy(iter, iter + vec.size(), vec.begin()); +//! \endcode +//! +//! This next example demonstrates how to use a \p counting_iterator with the \p thrust::copy_if function to compute the +//! indices of the non-zero elements of a \p device_vector. In this example, we use the \p make_counting_iterator +//! function to avoid specifying the type of the \p counting_iterator. +//! +//! \code +//! #include +//! #include +//! #include +//! #include +//! +//! int main() +//! { +//! // this example computes indices for all the nonzero values in a sequence +//! +//! // sequence of zero and nonzero values +//! thrust::device_vector stencil(8); +//! stencil[0] = 0; +//! stencil[1] = 1; +//! stencil[2] = 1; +//! stencil[3] = 0; +//! stencil[4] = 0; +//! stencil[5] = 1; +//! stencil[6] = 0; +//! stencil[7] = 1; +//! +//! // storage for the nonzero indices +//! thrust::device_vector indices(8); +//! +//! // compute indices of nonzero elements +//! using IndexIterator = thrust::device_vector::iterator; +//! +//! // use make_counting_iterator to define the sequence [0, 8) +//! IndexIterator indices_end = thrust::copy_if(thrust::make_counting_iterator(0), +//! thrust::make_counting_iterator(8), +//! stencil.begin(), +//! indices.begin(), +//! ::cuda::std::identity{}); +//! // indices now contains [1,2,5,7] +//! +//! return 0; +//! } +//! \endcode +//! +//! \see make_counting_iterator template class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator - : public detail::counting_iterator_base::type + : public detail::make_counting_iterator_base::type { - /*! \cond - */ - using super_t = typename detail::counting_iterator_base::type; - + //! \cond + using super_t = typename detail::make_counting_iterator_base::type; friend class iterator_core_access; public: using reference = typename super_t::reference; using difference_type = typename super_t::difference_type; + //! \endcond - /*! \endcond - */ - - /*! Default constructor initializes this \p counting_iterator's counter to - * `Incrementable{}`. - */ + //! Default constructor initializes this \p counting_iterator's counter to `Incrementable{}`. _CCCL_HOST_DEVICE counting_iterator() : super_t(Incrementable{}) {} - /*! Copy constructor copies the value of another counting_iterator - * with related System type. - * - * \param rhs The \p counting_iterator to copy. - */ - template < - class OtherSystem, - detail::enable_if_convertible_t< - typename thrust::iterator_system>::type, - typename thrust::iterator_system::type, - int> = 0> + //! Copy constructor copies the value of another counting_iterator with related System type. + //! + //! \param rhs The \p counting_iterator to copy. + template >::type, + typename iterator_system::type, + int> = 0> _CCCL_HOST_DEVICE counting_iterator(counting_iterator const& rhs) : super_t(rhs.base()) {} - /*! This \c explicit constructor copies the value of an \c Incrementable - * into a new \p counting_iterator's \c Incrementable counter. - * - * \param x The initial value of the new \p counting_iterator's \c Incrementable - * counter. - */ + //! This \c explicit constructor copies the value of an \c Incrementable into a new \p counting_iterator's \c + //! Incrementable counter. + //! + //! \param x The initial value of the new \p counting_iterator's \c Incrementable counter. _CCCL_HOST_DEVICE explicit counting_iterator(Incrementable x) : super_t(x) {} - /*! \cond - */ + //! \cond private: _CCCL_HOST_DEVICE reference dereference() const @@ -198,7 +322,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator _CCCL_HOST_DEVICE bool equal(counting_iterator const& y) const { - using e = thrust::detail::counting_iterator_equal; + using e = detail::counting_iterator_equal; return e::equal(this->base(), y.base()); } @@ -206,35 +330,29 @@ class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator _CCCL_HOST_DEVICE difference_type distance_to(counting_iterator const& y) const { - using d = typename thrust::detail::eval_if< - thrust::detail::is_numeric::value, - thrust::detail::identity_>, - thrust::detail::identity_< - thrust::detail::iterator_distance>>::type; + using d = typename detail::eval_if< + detail::is_numeric::value, + detail::identity_>, + detail::identity_>>::type; return d::distance(this->base(), y.base()); } - /*! \endcond - */ -}; // end counting_iterator + //! \endcond +}; -/*! \p make_counting_iterator creates a \p counting_iterator - * using an initial value for its \c Incrementable counter. - * - * \param x The initial value of the new \p counting_iterator's counter. - * \return A new \p counting_iterator whose counter has been initialized to \p x. - */ +//! \p make_counting_iterator creates a \p counting_iterator +//! using an initial value for its \c Incrementable counter. +//! +//! \param x The initial value of the new \p counting_iterator's counter. +//! \return A new \p counting_iterator whose counter has been initialized to \p x. template inline _CCCL_HOST_DEVICE counting_iterator make_counting_iterator(Incrementable x) { return counting_iterator(x); } -/*! \} // end fancyiterators - */ - -/*! \} // end iterators - */ +//! \} // end fancyiterators +//! \} // end iterators THRUST_NAMESPACE_END diff --git a/thrust/thrust/iterator/detail/constant_iterator_base.h b/thrust/thrust/iterator/detail/constant_iterator_base.h deleted file mode 100644 index 43ad83eed69..00000000000 --- a/thrust/thrust/iterator/detail/constant_iterator_base.h +++ /dev/null @@ -1,72 +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 - -#include - -THRUST_NAMESPACE_BEGIN - -// forward declaration of constant_iterator -template -class constant_iterator; - -namespace detail -{ - -template -struct constant_iterator_base -{ - using value_type = Value; - - // the reference type is the same as the value_type. - // we wish to avoid returning a reference to the internal state - // of the constant_iterator, which is prone to subtle bugs. - // consider the temporary iterator created in the expression - // *(iter + i) - using reference = value_type; - - // the incrementable type is int unless otherwise specified - using incrementable = - typename thrust::detail::ia_dflt_help>::type; - - using base_iterator = typename thrust::counting_iterator; - - using type = typename thrust::iterator_adaptor< - constant_iterator, - base_iterator, - value_type, - typename thrust::iterator_system::type, - typename thrust::iterator_traversal::type, - reference>; -}; // end constant_iterator_base - -} // namespace detail - -THRUST_NAMESPACE_END diff --git a/thrust/thrust/iterator/detail/counting_iterator.inl b/thrust/thrust/iterator/detail/counting_iterator.inl deleted file mode 100644 index 31fa54d71b4..00000000000 --- a/thrust/thrust/iterator/detail/counting_iterator.inl +++ /dev/null @@ -1,178 +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 -#include - -#include -#include -#include - -THRUST_NAMESPACE_BEGIN - -// forward declaration of counting_iterator -template -class counting_iterator; - -namespace detail -{ - -template -struct num_digits - : eval_if<::cuda::std::numeric_limits::is_specialized, - integral_constant::digits>, - integral_constant::digits - - (::cuda::std::numeric_limits::is_signed ? 1 : 0)>>::type -{}; // end num_digits - -template -struct integer_difference -//: eval_if< -// sizeof(Integer) >= sizeof(intmax_t), -// eval_if< -// is_signed::value, -// identity_, -// identity_ -// >, -// eval_if< -// sizeof(Integer) < sizeof(std::ptrdiff_t), -// identity_, -// identity_ -// > -// > -{ -private: - -public: - using type = - typename eval_if<::cuda::std::numeric_limits::is_signed - && (!::cuda::std::numeric_limits::is_bounded - || (int(::cuda::std::numeric_limits::digits) + 1 >= num_digits::value)), - identity_, - eval_if::digits) + 1 < num_digits::value, - identity_, - eval_if::digits) + 1 < num_digits::value, - identity_, - identity_>>>::type; -}; // end integer_difference - -template -struct numeric_difference - : eval_if<::cuda::std::is_integral::value, integer_difference, identity_> -{}; // end numeric_difference - -template -_CCCL_HOST_DEVICE typename numeric_difference::type numeric_distance(Number x, Number y) -{ - using difference_type = typename numeric_difference::type; - return difference_type(y) - difference_type(x); -} // end numeric_distance - -template -struct counting_iterator_base -{ - using system = typename thrust::detail::eval_if<::cuda::std::is_same::value, - thrust::detail::identity_, - thrust::detail::identity_>::type; - - using traversal = typename thrust::detail::ia_dflt_help< - Traversal, - thrust::detail::eval_if::value, - thrust::detail::identity_, - thrust::iterator_traversal>>::type; - - // unlike Boost, we explicitly use std::ptrdiff_t as the difference type - // for floating point counting_iterators - using difference = typename thrust::detail::ia_dflt_help< - Difference, - thrust::detail::eval_if::value, - thrust::detail::eval_if<::cuda::std::is_integral::value, - thrust::detail::numeric_difference, - thrust::detail::identity_<::cuda::std::ptrdiff_t>>, - thrust::iterator_difference>>::type; - - // our implementation departs from Boost's in that counting_iterator::dereference - // returns a copy of its counter, rather than a reference to it. returning a reference - // to the internal state of an iterator causes subtle bugs (consider the temporary - // iterator created in the expression *(iter + i)) and has no compelling use case - using type = - thrust::iterator_adaptor, - Incrementable, - Incrementable, - system, - traversal, - Incrementable, - difference>; -}; // end counting_iterator_base - -template -struct iterator_distance -{ - _CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y) - { - return y - x; - } -}; - -template -struct number_distance -{ - _CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y) - { - return static_cast(numeric_distance(x, y)); - } -}; - -template -struct counting_iterator_equal -{ - _CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y) - { - return x == y; - } -}; - -// specialization for floating point equality -template -struct counting_iterator_equal::value - || ::cuda::std::is_floating_point::value>> -{ - _CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y) - { - using d = number_distance; - return d::distance(x, y) == 0; - } -}; - -} // namespace detail -THRUST_NAMESPACE_END