来自线性化矩阵 CUDA 的唯一行

Question

我有一个存储为 thrust::device_vector<int>

的线性化矩阵

本质上就是一个维度为nc x nv的矩阵，存储在这个大小的线性数组中。

我想从此矩阵中获取唯一的行。如果至少有一个元素不同，则两行是唯一的。

我想使用 CUDA thrust::sort 和 thrust::unique 函数来执行此操作。我相信我需要构造一个对应于每一行的迭代器，然后使用一个按元素比较行的仿函数调用排序。但我不清楚这将如何完成。

使用跨步范围迭代器将允许我指定每一行的开始，但函子的实现尚不清楚。

这似乎是一个应该可以用推力解决的问题。有没有更好的方法？

Answer 1

我觉得你的方法可行。我建议对一个单独的行索引数组进行排序，而不是直接对矩阵进行排序，这样生成的行索引将按照矩阵行的排序顺序进行排序。

我们将创建一个带有两个行索引的排序仿函数，并使用它们索引矩阵的适当行。然后，该排序仿函数将使用 element-by-element 比较对指定的两个行进行排序。

我们将对传递给 thrust::unique 的 "equality" 仿函数使用类似的方法（传递两个行索引）。相等仿函数然后将测试两个指示的行是否相等。我本可以在这里使用 for-loop 作为排序仿函数，测试相等性 element-by-element，但为了多样性，我选择使用嵌套的 thrust::mismatch 算法。

这是一个有效的例子：

$ cat t1033.cu
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/unique.h>
#include <thrust/sequence.h>
#include <assert.h>
#include <iostream>
#include <thrust/execution_policy.h>
#include <thrust/mismatch.h>


typedef int mytype;

struct my_sort_func
{
  int cols;
  mytype *data;
  my_sort_func(int _cols, mytype *_data) : cols(_cols),data(_data) {};
  __host__ __device__
  bool operator()(int r1, int r2){
    for (int i = 0; i < cols; i++){
      if (data[cols*r1+i] < data[cols*r2+i])
        return true;
      else if (data[cols*r1+i] > data[cols*r2+i])
        return false;}
    return false;
    }
};

struct my_unique_func
{
  int cols;
  mytype *data;
  my_unique_func(int _cols, mytype *_data) : cols(_cols),data(_data) {};
  __device__
  bool operator()(int r1, int r2){
    thrust::pair<mytype *, mytype *> res = thrust::mismatch(thrust::seq, data+(r1*cols), data+(r1*cols)+cols, data+(r2*cols));
    return (res.first == data+(r1*cols)+cols);
    }
};

int main(){

  const int ncols = 3;
  mytype data[] = { 1, 2, 3, 1, 2, 3, 1, 3, 5, 2, 3, 4, 1, 2, 3, 1, 3, 5};
  size_t dsize = sizeof(data)/sizeof(mytype);
  assert ((dsize % ncols) == 0);
  int nrows = dsize/ncols;
  thrust::device_vector<mytype> d_data(data, data+dsize);
  thrust::device_vector<int> rowidx(nrows);  // reference rows by their index
  thrust::sequence(rowidx.begin(), rowidx.end());
  thrust::sort(rowidx.begin(), rowidx.end(), my_sort_func(ncols, thrust::raw_pointer_cast(d_data.data())));
  int rsize = thrust::unique(rowidx.begin(), rowidx.end(), my_unique_func(ncols, thrust::raw_pointer_cast(d_data.data()))) - rowidx.begin();
  thrust::host_vector<int> h_rowidx = rowidx;
  std::cout << "Unique rows: " << std::endl;
  for (int i = 0; i < rsize; i++){
    for (int j = 0; j < ncols; j++) std::cout << data[h_rowidx[i]*ncols+j] << ",";
    std::cout << std::endl;}
  return 0;
}
$ nvcc -o t1033 t1033.cu
$ ./t1033
Unique rows:
1,2,3,
1,3,5,
2,3,4,
$

备注：

我怀疑如果对输入矩阵进行转置，并且我们比较列（在转置矩阵中）而不是行，则整体性能会提高。它可能为排序操作提供一些好处，我怀疑它也可能为唯一操作提供一些好处。然而，给定的代码与您在问题中的描述相匹配，它应该是一个很好的路线图，说明如何在列案例中执行此操作，尽管必须为此进行重构。
这个方法实际上并不re-order矩阵行。为了提高效率，我想避免做很多 data-movement，因为问题陈述似乎并不依赖于它。如果您确实想要一个按排序顺序排列矩阵行的中间数据集，我仍然建议执行上述排序操作，然后使用结果 re-order 在单个操作中使用矩阵，使用以下两个之一可能的方法：scatter/gather 操作，或 thrust::permuation_iterator 结合 thrust::copy 操作。
有点聪明，嵌套的 thrust::mismatch 操作也可以用在排序仿函数中，代替 for-loop.

来自线性化矩阵 CUDA 的唯一行

Unique rows from Linearized Matrix CUDA

c++

matrix

cuda

thrust