Search code examples
cudathrust

Removing items during a large array transform in CUDA


Given a large array of values A that are transformed to an array B, so B = Transform(A). Where A and B are of different types and the transformation Transform() is reasonably expensive and the data size of B is larger than A. But also the results are to be filtered out based on a predicate Keep(B).

Is there a decent way to do this without writing an array of B out first then pruning the B entries to be kept?

I started in thrust to try:

typedef int A;
struct B { int a, b, c; };


struct FTransform : thrust::unary_function<A, B>
{
    __device__ B operator()(A a) const { return B{ a, a, a }; }
};

struct FKeep : thrust::unary_function<B, bool>
{
    __device__ bool operator()(B b) const { return (b.a & 1) == 0; }
};


thrust::device_vector<B> outputs(8);
thrust::device_vector<A> inputs(8);

std::generate(inputs.begin(), inputs.end(), rand);

auto first = thrust::make_transform_iterator(inputs.begin(), FTransform());
auto last = thrust::make_transform_iterator(inputs.end(), FTransform());

auto end = thrust::copy_if(first, last, outputs, FKeep());

However this gives compile errors (Cuda 9.2):

thrust/iterator/iterator_traits.h(49): error : class "thrust::device_vector<B, thrust::device_malloc_allocator<B>>" has no member "iterator_category"

thrust/detail/copy_if.inl(78): error : incomplete type is not allowed

thrust/detail/copy_if.inl(80): error : no instance of overloaded function "select_system" matches the argument list

thrust/detail/copy_if.inl(80): error : no instance of overloaded function "thrust::copy_if" matches the argument list

Solution

  • Here:

    auto end = thrust::copy_if(first, last, outputs, FKeep());
                                            ^^^^^^^
    

    outputs is not an iterator. You should be passing outputs.begin() there.

    With that change, your code compiles for me.