Fix compilation of async algorithms when RDC is enabled.
THRUST_RUNTIME_FUNCTION includes `__device__` when RDC is enabled. Since
the async algorithms use host-only futures and events, this causes builds
to fail. Making the async entry points regular `__host__` functions fixes
this.
This fixes
MODIFY
thrust/system/cuda/detail/async/copy.h
thrust/system/cuda/detail/async/copy.h
@@ -67,7 +67,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
-THRUST_RUNTIME_FUNCTION
auto async_copy_n(
FromPolicy& from_exec
, ToPolicy& to_exec
@@ -150,7 +149,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
-THRUST_RUNTIME_FUNCTION
auto async_copy_n(
thrust::cuda::execution_policy<FromPolicy>& from_exec
, thrust::cuda::execution_policy<ToPolicy>& to_exec
@@ -194,7 +192,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
-THRUST_RUNTIME_FUNCTION
auto async_copy_n(
FromPolicy& from_exec
, ToPolicy& to_exec
@@ -254,7 +251,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
-THRUST_RUNTIME_FUNCTION
auto async_copy_n(
FromPolicy& from_exec
, thrust::cuda::execution_policy<ToPolicy>& to_exec
@@ -358,7 +354,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
-THRUST_RUNTIME_FUNCTION
auto async_copy_n(
thrust::cuda::execution_policy<FromPolicy>& from_exec
, ToPolicy& to_exec
@@ -440,7 +435,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename OutputIt, typename Size
>
-THRUST_RUNTIME_FUNCTION
auto async_copy_n(
FromPolicy& from_exec
, ToPolicy& to_exec
@@ -486,7 +480,6 @@ template <
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
@@ -505,7 +498,6 @@ 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::cuda::execution_policy<ToPolicy>& to_exec
@@ -524,7 +516,6 @@ template <
typename FromPolicy, typename ToPolicy
, typename ForwardIt, typename Sentinel, typename OutputIt
>
-THRUST_RUNTIME_FUNCTION
auto async_copy(
thrust::cuda::execution_policy<FromPolicy>& from_exec
, thrust::cuda::execution_policy<ToPolicy>& to_exec

MODIFY
thrust/system/cuda/detail/async/for_each.h
thrust/system/cuda/detail/async/for_each.h
@@ -77,7 +77,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Size, typename UnaryFunction
>
-THRUST_RUNTIME_FUNCTION
auto async_for_each_n(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,
@@ -139,7 +138,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Sentinel, typename UnaryFunction
>
-THRUST_RUNTIME_FUNCTION
auto async_for_each(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,

MODIFY
thrust/system/cuda/detail/async/reduce.h
thrust/system/cuda/detail/async/reduce.h
@@ -60,7 +60,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Size, typename T, typename BinaryOp
>
-THRUST_RUNTIME_FUNCTION
auto async_reduce_n(
execution_policy<DerivedPolicy>& policy
, ForwardIt first
@@ -192,7 +191,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Sentinel, typename T, typename BinaryOp
>
-THRUST_RUNTIME_FUNCTION
auto async_reduce(
execution_policy<DerivedPolicy>& policy
, ForwardIt first
@@ -218,7 +216,6 @@ template <
, typename ForwardIt, typename Size, typename OutputIt
, typename T, typename BinaryOp
>
-THRUST_RUNTIME_FUNCTION
auto async_reduce_into_n(
execution_policy<DerivedPolicy>& policy
, ForwardIt first
@@ -330,7 +327,6 @@ template <
, typename ForwardIt, typename Sentinel, typename OutputIt
, typename T, typename BinaryOp
>
-THRUST_RUNTIME_FUNCTION
auto async_reduce_into(
execution_policy<DerivedPolicy>& policy
, ForwardIt first

MODIFY
thrust/system/cuda/detail/async/sort.h
thrust/system/cuda/detail/async/sort.h
@@ -65,7 +65,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Size, typename StrictWeakOrdering
>
-THRUST_RUNTIME_FUNCTION
auto async_stable_sort_n(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,
@@ -173,7 +172,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Size, typename StrictWeakOrdering
>
-THRUST_RUNTIME_FUNCTION
auto async_stable_sort_n(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,
@@ -289,7 +287,6 @@ auto async_stable_sort_n(
}
template <typename T, typename Size, typename StrictWeakOrdering>
-THRUST_RUNTIME_FUNCTION
typename std::enable_if<
is_operator_less_function_object<StrictWeakOrdering>::value
, cudaError_t
@@ -316,7 +313,6 @@ invoke_radix_sort(
}
template <typename T, typename Size, typename StrictWeakOrdering>
-THRUST_RUNTIME_FUNCTION
typename std::enable_if<
is_operator_greater_function_object<StrictWeakOrdering>::value
, cudaError_t
@@ -349,7 +345,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Size, typename StrictWeakOrdering
>
-THRUST_RUNTIME_FUNCTION
auto async_stable_sort_n(
execution_policy<DerivedPolicy>& policy
, ForwardIt first
@@ -504,7 +499,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Sentinel, typename StrictWeakOrdering
>
-THRUST_RUNTIME_FUNCTION
auto async_stable_sort(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,

MODIFY
thrust/system/cuda/detail/async/transform.h
thrust/system/cuda/detail/async/transform.h
@@ -78,7 +78,6 @@ template <
typename DerivedPolicy
, typename ForwardIt, typename Size, typename OutputIt, typename UnaryOperation
>
-THRUST_RUNTIME_FUNCTION
auto async_transform_n(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,
@@ -142,7 +141,6 @@ template <
, typename ForwardIt, typename Sentinel, typename OutputIt
, typename UnaryOperation
>
-THRUST_RUNTIME_FUNCTION
auto async_transform(
execution_policy<DerivedPolicy>& policy,
ForwardIt first,

Fix compilation of async algorithms when RDC is enabled.

THRUST_RUNTIME_FUNCTION includes `__device__` when RDC is enabled. Since
the async algorithms use host-only futures and events, this causes builds
to fail. Making the async entry points regular `__host__` functions fixes
this.

This fixes #1050.