CUDA Thrust - How can I write a function using multiple device vectors with different sizes?

I've been trying figure out how to perform a simple entropy calculation using four thrust device vectors.

I have four device vectors, representing two key-value pairs. The first pair of vectors contains the keys and the number of times that key appears. The second pair contains keys paired with the bins for calculating the entropy. In this second vector pair, keys appear multiple times, with each instance representing a different bin.

It looks something like this:

Device Vector Pair 1

KeyVal 6 8 9

Counts 1 3 2

Device Vector Pair 2

KeyVal 6 8 8 9 9

BinVal 1 1 2 1 1

Result Vector (contains the calculated entropy results)

KeyVal 8

Entropy 0.602

What I'm planning on doing is to use the first vector pair to check if a key appears enough times to calculate the entropy. If the count is large enough, the second vector pair will be used to calculate the entropy with the bin values for that key. I will need to use all of the bin values for that particular key. For example, if I wanted to calculate the entropy for keys that have appeared at least 3 times, I would find in the first vector pair that KeyVal 8 is ready. Then, I would search the second pair for all instances of KeyVal 8, and calculate the entropy using their corresponding BinVals. The entropy calculation is simple, it just involves adding together BinVal*Log(BinVal) for each relevant value. In my example, it would be entropy = 1*log(1) + 2*log(2).

However, I have no clue how to make this part work. I've tried using thrust::for_each to find all the keys that appear enough times to be tested, but I don't think it's possible to search for the keys in the second vector pair and perform the calculation within the for_each function.

Does anyone have suggestions on other ways to accomplish this?

Thank you for your help.

Solution

The two ideas I considered were:

Idea A:

Compute all the entropies
Select the ones that meet the criteria

Idea B:

Select the incoming data that meets the criteria
Compute the entropies.

Idea A seems to be doing unnecessary work - computing entropies that are or may be not needed. However as I worked through the process for Idea B, I ended up adding so many steps (such as computing prefix sums) to complete step 1 of Idea B, that it did not seem like it would be better. So I will present Idea A for now. Maybe m.s. or someone else will come along and post something better.

Step 1 of Idea A is handled by thrust::reduce_by_key along with an appropriate functor to compute the specific entropy function

Step 2 of Idea A is handled by thrust::copy_if

$ cat t827.cu
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/reduce.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <math.h>

// THRESH determines the minimum Counts value required for a KeyVal Entropy calculation to occur
#define THRESH 2

using namespace thrust::placeholders;


struct my_entropy : public thrust::unary_function<float, float>
{
  __host__ __device__
  float operator()(float val){
    return val*log10f(val);}  // if you want napierian log, change this to logf
};

int main(){

  int KeyVal1[]={6, 8, 9};
  int Counts[] ={1, 3, 2};
  int KeyVal2[]={6, 8, 8, 9, 9};
  float BinVal[] ={1, 1, 2, 1, 1};

  int dsize1 = sizeof(KeyVal1)/sizeof(int);
  int dsize2 = sizeof(KeyVal2)/sizeof(int);

  thrust::device_vector<int> d_KeyVal1(KeyVal1, KeyVal1+dsize1);
  thrust::device_vector<int> d_Counts(Counts, Counts+dsize1);
  thrust::device_vector<int> d_KeyVal2(KeyVal2, KeyVal2+dsize2);
  thrust::device_vector<float> d_BinVal(BinVal, BinVal+dsize2);


  // method 1 - just compute all entropies, then select the desired ones
  thrust::device_vector<float> entropies(dsize2);
  thrust::reduce_by_key(d_KeyVal2.begin(), d_KeyVal2.end(), thrust::make_transform_iterator(d_BinVal.begin(), my_entropy()), thrust::make_discard_iterator(), entropies.begin());
  thrust::device_vector<int> res_keys(dsize1);
  thrust::device_vector<float>res_ent(dsize1);
  int res_size = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(d_KeyVal1.begin(), entropies.begin())), thrust::make_zip_iterator(thrust::make_tuple(d_KeyVal1.end(), entropies.end())), d_Counts.begin(), thrust::make_zip_iterator(thrust::make_tuple(res_keys.begin(), res_ent.begin())), _1 >= THRESH) - thrust::make_zip_iterator(thrust::make_tuple(res_keys.begin(), res_ent.begin()));
  std::cout << "Counts threshold: " << THRESH << std::endl <<  "selected keys: " << std::endl;
  thrust::copy_n(res_keys.begin(), res_size, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl << "calculated entropies: " << std::endl;
  thrust::copy_n(res_ent.begin(), res_size, std::ostream_iterator<float>(std::cout, ","));
  std::cout << std::endl;

  return 0;
}
[bob@cluster1 misc]$ nvcc -o t827 t827.cu
$ ./t827
Counts threshold: 2
selected keys:
8,9,
calculated entropies:
0.60206,0,
$