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

Question

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

例如，假设我们有以下一对数组：

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 上执行的唯一操作，我不会推荐它。复制数据的成本 to/from GPU 可能会超过使用 GPU 可能带来的任何 efficiency/performance 好处。

编辑： 根据序列阈值可能比 2 长得多的评论，我会建议一种应该更有效的替代方法（方法 2）比 for 循环或强力方法（方法 1）。

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

如果我们查看推流压实区域，可用于解决此问题的算法是 thrust::copy_if()（特别是，为方便起见，采用模板阵列的版本）。

方法一：

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

对于一个给定的元素，对于序列length = 2的情况，我们可以构造一个完整的逻辑，如果我们知道：

元素
右边一位元素
左边一位的元素
左边两位的元素

基于以上内容，我们需要为那些未定义上述任何项目 2、3 或 4 的元素提出 "special case" 逻辑。

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

如果我左边的元素和我一样，但是左边两处的元素不一样，那么我属于输出
如果我左边的元素和我不一样，但是我右边的元素和我一样，我属于输出
否则我不属于输出

我将留给您为特殊情况构建必要的逻辑。（或者根据我提供的代码对其进行逆向工程）。

方法二：

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

正向和反向生成一个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 序列。

我们的 thrust::copy_if 仿函数的逻辑就变得相当简单了。对于给定的元素，如果该元素的 RS 和 FS 值的 sum 大于或等于所需的最小序列长度（-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

备注：

copy_func 为方法 1 的给定元素实现测试逻辑。它接收该元素的索引（通过模板）以及指向 ID 设备上的数据，以及数据的大小，通过函子初始化参数。变量r、m、l、l2指的是我右边的元素，我自己，我左边的元素，我左边两处的元素, 分别.
我们正在将指向 ID 数据的指针传递给仿函数。这允许仿函数检索（最多）4 个测试逻辑所需的元素。这避免了使用 thrust::zip_iterator 来提供所有这些值的混乱构造。请注意，仿函数中这些元素的读取应该很好地合并，因此相当高效，并且还可以从缓存中受益。
我不声称这是没有缺陷的。我认为我的测试逻辑是正确的，但也有可能我没有。您至少应该验证那部分代码的逻辑正确性。我的目的不是给你一段黑盒代码，而是展示你如何思考解决问题的方法。
对于长度超过 2 的键序列，这种方法可能会变得很麻烦。在这种情况下，我建议使用方法 2。（如果您已经有一个实现必要逻辑的顺序 for 循环，您可能能够将其修改后的版本放入方法 1 仿函数中以获得更长的键序列。这样的 for 循环可能仍然受益于缓存的合并访问和相邻访问。）

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

Making the number of key occurances equal using CUDA / Thrust

c++

arrays

cuda

gpu

thrust