如何使用推力将具有相同键的值从一个数组复制到另一个数组?

How to use thrust to copy values with the same keys from one array to another?

我有一个稀疏直方图,其中键作为 bin,值作为计数。我想使用直方图中的键和计数来创建另一个键值数组。这个数组中的键是 bin,但重复了 count 次,最后每个键的值都等于相应的计数。直方图和数组中的键均按升序排序。

例如,如果直方图如下所示:

    Key/Bins:       0 2 4 5 6 7
    Values/Counts:  3 2 1 4 2 3 

我想查找其值的键数组如下所示:

 Key:     0 0 0 2 2 4 5 5 5 5 6 6 7 7 7

使用直方图填充值后新的键值数组 看起来像这样:

   Key:     0 0 0 2 2 4 5 5 5 5 6 6 7 7 7
   Values:  3 3 3 2 2 1 4 4 4 4 2 2 3 3 3 

我可以使用循环并检查两个键是否相同来执行此操作,但是是否有使用 Thrust 执行此操作的有效方法?

谢谢!

这是一种可能的方法:

  1. 使用thrust::transform使用key数组的两个偏移版本来标记每个段的开始

    Key:     0 0 0 2 2 4 5 5 5 5 6 6 7 7 7
    seg:     0 0 0 1 0 1 1 0 0 0 1 0 1 0 0
    
  2. 标记每个段的开头,执行 thrust::inclusive_scan 将段标记转换为一组查找索引

    Key:     0 0 0 2 2 4 5 5 5 5 6 6 7 7 7
    seg:     0 0 0 1 0 1 1 0 0 0 1 0 1 0 0
    idx:     0 0 0 1 1 2 3 3 3 3 4 4 5 5 5
    
  3. 然后可以将这些查找索引与 thrust::permutation_iterator 一起使用,将索引值从 values/counts 数组复制到所需的输出

    Key:     0 0 0 2 2 4 5 5 5 5 6 6 7 7 7
    seg:     0 0 0 1 0 1 1 0 0 0 1 0 1 0 0
    idx:     0 0 0 1 1 2 3 3 3 3 4 4 5 5 5
    Val:     3 3 3 2 2 1 4 4 4 4 2 2 3 3 3
    

这是一个有效的例子:

$ cat t1299.cu
#include <thrust/device_vector.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/transform.h>
#include <thrust/remove.h>
#include <thrust/copy.h>
#include <iostream>

struct my_remove
{
  template <typename T>
  __host__ __device__
  bool operator()(const T &t)
  {
    return (thrust::get<0>(t) == 0);
  }
};

struct my_cmp
{
  template <typename T>
  __host__ __device__
  int operator()(const T &t)
  {
    if (thrust::get<0>(t) != thrust::get<1>(t)) return 1;
    return 0;
  }
};

using namespace thrust::placeholders;

int main(){

  int kb[] = {0,2,4,5,6,7};
  int vc[] = {3,2,1,4,2,3};
  int key[] = {0,0,0,2,2,4,5,5,5,5,6,6,7,7,7};
  int kbsize = sizeof(kb)/sizeof(int);
  int keysize = sizeof(key)/sizeof(int);

  thrust::device_vector<int> dkb(kb, kb+kbsize);
  thrust::device_vector<int> dvc(vc, vc+kbsize);
  thrust::device_vector<int> dkey(key, key+keysize);
  thrust::device_vector<int> dval(keysize);
  thrust::copy(dkey.begin(), dkey.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;

  //first remove any zero count values from kb/vc

//  thrust::remove_if(thrust::make_zip_iterator(thrust::make_tuple(dvc.begin(), dkb.begin())), thrust::make_zip_iterator(thrust::make_tuple(dvc.end(), dkb.end())), my_remove());

  // next, mark the segments in the key array

  thrust::transform(thrust::make_zip_iterator(thrust::make_tuple(dkey.begin(), dkey.begin()+1)), thrust::make_zip_iterator(thrust::make_tuple(dkey.end()-1, dkey.end())), dval.begin()+1, my_cmp());

  // prefix sum to generate lookup indices
  thrust::device_vector<int> didx(keysize);
  thrust::inclusive_scan(dval.begin(), dval.end(), didx.begin());

  // use lookup indices to populate values vector

  thrust::copy(thrust::make_permutation_iterator(dvc.begin(), didx.begin()), thrust::make_permutation_iterator(dvc.begin(), didx.end()), dval.begin());

  // display results

  thrust::copy(dval.begin(), dval.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -arch=sm_35 -o t1299 t1299.cu
$ ./t1299
0,0,0,2,2,4,5,5,5,5,6,6,7,7,7,
3,3,3,2,2,1,4,4,4,4,2,2,3,3,3,
$

推力优化的典型示例是寻找 "fusion" 操作。在这种情况下,由于我们有一个转换操作,紧接着是一个包容性扫描,一个简单的融合示例是将它们替换为 thrust::transform_inclusive_scan

您似乎正在 GPU 中进行 table 查找。这是一个在带有 CUDA 的 GPU 上使用布谷鸟哈希的实现。 https://github.com/shixing/cuckoo/tree/master

构建过程在CPU上实现:

create_hash_cpu(h_keys,h_values,h_key_1,h_key_2,h_value_1,h_value_2, N, EMPTY_KEY);

查找过程在 GPU 上:

cuckoo_lookup<<<(M+255)/256, 256>>>(d_keys_lookup, d_values_lookup, d_key_1, d_value_1, d_key_2, d_value_2, M, N, EMPTY_KEY, EMPTY_VALUE);