/* * 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 THRUST_NAMESPACE_BEGIN // forward declaration of counting_iterator template class counting_iterator; namespace detail { template struct counting_iterator_base { typedef typename thrust::detail::eval_if< // use any_system_tag if we are given use_default ::cuda::std::is_same::value, thrust::detail::identity_, thrust::detail::identity_>::type system; typedef typename thrust::detail::ia_dflt_help< Traversal, thrust::detail::eval_if::value, thrust::detail::identity_, thrust::iterator_traversal>>::type traversal; // unlike Boost, we explicitly use std::ptrdiff_t as the difference type // for floating point counting_iterators typedef 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_>, thrust::iterator_difference>>::type difference; // 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 typedef thrust::iterator_adaptor, // self Incrementable, // Base Incrementable, // XXX we may need to pass const here as Boost does system, traversal, Incrementable, difference> type; }; // 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) { typedef number_distance d; return d::distance(x, y) == 0; } }; } // namespace detail THRUST_NAMESPACE_END