使用CUDA/Thrust使关键出现次数相等

Question

有没有一种有效的方法来获取排序后的键/值数组对并确保每个键使用CUDA Thrust库都具有相同数量的元素？

例如，假设我们有下面的一对阵列：

ID: 1 2 2 3 3 3 
VN: 6 7 8 5 7 8 

如果我们想有两个每个键的出现，这将是结果：

ID: 2 2 3 3 
VN: 7 8 5 7 

实际阵列将会更大，包含数百万或更多的元素。我可以很容易地使用嵌套for循环来做到这一点，但我很想知道是否有更高效的方式来使用GPU转换数组。 Thrust似乎可能很有用，但我没有看到任何明显的功能。

谢谢你的帮助！

Answer 1

警告：如果这是您计划在GPU上执行的唯一操作，那么我不会推荐它。从GPU复制数据的成本可能会超过使用GPU的任何可能的效率/性能优势。

编辑：基于序列阈值可能会远远长于2的评论，我会建议一种替代方法（方法2），该方法应该比for-loop或brute-force更有效方法（方法1）。

一般来说，我会把这个问题放在一个名为stream compaction的类别中。流压缩通常是指采取一系列数据并将其减少为较小的数据序列。

如果我们查看逆冲流压实区域，可以对此问题进行处理的算法是thrust::copy_if()（特别是为了方便起见，需要模板阵列的版本）。

方法1：

要想想并行这个问题，我们必须问自己在什么条件下应该给定元素从输入到输出复制？如果我们能够形式化这个逻辑，我们可以构造一个推力函数，我们可以传递给thrust::copy_if来指示它复制哪些元素。

对于给定的元件，对于序列长度= 2的情况下，我们可以构建完整的逻辑，如果我们知道：

1. 元件
2. 元件一个位置向右
3. 元件一个地方向左
4. 元素两个地向左

基于以上，我们将需要拿出对于未定义上述第2,3或4项中的任何元素的“特例”逻辑。

忽略的特殊情况下，如果我们知道上面的4个项目，那么我们可以构造必要的逻辑如下：

1. 如果元素到我的左边是和我一样，但元素2地方左边是不同的，那么我在输出

2. 如果元素到我的左边是比我不同的归属，但该元素在我右边的是和我一样，我在输出属于

3. 否则，我不属于我n输出

我会留给你来构造特殊情况下的必要逻辑。 （或者从我提供的代码中反向工程）。

方法2：

对于长序列，方法1或方法1中的逻辑的一个for循环变体将产生至少1读出的数据每序列长度的元件设置的。对于长序列（例如2000），这将是低效的。因此另一种可能的方法将是如下：

1. 向前生成在两者exclusive_scan_by_key和反向方向上，使用ID值作为密钥，以及thrust::constant_iterator（值= 1）作为用于扫描的值。对于给定的数据集，创建中间结果是这样的：

   ID: 1 2 2 3 3 3 
   VN: 6 7 8 5 7 8 
   FS: 0 0 1 0 1 2 
   RS: 0 1 0 2 1 0 
   

其中FS和RS是向前的结果和反向扫描逐键。我们使用.rbegin()和.rend()reverse iterators生成反向扫描（RS）。请注意，为了生成如上所述的RS序列，必须对反向扫描输入和输出执行此操作。

1. 我们的thrust::copy_if仿函数的逻辑变得相当简单。对于给定元素，如果该元素的RS和FS值的总和大于或等于期望的最小序列长度（-1以考虑排他扫描操作）和，则FS值小于期望的最小序列长度，那么该元素属于输出。

下面是两种方法的整个例子，使用给定的数据，序列长度为2：

$ cat t1095.cu 
#include <thrust/device_vector.h> 
#include <thrust/copy.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/zip_iterator.h> 
#include <iostream> 

#include <thrust/scan.h> 
#include <thrust/iterator/constant_iterator.h> 

struct copy_func 
{ 
    int *d; 
    int dsize, r, l, m, l2; 
    copy_func(int *_d, int _dsize) : d(_d),dsize(_dsize) {}; 
    __host__ __device__ 
    bool operator()(int idx) 
    { 
    m = d[idx]; 
    // handle typical case 
    // this logic could be replaced by a for-loop for sequences of arbitrary length 
    if ((idx > 1) && (idx < dsize-1)){ 
     r = d[idx+1]; 
     l = d[idx-1]; 
     l2 = d[idx-2]; 
     if ((r == m) && (m != l)) return true; 
     if ((l == m) && (m != l2)) return true; 
     return false;} 
    // handle special cases 
    if (idx == 0){ 
     r = d[idx+1]; 
     return (r == m);} 
    if (idx == 1){ 
     r = d[idx+1]; 
     l = d[idx-1]; 
     if (l == m) return true; 
     else if (r == m) return true; 
     return false;} 
    if (idx == dsize-1){ 
     l = d[idx-1]; 
     l2 = d[idx-2]; 
     if ((m == l) && (m != l2)) return true; 
     return false;} 
    // could put assert(0) here, should never get here 
    return false; 
    } 
}; 

struct copy_func2 
{ 
    int thresh; 
    copy_func2(int _thresh) : thresh(_thresh) {}; 
    template <typename T> 
    __host__ __device__ 
    bool operator()(T t){ 
    return (((thrust::get<0>(t) + thrust::get<1>(t))>=(thresh-1)) && (thrust::get<0>(t) < thresh)); 
    } 
}; 

int main(){ 

    const int length_threshold = 2; 
    int ID[] = {1,2,2,3,3,3}; 
    int VN[] = {6,7,8,5,7,8}; 
    int dsize = sizeof(ID)/sizeof(int); 
    // we assume dsize > 3 
    thrust::device_vector<int> id(ID, ID+dsize); 
    thrust::device_vector<int> vn(VN, VN+dsize); 

    thrust::device_vector<int> res_id(dsize); 
    thrust::device_vector<int> res_vn(dsize); 
    thrust::counting_iterator<int> idx(0); 

    //method 1: sequence length threshold of 2 

    int rsize = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(id.begin(), vn.begin())), thrust::make_zip_iterator(thrust::make_tuple(id.end(), vn.end())), idx, thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())), copy_func(thrust::raw_pointer_cast(id.data()), dsize)) - thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())); 

    std::cout << "ID: "; 
    thrust::copy_n(res_id.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl << "VN: "; 
    thrust::copy_n(res_vn.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 

    //method 2: for arbitrary sequence length threshold 
    thrust::device_vector<int> res_fs(dsize); 
    thrust::device_vector<int> res_rs(dsize); 
    thrust::exclusive_scan_by_key(id.begin(), id.end(), thrust::constant_iterator<int>(1), res_fs.begin()); 
    thrust::exclusive_scan_by_key(id.rbegin(), id.rend(), thrust::constant_iterator<int>(1), res_rs.begin()); 
    rsize = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(id.begin(), vn.begin())), thrust::make_zip_iterator(thrust::make_tuple(id.end(), vn.end())), thrust::make_zip_iterator(thrust::make_tuple(res_fs.begin(), res_rs.rbegin())), thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())), copy_func2(length_threshold)) - thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())); 

    std::cout << "ID: "; 
    thrust::copy_n(res_id.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl << "VN: "; 
    thrust::copy_n(res_vn.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 
    return 0; 
} 

$ nvcc -o t1095 t1095.cu 
$ ./t1095 
ID: 2 2 3 3 
VN: 7 8 5 7 
ID: 2 2 3 3 
VN: 7 8 5 7 

注：

1. 的copy_func实现测试逻辑对于给定的元素。它通过函数初始化参数接收该元素的索引（通过模板）以及指向设备上的数据的指针以及数据的大小。变量r,m,l和l2分别指向右侧的元素，我自己，左侧的元素和左侧的两个元素。

2. 我们将一个指针传递给函数的ID数据。这允许仿函数检索测试逻辑的（最多）4个必要元素。这避免了一个棘手的建设推拉:: zip_iterator提供所有这些值。请注意，仿函数中这些元素的读取应该很好地融合在一起，因此效率相当高，并且也受益于缓存。

3. 我不认为这是无缺陷的。我认为我的测试逻辑是正确的，但有可能我没有。至少，您应该验证该部分代码的逻辑正确性。我的目的不是要给你一个黑盒子的代码，而是要演示如何思考你通过问题的方式。

4. 这种方法对于长于2的关键序列可能会很麻烦。在这种情况下，我会建议方法2.（如果您已经有了一个顺序的for循环来实现必要的逻辑，那么您可以放弃一个修改这种for-loop应该可能仍然受益于来自缓存的合并访问和相邻访问）。