Fix the one-policy overload of `thrust::async::copy` to not copy the policy,
because this leads to use-after-move in the future dependency extraction and
stream stealing code. Fixes
MODIFY
testing/async_copy.cu
testing/async_copy.cu
@@ -101,7 +101,7 @@ struct test_async_copy_device_to_host
void operator()(std::size_t n)
{
thrust::host_vector<T> h0(unittest::random_integers<T>(n));
- thrust::device_vector<T> h1(n);
+ thrust::host_vector<T> h1(n);
thrust::device_vector<T> d0(n);
thrust::copy(h0.begin(), h0.end(), d0.begin());
@@ -319,6 +319,84 @@ DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
///////////////////////////////////////////////////////////////////////////////
+template <typename T>
+struct test_async_copy_after
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0(unittest::random_integers<T>(n));
+ thrust::host_vector<T> h1(n);
+ thrust::device_vector<T> d0(n);
+ thrust::device_vector<T> d1(n);
+ thrust::device_vector<T> d2(n);
+
+ auto e0 = thrust::async::copy(
+ h0.begin(), h0.end(), d0.begin()
+ );
+
+ ASSERT_EQUAL(true, e0.valid_stream());
+
+ auto const e0_stream = e0.stream().native_handle();
+
+ auto e1 = thrust::async::copy(
+ thrust::device.after(e0), d0.begin(), d0.end(), d1.begin()
+ );
+
+ // Verify that double consumption of a future produces an exception.
+ ASSERT_THROWS_EQUAL(
+ auto x = thrust::async::copy(
+ thrust::device.after(e0), d0.begin(), d0.end(), d1.begin()
+ );
+ THRUST_UNUSED_VAR(x)
+ , thrust::event_error
+ , thrust::event_error(thrust::event_errc::no_state)
+ );
+
+ ASSERT_EQUAL_QUIET(e0_stream, e1.stream().native_handle());
+
+ auto after_policy2 = thrust::device.after(e1);
+
+ auto e2 = thrust::async::copy(
+ thrust::host, after_policy2
+ , h0.begin(), h0.end(), d2.begin()
+ );
+
+ // Verify that double consumption of a policy produces an exception.
+ ASSERT_THROWS_EQUAL(
+ auto x = thrust::async::copy(
+ thrust::host, after_policy2
+ , h0.begin(), h0.end(), d2.begin()
+ );
+ THRUST_UNUSED_VAR(x)
+ , thrust::event_error
+ , thrust::event_error(thrust::event_errc::no_state)
+ );
+
+ ASSERT_EQUAL_QUIET(e0_stream, e2.stream().native_handle());
+
+ auto e3 = thrust::async::copy(
+ thrust::device.after(e2), thrust::host
+ , d1.begin(), d1.end(), h1.begin()
+ );
+
+ ASSERT_EQUAL_QUIET(e0_stream, e3.stream().native_handle());
+
+ TEST_EVENT_WAIT(e3);
+
+ ASSERT_EQUAL(h0, h1);
+ ASSERT_EQUAL(h0, d0);
+ ASSERT_EQUAL(h0, d1);
+ ASSERT_EQUAL(h0, d2);
+ }
+};
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES(
+ test_async_copy_after
+, BuiltinNumericTypes
+);
+
+///////////////////////////////////////////////////////////////////////////////
+
// TODO: device_to_device NonContiguousIterator output (discard_iterator).
// TODO: host_to_device non trivially relocatable.

MODIFY
thrust/async/copy.h
thrust/async/copy.h
@@ -107,15 +107,20 @@ struct copy_fn final
, ForwardIt&& first, Sentinel&& last
, OutputIt&& output
)
- // ADL dispatch.
- THRUST_DECLTYPE_RETURNS(
- async_copy(
+// THRUST_DECLTYPE_RETURNS(
+ { return
+ copy_fn::call(
thrust::detail::derived_cast(thrust::detail::strip_const(exec))
- , 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>
__host__

MODIFY
thrust/iterator/detail/any_system_tag.h
thrust/iterator/detail/any_system_tag.h
@@ -30,8 +30,5 @@ struct any_system_tag
template<typename T> operator T () const {return T();}
};
-// TODO remove this in 1.7.0
-typedef THRUST_DEPRECATED any_system_tag any_space_tag;
-
} // end thrust

Fix the one-policy overload of `thrust::async::copy` to not copy the policy,
because this leads to use-after-move in the future dependency extraction and
stream stealing code. Fixes #1115.