No instance of overloaded function "thrust::remove_if" matches the argument list

Question

I have written program with remove_if. It uses array allocated by cudaMalloc and filled by previous part of the program(in device). After removing the array will be used by the next part (in device; without thrust). I want to avoid any copying device-host, host-device. I use this example: https://github.com/thrust/thrust/blob/master/examples/cuda/wrap_pointer.cu

Nvcc wrote: ** "remove_if.cu(19): error: no instance of overloaded function "thrust::remove_if" matches the argument list argument types are: (thrust::device_ptr, thrust::device_ptr, is_zero)". **

I have written simple program-example with the same error:

#include <stdio.h>
#include "book.h"
#include <thrust/remove.h>
#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
int main(void)
{
  int *dev_lsubpix;
  struct is_zero
  {
    __host__ __device__
    bool operator()(int x)
    {
      return (x<1);
    }
  };
HANDLE_ERROR( cudaMalloc((void**)&dev_lsubpix, 10*sizeof(int)));
thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(dev_lsubpix);
  int new_length=     thrust::remove_if(dev_ptr, dev_ptr+10, is_zero())-dev_ptr;
  cudaFree(dev_lsubpix);
}

Answer 1

Although it isn't very obvious why from the error, the problem lies with the scope of the predicate functor you are trying to use. Because you have declared the functor at main scope, it is not a valid type outside of main and the compiler doesn't know how to deal with an anonymous type.

If you refactor your code like this:

#include <stdio.h>
#include <thrust/remove.h>
#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
struct is_zero
{
    __host__ __device__
        bool operator()(int x)
        {
            return (x<1);
        }
};

int main(void)
{
    int *dev_lsubpix;
    cudaMalloc((void**)&dev_lsubpix, 10*sizeof(int));
    thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(dev_lsubpix);
    int new_length = thrust::remove_if(dev_ptr, dev_ptr+10, is_zero())-dev_ptr;
    cudaFree(dev_lsubpix);
}

so that the functor definition is at global scope, I think you will find that the code compiles correctly.