Search code examples
cudathrust

How to use thrust::remove_if to check and remove blocks 2×i and 2×i+1 together


I am using Cuda C++, and I have a big array Arr including 64-bit unsigned integers in a form like the below:

Arr = {a1, b1, a2, b2, ..., an, bn}

The number of items in Arr is 2n which is an even number. Now, given a boolean function f(int a, int b), I wonder if I can use thrust::remove_if to check f(a1,b1), f(a2, b2), ..., f(an, bn) and remove both consecutive numbers (ai, bi) together if needed?


Solution

  • Rather than zip_iterator, I think a simpler approach is just to reinterpret the array of 64-bit integers as an array of thrust::pair. here is an example:

    $ cat t2157.cu
    #include <iostream>
    #include <thrust/device_vector.h>
    #include <thrust/remove.h>
    #include <thrust/sequence.h>
    #include <thrust/copy.h>
    
    struct my_remove
    {
      template <typename T>
      __host__ __device__
      bool operator()(T t){
        return (thrust::get<0>(t) > thrust::get<1>(t));
      }
    };
    
    const size_t n = 32768;
    const size_t s = 2*n;
    using mt=unsigned long long;
    using dt=thrust::pair<mt,mt>;
    
    int main(){
    
      thrust::device_vector<mt> A(s);
      thrust::sequence(A.begin(), A.end());
      A[0] = 2;  // expecting removal of the first pair
      thrust::copy_n(A.begin(), 6, std::ostream_iterator<mt>(std::cout, ","));
      std::cout << std::endl;
      auto D = thrust::device_pointer_cast<dt>(reinterpret_cast<dt *>(thrust::raw_pointer_cast(A.data())));
      thrust::remove_if(D, D+n, my_remove());
      thrust::copy_n(A.begin(), 6, std::ostream_iterator<mt>(std::cout, ","));
      std::cout << std::endl;
    }
    $ nvcc -o t2157 t2157.cu
    $ compute-sanitizer ./t2157
    ========= COMPUTE-SANITIZER
    2,1,2,3,4,5,
    2,3,4,5,6,7,
    ========= ERROR SUMMARY: 0 errors
    $