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.
NVIDIA · May 25, 2020 · aaa6fb7 · aaa6fb7
1 parent 1b00acb
commit aaa6fb7
Show file tree

Hide file tree

Showing 5 changed files with 0 additions and 23 deletions.
diff --git a/thrust/system/cuda/detail/async/copy.h b/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

diff --git a/thrust/system/cuda/detail/async/for_each.h b/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,

diff --git a/thrust/system/cuda/detail/async/reduce.h b/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

diff --git a/thrust/system/cuda/detail/async/sort.h b/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,

diff --git a/thrust/system/cuda/detail/async/transform.h b/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,