如何使用推力将具有相同键的值从一个数组复制到另一个数组?
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 执行此操作的有效方法?
谢谢!
这是一种可能的方法:
使用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
标记每个段的开头,执行 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
然后可以将这些查找索引与 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);
我有一个稀疏直方图,其中键作为 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 执行此操作的有效方法?
谢谢!
这是一种可能的方法:
使用
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
标记每个段的开头,执行
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
然后可以将这些查找索引与
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);