thrust::counting_iterator complilation fail with relocatable device code

[hrvthzs]

Using CUDA Toolkit 10.0 with relocatable device code (-dc) the following code fails to compile and was working with CUDA Toolkit 8.0:

	using ci = thrust::counting_iterator<int>;
	thrust::reduce(thrust::device, ci(0), ci(1), 0);

cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch

[henryiii]

Dup of #924, the fix was going into the next-next version of CUDA, so that doesn't include CUDA 10.

PS: I've worked around it for the moment by always using the GitHub version of thrust.

[hrvthzs]

Sorry, yes it's a dup. I see it should be fixed by now. Where can I find the correspondig commit to this issue? Here on GitHub I have found only thrust 1.8.3.

[benjha]

Hi,

Any thoughts on this?

After testing thrust from the master branch and maintenance/cuda-10.1 and I am having a similar issue.

My code calls count_if in the next way and using CUDA 10:

   ```
   ...
    typedef thrust::device_vector<float>::iterator Iter;
    typedef thrust::device_vector<int>::iterator IterInt;
    thrust::device_vector<int>              d_clipFlag      ( size, -1  ); // num vertices
    strided_range<IterInt> d_clipX  ( d_clipFlag.begin(),   d_clipFlag.end(), 4);
    strided_range<IterInt> d_clipY  ( d_clipFlag.begin()+1, d_clipFlag.end(), 4);
    strided_range<IterInt> d_clipZ  ( d_clipFlag.begin()+2, d_clipFlag.end(), 4);
    strided_range<IterInt> d_clipW  ( d_clipFlag.begin()+3, d_clipFlag.end(), 4);

    compute_clipFlag();

    size_t numNotClipped = thrust::count_if(thrust::cuda::par,d_clipX.begin(), d_clipX.end(), 
    not_clipped<float>());

`../../../Programs/thrust/thrust/system/cuda/detail/cub/device/dispatch/dispatch_reduce.cuh(446): error: cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch
          detected during:
            instantiation of "cudaError_t thrust::cuda_cub::cub::DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, ReductionOpT, OutputT>::InvokeSingleTile<ActivePolicyT,SingleTileKernelT>(SingleTileKernelT) [with InputIteratorT=thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, OutputIteratorT=signed long *, OffsetT=int, ReductionOpT=thrust::plus<signed long>, OutputT=signed long, ActivePolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy130, SingleTileKernelT=void (*)(thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, signed long *, int, thrust::plus<signed long>, signed long)]" 
(599): here
            instantiation of "cudaError_t thrust::cuda_cub::cub::DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, ReductionOpT, OutputT>::Invoke<ActivePolicyT>() [with InputIteratorT=thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, OutputIteratorT=signed long *, OffsetT=int, ReductionOpT=thrust::plus<signed long>, OutputT=signed long, ActivePolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy130]" 
../../../Programs/thrust/thrust/system/cuda/detail/cub/block/../iterator/../util_device.cuh(332): here
            instantiation of "cudaError_t thrust::cuda_cub::cub::ChainedPolicy<PTX_VERSION, PolicyT, PolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=130, PolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy130, FunctorT=thrust::cuda_cub::cub::DispatchReduce<thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, signed long *, int, thrust::plus<signed long>, signed long>]" 
../../../Programs/thrust/thrust/system/cuda/detail/cub/block/../iterator/../util_device.cuh(315): here
            instantiation of "cudaError_t thrust::cuda_cub::cub::ChainedPolicy<PTX_VERSION, PolicyT, PrevPolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=200, PolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy200, PrevPolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy130, FunctorT=thrust::cuda_cub::cub::DispatchReduce<thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, signed long *, int, thrust::plus<signed long>, signed long>]" 
../../../Programs/thrust/thrust/system/cuda/detail/cub/block/../iterator/../util_device.cuh(315): here
            instantiation of "cudaError_t thrust::cuda_cub::cub::ChainedPolicy<PTX_VERSION, PolicyT, PrevPolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=300, PolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy300, PrevPolicyT=thrust::cuda_cub::cub::DeviceReducePolicy<signed long, int, thrust::plus<signed long>>::Policy200, FunctorT=thrust::cuda_cub::cub::DispatchReduce<thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, signed long *, int, thrust::plus<signed long>, signed long>]" 
../../../Programs/thrust/thrust/system/cuda/detail/cub/block/../iterator/../util_device.cuh(315): here
            [ 3 instantiation contexts not shown ]
            instantiation of "cudaError_t thrust::cuda_cub::cub::DeviceReduce::Reduce(void *, size_t &, InputIteratorT, OutputIteratorT, int, ReductionOpT, T, cudaStream_t, __nv_bool) [with InputIteratorT=thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, OutputIteratorT=signed long *, ReductionOpT=thrust::plus<signed long>, T=signed long]" 
../../../Programs/thrust/thrust/system/cuda/detail/reduce.h(948): here
            instantiation of "T thrust::cuda_cub::detail::reduce_n_impl(thrust::cuda_cub::execution_policy<Derived> &, InputIt, Size, T, BinaryOp) [with Derived=thrust::cuda_cub::par_t, InputIt=thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, Size=signed long, T=signed long, BinaryOp=thrust::plus<signed long>]" 
../../../Programs/thrust/thrust/system/cuda/detail/reduce.h(1018): here
            instantiation of "T thrust::cuda_cub::reduce_n(thrust::cuda_cub::execution_policy<Derived> &, InputIt, Size, T, BinaryOp) [with Derived=thrust::cuda_cub::par_t, InputIt=thrust::cuda_cub::transform_input_iterator_t<signed long, thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, not_clipped<float>>, Size=signed long, T=signed long, BinaryOp=thrust::plus<signed long>]" 
../../../Programs/thrust/thrust/system/cuda/detail/count.h(59): here
            instantiation of "thrust::iterator_traits<InputIt>::difference_type thrust::cuda_cub::count_if(thrust::cuda_cub::execution_policy<Derived> &, InputIt, InputIt, UnaryPred) [with Derived=thrust::cuda_cub::par_t, InputIt=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, UnaryPred=not_clipped<float>]" 
../../../Programs/thrust/thrust/detail/count.inl(51): here
            instantiation of "thrust::iterator_traits<InputIterator>::difference_type thrust::count_if(const thrust::detail::execution_policy_base<DerivedPolicy> &, InputIterator, InputIterator, Predicate) [with DerivedPolicy=thrust::cuda_cub::par_t, InputIterator=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::transform_iterator<strided_range<thrust::detail::normal_iterator<thrust::device_ptr<int>>>::stride_functor, thrust::counting_iterator<signed long, thrust::use_default, thrust::use_default, thrust::use_default>, thrust::use_default, thrust::use_default>>, Predicate=not_clipped<float>]" 
../source/cClipping.cu(218): here
`

[hrvthzs]

I have tried the latest GitHub version of thrust, but this issue is still not fixed.

[jrhemstad]

Looks like this was fixed as of CUDA 10.2 https://godbolt.org/z/ExWh9v15G