Search code examples
cudathrust

using std::bind2nd with thrust

I'm trying to use std::bind2nd with thrust. I have code that compiles with a host pointer, but not with a device pointer. I think they are identical and should work in both cases.

// this compiles fine
thrust::host_vector<unsigned int> h_vec(nwords);
thrust::host_vector<unsigned int> h_map(nwords);
thrust::host_vector<unsigned int> h_out1(nwords);
thrust::copy_if(h_vec.begin(), h_vec.end(), h_map.begin(), h_out1.begin(), 
                std::bind2nd(thrust::equal_to<int>(),1));

// this compilation fails with the error below
thrust::device_vector<unsigned int> d_map(nwords);
thrust::device_vector<unsigned int> d_vec(nwords);
thrust::device_vector<unsigned int> d_out1(nwords);
thrust::copy_if(d_vec.begin(), d_vec.end(), d_map.begin(), d_out1.begin(), 
                std::bind2nd(thrust::equal_to<int>(),1));

When I try to call the second copy_if with the bind2nd I get the error below:

 /opt/cuda/include/thrust/detail/internal_functional.h(99): warning: calling a
 __host__ function from a __host__ __device__ function is not allowed

Is there another way to using adapters for binary functions in thrust? I have seen some people using "thrust::bind2nd" in examples on the web but I can't find that in any of our header files.

Solution

  • It is possible to use adaptors in thrust. But if you want to use them on the GPU, the adaptor must be a __device__ function. This is why the first copy_if compiles and the second doesn't - your predicate is a host function, not a device function and the use of device_vector implies a device compilation trajectory.

    In short, if you want an adaptor function for use on the GPU, you will need to write one yourself, the standard library ones (bind1st, bind2nd, bind) can't be used.