/* * 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 #include #include #include #include #include #include THRUST_NAMESPACE_BEGIN // forward declare zip_iterator for zip_iterator_base template class zip_iterator; namespace detail { // Functors to be used with tuple algorithms // template class advance_iterator { public: inline _CCCL_HOST_DEVICE advance_iterator(DiffType step) : m_step(step) {} _CCCL_EXEC_CHECK_DISABLE template inline _CCCL_HOST_DEVICE void operator()(Iterator& it) const { thrust::advance(it, m_step); } private: DiffType m_step; }; // end advance_iterator struct increment_iterator { _CCCL_EXEC_CHECK_DISABLE template inline _CCCL_HOST_DEVICE void operator()(Iterator& it) { ++it; } }; // end increment_iterator struct decrement_iterator { _CCCL_EXEC_CHECK_DISABLE template inline _CCCL_HOST_DEVICE void operator()(Iterator& it) { --it; } }; // end decrement_iterator struct dereference_iterator { template struct apply { typedef typename iterator_traits::reference type; }; // end apply // XXX silence warnings of the form "calling a __host__ function from a __host__ __device__ function is not allowed _CCCL_EXEC_CHECK_DISABLE template _CCCL_HOST_DEVICE typename apply::type operator()(Iterator const& it) { return *it; } }; // end dereference_iterator // The namespace tuple_impl_specific provides two meta- // algorithms and two algorithms for tuples. namespace tuple_impl_specific { // define apply1 for tuple_meta_transform_impl template struct apply1 : UnaryMetaFunctionClass::template apply {}; // end apply1 // define apply2 for tuple_meta_accumulate_impl template struct apply2 : UnaryMetaFunctionClass::template apply {}; // end apply2 // Meta-accumulate algorithm for tuples. Note: The template // parameter StartType corresponds to the initial value in // ordinary accumulation. // template struct tuple_meta_accumulate; template struct tuple_meta_accumulate, BinaryMetaFun, StartType> { typedef typename thrust::detail::identity_::type type; }; template struct tuple_meta_accumulate, BinaryMetaFun, StartType> { typedef typename apply2, BinaryMetaFun, StartType>::type>::type type; }; template inline _CCCL_HOST_DEVICE Fun tuple_for_each_helper(Fun f) { return f; } template inline _CCCL_HOST_DEVICE Fun tuple_for_each_helper(Fun f, T& t, Ts&... ts) { f(t); return tuple_for_each_helper(f, ts...); } // for_each algorithm for tuples. template inline _CCCL_HOST_DEVICE Fun tuple_for_each(thrust::tuple& t, Fun f, thrust::index_sequence) { return tuple_for_each_helper(f, thrust::get(t)...); } // end tuple_for_each() // for_each algorithm for tuples. template inline _CCCL_HOST_DEVICE Fun tuple_for_each(thrust::tuple& t, Fun f) { return tuple_for_each(t, f, thrust::make_index_sequence>::value>{}); } } // namespace tuple_impl_specific // Metafunction to obtain the type of the tuple whose element types // are the value_types of an iterator tupel. // template struct tuple_of_value_types : tuple_meta_transform {}; // end tuple_of_value_types struct minimum_category_lambda { template struct apply : minimum_category {}; }; // Metafunction to obtain the minimal traversal tag in a tuple // of iterators. // template struct minimum_traversal_category_in_iterator_tuple { typedef typename tuple_meta_transform::type tuple_of_traversal_tags; typedef typename tuple_impl_specific:: tuple_meta_accumulate::type type; }; struct minimum_system_lambda { template struct apply : minimum_system {}; }; // Metafunction to obtain the minimal system tag in a tuple // of iterators. template struct minimum_system_in_iterator_tuple { typedef typename thrust::detail::tuple_meta_transform::type tuple_of_system_tags; typedef typename tuple_impl_specific:: tuple_meta_accumulate::type type; }; namespace zip_iterator_base_ns { template struct tuple_of_iterator_references_helper; template struct tuple_of_iterator_references_helper> { typedef thrust::detail::tuple_of_iterator_references::type...> type; }; template struct tuple_of_iterator_references { // get a thrust::tuple of the iterators' references typedef typename tuple_meta_transform::type tuple_of_references; // map thrust::tuple to tuple_of_iterator_references typedef typename tuple_of_iterator_references_helper< tuple_of_references, thrust::make_index_sequence::value>>::type type; }; } // namespace zip_iterator_base_ns /////////////////////////////////////////////////////////////////// // // Class zip_iterator_base // // Builds and exposes the iterator facade type from which the zip // iterator will be derived. // template struct zip_iterator_base { // private: // reference type is the type of the tuple obtained from the // iterators' reference types. typedef typename zip_iterator_base_ns::tuple_of_iterator_references::type reference; // Boost's Value type is the same as reference type. // typedef reference value_type; typedef typename tuple_of_value_types::type value_type; // Difference type is the first iterator's difference type typedef typename thrust::iterator_traits::type>::difference_type difference_type; // Iterator system is the minimum system tag in the // iterator tuple typedef typename minimum_system_in_iterator_tuple::type system; // Traversal category is the minimum traversal category in the // iterator tuple typedef typename minimum_traversal_category_in_iterator_tuple::type traversal_category; public: // The iterator facade type from which the zip iterator will // be derived. typedef thrust:: iterator_facade, value_type, system, traversal_category, reference, difference_type> type; }; // end zip_iterator_base } // namespace detail THRUST_NAMESPACE_END