Search code examples
cudathrust

Scan and generate a map (for use with thrust::scatter)


I'd like to use thrust to create an array which contains the all the indices that pass a test. I think I'm missing something here, as I can't seem to work out how. Is there a simple way to do this using thrust?

struct is_odd
{
  __device__
  bool operator()(int &x)
  {
    return (x % 2) == 1;
  }
};

int A[] = {1, 2, 1, 1, 4, 1};
int result[] = {-1, -1, -1, -1, -1};

thrust::find_map_if(A, A+5, result, is_odd()); // this function doesn't exist!

// result = {0, 2, 3, 5, -1} 

I need this map to scatter an arbitrary array of data (which is not A).


Solution

  • There are probably a lot of ways to tackle this. Here's one possible approach:

    1. Create an index array (or use a counting iterator)

      int  A[] = {1, 2, 1, 1, 4, 1};
      int iA[] = {0, 1, 2, 3, 4, 5};  
      

      you can use thrust::sequence to do this, for example. Or you can skip the explicit generation of iA and use a counting_iterator in the next step.

    2. Use thrust::remove_copy_if to take the index array and reduce it to the elements that correspond to the result of your test.

    Here's a fully worked example. Note that remove_copy_if copies elements for which the functor test is false:

    $ cat t596.cu
    #include <iostream>
    #include <thrust/device_vector.h>
    #include <thrust/iterator/counting_iterator.h>
    #include <thrust/remove.h>
    #include <thrust/copy.h>
    
    #define DSIZE 6
    
    struct is_odd
    {
      __device__
      bool operator()(int x) const
      {
        return !(x&1);
      }
    };
    
    int main(){
    
      int A[] = {1, 2, 1, 1, 4, 1};
      thrust::device_vector<int> d_A(A, A+DSIZE);
      thrust::device_vector<int> result(DSIZE, -1);
      thrust::counting_iterator<int> first(0);
      thrust::counting_iterator<int> last = first+DSIZE;
      thrust::device_vector<int>::iterator r_end = thrust::remove_copy_if(first, last, d_A.begin(), result.begin(), is_odd());
      thrust::copy(result.begin(), r_end, std::ostream_iterator<int>(std::cout, " "));
      std::cout << std::endl;
      thrust::copy(result.begin(), result.end(), std::ostream_iterator<int>(std::cout, " "));
      std::cout << std::endl;
    }
    
    $ nvcc -arch=sm_20 -o t596 t596.cu
    $ ./t596
    0 2 3 5
    0 2 3 5 -1 -1
    $