diff --git a/thrust/testing/pair.cu b/thrust/testing/pair.cu
index f5f6e92b5ca..27365781ea5 100644
--- a/thrust/testing/pair.cu
+++ b/thrust/testing/pair.cu
@@ -44,7 +44,7 @@ struct TestPairManipulation
     // test copy from pair
     p4.first  = T(2);
     p4.second = T(3);
-    
+
     P p5;
     p5 = p4;
     ASSERT_EQUAL(p4.first,  p5.first);
@@ -217,7 +217,7 @@ using PairConstVolatileTypes =
     unittest::type_list<thrust::pair<int, float>, thrust::pair<int, float> const,
                         thrust::pair<int, float> const volatile>;
 
-template <typename Pair> 
+template <typename Pair>
 struct TestPairTupleSize
 {
   void operator()()
@@ -289,3 +289,16 @@ void TestPairSwap(void)
 }
 DECLARE_UNITTEST(TestPairSwap);
 
+#if THRUST_CPP_DIALECT >= 2017
+void TestPairStructuredBindings(void)
+{
+  const int a = 42;
+  const int b = 1337;
+  thrust::pair<int,int> p(a,b);
+
+  auto [a2, b2] = p;
+  ASSERT_EQUAL(a, a2);
+  ASSERT_EQUAL(b, b2);
+}
+DECLARE_UNITTEST(TestPairStructuredBindings);
+#endif
diff --git a/thrust/testing/tuple.cu b/thrust/testing/tuple.cu
index 40dccbd2231..3efd351c284 100644
--- a/thrust/testing/tuple.cu
+++ b/thrust/testing/tuple.cu
@@ -491,4 +491,18 @@ void TestTupleSwap(void)
 }
 DECLARE_UNITTEST(TestTupleSwap);
 
-
+#if THRUST_CPP_DIALECT >= 2017
+void TestTupleStructuredBindings(void)
+{
+  const int a = 0;
+  const int b = 42;
+  const int c = 1337;
+  thrust::tuple<int,int,int> t(a,b,c);
+
+  auto [a2, b2, c2] = t;
+  ASSERT_EQUAL(a, a2);
+  ASSERT_EQUAL(b, b2);
+  ASSERT_EQUAL(c, c2);
+}
+DECLARE_UNITTEST(TestTupleStructuredBindings);
+#endif
diff --git a/thrust/testing/tuple_sort.cu b/thrust/testing/tuple_sort.cu
index 1684ba45e52..db51c118322 100644
--- a/thrust/testing/tuple_sort.cu
+++ b/thrust/testing/tuple_sort.cu
@@ -20,10 +20,7 @@ struct GetFunctor
 {
   template<typename Tuple>
   __host__ __device__
-  typename thrust::access_traits<
-    typename thrust::tuple_element<N, Tuple>::type
-  >::const_type
-  operator()(const Tuple &t)
+  typename thrust::tuple_element<N, Tuple>::type operator()(const Tuple &t)
   {
     return thrust::get<N>(t);
   }
@@ -64,7 +61,7 @@ struct TestTupleStableSort
 
      // select values
      transform(h_tuples.begin(), h_tuples.end(), h_values.begin(), GetFunctor<1>());
-     
+
      device_vector<T> d_values(h_values.size());
      transform(d_tuples.begin(), d_tuples.end(), d_values.begin(), GetFunctor<1>());
 
diff --git a/thrust/testing/tuple_transform.cu b/thrust/testing/tuple_transform.cu
index e77bd489be8..7aef96aa141 100644
--- a/thrust/testing/tuple_transform.cu
+++ b/thrust/testing/tuple_transform.cu
@@ -19,10 +19,7 @@ struct GetFunctor
 {
   template<typename Tuple>
   __host__ __device__
-  typename thrust::access_traits<
-    typename thrust::tuple_element<N, Tuple>::type
-  >::const_type
-  operator()(const Tuple &t)
+  typename thrust::tuple_element<N, Tuple>::type operator()(const Tuple &t)
   {
     return thrust::get<N>(t);
   }
diff --git a/thrust/testing/zip_iterator.cu b/thrust/testing/zip_iterator.cu
index c48ca21709c..b2493531cc3 100644
--- a/thrust/testing/zip_iterator.cu
+++ b/thrust/testing/zip_iterator.cu
@@ -285,9 +285,9 @@ void TestZipIteratorCopy(void)
   sequence(input0.begin(), input0.end(), T{0});
   sequence(input1.begin(), input1.end(), T{13});
 
-  copy( make_zip_iterator(make_tuple(input0.begin(),  input1.begin())),
-        make_zip_iterator(make_tuple(input0.end(),    input1.end())),
-        make_zip_iterator(make_tuple(output0.begin(), output1.begin())));
+  thrust::copy( make_zip_iterator(make_tuple(input0.begin(),  input1.begin())),
+                make_zip_iterator(make_tuple(input0.end(),    input1.end())),
+                make_zip_iterator(make_tuple(output0.begin(), output1.begin())));
 
   ASSERT_EQUAL(input0, output0);
   ASSERT_EQUAL(input1, output1);
diff --git a/thrust/thrust/detail/functional/actor.h b/thrust/thrust/detail/functional/actor.h
index cee0770a46b..208a8706180 100644
--- a/thrust/thrust/detail/functional/actor.h
+++ b/thrust/thrust/detail/functional/actor.h
@@ -65,10 +65,6 @@ template<typename Eval>
   __host__ __device__
   actor(const Eval &base);
 
-  __host__ __device__
-  typename apply_actor<eval_type, thrust::null_type >::type
-  operator()(void) const;
-
   template <typename... Ts>
   __host__ __device__
   typename apply_actor<eval_type, thrust::tuple<eval_ref<Ts>...>>::type
@@ -122,7 +118,7 @@ template<typename Eval>
 {
   typedef typename thrust::detail::functional::apply_actor<
     thrust::detail::functional::actor<Eval>,
-    thrust::null_type
+    thrust::tuple<>
   >::type type;
 }; // end result_of
 
diff --git a/thrust/thrust/detail/functional/actor.inl b/thrust/thrust/detail/functional/actor.inl
index e0bdebbbf3e..cd2ca66f374 100644
--- a/thrust/thrust/detail/functional/actor.inl
+++ b/thrust/thrust/detail/functional/actor.inl
@@ -54,18 +54,6 @@ template<typename Eval>
       : eval_type(base)
 {}
 
-template<typename Eval>
-  __host__ __device__
-  typename apply_actor<
-    typename actor<Eval>::eval_type,
-    typename thrust::null_type
-  >::type
-    actor<Eval>
-      ::operator()(void) const
-{
-  return eval_type::eval(thrust::null_type());
-} // end basic_environment::operator()
-
 // actor::operator() needs to construct a tuple of references to its
 // arguments. To make this work with thrust::reference<T>, we need to
 // detect thrust proxy references and store them as T rather than T&.
diff --git a/thrust/thrust/detail/functional/argument.h b/thrust/thrust/detail/functional/argument.h
index aac29f53713..b1ceb220555 100644
--- a/thrust/thrust/detail/functional/argument.h
+++ b/thrust/thrust/detail/functional/argument.h
@@ -41,9 +41,9 @@ template<unsigned int i, typename Env>
 };
 
 template<unsigned int i>
-  struct argument_helper<i,thrust::null_type>
+  struct argument_helper<i,thrust::tuple<>>
 {
-  typedef thrust::null_type type;
+  typedef thrust::tuple<> type;
 };
 
 
@@ -52,7 +52,7 @@ template<unsigned int i>
 {
   public:
     template<typename Env>
-      struct result
+    struct result
         : argument_helper<i,Env>
     {
     };
diff --git a/thrust/thrust/detail/functional/composite.h b/thrust/thrust/detail/functional/composite.h
index 41ee7473925..40466b3b8ea 100644
--- a/thrust/thrust/detail/functional/composite.h
+++ b/thrust/thrust/detail/functional/composite.h
@@ -36,33 +36,11 @@ namespace detail
 namespace functional
 {
 
-// XXX we should just take a single EvalTuple
-template<typename Eval0,
-         typename Eval1  = thrust::null_type,
-         typename Eval2  = thrust::null_type,
-         typename Eval3  = thrust::null_type,
-         typename Eval4  = thrust::null_type,
-         typename Eval5  = thrust::null_type,
-         typename Eval6  = thrust::null_type,
-         typename Eval7  = thrust::null_type,
-         typename Eval8  = thrust::null_type,
-         typename Eval9  = thrust::null_type,
-         typename Eval10 = thrust::null_type>
-  class composite;
+template <typename... Eval>
+class composite;
 
 template<typename Eval0, typename Eval1>
-  class composite<
-    Eval0,
-    Eval1,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type
-  >
+  class composite<Eval0, Eval1>
 {
   public:
     template<typename Env>
@@ -96,18 +74,7 @@ template<typename Eval0, typename Eval1>
 }; // end composite<Eval0,Eval1>
 
 template<typename Eval0, typename Eval1, typename Eval2>
-  class composite<
-    Eval0,
-    Eval1,
-    Eval2,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type,
-    thrust::null_type
-  >
+  class composite<Eval0, Eval1, Eval2>
 {
   public:
     template<typename Env>
diff --git a/thrust/thrust/detail/functional/operators/operator_adaptors.h b/thrust/thrust/detail/functional/operators/operator_adaptors.h
index 67326c5c1e5..addfbd3dce2 100644
--- a/thrust/thrust/detail/functional/operators/operator_adaptors.h
+++ b/thrust/thrust/detail/functional/operators/operator_adaptors.h
@@ -43,7 +43,7 @@ struct transparent_unary_operator
   using argument =
   typename thrust::detail::eval_if<
     thrust::tuple_size<Env>::value != 1,
-    thrust::detail::identity_<thrust::null_type>,
+    thrust::detail::identity_<thrust::tuple<>>,
     thrust::detail::functional::argument_helper<0, Env>
   >::type;
 
@@ -57,8 +57,8 @@ struct transparent_unary_operator
   template <typename Env>
   using result_type =
   typename thrust::detail::eval_if<
-    std::is_same<thrust::null_type, argument<Env>>::value,
-    thrust::detail::identity_<thrust::null_type>,
+    std::is_same<thrust::tuple<>, argument<Env>>::value,
+    thrust::detail::identity_<thrust::tuple<>>,
     result_type_impl<Env>
   >::type;
 
@@ -88,16 +88,16 @@ struct transparent_binary_operator
   using first_argument =
     typename thrust::detail::eval_if<
       thrust::tuple_size<Env>::value != 2,
-      thrust::detail::identity_<thrust::null_type>,
-      thrust::detail::functional::argument_helper<0, Env>
+      thrust::detail::identity_<thrust::tuple<>>,
+    thrust::detail::functional::argument_helper<0, Env>
     >::type;
 
   template <typename Env>
   using second_argument =
     typename thrust::detail::eval_if<
       thrust::tuple_size<Env>::value != 2,
-      thrust::detail::identity_<thrust::null_type>,
-      thrust::detail::functional::argument_helper<1, Env>
+      thrust::detail::identity_<thrust::tuple<>>,
+    thrust::detail::functional::argument_helper<1, Env>
     >::type;
 
   template <typename Env>
@@ -111,9 +111,9 @@ struct transparent_binary_operator
   template <typename Env>
   using result_type =
     typename thrust::detail::eval_if<
-      (std::is_same<thrust::null_type, first_argument<Env>>::value ||
-       std::is_same<thrust::null_type, second_argument<Env>>::value),
-      thrust::detail::identity_<thrust::null_type>,
+      (std::is_same<thrust::tuple<>, first_argument<Env>>::value ||
+       std::is_same<thrust::tuple<>, second_argument<Env>>::value),
+      thrust::detail::identity_<thrust::tuple<>>,
       result_type_impl<Env>
     >::type;
 
diff --git a/thrust/thrust/detail/pair.inl b/thrust/thrust/detail/pair.inl
deleted file mode 100644
index 4b7dd6eb0fe..00000000000
--- a/thrust/thrust/detail/pair.inl
+++ /dev/null
@@ -1,231 +0,0 @@
-/*
- *  Copyright 2008-2021 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 <thrust/detail/config.h>
-
-#include <thrust/pair.h>
-#include <thrust/detail/swap.h>
-#include <thrust/tuple.h>
-
-THRUST_NAMESPACE_BEGIN
-
-template <typename T1, typename T2>
-  __host__ __device__
-  pair<T1,T2>
-    ::pair(void)
-      :first(),second()
-{
-  ;
-} // end pair::pair()
-
-
-template <typename T1, typename T2>
-  __host__ __device__
-  pair<T1,T2>
-    ::pair(const T1 &x, const T2 &y)
-      :first(x),second(y)
-{
-  ;
-} // end pair::pair()
-
-
-template <typename T1, typename T2>
-  template <typename U1, typename U2>
-    __host__ __device__
-    pair<T1,T2>
-      ::pair(const pair<U1,U2> &p)
-        :first(p.first),second(p.second)
-{
-  ;
-} // end pair::pair()
-
-
-template <typename T1, typename T2>
-  template <typename U1, typename U2>
-    __host__ __device__
-    pair<T1,T2>
-      ::pair(const std::pair<U1,U2> &p)
-        :first(p.first),second(p.second)
-{
-  ;
-} // end pair::pair()
-
-
-template<typename T1, typename T2>
-  inline __host__ __device__
-    void pair<T1,T2>
-      ::swap(thrust::pair<T1,T2> &p)
-{
-  using thrust::swap;
-
-  swap(first, p.first);
-  swap(second, p.second);
-} // end pair::swap()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator==(const pair<T1,T2> &x, const pair<T1,T2> &y)
-{
-  return x.first == y.first && x.second == y.second;
-} // end operator==()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator<(const pair<T1,T2> &x, const pair<T1,T2> &y)
-{
-  return x.first < y.first || (!(y.first < x.first) && x.second < y.second);
-} // end operator<()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator!=(const pair<T1,T2> &x, const pair<T1,T2> &y)
-{
-  return !(x == y);
-} // end operator==()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator>(const pair<T1,T2> &x, const pair<T1,T2> &y)
-{
-  return y < x;
-} // end operator<()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator<=(const pair<T1,T2> &x, const pair<T1,T2> &y)
-{
-  return !(y < x);
-} // end operator<=()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator>=(const pair<T1,T2> &x, const pair<T1,T2> &y)
-{
-  return !(x < y);
-} // end operator>=()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    void swap(pair<T1,T2> &x, pair<T1,T2> &y)
-{
-  return x.swap(y);
-} // end swap()
-
-
-template <typename T1, typename T2>
-  inline __host__ __device__
-    pair<T1,T2> make_pair(T1 x, T2 y)
-{
-  return pair<T1,T2>(x,y);
-} // end make_pair()
-
-
-// specializations of tuple_element for pair
-template<typename T1, typename T2>
-  struct tuple_element<0, pair<T1,T2>>
-{
-  typedef T1 type;
-}; // end tuple_element
-
-template<typename T1, typename T2>
-  struct tuple_element<1, pair<T1,T2>>
-{
-  typedef T2 type;
-}; // end tuple_element
-
-
-// specialization of tuple_size for pair
-template<typename T1, typename T2>
-  struct tuple_size<pair<T1,T2>>
-{
-  static const unsigned int value = 2;
-}; // end tuple_size
-
-
-
-namespace detail
-{
-
-
-template<int N, typename Pair> struct pair_get {};
-
-template<typename Pair>
-  struct pair_get<0, Pair>
-{
-  inline __host__ __device__
-    const typename tuple_element<0, Pair>::type &
-      operator()(const Pair &p) const
-  {
-    return p.first;
-  } // end operator()()
-
-  inline __host__ __device__
-    typename tuple_element<0, Pair>::type &
-      operator()(Pair &p) const
-  {
-    return p.first;
-  } // end operator()()
-}; // end pair_get
-
-
-template<typename Pair>
-  struct pair_get<1, Pair>
-{
-  inline __host__ __device__
-    const typename tuple_element<1, Pair>::type &
-      operator()(const Pair &p) const
-  {
-    return p.second;
-  } // end operator()()
-
-  inline __host__ __device__
-    typename tuple_element<1, Pair>::type &
-      operator()(Pair &p) const
-  {
-    return p.second;
-  } // end operator()()
-}; // end pair_get
-
-} // end detail
-
-
-
-template<unsigned int N, typename T1, typename T2>
-  inline __host__ __device__
-    typename tuple_element<N, pair<T1,T2> >::type &
-      get(pair<T1,T2> &p)
-{
-  return detail::pair_get<N, pair<T1,T2> >()(p);
-} // end get()
-
-template<unsigned int N, typename T1, typename T2>
-  inline __host__ __device__
-    const typename tuple_element<N, pair<T1,T2> >::type &
-      get(const pair<T1,T2> &p)
-{
-  return detail::pair_get<N, pair<T1,T2> >()(p);
-} // end get()
-
-THRUST_NAMESPACE_END
diff --git a/thrust/thrust/detail/raw_reference_cast.h b/thrust/thrust/detail/raw_reference_cast.h
index eff45f0c273..99af0da183d 100644
--- a/thrust/thrust/detail/raw_reference_cast.h
+++ b/thrust/thrust/detail/raw_reference_cast.h
@@ -21,7 +21,6 @@
 #include <thrust/detail/type_traits/has_nested_type.h>
 #include <thrust/detail/type_traits.h>
 #include <thrust/detail/tuple_transform.h>
-#include <thrust/iterator/detail/tuple_of_iterator_references.h>
 
 
 // the order of declarations and definitions in this file is totally goofy
@@ -33,6 +32,8 @@ THRUST_NAMESPACE_BEGIN
 namespace detail
 {
 
+template<typename... Ts>
+class tuple_of_iterator_references;
 
 __THRUST_DEFINE_HAS_NESTED_TYPE(is_wrapped_reference, wrapped_reference_hint)
 
diff --git a/thrust/thrust/detail/tuple.inl b/thrust/thrust/detail/tuple.inl
deleted file mode 100644
index f4930bf4bd6..00000000000
--- a/thrust/thrust/detail/tuple.inl
+++ /dev/null
@@ -1,1004 +0,0 @@
-/*
- *  Copyright 2008-2021 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 <thrust/detail/config.h>
-
-#include <thrust/detail/type_traits.h>
-#include <thrust/detail/swap.h>
-
-THRUST_NAMESPACE_BEGIN
-
-// define null_type
-struct null_type {};
-
-// null_type comparisons
-__host__ __device__ inline
-bool operator==(const null_type&, const null_type&) { return true; }
-
-__host__ __device__ inline
-bool operator>=(const null_type&, const null_type&) { return true; }
-
-__host__ __device__ inline
-bool operator<=(const null_type&, const null_type&) { return true; }
-
-__host__ __device__ inline
-bool operator!=(const null_type&, const null_type&) { return false; }
-
-__host__ __device__ inline
-bool operator<(const null_type&, const null_type&) { return false; }
-
-__host__ __device__ inline
-bool operator>(const null_type&, const null_type&) { return false; }
-
-// forward declaration for tuple
-template <
-  class T0 = null_type, class T1 = null_type, class T2 = null_type,
-  class T3 = null_type, class T4 = null_type, class T5 = null_type,
-  class T6 = null_type, class T7 = null_type, class T8 = null_type,
-  class T9 = null_type>
-class tuple;
-
-
-template <size_t N, class T> struct tuple_element;
-
-template<size_t N, class T>
-  struct tuple_element_impl
-{
-  private:
-    typedef typename T::tail_type Next;
-
-  public:
-    /*! The result of this metafunction is returned in \c type.
-     */
-    typedef typename tuple_element_impl<N-1, Next>::type type;
-}; // end tuple_element
-
-template<class T>
-  struct tuple_element_impl<0,T>
-{
-  typedef typename T::head_type type;
-};
-
-template <size_t N, class T>
-  struct tuple_element<N, T const>
-{
-    using type = typename std::add_const<typename tuple_element<N, T>::type>::type;
-};
-
-template <size_t N, class T>
-struct tuple_element<N, T volatile>
-{
-    using type = typename std::add_volatile<typename tuple_element<N, T>::type>::type;
-};
-
-template <size_t N, class T>
-  struct tuple_element<N, T const volatile>
-{
-    using type = typename std::add_cv<typename tuple_element<N, T>::type>::type;
-};
-
-template <size_t N, class T>
-struct tuple_element{
-    using type = typename tuple_element_impl<N,T>::type;
-};
-
-// forward declaration of tuple_size
-template<class T> struct tuple_size;
-
-template<class T>
-  struct tuple_size<T const> : public tuple_size<T> {};
-
-template<class T>
-  struct tuple_size<T volatile> : public tuple_size<T> {};
-
-template<class T>
-  struct tuple_size<T const volatile> : public tuple_size<T> {};
-
-/*! This metafunction returns the number of elements
- *  of a \p tuple type of interest.
- *
- *  \tparam T A \c tuple type of interest.
- *
- *  \see pair
- *  \see tuple
- */
-template<class T>
-  struct tuple_size
-{
-  /*! The result of this metafunction is returned in \c value.
-   */
-  static const int value = 1 + tuple_size<typename T::tail_type>::value;
-}; // end tuple_size
-
-
-// specializations for tuple_size
-template<>
-  struct tuple_size< tuple<> >
-{
-  static const int value = 0;
-}; // end tuple_size< tuple<> >
-
-template<>
-  struct tuple_size<null_type>
-{
-  static const int value = 0;
-}; // end tuple_size<null_type>
-
-
-
-// forward declaration of detail::cons
-namespace detail
-{
-
-template <class HT, class TT> struct cons;
-
-} // end detail
-
-
-// -- some traits classes for get functions
-template <class T> struct access_traits
-{
-  typedef const T& const_type;
-  typedef T& non_const_type;
-
-  typedef const typename thrust::detail::remove_cv<T>::type& parameter_type;
-
-// used as the tuple constructors parameter types
-// Rationale: non-reference tuple element types can be cv-qualified.
-// It should be possible to initialize such types with temporaries,
-// and when binding temporaries to references, the reference must
-// be non-volatile and const. 8.5.3. (5)
-}; // end access_traits
-
-template <class T> struct access_traits<T&>
-{
-  typedef T& const_type;
-  typedef T& non_const_type;
-
-  typedef T& parameter_type;
-}; // end access_traits<T&>
-
-// forward declarations of get()
-template<int N, class HT, class TT>
-__host__ __device__
-inline typename access_traits<
-                  typename tuple_element<N, detail::cons<HT, TT> >::type
-                >::non_const_type
-// XXX we probably don't need to do this for any compiler we care about -jph
-//get(cons<HT, TT>& c BOOST_APPEND_EXPLICIT_TEMPLATE_NON_TYPE(int, N));
-get(detail::cons<HT, TT>& c);
-
-template<int N, class HT, class TT>
-__host__ __device__
-inline typename access_traits<
-                  typename tuple_element<N, detail::cons<HT, TT> >::type
-                >::const_type
-// XXX we probably don't need to do this for any compiler we care about -jph
-//get(const cons<HT, TT>& c BOOST_APPEND_EXPLICIT_TEMPLATE_NON_TYPE(int, N));
-get(const detail::cons<HT, TT>& c);
-
-namespace detail
-{
-
-// -- generate error template, referencing to non-existing members of this
-// template is used to produce compilation errors intentionally
-template<class T>
-class generate_error;
-
-// - cons getters --------------------------------------------------------
-// called: get_class<N>::get<RETURN_TYPE>(aTuple)
-
-template< int N >
-struct get_class
-{
-  template<class RET, class HT, class TT >
-  __host__ __device__
-  inline static RET get(const cons<HT, TT>& t)
-  {
-    // XXX we may not need to deal with this for any compiler we care about -jph
-    //return get_class<N-1>::BOOST_NESTED_TEMPLATE get<RET>(t.tail);
-    return get_class<N-1>::template get<RET>(t.tail);
-
-    // gcc 4.3 couldn't compile this:
-    //return get_class<N-1>::get<RET>(t.tail);
-  }
-
-  template<class RET, class HT, class TT >
-  __host__ __device__
-  inline static RET get(cons<HT, TT>& t)
-  {
-    // XXX we may not need to deal with this for any compiler we care about -jph
-    //return get_class<N-1>::BOOST_NESTED_TEMPLATE get<RET>(t.tail);
-    return get_class<N-1>::template get<RET>(t.tail);
-
-    // gcc 4.3 couldn't compile this:
-    //return get_class<N-1>::get<RET>(t.tail);
-  }
-}; // end get_class
-
-template<>
-struct get_class<0>
-{
-  template<class RET, class HT, class TT>
-  __host__ __device__
-  inline static RET get(const cons<HT, TT>& t)
-  {
-    return t.head;
-  }
-
-  template<class RET, class HT, class TT>
-  __host__ __device__
-  inline static RET get(cons<HT, TT>& t)
-  {
-    return t.head;
-  }
-}; // get get_class<0>
-
-
-template <bool If, class Then, class Else> struct IF
-{
-  typedef Then RET;
-};
-
-template <class Then, class Else> struct IF<false, Then, Else>
-{
-  typedef Else RET;
-};
-
-//  These helper templates wrap void types and plain function types.
-//  The rationale is to allow one to write tuple types with those types
-//  as elements, even though it is not possible to instantiate such object.
-//  E.g: typedef tuple<void> some_type; // ok
-//  but: some_type x; // fails
-
-template <class T> class non_storeable_type
-{
-  __host__ __device__
-  non_storeable_type();
-};
-
-template <class T> struct wrap_non_storeable_type
-{
-  // XXX is_function looks complicated; punt for now -jph
-  //typedef typename IF<
-  //  ::thrust::detail::is_function<T>::value, non_storeable_type<T>, T
-  //>::RET type;
-
-  typedef T type;
-};
-
-template <> struct wrap_non_storeable_type<void>
-{
-  typedef non_storeable_type<void> type;
-};
-
-
-template <class HT, class TT>
-  struct cons
-{
-  typedef HT head_type;
-  typedef TT tail_type;
-
-  typedef typename
-    wrap_non_storeable_type<head_type>::type stored_head_type;
-
-  stored_head_type head;
-  tail_type tail;
-
-  inline __host__ __device__
-  typename access_traits<stored_head_type>::non_const_type
-  get_head() { return head; }
-
-  inline __host__ __device__
-  typename access_traits<tail_type>::non_const_type
-  get_tail() { return tail; }
-
-  inline __host__ __device__
-  typename access_traits<stored_head_type>::const_type
-  get_head() const { return head; }
-
-  inline __host__ __device__
-  typename access_traits<tail_type>::const_type
-  get_tail() const { return tail; }
-
-  inline __host__ __device__
-  cons(void) : head(), tail() {}
-  //  cons() : head(detail::default_arg<HT>::f()), tail() {}
-
-  // the argument for head is not strictly needed, but it prevents
-  // array type elements. This is good, since array type elements
-  // cannot be supported properly in any case (no assignment,
-  // copy works only if the tails are exactly the same type, ...)
-
-  inline __host__ __device__
-  cons(typename access_traits<stored_head_type>::parameter_type h,
-       const tail_type& t)
-    : head (h), tail(t) {}
-
-  template <class T1, class T2, class T3, class T4, class T5,
-            class T6, class T7, class T8, class T9, class T10>
-  inline __host__ __device__
-  cons( T1& t1, T2& t2, T3& t3, T4& t4, T5& t5,
-        T6& t6, T7& t7, T8& t8, T9& t9, T10& t10 )
-    : head (t1),
-      tail (t2, t3, t4, t5, t6, t7, t8, t9, t10, static_cast<const null_type&>(null_type()))
-      {}
-
-  template <class T2, class T3, class T4, class T5,
-            class T6, class T7, class T8, class T9, class T10>
-  inline __host__ __device__
-  cons( const null_type& /*t1*/, T2& t2, T3& t3, T4& t4, T5& t5,
-        T6& t6, T7& t7, T8& t8, T9& t9, T10& t10 )
-    : head (),
-      tail (t2, t3, t4, t5, t6, t7, t8, t9, t10, static_cast<const null_type&>(null_type()))
-      {}
-
-
-  template <class HT2, class TT2>
-  inline __host__ __device__
-  cons( const cons<HT2, TT2>& u ) : head(u.head), tail(u.tail) {}
-
-#if THRUST_CPP_DIALECT >= 2011
-  cons(const cons &) = default;
-#endif
-
-  __thrust_exec_check_disable__
-  template <class HT2, class TT2>
-  inline __host__ __device__
-  cons& operator=( const cons<HT2, TT2>& u ) {
-    head=u.head; tail=u.tail; return *this;
-  }
-
-  // must define assignment operator explicitly, implicit version is
-  // illformed if HT is a reference (12.8. (12))
-  __thrust_exec_check_disable__
-  inline __host__ __device__
-  cons& operator=(const cons& u) {
-    head = u.head; tail = u.tail;  return *this;
-  }
-
-  // XXX enable when we support std::pair -jph
-  //template <class T1, class T2>
-  //__host__ __device__
-  //cons& operator=( const std::pair<T1, T2>& u ) {
-  //  //BOOST_STATIC_ASSERT(length<cons>::value == 2); // check length = 2
-  //  head = u.first; tail.head = u.second; return *this;
-  //}
-
-  // get member functions (non-const and const)
-  template <int N>
-  __host__ __device__
-  typename access_traits<
-             typename tuple_element<N, cons<HT, TT> >::type
-           >::non_const_type
-  get() {
-    return thrust::get<N>(*this); // delegate to non-member get
-  }
-
-  template <int N>
-  __host__ __device__
-  typename access_traits<
-             typename tuple_element<N, cons<HT, TT> >::type
-           >::const_type
-  get() const {
-    return thrust::get<N>(*this); // delegate to non-member get
-  }
-
-  inline __host__ __device__
-  void swap(cons &c)
-  {
-    using thrust::swap;
-
-    swap(head, c.head);
-    tail.swap(c.tail);
-  }
-};
-
-template <class HT>
-  struct cons<HT, null_type>
-{
-  typedef HT head_type;
-  typedef null_type tail_type;
-  typedef cons<HT, null_type> self_type;
-
-  typedef typename
-    wrap_non_storeable_type<head_type>::type stored_head_type;
-  stored_head_type head;
-
-  typename access_traits<stored_head_type>::non_const_type
-  inline __host__ __device__
-  get_head() { return head; }
-
-  inline __host__ __device__
-  null_type get_tail() { return null_type(); }
-
-  inline __host__ __device__
-  typename access_traits<stored_head_type>::const_type
-  get_head() const { return head; }
-
-  inline __host__ __device__
-  null_type get_tail() const { return null_type(); }
-
-  inline __host__ __device__
-  cons() : head() {}
-
-  inline __host__ __device__
-  cons(typename access_traits<stored_head_type>::parameter_type h,
-       const null_type& = null_type())
-    : head (h) {}
-
-  template<class T1>
-  inline __host__ __device__
-  cons(T1& t1, const null_type&, const null_type&, const null_type&,
-       const null_type&, const null_type&, const null_type&,
-       const null_type&, const null_type&, const null_type&)
-  : head (t1) {}
-
-  inline __host__ __device__
-  cons(const null_type&,
-       const null_type&, const null_type&, const null_type&,
-       const null_type&, const null_type&, const null_type&,
-       const null_type&, const null_type&, const null_type&)
-  : head () {}
-
-  template <class HT2>
-  inline __host__ __device__
-  cons( const cons<HT2, null_type>& u ) : head(u.head) {}
-
-#if THRUST_CPP_DIALECT >= 2011
-  cons(const cons &) = default;
-#endif
-
-  __thrust_exec_check_disable__
-  template <class HT2>
-  inline __host__ __device__
-  cons& operator=(const cons<HT2, null_type>& u )
-  {
-    head = u.head;
-    return *this;
-  }
-
-  // must define assignment operator explicitly, implicit version
-  // is illformed if HT is a reference
-  inline __host__ __device__
-  cons& operator=(const cons& u) { head = u.head; return *this; }
-
-  template <int N>
-  inline __host__ __device__
-  typename access_traits<
-             typename tuple_element<N, self_type>::type
-            >::non_const_type
-  // XXX we probably don't need this for the compilers we care about -jph
-  //get(BOOST_EXPLICIT_TEMPLATE_NON_TYPE(int, N))
-  get(void)
-  {
-    return thrust::get<N>(*this);
-  }
-
-  template <int N>
-  inline __host__ __device__
-  typename access_traits<
-             typename tuple_element<N, self_type>::type
-           >::const_type
-  // XXX we probably don't need this for the compilers we care about -jph
-  //get(BOOST_EXPLICIT_TEMPLATE_NON_TYPE(int, N)) const
-  get(void) const
-  {
-    return thrust::get<N>(*this);
-  }
-
-  inline __host__ __device__
-  void swap(cons &c)
-  {
-    using thrust::swap;
-
-    swap(head, c.head);
-  }
-}; // end cons
-
-template <class T0, class T1, class T2, class T3, class T4,
-          class T5, class T6, class T7, class T8, class T9>
-  struct map_tuple_to_cons
-{
-  typedef cons<T0,
-               typename map_tuple_to_cons<T1, T2, T3, T4, T5,
-                                          T6, T7, T8, T9, null_type>::type
-              > type;
-}; // end map_tuple_to_cons
-
-// The empty tuple is a null_type
-template <>
-  struct map_tuple_to_cons<null_type, null_type, null_type, null_type, null_type, null_type, null_type, null_type, null_type, null_type>
-{
-  typedef null_type type;
-}; // end map_tuple_to_cons<...>
-
-
-
-// ---------------------------------------------------------------------------
-// The call_traits for make_tuple
-
-// Must be instantiated with plain or const plain types (not with references)
-
-// from template<class T> foo(const T& t) : make_tuple_traits<const T>::type
-// from template<class T> foo(T& t) : make_tuple_traits<T>::type
-
-// Conversions:
-// T -> T,
-// references -> compile_time_error
-// array -> const ref array
-
-
-template<class T>
-struct make_tuple_traits {
-  typedef T type;
-
-  // commented away, see below  (JJ)
-  //  typedef typename IF<
-  //  boost::is_function<T>::value,
-  //  T&,
-  //  T>::RET type;
-
-};
-
-// The is_function test was there originally for plain function types,
-// which can't be stored as such (we must either store them as references or
-// pointers). Such a type could be formed if make_tuple was called with a
-// reference to a function.
-// But this would mean that a const qualified function type was formed in
-// the make_tuple function and hence make_tuple can't take a function
-// reference as a parameter, and thus T can't be a function type.
-// So is_function test was removed.
-// (14.8.3. says that type deduction fails if a cv-qualified function type
-// is created. (It only applies for the case of explicitly specifying template
-// args, though?)) (JJ)
-
-template<class T>
-struct make_tuple_traits<T&> {
-  typedef typename
-     detail::generate_error<T&>::
-       do_not_use_with_reference_type error;
-};
-
-// Arrays can't be stored as plain types; convert them to references.
-// All arrays are converted to const. This is because make_tuple takes its
-// parameters as const T& and thus the knowledge of the potential
-// non-constness of actual argument is lost.
-template<class T, int n>  struct make_tuple_traits <T[n]> {
-  typedef const T (&type)[n];
-};
-
-template<class T, int n>
-struct make_tuple_traits<const T[n]> {
-  typedef const T (&type)[n];
-};
-
-template<class T, int n>  struct make_tuple_traits<volatile T[n]> {
-  typedef const volatile T (&type)[n];
-};
-
-template<class T, int n>
-struct make_tuple_traits<const volatile T[n]> {
-  typedef const volatile T (&type)[n];
-};
-
-// XXX enable these if we ever care about reference_wrapper -jph
-//template<class T>
-//struct make_tuple_traits<reference_wrapper<T> >{
-//  typedef T& type;
-//};
-//
-//template<class T>
-//struct make_tuple_traits<const reference_wrapper<T> >{
-//  typedef T& type;
-//};
-
-
-// a helper traits to make the make_tuple functions shorter (Vesa Karvonen's
-// suggestion)
-template <
-  class T0 = null_type, class T1 = null_type, class T2 = null_type,
-  class T3 = null_type, class T4 = null_type, class T5 = null_type,
-  class T6 = null_type, class T7 = null_type, class T8 = null_type,
-  class T9 = null_type
->
-struct make_tuple_mapper {
-  typedef
-    tuple<typename make_tuple_traits<T0>::type,
-          typename make_tuple_traits<T1>::type,
-          typename make_tuple_traits<T2>::type,
-          typename make_tuple_traits<T3>::type,
-          typename make_tuple_traits<T4>::type,
-          typename make_tuple_traits<T5>::type,
-          typename make_tuple_traits<T6>::type,
-          typename make_tuple_traits<T7>::type,
-          typename make_tuple_traits<T8>::type,
-          typename make_tuple_traits<T9>::type> type;
-};
-
-} // end detail
-
-
-template<int N, class HT, class TT>
-__host__ __device__
-inline typename access_traits<
-                  typename tuple_element<N, detail::cons<HT, TT> >::type
-                >::non_const_type
-get(detail::cons<HT, TT>& c)
-{
-  //return detail::get_class<N>::BOOST_NESTED_TEMPLATE
-
-  // gcc 4.3 couldn't compile this:
-  //return detail::get_class<N>::
-
-  return detail::get_class<N>::template
-         get<
-           typename access_traits<
-             typename tuple_element<N, detail::cons<HT, TT> >::type
-           >::non_const_type,
-           HT,TT
-         >(c);
-}
-
-
-// get function for const cons-lists, returns a const reference to
-// the element. If the element is a reference, returns the reference
-// as such (that is, can return a non-const reference)
-template<int N, class HT, class TT>
-__host__ __device__
-inline typename access_traits<
-                  typename tuple_element<N, detail::cons<HT, TT> >::type
-                >::const_type
-get(const detail::cons<HT, TT>& c)
-{
-  //return detail::get_class<N>::BOOST_NESTED_TEMPLATE
-
-  // gcc 4.3 couldn't compile this:
-  //return detail::get_class<N>::
-
-  return detail::get_class<N>::template
-         get<
-           typename access_traits<
-             typename tuple_element<N, detail::cons<HT, TT> >::type
-           >::const_type,
-           HT,TT
-         >(c);
-}
-
-
-template<class T0>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0>::type
-    make_tuple(const T0& t0)
-{
-  typedef typename detail::make_tuple_mapper<T0>::type t;
-  return t(t0);
-} // end make_tuple()
-
-template<class T0, class T1>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1>::type
-    make_tuple(const T0& t0, const T1& t1)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1>::type t;
-  return t(t0,t1);
-} // end make_tuple()
-
-template<class T0, class T1, class T2>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2>::type t;
-  return t(t0,t1,t2);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3>::type t;
-  return t(t0,t1,t2,t3);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3, class T4>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3,T4>::type t;
-  return t(t0,t1,t2,t3,t4);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3, class T4, class T5>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3,T4,T5>::type t;
-  return t(t0,t1,t2,t3,t4,t5);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3,T4,T5,T6>::type t;
-  return t(t0,t1,t2,t3,t4,t5,t6);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6, class T7>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6, T7>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6, const T7& t7)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3,T4,T5,T6,T7>::type t;
-  return t(t0,t1,t2,t3,t4,t5,t6,t7);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6, class T7, class T8>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6, T7, T8>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6, const T7& t7, const T8& t8)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3,T4,T5,T6,T7,T8>::type t;
-  return t(t0,t1,t2,t3,t4,t5,t6,t7,t8);
-} // end make_tuple()
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6, class T7, class T8, class T9>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6, T7, T8, T9>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6, const T7& t7, const T8& t8, const T9& t9)
-{
-  typedef typename detail::make_tuple_mapper<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9>::type t;
-  return t(t0,t1,t2,t3,t4,t5,t6,t7,t8,t9);
-} // end make_tuple()
-
-
-template<typename T0>
-__host__ __device__ inline
-tuple<T0&> tie(T0 &t0)
-{
-  return tuple<T0&>(t0);
-}
-
-template<typename T0,typename T1>
-__host__ __device__ inline
-tuple<T0&,T1&> tie(T0 &t0, T1 &t1)
-{
-  return tuple<T0&,T1&>(t0,t1);
-}
-
-template<typename T0,typename T1, typename T2>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&> tie(T0 &t0, T1 &t1, T2 &t2)
-{
-  return tuple<T0&,T1&,T2&>(t0,t1,t2);
-}
-
-template<typename T0,typename T1, typename T2, typename T3>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3)
-{
-  return tuple<T0&,T1&,T2&,T3&>(t0,t1,t2,t3);
-}
-
-template<typename T0,typename T1, typename T2, typename T3, typename T4>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4)
-{
-  return tuple<T0&,T1&,T2&,T3&,T4&>(t0,t1,t2,t3,t4);
-}
-
-template<typename T0,typename T1, typename T2, typename T3, typename T4, typename T5>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5)
-{
-  return tuple<T0&,T1&,T2&,T3&,T4&,T5&>(t0,t1,t2,t3,t4,t5);
-}
-
-template<typename T0,typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6)
-{
-  return tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&>(t0,t1,t2,t3,t4,t5,t6);
-}
-
-template<typename T0,typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6, T7 &t7)
-{
-  return tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&>(t0,t1,t2,t3,t4,t5,t6,t7);
-}
-
-template<typename T0,typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6, T7 &t7, T8 &t8)
-{
-  return tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&>(t0,t1,t2,t3,t4,t5,t6,t7,t8);
-}
-
-template<typename T0,typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&,T9&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6, T7 &t7, T8 &t8, T9 &t9)
-{
-  return tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&,T9&>(t0,t1,t2,t3,t4,t5,t6,t7,t8,t9);
-}
-
-template<
-  typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9,
-  typename U0, typename U1, typename U2, typename U3, typename U4, typename U5, typename U6, typename U7, typename U8, typename U9
->
-__host__ __device__ inline
-void swap(thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> &x,
-          thrust::tuple<U0,U1,U2,U3,U4,U5,U6,U7,U8,U9> &y)
-{
-  return x.swap(y);
-}
-
-
-
-namespace detail
-{
-
-template<class T1, class T2>
-__host__ __device__
-inline bool eq(const T1& lhs, const T2& rhs) {
-  return lhs.get_head() == rhs.get_head() &&
-         eq(lhs.get_tail(), rhs.get_tail());
-}
-template<>
-__host__ __device__
-inline bool eq<null_type,null_type>(const null_type&, const null_type&) { return true; }
-
-template<class T1, class T2>
-__host__ __device__
-inline bool neq(const T1& lhs, const T2& rhs) {
-  return lhs.get_head() != rhs.get_head()  ||
-         neq(lhs.get_tail(), rhs.get_tail());
-}
-template<>
-__host__ __device__
-inline bool neq<null_type,null_type>(const null_type&, const null_type&) { return false; }
-
-template<class T1, class T2>
-__host__ __device__
-inline bool lt(const T1& lhs, const T2& rhs) {
-  return (lhs.get_head() < rhs.get_head())  ||
-            (!(rhs.get_head() < lhs.get_head()) &&
-             lt(lhs.get_tail(), rhs.get_tail()));
-}
-template<>
-__host__ __device__
-inline bool lt<null_type,null_type>(const null_type&, const null_type&) { return false; }
-
-template<class T1, class T2>
-__host__ __device__
-inline bool gt(const T1& lhs, const T2& rhs) {
-  return (lhs.get_head() > rhs.get_head())  ||
-            (!(rhs.get_head() > lhs.get_head()) &&
-             gt(lhs.get_tail(), rhs.get_tail()));
-}
-template<>
-__host__ __device__
-inline bool gt<null_type,null_type>(const null_type&, const null_type&) { return false; }
-
-template<class T1, class T2>
-__host__ __device__
-inline bool lte(const T1& lhs, const T2& rhs) {
-  return lhs.get_head() <= rhs.get_head()  &&
-          ( !(rhs.get_head() <= lhs.get_head()) ||
-            lte(lhs.get_tail(), rhs.get_tail()));
-}
-template<>
-__host__ __device__
-inline bool lte<null_type,null_type>(const null_type&, const null_type&) { return true; }
-
-template<class T1, class T2>
-__host__ __device__
-inline bool gte(const T1& lhs, const T2& rhs) {
-  return lhs.get_head() >= rhs.get_head()  &&
-          ( !(rhs.get_head() >= lhs.get_head()) ||
-            gte(lhs.get_tail(), rhs.get_tail()));
-}
-template<>
-__host__ __device__
-inline bool gte<null_type,null_type>(const null_type&, const null_type&) { return true; }
-
-} // end detail
-
-
-
-// equal ----
-
-template<class T1, class T2, class S1, class S2>
-__host__ __device__
-inline bool operator==(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
-{
-  // XXX support this eventually -jph
-  //// check that tuple lengths are equal
-  //BOOST_STATIC_ASSERT(tuple_size<T2>::value == tuple_size<S2>::value);
-
-  return  detail::eq(lhs, rhs);
-} // end operator==()
-
-// not equal -----
-
-template<class T1, class T2, class S1, class S2>
-__host__ __device__
-inline bool operator!=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
-{
-  // XXX support this eventually -jph
-  //// check that tuple lengths are equal
-  //BOOST_STATIC_ASSERT(tuple_size<T2>::value == tuple_size<S2>::value);
-
-  return detail::neq(lhs, rhs);
-} // end operator!=()
-
-// <
-template<class T1, class T2, class S1, class S2>
-__host__ __device__
-inline bool operator<(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
-{
-  // XXX support this eventually -jph
-  //// check that tuple lengths are equal
-  //BOOST_STATIC_ASSERT(tuple_size<T2>::value == tuple_size<S2>::value);
-
-  return detail::lt(lhs, rhs);
-} // end operator<()
-
-// >
-template<class T1, class T2, class S1, class S2>
-__host__ __device__
-inline bool operator>(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
-{
-  // XXX support this eventually -jph
-  //// check that tuple lengths are equal
-  //BOOST_STATIC_ASSERT(tuple_size<T2>::value == tuple_size<S2>::value);
-
-  return detail::gt(lhs, rhs);
-} // end operator>()
-
-// <=
-template<class T1, class T2, class S1, class S2>
-__host__ __device__
-inline bool operator<=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
-{
-  // XXX support this eventually -jph
-  //// check that tuple lengths are equal
-  //BOOST_STATIC_ASSERT(tuple_size<T2>::value == tuple_size<S2>::value);
-
-  return detail::lte(lhs, rhs);
-} // end operator<=()
-
-// >=
-template<class T1, class T2, class S1, class S2>
-__host__ __device__
-inline bool operator>=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
-{
-  // XXX support this eventually -jph
-  //// check that tuple lengths are equal
-  //BOOST_STATIC_ASSERT(tuple_size<T2>::value == tuple_size<S2>::value);
-
-  return detail::gte(lhs, rhs);
-} // end operator>=()
-
-THRUST_NAMESPACE_END
-
diff --git a/thrust/thrust/iterator/detail/tuple_of_iterator_references.h b/thrust/thrust/iterator/detail/tuple_of_iterator_references.h
index 78c5e8a28f4..b0a37b79e34 100644
--- a/thrust/thrust/iterator/detail/tuple_of_iterator_references.h
+++ b/thrust/thrust/iterator/detail/tuple_of_iterator_references.h
@@ -17,31 +17,40 @@
 #pragma once
 
 #include <thrust/detail/config.h>
+
+#include <cuda/std/type_traits>
+#include <cuda/std/tuple>
+
 #include <thrust/tuple.h>
 #include <thrust/pair.h>
 #include <thrust/detail/reference_forward_declaration.h>
+#include <thrust/detail/raw_reference_cast.h>
 
 THRUST_NAMESPACE_BEGIN
+
 namespace detail
 {
 
-  
 template<
   typename... Ts
 >
-  class tuple_of_iterator_references
-    : public thrust::tuple<Ts...>
+  class tuple_of_iterator_references : public thrust::tuple<Ts...>
 {
-  private:
-    typedef thrust::tuple<Ts...> super_t;
-
   public:
+    using super_t = thrust::tuple<Ts...>;
+    using super_t::super_t;
+
     // allow implicit construction from tuple<refs>
     inline __host__ __device__
-    tuple_of_iterator_references(const super_t &other)
+    tuple_of_iterator_references(const super_t& other)
       : super_t(other)
     {}
 
+    inline __host__ __device__
+    tuple_of_iterator_references(super_t&& other)
+      : super_t(::cuda::std::move(other))
+    {}
+
     // allow assignment from tuples
     // XXX might be worthwhile to guard this with an enable_if is_assignable
     __thrust_exec_check_disable__
@@ -68,79 +77,71 @@ template<
     // XXX perhaps we should generalize to reference<T>
     //     we could captures reference<pair> this way
     __thrust_exec_check_disable__
-    template<typename Pointer, typename Derived,
-             typename... Us>
+    template<typename Pointer, typename Derived, typename... Us>
     inline __host__ __device__
-// XXX gcc-4.2 crashes on is_assignable
-//    typename thrust::detail::enable_if<
-//      thrust::detail::is_assignable<
-//        super_t,
-//        const thrust::tuple<Us...>
-//      >::value,
-//      tuple_of_iterator_references &
-//    >::type
-    tuple_of_iterator_references &
+    tuple_of_iterator_references&
     operator=(const thrust::reference<thrust::tuple<Us...>, Pointer, Derived> &other)
     {
       typedef thrust::tuple<Us...> tuple_type;
 
       // XXX perhaps this could be accelerated
-      tuple_type other_tuple = other;
-      super_t::operator=(other_tuple);
+      super_t::operator=(tuple_type{other});
       return *this;
     }
 
+    template<class... Us, ::cuda::std::__enable_if_t<sizeof...(Us) == sizeof...(Ts), int> = 0>
+    inline __host__ __device__
+    constexpr operator thrust::tuple<Us...>() const {
+      return to_tuple<Us...>(typename ::cuda::std::__make_tuple_indices<sizeof...(Ts)>::type{});
+    }
 
-    // duplicate thrust::tuple's constructors
+    // this overload of swap() permits swapping tuple_of_iterator_references returned as temporaries from
+    // iterator dereferences
+    template<class... Us>
     inline __host__ __device__
-    tuple_of_iterator_references() {}
+    friend void swap(tuple_of_iterator_references x, tuple_of_iterator_references<Us...> y)
+    {
+      x.swap(y);
+    }
 
+private:
+    template<class... Us, size_t... Id>
     inline __host__ __device__
-    tuple_of_iterator_references(typename access_traits<Ts>::parameter_type... ts)
-      : super_t(ts...)
-    {}
+    constexpr thrust::tuple<Us...> to_tuple(::cuda::std::__tuple_indices<Id...>) const {
+      return {get<Id>(*this)...};
+    }
 };
 
+} // end detail
 
-// this overload of swap() permits swapping tuple_of_iterator_references returned as temporaries from
-// iterator dereferences
-template<
-  typename... Ts,
-  typename... Us
->
-inline __host__ __device__
-void swap(tuple_of_iterator_references<Ts...> x,
-          tuple_of_iterator_references<Us...> y)
-{
-  x.swap(y);
-}
-
+THRUST_NAMESPACE_END
 
-} // end detail
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
 
 // define tuple_size, tuple_element, etc.
-template<class... Ts>
-struct tuple_size<detail::tuple_of_iterator_references<Ts...>>
-  : std::integral_constant<size_t, sizeof...(Ts)>
+template <class... Ts>
+struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
+    : integral_constant<size_t, sizeof...(Ts)>
 {};
 
-template<size_t i>
-struct tuple_element<i, detail::tuple_of_iterator_references<>> {};
-
-
-template<class T, class... Ts>
-struct tuple_element<0, detail::tuple_of_iterator_references<T,Ts...>>
-{
-  using type = T;
-};
+template <size_t Id, class... Ts>
+struct tuple_element<Id, THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
+    : _CUDA_VSTD::tuple_element<Id, _CUDA_VSTD::tuple<Ts...>>
+{};
 
+_LIBCUDACXX_END_NAMESPACE_STD
 
-template<size_t i, class T, class... Ts>
-struct tuple_element<i, detail::tuple_of_iterator_references<T,Ts...>>
-{
-  using type = typename tuple_element<i - 1, detail::tuple_of_iterator_references<Ts...>>::type;
-};
+// structured bindings suppport
+namespace std {
 
+template <class... Ts>
+struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
+    : integral_constant<size_t, sizeof...(Ts)>
+{};
 
-THRUST_NAMESPACE_END
+template <size_t Id, class... Ts>
+struct tuple_element<Id, THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
+    : _CUDA_VSTD::tuple_element<Id, _CUDA_VSTD::tuple<Ts...>>
+{};
 
+} // namespace std
diff --git a/thrust/thrust/pair.h b/thrust/thrust/pair.h
index eb2138aaf2c..8b1b6b91f87 100644
--- a/thrust/thrust/pair.h
+++ b/thrust/thrust/pair.h
@@ -21,7 +21,8 @@
 #pragma once
 
 #include <thrust/detail/config.h>
-#include <utility>
+
+#include <cuda/std/utility>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -33,193 +34,6 @@ THRUST_NAMESPACE_BEGIN
  *  \{
  */
 
-/*! \p pair is a generic data structure encapsulating a heterogeneous
- *  pair of values.
- *
- *  \tparam T1 The type of \p pair's first object type.  There are no
- *          requirements on the type of \p T1. <tt>T1</tt>'s type is
- *          provided by <tt>pair::first_type</tt>.
- *
- *  \tparam T2 The type of \p pair's second object type.  There are no
- *          requirements on the type of \p T2. <tt>T2</tt>'s type is
- *          provided by <tt>pair::second_type</tt>.
- */
-template <typename T1, typename T2>
-  struct pair
-{
-  /*! \p first_type is the type of \p pair's first object type.
-   */
-  typedef T1 first_type;
-
-  /*! \p second_type is the type of \p pair's second object type.
-   */
-  typedef T2 second_type;
-
-  /*! The \p pair's first object.
-   */
-  first_type first;
-
-  /*! The \p pair's second object.
-   */
-  second_type second;
-
-  /*! \p pair's default constructor constructs \p first
-   *  and \p second using \c first_type & \c second_type's
-   *  default constructors, respectively.
-   */
-  __host__ __device__ pair(void);
-
-  /*! This constructor accepts two objects to copy into this \p pair.
-   *
-   *  \param x The object to copy into \p first.
-   *  \param y The object to copy into \p second.
-   */
-  inline __host__ __device__
-  pair(const T1 &x, const T2 &y);
-
-  /*! This copy constructor copies from a \p pair whose types are
-   *  convertible to this \p pair's \c first_type and \c second_type,
-   *  respectively.
-   *
-   *  \param p The \p pair to copy from.
-   *
-   *  \tparam U1 is convertible to \c first_type.
-   *  \tparam U2 is convertible to \c second_type.
-   */
-  template <typename U1, typename U2>
-  inline __host__ __device__
-  pair(const pair<U1,U2> &p);
-
-  /*! This copy constructor copies from a <tt>std::pair</tt> whose types are
-   *  convertible to this \p pair's \c first_type and \c second_type,
-   *  respectively.
-   *
-   *  \param p The <tt>std::pair</tt> to copy from.
-   *
-   *  \tparam U1 is convertible to \c first_type.
-   *  \tparam U2 is convertible to \c second_type.
-   */
-  template <typename U1, typename U2>
-  inline __host__ __device__
-  pair(const std::pair<U1,U2> &p);
-
-  /*! \p swap swaps the elements of two <tt>pair</tt>s.
-   *  
-   *  \param p The other <tt>pair</tt> with which to swap.
-   */
-  inline __host__ __device__
-  void swap(pair &p);
-}; // end pair
-
-
-/*! This operator tests two \p pairs for equality.
- *
- *  \param x The first \p pair to compare.
- *  \param y The second \p pair to compare.
- *  \return \c true if and only if <tt>x.first == y.first && x.second == y.second</tt>.
- *  
- *  \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
- *  \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator==(const pair<T1,T2> &x, const pair<T1,T2> &y);
-
-
-/*! This operator tests two pairs for ascending ordering.
- *
- *  \param x The first \p pair to compare.
- *  \param y The second \p pair to compare.
- *  \return \c true if and only if <tt>x.first < y.first || (!(y.first < x.first) && x.second < y.second)</tt>.
- *
- *  \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- *  \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator<(const pair<T1,T2> &x, const pair<T1,T2> &y);
-
-
-/*! This operator tests two pairs for inequality.
- *
- *  \param x The first \p pair to compare.
- *  \param y The second \p pair to compare.
- *  \return \c true if and only if <tt>!(x == y)</tt>.
- *
- *  \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
- *  \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator!=(const pair<T1,T2> &x, const pair<T1,T2> &y);
-
-
-/*! This operator tests two pairs for descending ordering.
- *
- *  \param x The first \p pair to compare.
- *  \param y The second \p pair to compare.
- *  \return \c true if and only if <tt>y < x</tt>.
- *
- *  \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- *  \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator>(const pair<T1,T2> &x, const pair<T1,T2> &y);
-
-
-/*! This operator tests two pairs for ascending ordering or equivalence.
- *
- *  \param x The first \p pair to compare.
- *  \param y The second \p pair to compare.
- *  \return \c true if and only if <tt>!(y < x)</tt>.
- *
- *  \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- *  \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator<=(const pair<T1,T2> &x, const pair<T1,T2> &y);
-
-
-/*! This operator tests two pairs for descending ordering or equivalence.
- *
- *  \param x The first \p pair to compare.
- *  \param y The second \p pair to compare.
- *  \return \c true if and only if <tt>!(x < y)</tt>.
- *
- *  \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- *  \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    bool operator>=(const pair<T1,T2> &x, const pair<T1,T2> &y);
-
-
-/*! \p swap swaps the contents of two <tt>pair</tt>s.
- *
- *  \param x The first \p pair to swap.
- *  \param y The second \p pair to swap.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    void swap(pair<T1,T2> &x, pair<T1,T2> &y);
-
-
-/*! This convenience function creates a \p pair from two objects.
- *
- *  \param x The first object to copy from.
- *  \param y The second object to copy from.
- *  \return A newly-constructed \p pair copied from \p a and \p b.
- *
- *  \tparam T1 There are no requirements on the type of \p T1.
- *  \tparam T2 There are no requirements on the type of \p T2.
- */
-template <typename T1, typename T2>
-  inline __host__ __device__
-    pair<T1,T2> make_pair(T1 x, T2 y);
-
-
 /*! This convenience metafunction is included for compatibility with
  *  \p tuple. It returns either the type of a \p pair's
  *  \c first_type or \c second_type in its nested type, \c type.
@@ -227,8 +41,8 @@ template <typename T1, typename T2>
  *  \tparam N This parameter selects the member of interest.
  *  \tparam T A \c pair type of interest.
  */
-template<size_t N, class T> struct tuple_element;
-
+template <size_t N, class T>
+using tuple_element = ::cuda::std::tuple_element<N, T>;
 
 /*! This convenience metafunction is included for compatibility with
  *  \p tuple. It returns \c 2, the number of elements of a \p pair,
@@ -236,39 +50,28 @@ template<size_t N, class T> struct tuple_element;
  *
  *  \tparam Pair A \c pair type of interest.
  */
-template<typename Pair> struct tuple_size;
+template <class T>
+using tuple_size = ::cuda::std::tuple_size<T>;
 
-
-/*! This convenience function returns a reference to either the first or
- *  second member of a \p pair.
+/*! \p pair is a generic data structure encapsulating a heterogeneous
+ *  pair of values.
  *
- *  \param p The \p pair of interest.
- *  \return \c p.first or \c p.second, depending on the template
- *          parameter.
+ *  \tparam T1 The type of \p pair's first object type.  There are no
+ *          requirements on the type of \p T1. <tt>T1</tt>'s type is
+ *          provided by <tt>pair::first_type</tt>.
  *
- *  \tparam N This parameter selects the member of interest.
+ *  \tparam T2 The type of \p pair's second object type.  There are no
+ *          requirements on the type of \p T2. <tt>T2</tt>'s type is
+ *          provided by <tt>pair::second_type</tt>.
  */
-// XXX comment out these prototypes as a WAR to a problem on MSVC 2005
-//template<unsigned int N, typename T1, typename T2>
-//  inline __host__ __device__
-//    typename tuple_element<N, pair<T1,T2> >::type &
-//      get(pair<T1,T2> &p);
+template <class T, class U>
+using pair = ::cuda::std::pair<T, U>;
 
+using ::cuda::std::get;
+using ::cuda::std::make_pair;
 
-/*! This convenience function returns a const reference to either the
- *  first or second member of a \p pair.
- *
- *  \param p The \p pair of interest.
- *  \return \c p.first or \c p.second, depending on the template
- *          parameter.
- *
- *  \tparam i This parameter selects the member of interest.
+/*! \endcond
  */
-// XXX comment out these prototypes as a WAR to a problem on MSVC 2005
-//template<int N, typename T1, typename T2>
-//  inline __host__ __device__
-//    const typename tuple_element<N, pair<T1,T2> >::type &
-//      get(const pair<T1,T2> &p);
 
 /*! \} // pair
  */
@@ -277,5 +80,3 @@ template<typename Pair> struct tuple_size;
  */
 
 THRUST_NAMESPACE_END
-
-#include <thrust/detail/pair.inl>
diff --git a/thrust/thrust/tuple.h b/thrust/thrust/tuple.h
index 04f3154a3b3..cfea012ee80 100644
--- a/thrust/thrust/tuple.h
+++ b/thrust/thrust/tuple.h
@@ -14,7 +14,6 @@
  *  limitations under the License.
  */
 
-
 /*! \file tuple.h
  *  \brief A type encapsulating a heterogeneous collection of elements.
  */
@@ -31,8 +30,11 @@
 #pragma once
 
 #include <thrust/detail/config.h>
-#include <thrust/detail/tuple.inl>
-#include <thrust/pair.h>
+
+#include <tuple>
+
+#include <cuda/std/tuple>
+#include <cuda/std/utility>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -44,14 +46,6 @@ THRUST_NAMESPACE_BEGIN
  *  \{
  */
 
-/*! \cond
- */
-
-struct null_type;
-
-/*! \endcond
- */
-
 /*! This metafunction returns the type of a
  *  \p tuple's <tt>N</tt>th element.
  *
@@ -61,7 +55,8 @@ struct null_type;
  *  \see pair
  *  \see tuple
  */
-template <size_t N, class T> struct tuple_element;
+template <size_t N, class T>
+using tuple_element = ::cuda::std::tuple_element<N, T>;
 
 /*! This metafunction returns the number of elements
  *  of a \p tuple type of interest.
@@ -71,73 +66,8 @@ template <size_t N, class T> struct tuple_element;
  *  \see pair
  *  \see tuple
  */
-template <class T> struct tuple_size;
-
-
-// get function for non-const cons-lists, returns a reference to the element
-
-/*! The \p get function returns a reference to a \p tuple element of
- *  interest.
- *
- *  \param t A reference to a \p tuple of interest.
- *  \return A reference to \p t's <tt>N</tt>th element.
- *
- *  \tparam N The index of the element of interest.
- *
- *  The following code snippet demonstrates how to use \p get to print
- *  the value of a \p tuple element.
- *
- *  \code
- *  #include <thrust/tuple.h>
- *  #include <iostream>
- *  ...
- *  thrust::tuple<int, const char *> t(13, "thrust");
- *
- *  std::cout << "The 1st value of t is " << thrust::get<0>(t) << std::endl;
- *  \endcode
- *
- *  \see pair
- *  \see tuple
- */
-template<int N, class HT, class TT>
-__host__ __device__
-inline typename access_traits<
-                  typename tuple_element<N, detail::cons<HT, TT> >::type
-                >::non_const_type
-get(detail::cons<HT, TT>& t);
-
-
-/*! The \p get function returns a \c const reference to a \p tuple element of
- *  interest.
- *
- *  \param t A reference to a \p tuple of interest.
- *  \return A \c const reference to \p t's <tt>N</tt>th element.
- *
- *  \tparam N The index of the element of interest.
- *
- *  The following code snippet demonstrates how to use \p get to print
- *  the value of a \p tuple element.
- *
- *  \code
- *  #include <thrust/tuple.h>
- *  #include <iostream>
- *  ...
- *  thrust::tuple<int, const char *> t(13, "thrust");
- *
- *  std::cout << "The 1st value of t is " << thrust::get<0>(t) << std::endl;
- *  \endcode
- *
- *  \see pair
- *  \see tuple
- */
-template<int N, class HT, class TT>
-__host__ __device__
-inline typename access_traits<
-                  typename tuple_element<N, detail::cons<HT, TT> >::type
-                >::const_type
-get(const detail::cons<HT, TT>& t);
-
-
+template <class T>
+using tuple_size = ::cuda::std::tuple_size<T>;
 
 /*! \brief \p tuple is a class template that can be instantiated with up to ten
  *  arguments. Each template argument specifies the type of element in the \p
@@ -155,7 +85,7 @@ get(const detail::cons<HT, TT>& t);
  *  \code
  *  #include <thrust/tuple.h>
  *  #include <iostream>
- *  
+ *
  *  int main() {
  *    // Create a tuple containing an `int`, a `float`, and a string.
  *    thrust::tuple<int, float, const char*> t(13, 0.1f, "thrust");
@@ -178,390 +108,12 @@ get(const detail::cons<HT, TT>& t);
  *  \see tuple_size
  *  \see tie
  */
-template <class T0, class T1, class T2, class T3, class T4,
-          class T5, class T6, class T7, class T8, class T9>
-  class tuple
-  /*! \cond
-   */
-    : public detail::map_tuple_to_cons<T0, T1, T2, T3, T4, T5, T6, T7, T8, T9>::type
-  /*! \endcond
-   */
-{
-  /*! \cond
-   */
-
-  private:
-  typedef typename detail::map_tuple_to_cons<T0, T1, T2, T3, T4, T5, T6, T7, T8, T9>::type inherited;
-
-  /*! \endcond
-   */
-
-  public:
-
-  /*! \p tuple's no-argument constructor initializes each element.
-   */
-  inline __host__ __device__
-  tuple(void) {}
-
-  /*! \p tuple's one-argument constructor copy constructs the first element from the given parameter
-   *     and intializes all other elements.
-   *  \param t0 The value to assign to this \p tuple's first element.
-   */
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0)
-    : inherited(t0,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  /*! \p tuple's one-argument constructor copy constructs the first two elements from the given parameters
-   *     and intializes all other elements.
-   *  \param t0 The value to assign to this \p tuple's first element.
-   *  \param t1 The value to assign to this \p tuple's second element.
-   *  \note \p tuple's constructor has ten variants of this form, the rest of which are ommitted here for brevity.
-   */
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1)
-    : inherited(t0, t1,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  /*! \cond
-   */
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2)
-    : inherited(t0, t1, t2,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3)
-    : inherited(t0, t1, t2, t3,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3,
-        typename access_traits<T4>::parameter_type t4)
-    : inherited(t0, t1, t2, t3, t4,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3,
-        typename access_traits<T4>::parameter_type t4,
-        typename access_traits<T5>::parameter_type t5)
-    : inherited(t0, t1, t2, t3, t4, t5,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3,
-        typename access_traits<T4>::parameter_type t4,
-        typename access_traits<T5>::parameter_type t5,
-        typename access_traits<T6>::parameter_type t6)
-    : inherited(t0, t1, t2, t3, t4, t5, t6,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3,
-        typename access_traits<T4>::parameter_type t4,
-        typename access_traits<T5>::parameter_type t5,
-        typename access_traits<T6>::parameter_type t6,
-        typename access_traits<T7>::parameter_type t7)
-    : inherited(t0, t1, t2, t3, t4, t5, t6, t7,
-                static_cast<const null_type&>(null_type()),
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3,
-        typename access_traits<T4>::parameter_type t4,
-        typename access_traits<T5>::parameter_type t5,
-        typename access_traits<T6>::parameter_type t6,
-        typename access_traits<T7>::parameter_type t7,
-        typename access_traits<T8>::parameter_type t8)
-    : inherited(t0, t1, t2, t3, t4, t5, t6, t7, t8,
-                static_cast<const null_type&>(null_type())) {}
-
-  inline __host__ __device__
-  tuple(typename access_traits<T0>::parameter_type t0,
-        typename access_traits<T1>::parameter_type t1,
-        typename access_traits<T2>::parameter_type t2,
-        typename access_traits<T3>::parameter_type t3,
-        typename access_traits<T4>::parameter_type t4,
-        typename access_traits<T5>::parameter_type t5,
-        typename access_traits<T6>::parameter_type t6,
-        typename access_traits<T7>::parameter_type t7,
-        typename access_traits<T8>::parameter_type t8,
-        typename access_traits<T9>::parameter_type t9)
-    : inherited(t0, t1, t2, t3, t4, t5, t6, t7, t8, t9) {}
-
-
-  template<class U1, class U2>
-  inline __host__ __device__
-  tuple(const detail::cons<U1, U2>& p) : inherited(p) {}
-
-  __thrust_exec_check_disable__
-  template <class U1, class U2>
-  inline __host__ __device__
-  tuple& operator=(const detail::cons<U1, U2>& k)
-  {
-    inherited::operator=(k);
-    return *this;
-  }
-
-  /*! \endcond
-   */
-
-  /*! This assignment operator allows assigning the first two elements of this \p tuple from a \p pair.
-   *  \param k A \p pair to assign from.
-   */
-  __thrust_exec_check_disable__
-  template <class U1, class U2>
-  __host__ __device__ inline
-  tuple& operator=(const thrust::pair<U1, U2>& k) {
-    //BOOST_STATIC_ASSERT(length<tuple>::value == 2);// check_length = 2
-    this->head = k.first;
-    this->tail.head = k.second;
-    return *this;
-  }
-
-  /*! \p swap swaps the elements of two <tt>tuple</tt>s.
-   *
-   *  \param t The other <tt>tuple</tt> with which to swap.
-   */
-  inline __host__ __device__
-  void swap(tuple &t)
-  {
-    inherited::swap(t);
-  }
-};
-
-/*! \cond
- */
-
-template <>
-class tuple<null_type, null_type, null_type, null_type, null_type, null_type, null_type, null_type, null_type, null_type>  :
-  public null_type
-{
-public:
-  typedef null_type inherited;
-};
-
-/*! \endcond
- */
-
-
-/*! This version of \p make_tuple creates a new \c tuple object from a
- *  single object.
- *
- *  \param t0 The object to copy from.
- *  \return A \p tuple object with a single member which is a copy of \p t0.
- */
-template<class T0>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0>::type
-    make_tuple(const T0& t0);
-
-/*! This version of \p make_tuple creates a new \c tuple object from two
- *  objects.
- *
- *  \param t0 The first object to copy from.
- *  \param t1 The second object to copy from.
- *  \return A \p tuple object with two members which are copies of \p t0
- *          and \p t1.
- *
- *  \note \p make_tuple has ten variants, the rest of which are omitted here
- *        for brevity.
- */
-template<class T0, class T1>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1>::type
-    make_tuple(const T0& t0, const T1& t1);
-
-/*! This version of \p tie creates a new \c tuple whose single element is
- *  a reference which refers to this function's argument.
- *
- *  \param t0 The object to reference.
- *  \return A \p tuple object with one member which is a reference to \p t0.
- */
-template<typename T0>
-__host__ __device__ inline
-tuple<T0&> tie(T0& t0);
-
-/*! This version of \p tie creates a new \c tuple of references object which
- *  refers to this function's arguments.
- *
- *  \param t0 The first object to reference.
- *  \param t1 The second object to reference.
- *  \return A \p tuple object with two members which are references to \p t0
- *          and \p t1.
- *
- *  \note \p tie has ten variants, the rest of which are omitted here for
- *           brevity.
- */
-template<typename T0, typename T1>
-__host__ __device__ inline
-tuple<T0&,T1&> tie(T0& t0, T1& t1);
-
-/*! \p swap swaps the contents of two <tt>tuple</tt>s.
- *
- *  \param x The first \p tuple to swap.
- *  \param y The second \p tuple to swap.
- */
-template<
-  typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9,
-  typename U0, typename U1, typename U2, typename U3, typename U4, typename U5, typename U6, typename U7, typename U8, typename U9
->
-inline __host__ __device__
-void swap(tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> &x,
-          tuple<U0,U1,U2,U3,U4,U5,U6,U7,U8,U9> &y);
-
-
-
-/*! \cond
- */
-
-template<class T0, class T1, class T2>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2);
-
-template<class T0, class T1, class T2, class T3>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3);
-
-template<class T0, class T1, class T2, class T3, class T4>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4);
-
-template<class T0, class T1, class T2, class T3, class T4, class T5>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5);
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6);
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6, class T7>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6, T7>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6, const T7& t7);
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6, class T7, class T8>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6, T7, T8>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6, const T7& t7, const T8& t8);
-
-template<class T0, class T1, class T2, class T3, class T4, class T5, class T6, class T7, class T8, class T9>
-__host__ __device__ inline
-  typename detail::make_tuple_mapper<T0, T1, T2, T3, T4, T5, T6, T7, T8, T9>::type
-    make_tuple(const T0& t0, const T1& t1, const T2& t2, const T3& t3, const T4& t4, const T5& t5, const T6& t6, const T7& t7, const T8& t8, const T9& t9);
-
-template<typename T0, typename T1, typename T2>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&> tie(T0 &t0, T1 &t1, T2 &t2);
-
-template<typename T0, typename T1, typename T2, typename T3>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3);
-
-template<typename T0, typename T1, typename T2, typename T3, typename T4>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4);
-
-template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5);
-
-template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6);
-
-template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6, T7 &t7);
-
-template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6, T7 &t7, T8 &t8);
-
-template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
-__host__ __device__ inline
-tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&,T9&> tie(T0 &t0, T1 &t1, T2 &t2, T3 &t3, T4 &t4, T5 &t5, T6 &t6, T7 &t7, T8 &t8, T9 &t9);
-
-
-__host__ __device__ inline
-bool operator==(const null_type&, const null_type&);
-
-__host__ __device__ inline
-bool operator>=(const null_type&, const null_type&);
-
-__host__ __device__ inline
-bool operator<=(const null_type&, const null_type&);
-
-__host__ __device__ inline
-bool operator!=(const null_type&, const null_type&);
-
-__host__ __device__ inline
-bool operator<(const null_type&, const null_type&);
+template <class... T>
+using tuple = ::cuda::std::tuple<T...>;
 
-__host__ __device__ inline
-bool operator>(const null_type&, const null_type&);
+using ::cuda::std::get;
+using ::cuda::std::make_tuple;
+using ::cuda::std::tie;
 
 /*! \endcond
  */