`thrust::async::copy`: Add an overload that takes two policies to allow explicit
cross system copies.
Bug
MODIFY
testing/async_copy.cu
testing/async_copy.cu
@@ -30,6 +30,29 @@ DECLARE_VARIABLE_UNITTEST(
test_async_copy_host_to_device_trivially_relocatable
);
+template <typename T>
+__host__
+void
+test_async_copy_host_to_device_trivially_relocatable_with_policies(
+ std::size_t n
+)
+{
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(n);
+
+ auto f0 = thrust::async::copy(
+ thrust::host, thrust::device
+ , h0_data.begin(), h0_data.end(), d0_data.begin()
+ );
+
+ std::move(f0).get();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+}
+DECLARE_VARIABLE_UNITTEST(
+ test_async_copy_host_to_device_trivially_relocatable_with_policies
+);
+
template <typename T>
__host__
void
@@ -58,6 +81,35 @@ DECLARE_VARIABLE_UNITTEST(
test_async_copy_device_to_host_trivially_relocatable
);
+template <typename T>
+__host__
+void
+test_async_copy_device_to_host_trivially_relocatable_with_policies(
+ std::size_t n
+)
+{
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> h1_data(n);
+ thrust::device_vector<T> d0_data(n);
+
+ thrust::copy(h0_data.begin(), h0_data.end(), d0_data.begin());
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ auto f0 = thrust::async::copy(
+ thrust::device, thrust::host
+ , d0_data.begin(), d0_data.end(), h1_data.begin()
+ );
+
+ std::move(f0).get();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ ASSERT_EQUAL(d0_data, h1_data);
+}
+DECLARE_VARIABLE_UNITTEST(
+ test_async_copy_device_to_host_trivially_relocatable_with_policies
+);
+
template <typename T>
struct test_async_copy_device_to_device
{

MODIFY
thrust/async/copy.h
thrust/async/copy.h
@@ -40,20 +40,22 @@ namespace unimplemented
{
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename Sentinel, typename OutputIt
>
__host__ __device__
+auto async_copy(
+ thrust::execution_policy<FromPolicy>& from_exec
+, thrust::execution_policy<ToPolicy>& to_exec
+, ForwardIt first, Sentinel last, OutputIt output
+) ->
future<
- OutputIt, DerivedPolicy
+ OutputIt
+ , decltype(thrust::detail::select_system(from_exec, to_exec))
, typename thrust::detail::pointer_traits<
thrust::host_memory_resource::pointer
>::template rebind<OutputIt>::other
>
-async_copy(
- thrust::execution_policy<DerivedPolicy>& exec
-, ForwardIt first, Sentinel last, OutputIt output
-)
{
THRUST_STATIC_ASSERT_MSG(
(thrust::detail::depend_on_instantiation<ForwardIt, false>::value)
@@ -64,34 +66,53 @@ async_copy(
} // namespace unimplemented
+namespace copy_detail
+{
+
struct copy_fn final
{
__thrust_exec_check_disable__
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename Sentinel, typename OutputIt
>
__host__ __device__
- future<
- OutputIt, DerivedPolicy
- , typename thrust::detail::pointer_traits<
- thrust::host_memory_resource::pointer
- >::template rebind<OutputIt>::other
+ static auto call(
+ thrust::detail::execution_policy_base<FromPolicy> const& from_exec
+ , thrust::detail::execution_policy_base<ToPolicy> const& to_exec
+ , ForwardIt&& first, Sentinel&& last
+ , OutputIt&& output
+ )
+ // ADL dispatch.
+ THRUST_DECLTYPE_RETURNS(
+ async_copy(
+ thrust::detail::derived_cast(thrust::detail::strip_const(from_exec))
+ , thrust::detail::derived_cast(thrust::detail::strip_const(to_exec))
+ , THRUST_FWD(first), THRUST_FWD(last)
+ , THRUST_FWD(output)
+ )
+ )
+
+ __thrust_exec_check_disable__
+ template <
+ typename DerivedPolicy
+ , typename ForwardIt, typename Sentinel, typename OutputIt
>
- static call(
+ __host__ __device__
+ static auto call(
thrust::detail::execution_policy_base<DerivedPolicy> const& exec
, ForwardIt&& first, Sentinel&& last
, OutputIt&& output
)
- {
// ADL dispatch.
- using thrust::async::unimplemented::async_copy;
- return async_copy(
+ THRUST_DECLTYPE_RETURNS(
+ async_copy(
thrust::detail::derived_cast(thrust::detail::strip_const(exec))
+ , thrust::detail::derived_cast(thrust::detail::strip_const(exec))
, THRUST_FWD(first), THRUST_FWD(last)
, THRUST_FWD(output)
- );
- }
+ )
+ )
__thrust_exec_check_disable__
template <typename ForwardIt, typename Sentinel, typename OutputIt>
@@ -99,10 +120,8 @@ struct copy_fn final
static auto call(ForwardIt&& first, Sentinel&& last, OutputIt&& output)
THRUST_DECLTYPE_RETURNS(
copy_fn::call(
- thrust::detail::select_system(
typename thrust::iterator_system<ForwardIt>::type{}
, typename thrust::iterator_system<OutputIt>::type{}
- )
, THRUST_FWD(first), THRUST_FWD(last)
, THRUST_FWD(output)
)
@@ -116,7 +135,9 @@ struct copy_fn final
)
};
-THRUST_INLINE_CONSTANT copy_fn copy{};
+} // namespace copy_detail
+
+THRUST_INLINE_CONSTANT copy_detail::copy_fn copy{};
} // namespace async

MODIFY
thrust/system/cuda/detail/async/copy.h
thrust/system/cuda/detail/async/copy.h
@@ -60,12 +60,13 @@ namespace system { namespace cuda { namespace detail
// TriviallyRelocatable value type
// Device to host, host to device
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
THRUST_RUNTIME_FUNCTION
auto async_copy_n(
- execution_policy<DerivedPolicy>& policy
+ FromPolicy& from_exec
+, ToPolicy& to_exec
, ForwardIt first
, Size n
, OutputIt output
@@ -78,14 +79,16 @@ auto async_copy_n(
, typename iterator_traits<OutputIt>::value_type
>
, disjunction<
- decltype(is_host_to_device_copy(policy))
- , decltype(is_device_to_host_copy(policy))
+ decltype(is_host_to_device_copy(from_exec, to_exec))
+ , decltype(is_device_to_host_copy(from_exec, to_exec))
>
>::value
, unique_eager_future<
OutputIt
, typename thrust::detail::allocator_traits<
- decltype(get_async_universal_host_pinned_allocator(policy))
+ decltype(get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ ))
>::template rebind_traits<OutputIt>::pointer
>
>::type
@@ -102,7 +105,10 @@ auto async_copy_n(
// Workaround for an NVCC bug; when two SFINAE-enabled overloads are only
// distinguishable by a part of a SFINAE condition that is in a `decltype`,
// NVCC thinks they are the same overload and emits an error.
-template <typename ExecutionPolicy, typename ForwardIt, typename OutputIt>
+template <
+ typename FromPolicy, typename ToPolicy
+, typename ForwardIt, typename OutputIt
+>
struct is_buffered_trivially_relocatable_host_to_device_copy
: thrust::integral_constant<
bool
@@ -112,7 +118,12 @@ struct is_buffered_trivially_relocatable_host_to_device_copy
typename iterator_traits<ForwardIt>::value_type
, typename iterator_traits<OutputIt>::value_type
>::value
- && decltype(is_host_to_device_copy(std::declval<ExecutionPolicy>()))::value
+ && decltype(
+ is_host_to_device_copy(
+ std::declval<FromPolicy const&>()
+ , std::declval<ToPolicy const&>()
+ )
+ )::value
>
{};
@@ -120,24 +131,29 @@ struct is_buffered_trivially_relocatable_host_to_device_copy
// TriviallyRelocatable value type
// Host to device
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
THRUST_RUNTIME_FUNCTION
auto async_copy_n(
- execution_policy<DerivedPolicy>& policy
+ FromPolicy& from_exec
+, thrust::system::cuda::execution_policy<ToPolicy>& to_exec
, ForwardIt first
, Size n
, OutputIt output
) ->
typename std::enable_if<
is_buffered_trivially_relocatable_host_to_device_copy<
- execution_policy<DerivedPolicy>, ForwardIt, OutputIt
+ FromPolicy
+ , thrust::system::cuda::execution_policy<ToPolicy>
+ , ForwardIt, OutputIt
>::value
, unique_eager_future<
OutputIt
, typename thrust::detail::allocator_traits<
- decltype(get_async_universal_host_pinned_allocator(policy))
+ decltype(get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ ))
>::template rebind_traits<OutputIt>::pointer
>
>::type
@@ -157,7 +173,10 @@ auto async_copy_n(
// Workaround for an NVCC bug; when two SFINAE-enabled overloads are only
// distinguishable by a part of a SFINAE condition that is in a `decltype`,
// NVCC thinks they are the same overload and emits an error.
-template <typename ExecutionPolicy, typename ForwardIt, typename OutputIt>
+template <
+ typename FromPolicy, typename ToPolicy
+, typename ForwardIt, typename OutputIt
+>
struct is_buffered_trivially_relocatable_device_to_host_copy
: thrust::integral_constant<
bool
@@ -167,7 +186,12 @@ struct is_buffered_trivially_relocatable_device_to_host_copy
typename iterator_traits<ForwardIt>::value_type
, typename iterator_traits<OutputIt>::value_type
>::value
- && decltype(is_device_to_host_copy(std::declval<ExecutionPolicy>()))::value
+ && decltype(
+ is_device_to_host_copy(
+ std::declval<FromPolicy const&>()
+ , std::declval<ToPolicy const&>()
+ )
+ )::value
>
{};
@@ -175,24 +199,29 @@ struct is_buffered_trivially_relocatable_device_to_host_copy
// TriviallyRelocatable value type
// Device to host
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
THRUST_RUNTIME_FUNCTION
auto async_copy_n(
- execution_policy<DerivedPolicy>& policy
+ thrust::system::cuda::execution_policy<FromPolicy>& from_exec
+, ToPolicy& to_exec
, ForwardIt first
, Size n
, OutputIt output
) ->
typename std::enable_if<
is_buffered_trivially_relocatable_device_to_host_copy<
- execution_policy<DerivedPolicy>, ForwardIt, OutputIt
+ thrust::system::cuda::execution_policy<FromPolicy>
+ , ToPolicy
+ , ForwardIt, OutputIt
>::value
, unique_eager_future<
OutputIt
, typename thrust::detail::allocator_traits<
- decltype(get_async_universal_host_pinned_allocator(policy))
+ decltype(get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ ))
>::template rebind_traits<OutputIt>::pointer
>
>::type
@@ -222,12 +251,13 @@ void async_copy_n_compile_failure_non_trivially_relocatable_elements()
// Non-TriviallyRelocatable value type
// Host to device, device to host
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
THRUST_RUNTIME_FUNCTION
auto async_copy_n(
- execution_policy<DerivedPolicy>& policy
+ FromPolicy& from_exec
+, ToPolicy& to_exec
, ForwardIt first
, Size n
, OutputIt output
@@ -241,14 +271,16 @@ auto async_copy_n(
>
>
, disjunction<
- decltype(is_host_to_device_copy(policy))
- , decltype(is_device_to_host_copy(policy))
+ decltype(is_host_to_device_copy(from_exec, to_exec))
+ , decltype(is_device_to_host_copy(from_exec, to_exec))
>
>::value
, unique_eager_future<
OutputIt
, typename thrust::detail::allocator_traits<
- decltype(get_async_universal_host_pinned_allocator(policy))
+ decltype(get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ ))
>::template rebind_traits<OutputIt>::pointer
>
>::type
@@ -268,12 +300,13 @@ auto async_copy_n(
// Non-ContiguousIterator input or output iterator, or non-TriviallyRelocatable value type
// Device to device
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
THRUST_RUNTIME_FUNCTION
auto async_copy_n(
- execution_policy<DerivedPolicy>& policy
+ thrust::system::cuda::execution_policy<FromPolicy>& from_exec
+, thrust::system::cuda::execution_policy<ToPolicy>& to_exec
, ForwardIt first
, Size n
, OutputIt output
@@ -283,31 +316,37 @@ auto async_copy_n(
negation<
thrust::is_trivially_relocatable_sequence_copy<ForwardIt, OutputIt>
>
- , decltype(is_device_to_device_copy(policy))
+ , decltype(is_device_to_device_copy(from_exec, to_exec))
>::value
, unique_eager_future<
OutputIt
, typename thrust::detail::allocator_traits<
- decltype(get_async_universal_host_pinned_allocator(policy))
+ decltype(get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ ))
>::template rebind_traits<OutputIt>::pointer
>
>::type
{
using T = typename thrust::iterator_traits<ForwardIt>::value_type;
- return async_transform_n(policy, first, n, output, thrust::identity<T>());
+ return async_transform_n(
+ select_device_system(from_exec, to_exec)
+ , first, n, output, thrust::identity<T>()
+ );
}
// ContiguousIterator input and output iterators
// TriviallyCopyable elements
// Host to device, device to host, device to device
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
THRUST_RUNTIME_FUNCTION
auto async_copy_n(
- execution_policy<DerivedPolicy>& policy
+ FromPolicy& from_exec
+, ToPolicy& to_exec
, ForwardIt first
, Size n
, OutputIt output
@@ -317,14 +356,18 @@ auto async_copy_n(
, unique_eager_future<
OutputIt
, typename thrust::detail::allocator_traits<
- decltype(get_async_universal_host_pinned_allocator(policy))
+ decltype(get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ ))
>::template rebind_traits<OutputIt>::pointer
>
>::type
{
using T = typename thrust::iterator_traits<ForwardIt>::value_type;
- auto const uhp_alloc = get_async_universal_host_pinned_allocator(policy);
+ auto const uhp_alloc = get_async_universal_host_pinned_allocator(
+ select_device_system(from_exec, to_exec)
+ );
using return_type = OutputIt;
@@ -340,7 +383,9 @@ auto async_copy_n(
// Set up stream with dependencies.
- cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(policy);
+ cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(
+ select_device_system(from_exec, to_exec)
+ );
if (thrust::cuda_cub::default_stream() != user_raw_stream)
{
@@ -371,7 +416,7 @@ auto async_copy_n(
thrust::raw_pointer_cast(&*output)
, thrust::raw_pointer_cast(&*first)
, sizeof(T) * n
- , direction_of_copy(policy)
+ , direction_of_copy(from_exec, to_exec)
, fp.future.stream()
)
, "after copy launch"
@@ -387,19 +432,58 @@ namespace cuda_cub
// ADL entry point.
template <
- typename DerivedPolicy
+ typename FromPolicy, typename ToPolicy
+, typename ForwardIt, typename Sentinel, typename OutputIt
+>
+THRUST_RUNTIME_FUNCTION
+auto async_copy(
+ thrust::cuda::execution_policy<FromPolicy>& from_exec
+, thrust::cpp::execution_policy<ToPolicy>& to_exec
+, ForwardIt first
+, Sentinel last
+, OutputIt output
+)
+THRUST_DECLTYPE_RETURNS(
+ thrust::system::cuda::detail::async_copy_n(
+ from_exec, to_exec, first, thrust::distance(first, last), output
+ )
+)
+
+// ADL entry point.
+template <
+ typename FromPolicy, typename ToPolicy
+, typename ForwardIt, typename Sentinel, typename OutputIt
+>
+THRUST_RUNTIME_FUNCTION
+auto async_copy(
+ thrust::cpp::execution_policy<FromPolicy>& from_exec
+, thrust::system::cuda::execution_policy<ToPolicy>& to_exec
+, ForwardIt first
+, Sentinel last
+, OutputIt output
+)
+THRUST_DECLTYPE_RETURNS(
+ thrust::system::cuda::detail::async_copy_n(
+ from_exec, to_exec, first, thrust::distance(first, last), output
+ )
+)
+
+// ADL entry point.
+template <
+ typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename Sentinel, typename OutputIt
>
THRUST_RUNTIME_FUNCTION
auto async_copy(
- execution_policy<DerivedPolicy>& policy
+ thrust::system::cuda::execution_policy<FromPolicy>& from_exec
+, thrust::system::cuda::execution_policy<ToPolicy>& to_exec
, ForwardIt first
, Sentinel last
, OutputIt output
)
THRUST_DECLTYPE_RETURNS(
thrust::system::cuda::detail::async_copy_n(
- policy, first, thrust::distance(first, last), output
+ from_exec, to_exec, first, thrust::distance(first, last), output
)
)

MODIFY
thrust/system/cuda/detail/async/sort.h
thrust/system/cuda/detail/async/sort.h
@@ -102,24 +102,12 @@ auto async_stable_sort_n(
Size n,
StrictWeakOrdering comp
) ->
-/* typename std::enable_if<
- conjunction<
- is_contiguous_iterator<ForwardIt>
- , negation<
- std::is_scalar<
- typename thrust::iterator_traits<ForwardIt>::value_type
- >
- >
- >::value
- ,
-*/
unique_eager_future<
void
, typename thrust::detail::allocator_traits<
decltype(get_async_device_allocator(policy))
>::template rebind_traits<void>::pointer
>
-// >::type
{
using T = typename thrust::iterator_traits<ForwardIt>::value_type;

MODIFY
thrust/system/cuda/detail/cross_system.h
thrust/system/cuda/detail/cross_system.h
@@ -56,8 +56,10 @@ namespace cuda_cub {
// Device to host.
template <class Sys1, class Sys2>
THRUST_CONSTEXPR __host__ __device__
- auto direction_of_copy(execution_policy<Sys1> const &,
- thrust::cpp::execution_policy<Sys2> const &)
+ auto direction_of_copy(
+ thrust::system::cuda::execution_policy<Sys1> const&
+ , thrust::cpp::execution_policy<Sys2> const&
+ )
THRUST_DECLTYPE_RETURNS(
thrust::detail::integral_constant<
cudaMemcpyKind, cudaMemcpyDeviceToHost
@@ -67,14 +69,29 @@ namespace cuda_cub {
// Host to device.
template <class Sys1, class Sys2>
THRUST_CONSTEXPR __host__ __device__
- auto direction_of_copy(thrust::cpp::execution_policy<Sys1> const &,
- execution_policy<Sys2> const &)
+ auto direction_of_copy(
+ thrust::cpp::execution_policy<Sys1> const&
+ , thrust::system::cuda::execution_policy<Sys2> const&
+ )
THRUST_DECLTYPE_RETURNS(
thrust::detail::integral_constant<
cudaMemcpyKind, cudaMemcpyHostToDevice
>{}
)
+ // Device to device.
+ template <class Sys1, class Sys2>
+ THRUST_CONSTEXPR __host__ __device__
+ auto direction_of_copy(
+ thrust::system::cuda::execution_policy<Sys1> const&
+ , thrust::system::cuda::execution_policy<Sys2> const&
+ )
+ THRUST_DECLTYPE_RETURNS(
+ thrust::detail::integral_constant<
+ cudaMemcpyKind, cudaMemcpyDeviceToDevice
+ >{}
+ )
+
// Device to device.
template <class DerivedPolicy>
THRUST_CONSTEXPR __host__ __device__
@@ -97,12 +114,46 @@ namespace cuda_cub {
)
)
+ template <typename ExecutionPolicy0, typename ExecutionPolicy1>
+ THRUST_CONSTEXPR __host__ __device__
+ auto is_device_to_host_copy(
+ ExecutionPolicy0 const& exec0
+ , ExecutionPolicy1 const& exec1
+ )
+ noexcept ->
+ thrust::detail::integral_constant<
+ bool
+ , cudaMemcpyDeviceToHost
+ == decltype(direction_of_copy(exec0, exec1))::value
+ >
+ {
+ return {};
+ }
+
template <typename ExecutionPolicy>
THRUST_CONSTEXPR __host__ __device__
auto is_device_to_host_copy(ExecutionPolicy const& exec)
- THRUST_NOEXCEPT ->
+ noexcept ->
+ thrust::detail::integral_constant<
+ bool
+ , cudaMemcpyDeviceToHost
+ == decltype(direction_of_copy(exec))::value
+ >
+ {
+ return {};
+ }
+
+ template <typename ExecutionPolicy0, typename ExecutionPolicy1>
+ THRUST_CONSTEXPR __host__ __device__
+ auto is_host_to_device_copy(
+ ExecutionPolicy0 const& exec0
+ , ExecutionPolicy1 const& exec1
+ )
+ noexcept ->
thrust::detail::integral_constant<
- bool, cudaMemcpyDeviceToHost == decltype(direction_of_copy(exec))::value
+ bool
+ , cudaMemcpyHostToDevice
+ == decltype(direction_of_copy(exec0, exec1))::value
>
{
return {};
@@ -111,9 +162,27 @@ namespace cuda_cub {
template <typename ExecutionPolicy>
THRUST_CONSTEXPR __host__ __device__
auto is_host_to_device_copy(ExecutionPolicy const& exec)
- THRUST_NOEXCEPT ->
+ noexcept ->
thrust::detail::integral_constant<
- bool, cudaMemcpyHostToDevice == decltype(direction_of_copy(exec))::value
+ bool
+ , cudaMemcpyHostToDevice
+ == decltype(direction_of_copy(exec))::value
+ >
+ {
+ return {};
+ }
+
+ template <typename ExecutionPolicy0, typename ExecutionPolicy1>
+ THRUST_CONSTEXPR __host__ __device__
+ auto is_device_to_device_copy(
+ ExecutionPolicy0 const& exec0
+ , ExecutionPolicy1 const& exec1
+ )
+ noexcept ->
+ thrust::detail::integral_constant<
+ bool
+ , cudaMemcpyDeviceToDevice
+ == decltype(direction_of_copy(exec0, exec1))::value
>
{
return {};
@@ -124,11 +193,61 @@ namespace cuda_cub {
auto is_device_to_device_copy(ExecutionPolicy const& exec)
noexcept ->
thrust::detail::integral_constant<
- bool, cudaMemcpyDeviceToDevice == decltype(direction_of_copy(exec))::value
+ bool
+ , cudaMemcpyDeviceToDevice
+ == decltype(direction_of_copy(exec))::value
>
{
return {};
}
+
+ // Device to host.
+ template <class Sys1, class Sys2>
+ __host__ __device__
+ auto
+ select_device_system(execution_policy<Sys1> & sys1,
+ thrust::cpp::execution_policy<Sys2> &)
+ THRUST_DECLTYPE_RETURNS(sys1)
+
+ // Device to host.
+ template <class Sys1, class Sys2>
+ __host__ __device__
+ auto
+ select_device_system(execution_policy<Sys1> const & sys1,
+ thrust::cpp::execution_policy<Sys2> const &)
+ THRUST_DECLTYPE_RETURNS(sys1)
+
+ // Host to device.
+ template <class Sys1, class Sys2>
+ __host__ __device__
+ auto
+ select_device_system(thrust::cpp::execution_policy<Sys1> &,
+ execution_policy<Sys2> & sys2)
+ THRUST_DECLTYPE_RETURNS(sys2)
+
+ // Host to device.
+ template <class Sys1, class Sys2>
+ __host__ __device__
+ auto
+ select_device_system(thrust::cpp::execution_policy<Sys1> const &,
+ execution_policy<Sys2> const & sys2)
+ THRUST_DECLTYPE_RETURNS(sys2)
+
+ // Device to device.
+ template <class Sys1, class Sys2>
+ __host__ __device__
+ auto
+ select_device_system(execution_policy<Sys1> &sys1,
+ execution_policy<Sys2> &)
+ THRUST_DECLTYPE_RETURNS(sys1)
+
+ // Device to device.
+ template <class Sys1, class Sys2>
+ __host__ __device__
+ auto
+ select_device_system(execution_policy<Sys1> const &sys1,
+ execution_policy<Sys2> const &)
+ THRUST_DECLTYPE_RETURNS(sys1)
#endif
// Device to host.
@@ -147,7 +266,7 @@ namespace cuda_cub {
template <class Sys1, class Sys2>
__host__ __device__
cross_system<Sys1, Sys2>
- select_system(const thrust::cpp::execution_policy<Sys1> &sys1,
+ select_system(thrust::cpp::execution_policy<Sys1> const &sys1,
execution_policy<Sys2> const & sys2)
{
thrust::cpp::execution_policy<Sys1> &non_const_sys1 = const_cast<thrust::cpp::execution_policy<Sys1> &>(sys1);

MODIFY
thrust/type_traits/remove_cvref.h
thrust/type_traits/remove_cvref.h
#pragma once
#include <thrust/detail/config.h>
-
-#if THRUST_CPP_DIALECT >= 2011
-# include <type_traits>
-#else
#include <thrust/detail/type_traits.h>
-#endif
THRUST_BEGIN_NS

`thrust::async::copy`: Add an overload that takes two policies to allow explicit
cross system copies.

Bug 2379510