Warnings and errors when enabling RDC

[alliepiper]

There are a quite a few of these when compiling with RDC enabled.

The warning at reduce.h:275 is caused by async_reduce_into_n being marked up with THRUST_RUNTIME_FUNCTION (includes __host__ __device__ when RDC is enabled), which calls the __host__ only function make_dependent_event. The others are probably similar.

[alliepiper]

The build will eventually error out as well:

/home/allie/code/src/thrust/cub/device/dispatch/dispatch_reduce.cuh(452): error: cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch
          detected during:
            instantiation of "cudaError_t cub::DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, ReductionOpT, OutputT>::InvokeSingleTile<ActivePolicyT,SingleTileKernelT>(SingleTileKernelT) [with InputIteratorT=thrust::zip_iterator<thrust::tuple<thrust::cuda_cub::transform_input_iterator_t<__nv_bool, thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, equal_to_value_pred<custom_numeric>>, thrust::cuda_cub::counting_iterator_t<signed long>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIteratorT=thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> *, OffsetT=int, ReductionOpT=thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputT=thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, ActivePolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy130, SingleTileKernelT=void (*)(thrust::zip_iterator<thrust::tuple<thrust::cuda_cub::transform_input_iterator_t<__nv_bool, thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, equal_to_value_pred<custom_numeric>>, thrust::cuda_cub::counting_iterator_t<signed long>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> *, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>)]" 
(605): here
            instantiation of "cudaError_t cub::DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, ReductionOpT, OutputT>::Invoke<ActivePolicyT>() [with InputIteratorT=thrust::zip_iterator<thrust::tuple<thrust::cuda_cub::transform_input_iterator_t<__nv_bool, thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, equal_to_value_pred<custom_numeric>>, thrust::cuda_cub::counting_iterator_t<signed long>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIteratorT=thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> *, OffsetT=int, ReductionOpT=thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputT=thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, ActivePolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy130]" 
/home/allie/code/src/thrust/cub/block/../iterator/../util_device.cuh(500): here
            instantiation of "cudaError_t cub::ChainedPolicy<PTX_VERSION, PolicyT, PolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=130, PolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy130, FunctorT=cub::DispatchReduce<thrust::zip_iterator<thrust::tuple<thrust::cuda_cub::transform_input_iterator_t<__nv_bool, thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, equal_to_value_pred<custom_numeric>>, thrust::cuda_cub::counting_iterator_t<signed long>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> *, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>]" 
/home/allie/code/src/thrust/cub/block/../iterator/../util_device.cuh(483): here
            instantiation of "cudaError_t cub::ChainedPolicy<PTX_VERSION, PolicyT, PrevPolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=200, PolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy200, PrevPolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy130, FunctorT=cub::DispatchReduce<thrust::zip_iterator<thrust::tuple<thrust::cuda_cub::transform_input_iterator_t<__nv_bool, thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, equal_to_value_pred<custom_numeric>>, thrust::cuda_cub::counting_iterator_t<signed long>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> *, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>]" 
/home/allie/code/src/thrust/cub/block/../iterator/../util_device.cuh(483): here
            instantiation of "cudaError_t cub::ChainedPolicy<PTX_VERSION, PolicyT, PrevPolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=300, PolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy300, PrevPolicyT=cub::DeviceReducePolicy<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>>::Policy200, FunctorT=cub::DispatchReduce<thrust::zip_iterator<thrust::tuple<thrust::cuda_cub::transform_input_iterator_t<__nv_bool, thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, equal_to_value_pred<custom_numeric>>, thrust::cuda_cub::counting_iterator_t<signed long>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> *, int, thrust::cuda_cub::__find_if::functor<thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::tuple<__nv_bool, signed long, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>]" 
/home/allie/code/src/thrust/cub/block/../iterator/../util_device.cuh(483): here
            [ 7 instantiation contexts not shown ]
            instantiation of "InputIt thrust::cuda_cub::find_if_n(thrust::cuda_cub::execution_policy<Derived> &, InputIt, Size, Predicate) [with Derived=thrust::cuda_cub::tag, InputIt=thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, Size=signed long, Predicate=equal_to_value_pred<custom_numeric>]" 
/home/allie/code/src/thrust/thrust/system/cuda/detail/find.h(181): here
            instantiation of "InputIt thrust::cuda_cub::find_if(thrust::cuda_cub::execution_policy<Derived> &, InputIt, InputIt, Predicate) [with Derived=thrust::cuda_cub::tag, InputIt=thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, Predicate=equal_to_value_pred<custom_numeric>]" 
/home/allie/code/src/thrust/thrust/detail/find.inl(54): here
            instantiation of "InputIterator thrust::find_if(const thrust::detail::execution_policy_base<DerivedPolicy> &, InputIterator, InputIterator, Predicate) [with DerivedPolicy=thrust::cuda_cub::tag, InputIterator=thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, Predicate=equal_to_value_pred<custom_numeric>]" 
/home/allie/code/src/thrust/thrust/detail/find.inl(96): here
            instantiation of "InputIterator thrust::find_if(InputIterator, InputIterator, Predicate) [with InputIterator=thrust::detail::normal_iterator<thrust::device_ptr<custom_numeric>>, Predicate=equal_to_value_pred<custom_numeric>]" 
/home/allie/code/src/thrust/testing/find.cu(112): here
            instantiation of "void TestFindIfSimple<Vector>() [with Vector=thrust::device_vector<custom_numeric, thrust::device_allocator<custom_numeric>>]" 
/home/allie/code/src/thrust/testing/find.cu(119): here