Skip to content

HTTPS clone URL

Subversion checkout URL

You can clone with HTTPS or Subversion.

Download ZIP

Loading…

multi_permutation_iterator (sorry for the dupe) #247

Open
wants to merge 8 commits into from

1 participant

Andrew Corrigan
Andrew Corrigan

multi_permutation_iterator resolves issue 70 [1], which was opened based on a request [2] I made over two years ago.

This iterator is general purpose, but is of particular significance due to its application in efficiently implementing stencil operations and multi-dimensional arrays, both of which are areas in which Thrust is currently lacking.

If a proliferation of new iterators is to be avoided then perhaps the functionality provided by multi_permutation_iterator could be integrated into permutation_iterator, to directly generalize permutation_iterators present capabilities? As it stands they have analogous functionality, both index into another iterator. Their difference only lies in the types of indexes allowed: in the case of permutation iterator only types such as int can be used as indexes and therefore permutation_iterator dereferences to produces a single value, whereas multi_permutation_iterator allows for tuples to be used as input and therefore dereferences to produce multiple values.

[1] thrust/thrust#70
[2] https://groups.google.com/d/msg/thrust-users/125vfWIGeQ4/LnTFRUu4eTwJ

PS: My apologies for the duplicate pull request. I accidentally closed the previous one, and couldn't figure out how to re-open it.

Andrew Corrigan

I just wanted to mention that I've made multi_permutation_iterator compatible with the latest improvements to iterator_facade and iterator_adapter.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
This page is out of date. Refresh to see the latest.
37 thrust/iterator/detail/multi_permutation_iterator.inl
View
@@ -0,0 +1,37 @@
+/*
+ * Copyright 2008-2012 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.
+ */
+
+#include <thrust/iterator/multi_permutation_iterator.h>
+
+namespace thrust
+{
+
+template <typename ElementIterator, typename IndexTupleIterator>
+ typename multi_permutation_iterator<ElementIterator, IndexTupleIterator>::super_t::reference
+ multi_permutation_iterator<ElementIterator, IndexTupleIterator>
+ ::dereference(void) const
+{
+ typename thrust::iterator_traits<IndexTupleIterator>::value_type index_tuple = *(this->base()); // this converts from reference to value, and avoids issues with device_reference not being a tuple when dispatcing the host_device_transform below
+
+#ifndef WAR_NVCC_CANNOT_HANDLE_DEPENDENT_TEMPLATE_TEMPLATE_ARGUMENT
+ return thrust::detail::tuple_host_device_transform<detail::tuple_dereference_iterator<ElementIterator>::template apply>(index_tuple, detail::tuple_dereference_iterator<ElementIterator>(m_element_iterator));
+#else
+ return thrust::detail::multi_permutation_iterator_tuple_transform_ns::tuple_host_device_transform< typename multi_permutation_iterator<ElementIterator, IndexTupleIterator>::super_t::reference >(index_tuple, detail::tuple_dereference_iterator<ElementIterator>(m_element_iterator));
+#endif
+} // end multi_permutation_iterator::dereference()
+
+} // end thrust
+
210 thrust/iterator/detail/multi_permutation_iterator_base.h
View
@@ -0,0 +1,210 @@
+/*
+ * Copyright 2008-2012 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
+
+// this iterator's implementation is based on zip_iterator, so we grab all of its includes
+#include <thrust/iterator/detail/zip_iterator_base.h>
+
+//
+// nvcc cannot compile the line from below:
+//
+// return thrust::detail::tuple_host_device_transform<detail::tuple_dereference_iterator<Iterator>::template apply>(i, detail::tuple_dereference_iterator<Iterator>(this->_it));
+//
+// it incorrectly reports:
+//
+// error: dependent-name 'thrust::detail::tuple_dereference_iterator::apply' is parsed as a non-type, but instantiation yields a type
+// note: say 'typename thrust::detail::tuple_dereference_iterator::apply' if a type is meant
+//
+// To see that this a compiler bug, compare the analogous line from zip_iterator_base.h:
+// since detail::dereference_iterator is not a template, and thus thrust::detail::tuple_dereference_iterator::apply
+// is not a dependent-name, nvcc is able to compile it correctly
+//
+// to workaround this, we reimplement tuple_host_device_transform
+// in a way that avoids passing this dependent-name-template-template apply in (TupleMetaFunction)
+// and instead just specify XfrmTuple directly
+//
+#ifdef __CUDACC__
+# define WAR_NVCC_CANNOT_HANDLE_DEPENDENT_TEMPLATE_TEMPLATE_ARGUMENT
+#endif
+
+
+#ifdef WAR_NVCC_CANNOT_HANDLE_DEPENDENT_TEMPLATE_TEMPLATE_ARGUMENT
+# include <thrust/iterator/detail/multi_permutation_iterator_tuple_transform.h>
+#endif
+
+namespace thrust
+{
+
+template<typename,typename> class multi_permutation_iterator;
+
+namespace detail
+{
+
+template<typename Substitute, typename T>
+struct tuple_substitution
+{
+ typedef Substitute type;
+};
+
+template<typename Substitute>
+ struct tuple_substitution<Substitute, thrust::null_type>
+{
+ typedef thrust::null_type type;
+};
+
+template<typename Substitute, typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
+struct tuple_substitution <Substitute, thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> >
+{
+ typedef thrust::tuple<
+ typename tuple_substitution<Substitute,T0>::type,
+ typename tuple_substitution<Substitute,T1>::type,
+ typename tuple_substitution<Substitute,T2>::type,
+ typename tuple_substitution<Substitute,T3>::type,
+ typename tuple_substitution<Substitute,T4>::type,
+ typename tuple_substitution<Substitute,T5>::type,
+ typename tuple_substitution<Substitute,T6>::type,
+ typename tuple_substitution<Substitute,T7>::type,
+ typename tuple_substitution<Substitute,T8>::type,
+ typename tuple_substitution<Substitute,T9>::type
+ > type;
+}; // end of tuple_substitution
+
+template<typename Substitute, typename T>
+struct tuple_reference_substitution
+{
+ typedef Substitute type;
+};
+
+template<typename Substitute>
+ struct tuple_reference_substitution<Substitute, thrust::null_type>
+{
+ typedef thrust::null_type type;
+};
+
+template<typename Substitute, typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
+struct tuple_reference_substitution <Substitute, thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> >
+{
+ typedef thrust::detail::tuple_of_iterator_references<
+ typename tuple_reference_substitution<Substitute,T0>::type,
+ typename tuple_reference_substitution<Substitute,T1>::type,
+ typename tuple_reference_substitution<Substitute,T2>::type,
+ typename tuple_reference_substitution<Substitute,T3>::type,
+ typename tuple_reference_substitution<Substitute,T4>::type,
+ typename tuple_reference_substitution<Substitute,T5>::type,
+ typename tuple_reference_substitution<Substitute,T6>::type,
+ typename tuple_reference_substitution<Substitute,T7>::type,
+ typename tuple_reference_substitution<Substitute,T8>::type,
+ typename tuple_reference_substitution<Substitute,T9>::type
+ > type;
+}; // end of tuple_reference_substitution
+
+// Metafunction to obtain a tuple whose element types
+// are all the reference type of ElementIterator,
+// where the tuple is the same size as IndexTupleIterator
+template<typename ElementIterator, typename IndexTupleIterator>
+ struct indexed_tuple_of_references
+ : tuple_reference_substitution<typename iterator_traits<ElementIterator>::reference, typename thrust::iterator_traits<IndexTupleIterator>::value_type>
+{
+}; // end indexed_tuple_of_references
+
+
+// Metafunction to obtain a tuple whose element types
+// are all the value type of ElementIterator,
+// where the tuple is the same size as IndexTupleIterator
+template<typename ElementIterator, typename IndexTupleIterator>
+ struct indexed_tuple_of_value_types
+ : tuple_substitution<typename iterator_traits<ElementIterator>::value_type, typename thrust::iterator_traits<IndexTupleIterator>::value_type>
+{
+}; // end indexed_tuple_of_value_types
+
+template<typename ElementIterator,
+ typename IndexTupleIterator>
+struct multi_permutation_iterator_base
+{
+ typedef typename thrust::iterator_system< ElementIterator>::type System1;
+ typedef typename thrust::iterator_system<IndexTupleIterator>::type System2;
+
+ //private:
+ // reference type is the type of the tuple obtained from the
+ // iterators' reference types.
+ typedef typename indexed_tuple_of_references<ElementIterator,IndexTupleIterator>::type reference;
+
+ // Boost's Value type is the same as reference type.
+ //typedef reference value_type;
+ typedef typename indexed_tuple_of_value_types<ElementIterator,IndexTupleIterator>::type value_type;
+
+ // Difference type is the IndexTuple iterator's difference type
+ typedef typename thrust::iterator_traits<
+ IndexTupleIterator
+ >::difference_type difference_type;
+
+ // Iterator system is the minimum system tag in the
+ // iterator tuple
+ typedef typename
+ detail::minimum_system<System1,System2>::type system;
+
+public:
+
+ typedef thrust::iterator_adaptor<
+ multi_permutation_iterator<ElementIterator,IndexTupleIterator>,
+ IndexTupleIterator,
+ value_type,
+ system,
+ thrust::use_default,
+ reference
+ > type;
+}; // end multi_permutation_iterator_base
+
+template<typename Iterator>
+struct tuple_dereference_iterator
+{
+ Iterator _it;
+ __host__ __device__ tuple_dereference_iterator(Iterator it) : _it(it) {}
+
+ template<typename Offset>
+ struct apply
+ {
+ typedef typename
+ tuple_reference_substitution<typename thrust::iterator_traits<Iterator>::reference, Offset>::type
+ type;
+ }; // end apply
+
+ // XXX silence warnings of the form "calling a __host__ function from a __host__ __device__ function is not allowed
+ __thrust_hd_warning_disable__
+ template<typename Offset>
+ __host__ __device__
+ typename apply<Offset>::type operator()(Offset const& i)
+ { return *(this->_it + i); }
+
+ // XXX silence warnings of the form "calling a __host__ function from a __host__ __device__ function is not allowed
+ __thrust_hd_warning_disable__
+ template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
+ __host__ __device__
+ typename apply<thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> >::type operator()(thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> const& i)
+ {
+#ifndef WAR_NVCC_CANNOT_HANDLE_DEPENDENT_TEMPLATE_TEMPLATE_ARGUMENT
+ return thrust::detail:: tuple_host_device_transform<detail::tuple_dereference_iterator<Iterator>::template apply >(i, detail::tuple_dereference_iterator<Iterator>(this->_it));
+#else
+ return thrust::detail::multi_permutation_iterator_tuple_transform_ns::tuple_host_device_transform<typename apply<thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> >::type>(i, detail::tuple_dereference_iterator<Iterator>(this->_it));
+#endif
+ }
+}; // end tuple_dereference_iterator
+
+} // end detail
+
+} // end thrust
+
260 thrust/iterator/detail/multi_permutation_iterator_tuple_transform.h
View
@@ -0,0 +1,260 @@
+/*
+ * Copyright 2008-2012 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
+
+//
+// this code is a workaround for a bug in nvcc
+// see multi_permutation_iterator_base.h for an explanation
+//
+// note: we do not re-implement do_it_on_the_host / tuple_host_transform
+// since it is no longer needed
+//
+#ifdef WAR_NVCC_CANNOT_HANDLE_DEPENDENT_TEMPLATE_TEMPLATE_ARGUMENT
+
+#include <thrust/tuple.h>
+
+namespace thrust
+{
+
+namespace detail
+{
+
+namespace multi_permutation_iterator_tuple_transform_ns
+{
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction,
+ unsigned int sz = thrust::tuple_size<Tuple>::value>
+ struct tuple_transform_functor;
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,0>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ { return thrust::null_type();
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,1>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,2>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,3>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,4>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,5>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)),
+ f(thrust::get<4>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,6>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)),
+ f(thrust::get<4>(t)),
+ f(thrust::get<5>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,7>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)),
+ f(thrust::get<4>(t)),
+ f(thrust::get<5>(t)),
+ f(thrust::get<6>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,8>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)),
+ f(thrust::get<4>(t)),
+ f(thrust::get<5>(t)),
+ f(thrust::get<6>(t)),
+ f(thrust::get<7>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,9>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)),
+ f(thrust::get<4>(t)),
+ f(thrust::get<5>(t)),
+ f(thrust::get<6>(t)),
+ f(thrust::get<7>(t)),
+ f(thrust::get<8>(t)));
+ }
+};
+
+
+template<typename Tuple,
+ typename XfrmTuple,
+ typename UnaryFunction>
+ struct tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction,10>
+{
+ static __host__ __device__
+ XfrmTuple
+ do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
+ {
+ return XfrmTuple(f(thrust::get<0>(t)),
+ f(thrust::get<1>(t)),
+ f(thrust::get<2>(t)),
+ f(thrust::get<3>(t)),
+ f(thrust::get<4>(t)),
+ f(thrust::get<5>(t)),
+ f(thrust::get<6>(t)),
+ f(thrust::get<7>(t)),
+ f(thrust::get<8>(t)),
+ f(thrust::get<9>(t)));
+ }
+};
+
+template<typename XfrmTuple,
+ typename Tuple,
+ typename UnaryFunction>
+XfrmTuple
+__host__ __device__
+tuple_host_device_transform(const Tuple &t, UnaryFunction f)
+{
+ return tuple_transform_functor<Tuple,XfrmTuple,UnaryFunction>::do_it_on_the_host_or_device(t,f);
+}
+
+} // end multi_permutation_iterator_tuple_transform_ns
+
+} // end detail
+
+} // end thrust
+
+#endif
+
208 thrust/iterator/multi_permutation_iterator.h
View
@@ -0,0 +1,208 @@
+/*
+ * Copyright 2008-2012 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.
+ */
+
+/*! \file multi_permutation_iterator.h
+ * \brief An iterator which iterates over a permutation of a range.
+ */
+
+/*
+ * (C) Copyright Toon Knapen 2001.
+ * (C) Copyright David Abrahams 2003.
+ * (C) Copyright Roland Richter 2003.
+ *
+ * Distributed under the Boost Software License, Version 1.0.
+ * (See accompanying NOTICE file for the complete license)
+ *
+ * For more information, see http://www.boost.org
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/type_traits.h>
+#include <thrust/iterator/detail/multi_permutation_iterator_base.h>
+#include <thrust/iterator/iterator_facade.h>
+#include <thrust/iterator/iterator_traits.h>
+
+namespace thrust
+{
+
+
+/*! \addtogroup iterators
+ * \{
+ */
+
+/*! \addtogroup fancyiterator Fancy Iterators
+ * \ingroup iterators
+ * \{
+ */
+
+/*! \p multi_permutation_iterator is an iterator which represents a pointer into a
+ * reordered view of a given range. \p multi_permutation_iterator is an imprecise name;
+ * the reordered view need not be a strict permutation. This iterator is useful
+ * for fusing a scatter or gather operation with other algorithms.
+ *
+ * This iterator takes two arguments:
+ *
+ * - an iterator to the range \c V on which the "permutation" will be applied
+ * - the reindexing scheme that defines how the elements of \c V will be permuted.
+ *
+ * Note that \p multi_permutation_iterator is not limited to strict permutations of the
+ * given range \c V. The distance between begin and end of the reindexing iterators
+ * is allowed to be smaller compared to the size of the range \c V, in which case
+ * the \p multi_permutation_iterator only provides a "permutation" of a subrange of \c V.
+ * The indices neither need to be unique. In this same context, it must be noted
+ * that the past-the-end \p multi_permutation_iterator is completely defined by means of
+ * the past-the-end iterator to the indices.
+ *
+ * The following code snippet demonstrates how to create a \p multi_permutation_iterator
+ * which represents a reordering of the contents of a \p device_vector.
+ *
+ * \code
+ * #include <thrust/iterator/multi_permutation_iterator.h>
+ * #include <thrust/device_vector.h>
+ * ...
+ * thrust::device_vector<float> values(4);
+ * values[0] = 10.0f;
+ * values[1] = 20.0f;
+ * values[2] = 30.0f;
+ * values[3] = 40.0f;
+ * values[4] = 50.0f;
+ * values[5] = 60.0f;
+ * values[6] = 70.0f;
+ * values[7] = 80.0f;
+ *
+ * thrust::device_vector<int> indices(4);
+ * indices[0] = 2;
+ * indices[1] = 6;
+ * indices[2] = 1;
+ * indices[3] = 3;
+ *
+ * typedef thrust::device_vector<float>::iterator ElementIterator;
+ * typedef thrust::device_vector<int>::iterator IndexTupleIterator;
+ *
+ * thrust::multi_permutation_iterator<ElementIterator,IndexTupleIterator> iter(values.begin(), indices.begin());
+ *
+ * *iter; // returns 30.0f;
+ * iter[0]; // returns 30.0f;
+ * iter[1]; // returns 70.0f;
+ * iter[2]; // returns 20.0f;
+ * iter[3]; // returns 40.0f;
+ *
+ * // iter[4] is an out-of-bounds error
+ *
+ * *iter = -1.0f; // sets values[2] to -1.0f;
+ * iter[0] = -1.0f; // sets values[2] to -1.0f;
+ * iter[1] = -1.0f; // sets values[6] to -1.0f;
+ * iter[2] = -1.0f; // sets values[1] to -1.0f;
+ * iter[3] = -1.0f; // sets values[3] to -1.0f;
+ *
+ * // values is now {10, -1, -1, -1, 50, 60, -1, 80}
+ * \endcode
+ *
+ * \see make_multi_permutation_iterator
+ */
+template <typename ElementIterator,
+ typename IndexTupleIterator>
+ class multi_permutation_iterator
+ : public thrust::detail::multi_permutation_iterator_base<
+ ElementIterator,
+ IndexTupleIterator
+ >::type
+{
+ /*! \cond
+ */
+ private:
+ typedef typename detail::multi_permutation_iterator_base<ElementIterator,IndexTupleIterator>::type super_t;
+
+ friend class thrust::iterator_core_access;
+ /*! \endcond
+ */
+
+ public:
+ /*! Null constructor calls the null constructor of this \p multi_permutation_iterator's
+ * element iterator.
+ */
+ __host__ __device__
+ multi_permutation_iterator()
+ : m_element_iterator() {}
+
+ /*! Constructor accepts an \c ElementIterator into a range of values and an
+ * \c IndexTupleIterator into a range of indices defining the indexing scheme on the
+ * values.
+ *
+ * \param x An \c ElementIterator pointing this \p multi_permutation_iterator's range of values.
+ * \param y An \c IndexTupleIterator pointing to an indexing scheme to use on \p x.
+ */
+ __host__ __device__
+ explicit multi_permutation_iterator(ElementIterator x, IndexTupleIterator y)
+ : super_t(y), m_element_iterator(x) {}
+
+ /*! Copy constructor accepts a related \p multi_permutation_iterator.
+ * \param r A compatible \p multi_permutation_iterator to copy from.
+ */
+ template<typename OtherElementIterator, typename OtherIndexTupleIterator>
+ __host__ __device__
+ multi_permutation_iterator(multi_permutation_iterator<OtherElementIterator,OtherIndexTupleIterator> const &r
+ // XXX remove these guards when we have static_assert
+ , typename detail::enable_if_convertible<OtherElementIterator, ElementIterator>::type* = 0
+ , typename detail::enable_if_convertible<OtherIndexTupleIterator, IndexTupleIterator>::type* = 0
+ )
+ : super_t(r.base()), m_element_iterator(r.m_element_iterator)
+ {}
+
+ /*! \cond
+ */
+ private:
+ __host__ __device__
+ typename super_t::reference dereference() const;
+
+ // make friends for the copy constructor
+ template<typename,typename> friend class multi_permutation_iterator;
+
+ ElementIterator m_element_iterator;
+ /*! \endcond
+ */
+}; // end multi_permutation_iterator
+
+
+/*! \p make_multi_permutation_iterator creates a \p multi_permutation_iterator
+ * from an \c ElementIterator pointing to a range of elements to "permute"
+ * and an \c IndexTupleIterator pointing to a range of indices defining an indexing
+ * scheme on the values.
+ *
+ * \param e An \c ElementIterator pointing to a range of values.
+ * \param i An \c IndexTupleIterator pointing to an indexing scheme to use on \p e.
+ * \return A new \p multi_permutation_iterator which permutes the range \p e by \p i.
+ * \see multi_permutation_iterator
+ */
+template<typename ElementIterator, typename IndexTupleIterator>
+__host__ __device__
+multi_permutation_iterator<ElementIterator,IndexTupleIterator> make_multi_permutation_iterator(ElementIterator e, IndexTupleIterator i)
+{
+ return multi_permutation_iterator<ElementIterator,IndexTupleIterator>(e,i);
+}
+
+/*! \} // end fancyiterators
+ */
+
+/*! \} // end iterators
+ */
+
+} // end thrust
+
+#include <thrust/iterator/detail/multi_permutation_iterator.inl>
+
Something went wrong with that request. Please try again.