cuda Thrust how to get the values associated to key

Question

I got a list of values/keys in one side and in the other side i got a list of keys.

I want to get the values associated at this list of keys in a list.

here my pseudocode

I don't know how to make my predicate and by the way if it's the good way to reach my goal.

void pseudoCode()
{
    const int N = 7;
    thrust::device_vector<vec3> keys1(N);

    values[0] = vec3(1.01,1.01,1.0156); 
    values[1] = vec3(1.01,1.01,1.01561);
    values[2] = vec3(1.02,1.52,1.02); 
    values[3] = vec3(1.02,1.52,1.02); 
    values[4] = vec3(1.0,1.0,1.0); 
    values[5] = vec3(5.0,1.0,1.0); 
    values[6] = vec3(5.0,1.5,1.0); 
  
    thrust::device_vector<long> values(N);
    keys1[0] = 0; 
    keys1[1] = 1;
    keys1[2] = 2;
    keys1[3] = 5;
    keys1[4] = 9;
    keys1[5] = 19;
    keys1[6] = 22;

    thrust::device_vector<long> to_find(6);
    to_find[0] = 1; 
    to_find[1] = 5;
    to_find[2] = 9;
    
    thrust::device_vector<vec3> output(6);

    auto begin = thrust::make_zip_iterator(thrust::make_tuple(keys1.begin(), values.begin()));
    auto end = thrust::make_zip_iterator(thrust::make_tuple(keys1.end(), values.end()));


    thrust::copy_if(begin, end, output, pred(to_find));

    // result 

    output[0]=values[1];
    output[1]=values[3];
    output[2]=values[4];
    

}

I try to make my predicate here is my code

struct FindValue
{
    thrust::device_vector<long>  ToFind;
    explicit FindValue(thrust::device_vector<long> toFind) : ToFind{toFind} {}

    __host__ __device__
    bool operator()(thrust::tuple<vec3, long> x)
    {
        long v = thrust::get<1>(x);
        size_t N = ToFind.size(); 
        for (size_t index=0; index<N; index++)
        {
            if (v==ToFind[index]) return true;
        }   
        return false;
    }
};

and change the call of copy_if

thrust::copy_if(begin, end, output, FindValue(to_find));

but i got the errors

/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/copy_if.inl(75): error: class "thrust::iterator_system<thrust::device_vector<vec3, thrust::device_allocator<vec3>>>" has no member "type"
          detected during instantiation of "OutputIterator thrust::copy_if(InputIterator, InputIterator, OutputIterator, Predicate) [with InputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<vec3>>, thrust::detail::normal_iterator<thrust::device_ptr<long>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIterator=thrust::device_vector<vec3, thrust::device_allocator<vec3>>, Predicate=FindValue]" 
cpp_addin_new2.cu(768): here

/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/copy_if.inl(80): error: no instance of overloaded function "select_system" matches the argument list
            argument types are: (System1, System2)
          detected during instantiation of "OutputIterator thrust::copy_if(InputIterator, InputIterator, OutputIterator, Predicate) [with InputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<vec3>>, thrust::detail::normal_iterator<thrust::device_ptr<long>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIterator=thrust::device_vector<vec3, thrust::device_allocator<vec3>>, Predicate=FindValue]" 
cpp_addin_new2.cu(768): here

/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/copy_if.inl(80): error: no instance of overloaded function "thrust::copy_if" matches the argument list
            argument types are: (<error-type>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<vec3>>, thrust::detail::normal_iterator<thrust::device_ptr<long>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<vec3>>, thrust::detail::normal_iterator<thrust::device_ptr<long>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::device_vector<vec3, thrust::device_allocator<vec3>>, FindValue)
          detected during instantiation of "OutputIterator thrust::copy_if(InputIterator, InputIterator, OutputIterator, Predicate) [with InputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<vec3>>, thrust::detail::normal_iterator<thrust::device_ptr<long>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIterator=thrust::device_vector<vec3, thrust::device_allocator<vec3>>, Predicate=FindValue]" 
cpp_addin_new2.cu(768): here

3 errors detected in the compilation of "cpp_addin_new2.cu".

Answer 1

The main problem in your current approach is that thrust::device_vector cannot be used in device code. The functor needs to store pointer and size instead. There are also issues with datatypes, for example, you have long values, and want to copy values, but the output vector contains vec3 i.e. keys. The current functor also searches to_find for values instead of keys.

Anyways, a correct functor which searches for keys could look like this.

template<class KeyType>
struct CopyIfPred{
    int numToFind;
    KeyType* to_find;

    __host__ __device__
    CopyIfPred(KeyType* to_find_, int numToFind_) : numToFind(numToFind_), to_find(to_find_){}

    template<class Tup>
    __host__ __device__
    bool operator()(const Tup& tup) const noexcept{
        const auto key = thrust::get<0>(tup);
        const bool found = thrust::binary_search(
            thrust::seq, 
            to_find, 
            to_find + numToFind, 
            key);
        return found;
    }
};

It can be used like this:

        auto begin = thrust::make_zip_iterator(thrust::make_tuple(
            keys1.begin(), 
            values.begin()
        ));
        auto end = begin + N;
        auto outputbegin = thrust::make_zip_iterator(thrust::make_tuple(
            thrust::make_discard_iterator(), // discard copied keys
            output.begin()
        ));

        CopyIfPred<KeyType> pred(
            thrust::raw_pointer_cast(to_find.data()), 
            to_find.size());


        auto endIters = thrust::copy_if(begin, end, outputbegin, pred);
        std::size_t outputsize = thrust::distance(outputbegin, endIters);

        for(std::size_t i = 0; i < outputsize; i++){
            std::cout << output[i] << "\n";
        }