Type of return value of thrust::remove_if

2019-02-15 06:06发布

问题:

I have two arrays of integers dmap and dflag on the device of the same length and I have wrapped them with thrust device pointers, dmapt and dflagt

There are some elements in the dmap array with value -1. I want to remove these -1's and the corresponding values from the dflag array.

I am using the remove_if function to do this, but I cannot figure out what the return value of this call is or how I should use this returned value to get .

( I want to pass these reduced arrays to the reduce_by_key function where dflagt will be used as the keys. )

I am using the following call for doing the reduction. Please let me know how I can store the returned value in a variable and use it to address the individual arrays dflag and dmap

thrust::remove_if( 
    thrust::make_zip_iterator(thrust::make_tuple(dmapt, dflagt)), 
    thrust::make_zip_iterator(thrust::make_tuple(dmapt+numindices, dflagt+numindices)), 
    minus_one_equality_test() 
); 

where the predicate functor used above is defined as

struct minus_one_equality_test
{ 
    typedef typename thrust::tuple<int,int> Tuple; 
    __host__ __device__ 
    bool operator()(const Tuple& a ) 
    { 
        return  thrust::get<0>(a) ==  (-1); 
    } 
} 

回答1:

The return value is a zip_iterator which marks the new end of the sequence of tuples for which your functor returned true during the remove_if call. To access the new end iterator of the underlying array you will need to retrieve a tuple iterator from the zip_iterator; the contents of that tuple are then the new end iterators of the original arrays you used to build the zip_iterator. It is a lot more convoluted in words than in code:

#include <thrust/tuple.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/remove.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>

#include <iostream>

struct minus_one_equality_test
{ 
    typedef thrust::tuple<int,int> Tuple; 
    __host__ __device__ 
    bool operator()(const Tuple& a ) 
    { 
        return  thrust::get<0>(a) ==  (-1); 
    }; 
}; 


int main(void)
{
    const int numindices = 10;

    int mapt[numindices] = { 1, 2, -1, 4, 5, -1, 7, 8, -1, 10 };
    int flagt[numindices] = { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10 };

    thrust::device_vector<int> vmapt(10);
    thrust::device_vector<int> vflagt(10);

    thrust::copy(mapt, mapt+numindices, vmapt.begin());
    thrust::copy(flagt, flagt+numindices, vflagt.begin());

    thrust::device_ptr<int> dmapt = vmapt.data();
    thrust::device_ptr<int> dflagt = vflagt.data();

    typedef thrust::device_vector< int >::iterator  VIt;
    typedef thrust::tuple< VIt, VIt > TupleIt;
    typedef thrust::zip_iterator< TupleIt >  ZipIt;

    ZipIt Zend = thrust::remove_if(  
        thrust::make_zip_iterator(thrust::make_tuple(dmapt, dflagt)), 
        thrust::make_zip_iterator(thrust::make_tuple(dmapt+numindices, dflagt+numindices)), 
        minus_one_equality_test() 
    ); 

    TupleIt Tend = Zend.get_iterator_tuple();
    VIt vmapt_end = thrust::get<0>(Tend);

    for(VIt x = vmapt.begin(); x != vmapt_end; x++) {
        std::cout << *x << std::endl;
    }

    return 0;
}

If you compile this and run it, you should see something like this:

$ nvcc -arch=sm_12 remove_if.cu 
$ ./a.out
1
2
4
5
7
8
10

In this example I only "retrieve" the shorted contents of the first element of the tuple, the second is accessed in the same way, ie. the iterator marking the new end of the vector is thrust::get<1>(Tend).



标签: cuda thrust