Backport to 2.8: Deprecate thrust::async (#3324) (#3388)

bernhardmgruber · web-flow · commit aa71d81a73ed · 2025-01-21T15:01:07.000+01:00
Fixes: #100
diff --git a/thrust/testing/async/exclusive_scan/mixin.h b/thrust/testing/async/exclusive_scan/mixin.h
@@ -93,12 +93,14 @@ struct simple
     PostfixArgTuple&& postfix_tuple,
     std::index_sequence<PostfixArgIndices...>)
   {
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     auto e = thrust::async::exclusive_scan(
       std::get<PrefixArgIndices>(THRUST_FWD(prefix_tuple))...,
       input.cbegin(),
       input.cend(),
       output.begin(),
       std::get<PostfixArgIndices>(THRUST_FWD(postfix_tuple))...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
     return e;
   }
 };
diff --git a/thrust/testing/async/exclusive_scan/using_vs_adl.cu b/thrust/testing/async/exclusive_scan/using_vs_adl.cu
@@ -69,13 +69,15 @@ struct using_namespace
     // Importing the CPO into the current namespace should unambiguously resolve
     // this call to the CPO, as opposed to resolving to the thrust:: algorithm
     // via ADL. This is verified by checking that an event is returned.
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     using namespace thrust::async;
     thrust::device_event e = exclusive_scan(
       std::get<PrefixArgIndices>(THRUST_FWD(prefix_tuple))...,
       input.cbegin(),
       input.cend(),
       output.begin(),
       std::get<PostfixArgIndices>(THRUST_FWD(postfix_tuple))...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
     return e;
   }
 };
@@ -100,12 +102,14 @@ struct using_cpo
     // this call to the CPO, as opposed to resolving to the thrust:: algorithm
     // via ADL. This is verified by checking that an event is returned.
     using thrust::async::exclusive_scan;
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     thrust::device_event e = exclusive_scan(
       std::get<PrefixArgIndices>(THRUST_FWD(prefix_tuple))...,
       input.cbegin(),
       input.cend(),
       output.begin(),
       std::get<PostfixArgIndices>(THRUST_FWD(postfix_tuple))...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
     return e;
   }
 };
diff --git a/thrust/testing/async/inclusive_scan/mixin.h b/thrust/testing/async/inclusive_scan/mixin.h
@@ -109,12 +109,14 @@ struct simple
     PostfixArgTuple&& postfix_tuple,
     std::index_sequence<PostfixArgIndices...>)
   {
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     auto e = thrust::async::inclusive_scan(
       std::get<PrefixArgIndices>(THRUST_FWD(prefix_tuple))...,
       input.cbegin(),
       input.cend(),
       output.begin(),
       std::get<PostfixArgIndices>(THRUST_FWD(postfix_tuple))...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
     return e;
   }
 };
diff --git a/thrust/testing/async/inclusive_scan/using_vs_adl.cu b/thrust/testing/async/inclusive_scan/using_vs_adl.cu
@@ -69,13 +69,15 @@ struct using_namespace
     // Importing the CPO into the current namespace should unambiguously resolve
     // this call to the CPO, as opposed to resolving to the thrust:: algorithm
     // via ADL. This is verified by checking that an event is returned.
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     using namespace thrust::async;
     thrust::device_event e = inclusive_scan(
       std::get<PrefixArgIndices>(THRUST_FWD(prefix_tuple))...,
       input.cbegin(),
       input.cend(),
       output.begin(),
       std::get<PostfixArgIndices>(THRUST_FWD(postfix_tuple))...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
     return e;
   }
 };
@@ -100,12 +102,14 @@ struct using_cpo
     // this call to the CPO, as opposed to resolving to the thrust:: algorithm
     // via ADL. This is verified by checking that an event is returned.
     using thrust::async::inclusive_scan;
+    _CCCL_SUPPRESS_DEPRECATED_PUSH
     thrust::device_event e = inclusive_scan(
       std::get<PrefixArgIndices>(THRUST_FWD(prefix_tuple))...,
       input.cbegin(),
       input.cend(),
       output.begin(),
       std::get<PostfixArgIndices>(THRUST_FWD(postfix_tuple))...);
+    _CCCL_SUPPRESS_DEPRECATED_POP
     return e;
   }
 };
diff --git a/thrust/testing/async_copy.cu b/thrust/testing/async_copy.cu
@@ -10,6 +10,8 @@
 #  include <unittest/unittest.h>
 #  include <unittest/util_async.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 #  define DEFINE_ASYNC_COPY_CALLABLE(name, ...)                                               \
     struct THRUST_PP_CAT2(name, _fn)                                                          \
     {                                                                                         \
diff --git a/thrust/testing/async_for_each.cu b/thrust/testing/async_for_each.cu
@@ -8,6 +8,8 @@
 
 #  include <unittest/unittest.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 #  define DEFINE_ASYNC_FOR_EACH_CALLABLE(name, ...)                                           \
     struct THRUST_PP_CAT2(name, _fn)                                                          \
     {                                                                                         \
@@ -21,7 +23,6 @@
     /**/
 
 DEFINE_ASYNC_FOR_EACH_CALLABLE(invoke_async_for_each);
-
 DEFINE_ASYNC_FOR_EACH_CALLABLE(invoke_async_for_each_device, thrust::device);
 
 #  undef DEFINE_ASYNC_FOR_EACH_CALLABLE
diff --git a/thrust/testing/async_reduce.cu b/thrust/testing/async_reduce.cu
@@ -12,6 +12,8 @@
 #  include <unittest/unittest.h>
 #  include <unittest/util_async.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 template <typename T>
 struct custom_plus
 {
@@ -539,12 +541,16 @@ struct test_async_reduce_using
     // When you import the customization points into the global namespace,
     // they should be selected instead of the synchronous algorithms.
     {
+      _CCCL_SUPPRESS_DEPRECATED_PUSH
       using namespace thrust::async;
       f0a = reduce(d0a.begin(), d0a.end());
+      _CCCL_SUPPRESS_DEPRECATED_POP
     }
     {
+      _CCCL_SUPPRESS_DEPRECATED_PUSH
       using thrust::async::reduce;
       f0b = reduce(d0b.begin(), d0b.end());
+      _CCCL_SUPPRESS_DEPRECATED_POP
     }
 
     // ADL should find the synchronous algorithms.
diff --git a/thrust/testing/async_reduce_into.cu b/thrust/testing/async_reduce_into.cu
@@ -13,6 +13,8 @@
 #  include <unittest/unittest.h>
 #  include <unittest/util_async.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 template <typename T>
 struct custom_plus
 {
diff --git a/thrust/testing/async_sort.cu b/thrust/testing/async_sort.cu
@@ -15,6 +15,8 @@
 
 #  include <unittest/unittest.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 enum wait_policy
 {
   wait_for_futures,
diff --git a/thrust/testing/async_transform.cu b/thrust/testing/async_transform.cu
@@ -10,6 +10,8 @@
 #  include <unittest/unittest.h>
 #  include <unittest/util_async.h>
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+
 template <typename T>
 struct divide_by_2
 {
@@ -401,12 +403,16 @@ struct test_async_transform_using
       // When you import the customization points into the global namespace,
       // they should be selected instead of the synchronous algorithms.
       {
+        _CCCL_SUPPRESS_DEPRECATED_PUSH
         using namespace thrust::async;
         f0a = transform(d0a.begin(), d0a.end(), d1a.begin(), op);
+        _CCCL_SUPPRESS_DEPRECATED_POP
       }
       {
+        _CCCL_SUPPRESS_DEPRECATED_PUSH
         using thrust::async::transform;
         f0b = transform(d0b.begin(), d0b.end(), d1b.begin(), op);
+        _CCCL_SUPPRESS_DEPRECATED_POP
       }
 
       // ADL should find the synchronous algorithms.
diff --git a/thrust/thrust/async/copy.h b/thrust/thrust/async/copy.h
@@ -52,7 +52,7 @@ namespace unimplemented
 {
 
 template <typename FromPolicy, typename ToPolicy, typename ForwardIt, typename Sentinel, typename OutputIt>
-_CCCL_HOST event<FromPolicy> async_copy(
+CCCL_DEPRECATED _CCCL_HOST event<FromPolicy> async_copy(
   thrust::execution_policy<FromPolicy>& from_exec,
   thrust::execution_policy<ToPolicy>& to_exec,
   ForwardIt first,
@@ -73,6 +73,7 @@ using thrust::async::unimplemented::async_copy;
 
 struct copy_fn final
 {
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   template <typename FromPolicy, typename ToPolicy, typename ForwardIt, typename Sentinel, typename OutputIt>
   _CCCL_HOST static auto
   call(thrust::detail::execution_policy_base<FromPolicy> const& from_exec,
@@ -86,41 +87,43 @@ struct copy_fn final
       thrust::detail::derived_cast(thrust::detail::strip_const(to_exec)),
       THRUST_FWD(first),
       THRUST_FWD(last),
-      THRUST_FWD(output)))
-
-      template <typename DerivedPolicy, typename ForwardIt, typename Sentinel, typename OutputIt>
-      _CCCL_HOST static auto call(thrust::detail::execution_policy_base<DerivedPolicy> const& exec,
-                                  ForwardIt&& first,
-                                  Sentinel&& last,
-                                  OutputIt&& output)
-        THRUST_RETURNS(copy_fn::call(
-          thrust::detail::derived_cast(thrust::detail::strip_const(exec))
-          // Synthesize a suitable new execution policy, because we don't want to
-          // try and extract twice from the one we were passed.
-          ,
-          typename ::cuda::std::remove_cvref_t<
-            decltype(thrust::detail::derived_cast(thrust::detail::strip_const(exec)))>::tag_type{},
+      THRUST_FWD(output))) _CCCL_SUPPRESS_DEPRECATED_POP
+
+    template <typename DerivedPolicy, typename ForwardIt, typename Sentinel, typename OutputIt>
+    _CCCL_HOST static auto call(thrust::detail::execution_policy_base<DerivedPolicy> const& exec,
+                                ForwardIt&& first,
+                                Sentinel&& last,
+                                OutputIt&& output)
+      THRUST_RETURNS(copy_fn::call(
+        thrust::detail::derived_cast(thrust::detail::strip_const(exec))
+        // Synthesize a suitable new execution policy, because we don't want to
+        // try and extract twice from the one we were passed.
+        ,
+        typename ::cuda::std::remove_cvref_t<
+          decltype(thrust::detail::derived_cast(thrust::detail::strip_const(exec)))>::tag_type{},
+        THRUST_FWD(first),
+        THRUST_FWD(last),
+        THRUST_FWD(output)))
+
+        template <typename ForwardIt, typename Sentinel, typename OutputIt>
+        _CCCL_HOST static auto call(ForwardIt&& first, Sentinel&& last, OutputIt&& output) THRUST_RETURNS(copy_fn::call(
+          thrust::detail::select_system(
+            typename thrust::iterator_system<::cuda::std::remove_cvref_t<ForwardIt>>::type{}),
+          thrust::detail::select_system(typename thrust::iterator_system<::cuda::std::remove_cvref_t<OutputIt>>::type{}),
           THRUST_FWD(first),
           THRUST_FWD(last),
           THRUST_FWD(output)))
 
-          template <typename ForwardIt, typename Sentinel, typename OutputIt>
-          _CCCL_HOST static auto call(ForwardIt&& first, Sentinel&& last, OutputIt&& output)
-            THRUST_RETURNS(copy_fn::call(
-              thrust::detail::select_system(
-                typename thrust::iterator_system<::cuda::std::remove_cvref_t<ForwardIt>>::type{}),
-              thrust::detail::select_system(
-                typename thrust::iterator_system<::cuda::std::remove_cvref_t<OutputIt>>::type{}),
-              THRUST_FWD(first),
-              THRUST_FWD(last),
-              THRUST_FWD(output)))
-
-              template <typename... Args>
-              _CCCL_NODISCARD _CCCL_HOST auto operator()(Args&&... args) const THRUST_RETURNS(call(THRUST_FWD(args)...))
+          template <typename... Args>
+          CCCL_DEPRECATED _CCCL_NODISCARD _CCCL_HOST auto operator()(Args&&... args) const
+    THRUST_RETURNS(call(THRUST_FWD(args)...))
 };
 
 } // namespace copy_detail
 
+// note: cannot add a CCCL_DEPRECATED here because the global variable is emitted into cudafe1.stub.c and we cannot
+// suppress the warning there
+//! deprecated [Since 2.8.0]
 _CCCL_GLOBAL_CONSTANT copy_detail::copy_fn copy{};
 
 /*! \endcond
diff --git a/thrust/thrust/async/for_each.h b/thrust/thrust/async/for_each.h
@@ -52,7 +52,7 @@ namespace unimplemented
 {
 
 template <typename DerivedPolicy, typename ForwardIt, typename Sentinel, typename UnaryFunction>
-_CCCL_HOST event<DerivedPolicy>
+CCCL_DEPRECATED _CCCL_HOST event<DerivedPolicy>
 async_for_each(thrust::execution_policy<DerivedPolicy>&, ForwardIt, Sentinel, UnaryFunction)
 {
   THRUST_STATIC_ASSERT_MSG((thrust::detail::depend_on_instantiation<ForwardIt, false>::value),
@@ -69,6 +69,7 @@ using thrust::async::unimplemented::async_for_each;
 
 struct for_each_fn final
 {
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   template <typename DerivedPolicy, typename ForwardIt, typename Sentinel, typename UnaryFunction>
   _CCCL_HOST static auto
   call(thrust::detail::execution_policy_base<DerivedPolicy> const& exec,
@@ -79,22 +80,25 @@ struct for_each_fn final
     THRUST_RETURNS(async_for_each(thrust::detail::derived_cast(thrust::detail::strip_const(exec)),
                                   THRUST_FWD(first),
                                   THRUST_FWD(last),
-                                  THRUST_FWD(f)))
-
-      template <typename ForwardIt, typename Sentinel, typename UnaryFunction>
-      _CCCL_HOST static auto call(ForwardIt&& first, Sentinel&& last, UnaryFunction&& f)
-        THRUST_RETURNS(for_each_fn::call(
-          thrust::detail::select_system(typename iterator_system<::cuda::std::remove_cvref_t<ForwardIt>>::type{}),
-          THRUST_FWD(first),
-          THRUST_FWD(last),
-          THRUST_FWD(f)))
-
-          template <typename... Args>
-          _CCCL_NODISCARD _CCCL_HOST auto operator()(Args&&... args) const THRUST_RETURNS(call(THRUST_FWD(args)...))
+                                  THRUST_FWD(f))) _CCCL_SUPPRESS_DEPRECATED_POP
+
+    template <typename ForwardIt, typename Sentinel, typename UnaryFunction>
+    _CCCL_HOST static auto call(ForwardIt&& first, Sentinel&& last, UnaryFunction&& f) THRUST_RETURNS(for_each_fn::call(
+      thrust::detail::select_system(typename iterator_system<::cuda::std::remove_cvref_t<ForwardIt>>::type{}),
+      THRUST_FWD(first),
+      THRUST_FWD(last),
+      THRUST_FWD(f)))
+
+      template <typename... Args>
+      CCCL_DEPRECATED _CCCL_NODISCARD _CCCL_HOST auto operator()(Args&&... args) const
+    THRUST_RETURNS(call(THRUST_FWD(args)...))
 };
 
 } // namespace for_each_detail
 
+// note: cannot add a CCCL_DEPRECATED here because the global variable is emitted into cudafe1.stub.c and we cannot
+// suppress the warning there
+//! deprecated [Since 2.8.0]
 _CCCL_GLOBAL_CONSTANT for_each_detail::for_each_fn for_each{};
 
 /*! \endcond
diff --git a/thrust/thrust/async/reduce.h b/thrust/thrust/async/reduce.h
diff --git a/thrust/thrust/async/scan.h b/thrust/thrust/async/scan.h
diff --git a/thrust/thrust/async/sort.h b/thrust/thrust/async/sort.h
diff --git a/thrust/thrust/async/transform.h b/thrust/thrust/async/transform.h

Original file line number	Diff line number	Diff line change
`@@ -12,6 +12,8 @@`
`12`	`12`	`# include <unittest/unittest.h>`
`13`	`13`	`# include <unittest/util_async.h>`
`14`	`14`
	`15`	`+_CCCL_SUPPRESS_DEPRECATED_PUSH`
	`16`	`+`
`15`	`17`	`template <typename T>`
`16`	`18`	`struct custom_plus`
`17`	`19`	`{`
`@@ -539,12 +541,16 @@ struct test_async_reduce_using`
`539`	`541`	`// When you import the customization points into the global namespace,`
`540`	`542`	`// they should be selected instead of the synchronous algorithms.`
`541`	`543`	`{`
	`544`	`+ _CCCL_SUPPRESS_DEPRECATED_PUSH`
`542`	`545`	`using namespace thrust::async;`
`543`	`546`	`f0a = reduce(d0a.begin(), d0a.end());`
	`547`	`+ _CCCL_SUPPRESS_DEPRECATED_POP`
`544`	`548`	`}`
`545`	`549`	`{`
	`550`	`+ _CCCL_SUPPRESS_DEPRECATED_PUSH`
`546`	`551`	`using thrust::async::reduce;`
`547`	`552`	`f0b = reduce(d0b.begin(), d0b.end());`
	`553`	`+ _CCCL_SUPPRESS_DEPRECATED_POP`
`548`	`554`	`}`
`549`	`555`
`550`	`556`	`// ADL should find the synchronous algorithms.`
Original file line number	Diff line number	Diff line change
`@@ -13,6 +13,8 @@`
`13`	`13`	`# include <unittest/unittest.h>`
`14`	`14`	`# include <unittest/util_async.h>`
`15`	`15`
	`16`	`+_CCCL_SUPPRESS_DEPRECATED_PUSH`
	`17`	`+`
`16`	`18`	`template <typename T>`
`17`	`19`	`struct custom_plus`
`18`	`20`	`{`
Original file line number	Diff line number	Diff line change
`@@ -15,6 +15,8 @@`
`15`	`15`
`16`	`16`	`# include <unittest/unittest.h>`
`17`	`17`
	`18`	`+_CCCL_SUPPRESS_DEPRECATED_PUSH`
	`19`	`+`
`18`	`20`	`enum wait_policy`
`19`	`21`	`{`
`20`	`22`	`wait_for_futures,`
Original file line number	Diff line number	Diff line change
`@@ -10,6 +10,8 @@`
`10`	`10`	`# include <unittest/unittest.h>`
`11`	`11`	`# include <unittest/util_async.h>`
`12`	`12`
	`13`	`+_CCCL_SUPPRESS_DEPRECATED_PUSH`
	`14`	`+`
`13`	`15`	`template <typename T>`
`14`	`16`	`struct divide_by_2`
`15`	`17`	`{`
`@@ -401,12 +403,16 @@ struct test_async_transform_using`
`401`	`403`	`// When you import the customization points into the global namespace,`
`402`	`404`	`// they should be selected instead of the synchronous algorithms.`
`403`	`405`	`{`
	`406`	`+ _CCCL_SUPPRESS_DEPRECATED_PUSH`
`404`	`407`	`using namespace thrust::async;`
`405`	`408`	`f0a = transform(d0a.begin(), d0a.end(), d1a.begin(), op);`
	`409`	`+ _CCCL_SUPPRESS_DEPRECATED_POP`
`406`	`410`	`}`
`407`	`411`	`{`
	`412`	`+ _CCCL_SUPPRESS_DEPRECATED_PUSH`
`408`	`413`	`using thrust::async::transform;`
`409`	`414`	`f0b = transform(d0b.begin(), d0b.end(), d1b.begin(), op);`
	`415`	`+ _CCCL_SUPPRESS_DEPRECATED_POP`
`410`	`416`	`}`
`411`	`417`
`412`	`418`	`// ADL should find the synchronous algorithms.`