Deprecate thrust::async (NVIDIA#3324)

Fixes: NVIDIA#100
davebayer · Jan 18, 2025 · 9224888 · 9224888
1 parent 5ff9e06
commit 9224888
Show file tree

Hide file tree

Showing 16 changed files with 321 additions and 233 deletions.
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
 {
@@ -395,12 +397,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
@@ -50,7 +50,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,
@@ -72,7 +72,7 @@ using thrust::async::unimplemented::async_copy;
 struct copy_fn final
 {
   template <typename FromPolicy, typename ToPolicy, typename ForwardIt, typename Sentinel, typename OutputIt>
-  _CCCL_HOST static auto
+  _CCCL_HOST _CCCL_SUPPRESS_DEPRECATED_PUSH static auto
   call(thrust::detail::execution_policy_base<FromPolicy> const& from_exec,
        thrust::detail::execution_policy_base<ToPolicy> const& to_exec,
        ForwardIt&& first,
@@ -84,38 +84,41 @@ 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 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 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<remove_cvref_t<ForwardIt>>::type{}),
+          thrust::detail::select_system(typename thrust::iterator_system<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<remove_cvref_t<ForwardIt>>::type{}),
-              thrust::detail::select_system(typename thrust::iterator_system<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
@@ -50,7 +50,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),
@@ -68,7 +68,7 @@ using thrust::async::unimplemented::async_for_each;
 struct for_each_fn final
 {
   template <typename DerivedPolicy, typename ForwardIt, typename Sentinel, typename UnaryFunction>
-  _CCCL_HOST static auto
+  _CCCL_HOST _CCCL_SUPPRESS_DEPRECATED_PUSH static auto
   call(thrust::detail::execution_policy_base<DerivedPolicy> const& exec,
        ForwardIt&& first,
        Sentinel&& last,
@@ -77,22 +77,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<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<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