cudathruststream-compaction

thrust remove copy unique by key


I'm a bit confused on the best way to do the following:

Say I have the following sorted key value pairs

(K:V) (0 : .5)(0 : .7)(0 : .9) (1 : .2) (1 : .6) (1 : .8)

and so on..

I want to remove copy the minimum value of each key so I'll have 2 results

minimum by key

(0 : .5)(1 : .2)

remaining

(0 : .7)(0 : .9)(1 : .6)(1 : .8)

thrust::unique_by_key_copy would appear to be able to give me the minimum by key since (K:V) is sorted. However I'm not sure how to remove those that are chosen from original to get the remaining.

Any thoughts or advice greatly appreciated


Solution

  • I'm sure there are a number of ways to do this.

    One possible approach is to use thrust::adjacent_difference to create a "flag array".

    Then use thrust::copy_if (the stencil version) using the flag array (as the stencil) to select first the elements you want in result 1, then the elements you want in result 2 (using the inverse of the flag array, logically).

    Here's a worked example:

    $ cat t671.cu
    #include <thrust/copy.h>
    #include <thrust/adjacent_difference.h>
    #include <thrust/device_vector.h>
    #include <thrust/host_vector.h>
    #include <thrust/iterator/zip_iterator.h>
    #include <iostream>
    
    
    struct is_zero
    {
      __host__ __device__
      bool operator()(const int s_val) const {
        return (!s_val);}
    };
    struct is_not_zero
    {
      __host__ __device__
      bool operator()(const int s_val) const {
        return (s_val);}
    };
    
    int main(){
    
      int keys[] = {0, 0, 0, 1, 1, 1 };
      float vals[] = {0.5f, 0.7f, 0.9f, 0.2f, 0.6f, 0.8f };
      int my_size = sizeof(keys)/sizeof(keys[0]);
      thrust::host_vector<int>     h_keys(keys, keys+my_size);
      thrust::host_vector<float>   h_vals(vals, vals+my_size);
      thrust::device_vector<int>   d_keys = h_keys;
      thrust::device_vector<float> d_vals = h_vals;
      thrust::device_vector<int>   flags(my_size);
      thrust::device_vector<int>   d_keys1(my_size);
      thrust::device_vector<float> d_vals1(my_size);
      thrust::device_vector<int>   d_keys2(my_size);
      thrust::device_vector<float> d_vals2(my_size);
      thrust::adjacent_difference(d_keys.begin(), d_keys.end(), flags.begin());
      flags[0] = 1;  // first element is always included
      int len1 = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(d_keys.begin(), d_vals.begin())), thrust::make_zip_iterator(thrust::make_tuple(d_keys.end(), d_vals.end())), flags.begin(), thrust::make_zip_iterator(thrust::make_tuple(d_keys1.begin(), d_vals1.begin())), is_zero()) - thrust::make_zip_iterator(thrust::make_tuple(d_keys1.begin(), d_vals1.begin()));
      int len2 = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(d_keys.begin(), d_vals.begin())), thrust::make_zip_iterator(thrust::make_tuple(d_keys.end(), d_vals.end())), flags.begin(), thrust::make_zip_iterator(thrust::make_tuple(d_keys2.begin(), d_vals2.begin())), is_not_zero()) - thrust::make_zip_iterator(thrust::make_tuple(d_keys2.begin(), d_vals2.begin()));
    
      thrust::host_vector<int>   h_keys1(len1);
      thrust::host_vector<float> h_vals1(len1);
      thrust::copy_n(d_keys1.begin(), len1, h_keys1.begin());
      thrust::copy_n(d_vals1.begin(), len1, h_vals1.begin());
      thrust::host_vector<int>   h_keys2(len2);
      thrust::host_vector<float> h_vals2(len2);
      thrust::copy_n(d_keys2.begin(), len2, h_keys2.begin());
      thrust::copy_n(d_vals2.begin(), len2, h_vals2.begin());
    
      std::cout << std::endl  << "Input keys:" << std::endl;
      thrust::copy(h_keys.begin(), h_keys.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl  << "Input values:" << std::endl;
      thrust::copy(h_vals.begin(), h_vals.end(), std::ostream_iterator<float>(std::cout, ","));
    
    
    
      std::cout << std::endl  << "Output keys 1:" << std::endl;
      thrust::copy(h_keys1.begin(), h_keys1.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl  << "Output values 1:" << std::endl;
      thrust::copy(h_vals1.begin(), h_vals1.end(), std::ostream_iterator<float>(std::cout, ","));
      std::cout << std::endl  << "Output keys 2:" << std::endl;
      thrust::copy(h_keys2.begin(), h_keys2.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl  << "Output values 2:" << std::endl;
      thrust::copy(h_vals2.begin(), h_vals2.end(), std::ostream_iterator<float>(std::cout, ","));
      std::cout << std::endl;
    
      return 0;
    }
    
    $ nvcc -arch=sm_20 -o t671 t671.cu
    $ ./t671
    
    Input keys:
    0,0,0,1,1,1,
    Input values:
    0.5,0.7,0.9,0.2,0.6,0.8,
    Output keys 1:
    0,0,1,1,
    Output values 1:
    0.7,0.9,0.6,0.8,
    Output keys 2:
    0,1,
    Output values 2:
    0.5,0.2,
    $