Как запустить thrust :: count_if на устройстве? (Куда)

Я хотел бы реализовать RANSAC. Я генерирую 60 тысяч очков и 500 самолетов, и я хотел бы посчитать для каждой плоскости, сколько точек находится рядом с ними. Затем выберите тот, который имеет максимальное значение.

После того как я сгенерировал векторы (d_vec) и самолеты (d_pl) и передал их в графический процессор, я использую thrust::transform и внутри этого thrust:count_if посчитать количество близких точек.

К сожалению, я получаю эту ошибку:

1>D:\Projects\cuda\CudaTest\CudaTest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\bin\nvcc.exe" -gencode=arch=compute_30,code=\"sm_30,compute_30\" --use-local-env --cl-version 2015 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\bin\x86_amd64" -x cu  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\include"     --keep-dir x64\Release -maxrregcount=0  --machine 64 --compile -cudart static     -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /O2 /FS /Zi  /MD " -o x64\Release\kernel.cu.obj "D:\Projects\cuda\CudaTest\CudaTest\kernel.cu"1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\include\thrust/detail/type_traits/pointer_traits.h(201): error : calling a __host__ function("thrust::detail::vector_base< ::Vec3,  ::thrust::device_malloc_allocator< ::Vec3> > ::begin") from a __device__ function("thrust::cuda_cub::__transform::unary_transform_f< ::thrust::detail::normal_iterator< ::thrust::device_ptr< ::Plane> > ,  ::thrust::detail::normal_iterator< ::thrust::device_ptr<int> > ,  ::thrust::cuda_cub::__transform::no_stencil_tag,  ::plane_functor,  ::thrust::cuda_cub::__transform::always_true_predicate> ::operator ()<long long> ") is not allowed
1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\include\thrust/detail/type_traits/pointer_traits.h(201): error : identifier "thrust::detail::vector_base< ::Vec3,  ::thrust::device_malloc_allocator< ::Vec3> > ::begin" is undefined in device code
1>D:/Projects/cuda/CudaTest/CudaTest/kernel.cu(84): error : calling a __host__ function("thrust::detail::vector_base< ::Vec3,  ::thrust::device_malloc_allocator< ::Vec3> > ::end") from a __device__ function("thrust::cuda_cub::__transform::unary_transform_f< ::thrust::detail::normal_iterator< ::thrust::device_ptr< ::Plane> > ,  ::thrust::detail::normal_iterator< ::thrust::device_ptr<int> > ,  ::thrust::cuda_cub::__transform::no_stencil_tag,  ::plane_functor,  ::thrust::cuda_cub::__transform::always_true_predicate> ::operator ()<long long> ") is not allowed
1>D:/Projects/cuda/CudaTest/CudaTest/kernel.cu(84): error : identifier "thrust::detail::vector_base< ::Vec3,  ::thrust::device_malloc_allocator< ::Vec3> > ::end" is undefined in device code

Как можно вызвать thrust :: count_if из кода устройства? Что я не прав?
Это полный код:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <thrust/execution_policy.h>
#include <algorithm>
#include <iostream>
#include <cstdlib>
#include <time.h>
#include <thrust/count.h>
#include <thrust/extrema.h>

struct Vec3 {
float x;
float y;
float z;friend std::ostream& operator<<(std::ostream& os, const Vec3& dt);
};

std::ostream& operator<<(std::ostream& os, const Vec3& dt)
{
os << dt.x << ", " << dt.y << ", " << dt.z;
return os;
}

struct Plane {

float a;
float b;
float c;
float d;

// https://keisan.casio.com/exec/system/1223596129
static Plane FromPoints(Vec3 A, Vec3 B, Vec3 C) {
Plane ret;

ret.a = (B.y - A.y)*(C.z - A.z) - (C.y - A.y)*(B.z - A.z);
ret.b = (B.z - A.z)*(C.x - A.x) - (C.z - A.z)*(B.x - A.x);
ret.c = (B.x - A.x)*(C.y - A.y) - (C.x - A.x)*(B.y - A.y);

ret.d = -(ret.a*A.x + ret.b*A.y + ret.c*A.z);

return ret;

}

};

Vec3 generator() {
return {
float(rand()) / float(RAND_MAX) * 1000.f,
float(rand()) / float(RAND_MAX) * 1000.f,
float(rand()) / float(RAND_MAX) * 1000.f
};
}

int index_generator() {
return rand() % 69632;
}

struct plane_distance {

const Plane pl;

__device__ plane_distance(const Plane pl) : pl(pl) {}

__device__ bool operator()(const Vec3& vv) const {
return fabsf(pl.a*vv.x + pl.b*vv.y + pl.c*vv.z + pl.d) / sqrtf(pl.a*pl.a + pl.b*pl.b + pl.c*pl.c) > 0.128f;
}

};

struct plane_functor
{
thrust::device_vector<Vec3>& d_vec;

plane_functor(thrust::device_vector<Vec3>& d_vec) : d_vec(d_vec) {}

__device__ int operator()(const Plane& pl) const {

return thrust::count_if(thrust::device, d_vec.begin(), d_vec.end(), plane_distance(pl));

}
};

int main(void)
{// Generate random points for testing

std::cout << "Generating..." << std::endl;

// generate random vectors serially
thrust::host_vector<Vec3> h_vec(65536);
std::generate(h_vec.begin(), h_vec.end(), generator);

// Generate random planes
thrust::host_vector<Plane> h_pl(512);
std::generate(h_pl.begin(), h_pl.end(), [&h_vec]() {

return Plane::FromPoints(
h_vec[index_generator()],
h_vec[index_generator()],
h_vec[index_generator()]
);

});

std::cout << "Transfer" << std::endl;

// transfer data to the device
thrust::device_vector<Vec3> d_vec = h_vec;
thrust::device_vector<Plane> d_pl = h_pl;
thrust::device_vector<int> counts(512);

std::cout << "Searching" << std::endl;

thrust::transform(thrust::device, d_pl.begin(), d_pl.end(), counts.begin(), plane_functor(d_vec));

auto result = thrust::max_element(thrust::device, counts.begin(), counts.end());

std::cout << "Press any key to exit" << std::endl;
std::cin.get();

return 0;
}