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

Question

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

1> detected during: 1> instantiation of "thrust::iterator_adaptor::difference_type thrust::iterator_adaptor::distance_to(const thrust::iterator_adaptor &) const [with Derived=thrust::detail::normal_iterator>, Base=thrust::device_ptr, Value=thrust::use_default, System=thrust::use_default, Traversal=thrust::use_default, Reference=thrust::use_default, Difference=thrust::use_default, OtherDerived=thrust::device_ptr, OtherIterator=int *, V=signed int, S=thrust::device_system_tag, T=thrust::random_access_traversal_tag, R=thrust::device_reference, 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>, Facade2=thrust::device_ptr]" 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::type thrust::iterator_core_access::distance_from(const Facade1 &, const Facade2 &) [with Facade1=thrust::detail::normal_iterator>, Facade2=thrust::device_ptr]" 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>::type thrust::operator-(const thrust::iterator_facade &, const thrust::iterator_facade &) [with Derived1=thrust::detail::normal_iterator>, Value1=signed int, System1=thrust::device_system_tag, Traversal1=thrust::random_access_traversal_tag, Reference1=thrust::device_reference, Difference1=signed int, Derived2=thrust::device_ptr, Value2=signed int, System2=thrust::device_system_tag, Traversal2=thrust::random_access_traversal_tag, Reference2=thrust::device_reference, 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.

Answer 1

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);