Replace thrust::tuple implementation with cuda::std::tuple

Our old implementation of `tuple` has a lot of limitations. So rather than using the old homegrown version reuse `cuda::std::tuple`
NVIDIA · Mar 6, 2023 · 57adb2b · 57adb2b
1 parent 70ff353
commit 57adb2b
Show file tree

Hide file tree

Showing 16 changed files with 114 additions and 2,051 deletions.
diff --git a/testing/pair.cu b/testing/pair.cu
@@ -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);
+
+  [a2, b2] = p;
+  ASSERT_EQUAL(a, a2);
+  ASSERT_EQUAL(b, b2);
+}
+DECLARE_UNITTEST(TestPairStructuredBindings);
+#endif
diff --git a/testing/tuple.cu b/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);
+
+  [a2, b2, c2] = t;
+  ASSERT_EQUAL(a, a2);
+  ASSERT_EQUAL(b, b2);
+  ASSERT_EQUAL(c, c2);
+}
+DECLARE_UNITTEST(TestTupleStructuredBindings);
+#endif
diff --git a/testing/tuple_sort.cu b/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)
+  decltype(auto) operator()(const Tuple &t)
   {
     return thrust::get<N>(t);
   }

diff --git a/testing/tuple_transform.cu b/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)
+  decltype(auto) operator()(const Tuple &t)
   {
     return thrust::get<N>(t);
   }

diff --git a/testing/zip_iterator.cu b/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/detail/functional/actor.h b/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/detail/functional/actor.inl b/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/detail/functional/argument.h b/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/detail/functional/composite.h b/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/detail/functional/operators/operator_adaptors.h b/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;