[SOLVED] Thrust: how to get the number of elements copied by the copy_if function when using device

Thrust: how to get the number of elements copied by the copy_if function when using device_ptr

I am using the thrust::copy_if function of the Thrust library coupled with counting iterators to get the indices of nonzero elements in an array. I also need to get the number of copied elements.

I am using the code from the 'counting_iterator.cu' example, except that in my application I need to reuse pre-allocated arrays, so I wrap them with thrust::device_ptr and then pass them to the thrust::copy_if function. This is the code:

using namespace thrust;

int output[5];
thrust::device_ptr<int> tp_output = device_pointer_cast(output);

float stencil[5];
stencil[0] = 0;
stencil[1] = 0;
stencil[2] = 1;
stencil[3] = 0;
stencil[4] = 1;
device_ptr<float> tp_stencil = device_pointer_cast(stencil);

device_vector<int>::iterator output_end = copy_if(make_counting_iterator<int>(0), 
     make_counting_iterator<int>(5), 
     tp_stencil, 
     tp_output, 
     _1 == 1);

int number_of_ones = output_end - tp_output;

If I comment the last line of code, the function fills correctly the output array. However, when I uncomment it, I get the following compilation error:

1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include\thrust/iterator/iterator_adaptor.h(223): error : no operator "-" matches these operands

1>              operand types are: int *const - const thrust::device_ptr<int>

1>            detected during:
1>              instantiation of "thrust::iterator_adaptor<Derived, Base, Value, System, Traversal, Reference, Difference>::difference_type thrust::iterator_adaptor<Derived, Base, Value, System, Traversal, Reference, Difference>::distance_to(const thrust::iterator_adaptor<OtherDerived, OtherIterator, V, S, T, R, D> &) const [with Derived=thrust::detail::normal_iterator<thrust::device_ptr<int>>, Base=thrust::device_ptr<int>, Value=thrust::use_default, System=thrust::use_default, Traversal=thrust::use_default, Reference=thrust::use_default, Difference=thrust::use_default, OtherDerived=thrust::device_ptr<int>, OtherIterator=int *, V=signed int, S=thrust::device_system_tag, T=thrust::random_access_traversal_tag, R=thrust::device_reference<int>, D=ptrdiff_t]" 
1>  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include\thrust/iterator/iterator_facade.h(181): here
1>              instantiation of "Facade1::difference_type thrust::iterator_core_access::distance_from(const Facade1 &, const Facade2 &, thrust::detail::true_type) [with Facade1=thrust::detail::normal_iterator<thrust::device_ptr<int>>, Facade2=thrust::device_ptr<int>]" 
1>  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include\thrust/iterator/iterator_facade.h(202): here
1>              instantiation of "thrust::detail::distance_from_result<Facade1, Facade2>::type thrust::iterator_core_access::distance_from(const Facade1 &, const Facade2 &) [with Facade1=thrust::detail::normal_iterator<thrust::device_ptr<int>>, Facade2=thrust::device_ptr<int>]" 
1>  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include\thrust/iterator/iterator_facade.h(506): here
1>              instantiation of "thrust::detail::distance_from_result<thrust::iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1>, thrust::iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2>>::type thrust::operator-(const thrust::iterator_facade<Derived1, Value1, System1, Traversal1, Reference1, Difference1> &, const thrust::iterator_facade<Derived2, Value2, System2, Traversal2, Reference2, Difference2> &) [with Derived1=thrust::detail::normal_iterator<thrust::device_ptr<int>>, Value1=signed int, System1=thrust::device_system_tag, Traversal1=thrust::random_access_traversal_tag, Reference1=thrust::device_reference<int>, Difference1=signed int, Derived2=thrust::device_ptr<int>, Value2=signed int, System2=thrust::device_system_tag, Traversal2=thrust::random_access_traversal_tag, Reference2=thrust::device_reference<int>, Difference2=signed int]" 
1>  C:/ProgramData/NVIDIA Corporation/CUDA Samples/v5.5/7_CUDALibraries/nsgaIIparallelo_23ott/rank_cuda.cu(70): here

If I use thrust::device_vector for the output array instead, everything is okay:

using namespace thrust;

thrust::device_vector<int> output(5);

float stencil[5];
stencil[0] = 0;
stencil[1] = 0;
stencil[2] = 1;
stencil[3] = 0;
stencil[4] = 1;
device_ptr<float> tp_stencil = device_pointer_cast(stencil);

device_vector<int>::iterator output_end = copy_if(make_counting_iterator<int>(0), 
     make_counting_iterator<int>(5), 
     tp_stencil, 
     output.begin(), 
     _1 == 1);

int number_of_ones = output_end - output.begin();

Can you suggest any solution to this problem? Thank you.

Solution

Try using device_ptr instead of device_vector::iterator in your copy_if call:

thrust::device_ptr<int> output_end = copy_if(make_counting_iterator<int>(0),
 make_counting_iterator<int>(5),
 tp_stencil,
 tp_output,
 _1 == 1);