diff --git a/thrust/iterator/detail/multi_permutation_iterator.inl b/thrust/iterator/detail/multi_permutation_iterator.inl
new file mode 100644
index 000000000..c2cedc801
--- /dev/null
+++ b/thrust/iterator/detail/multi_permutation_iterator.inl
@@ -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
+
diff --git a/thrust/iterator/detail/multi_permutation_iterator_base.h b/thrust/iterator/detail/multi_permutation_iterator_base.h
new file mode 100644
index 000000000..c0dbadd45
--- /dev/null
+++ b/thrust/iterator/detail/multi_permutation_iterator_base.h
@@ -0,0 +1,211 @@
+/*
+ *  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;
+    
+    // Boost's Pointer type is just value_type *
+    //typedef value_type * pointer;
+    typedef reference * pointer;
+    
+    // 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::experimental::iterator_adaptor<
+    multi_permutation_iterator<ElementIterator,IndexTupleIterator>,
+    IndexTupleIterator,
+    pointer,
+    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
+
+  template<typename Offset>
+  __host__ __device__
+  typename apply<Offset>::type operator()(Offset const& i)
+  { return *(this->_it + i); }
+
+  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
+
diff --git a/thrust/iterator/detail/multi_permutation_iterator_tuple_transform.h b/thrust/iterator/detail/multi_permutation_iterator_tuple_transform.h
new file mode 100644
index 000000000..94d9a01e8
--- /dev/null
+++ b/thrust/iterator/detail/multi_permutation_iterator_tuple_transform.h
@@ -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
+
diff --git a/thrust/iterator/multi_permutation_iterator.h b/thrust/iterator/multi_permutation_iterator.h
new file mode 100644
index 000000000..f36af7dec
--- /dev/null
+++ b/thrust/iterator/multi_permutation_iterator.h
@@ -0,0 +1,207 @@
+/*
+ *  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_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 experimental::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>
+