如何在不排序的情况下使用推力将相等的元素放在一起

Question

我有一个元素数组，每个元素只定义 "equal to" 运算符。换句话说，没有为这种类型的元素定义排序。

由于我不能像在推力直方图示例中那样使用 thrust::sort 如何使用推力将相等的元素放在一起？

例如：

我的数组最初是

a e t b c a c e t a

其中相同的字符表示相同的元素。

经过细化，数组应该是

a a a t t b c c e e

但也可以

a a a c c t t e e b

或任何其他排列

谢谢

Answer 1

In the discussion 我们发现您的真正目标是消除 float4 元素向量中的重复项。为了应用 thrust::unique 需要对元素进行排序。

所以你需要一个4维数据的排序方法。这可以使用 space 填充曲线来完成。我以前使用过 z-order curve (aka morton code) to sort 3D data. There are efficient CUDA implementations for the 3D case available，但是快速谷歌搜索并没有 return 4D 案例的现成可用实现。

我找到一篇论文，其中列出了使用 z 顺序曲线对 n 维数据点进行排序的通用算法： Fast construction of k-Nearest Neighbor Graphs for Point Clouds （参见算法 1：浮点莫顿阶算法）。这个算法还有一个C++ implementation available。

对于 4D 数据，循环可以展开，但可能有更简单、更有效的算法可用。

所以（未完全实现的）操作序列将如下所示：

#include <thrust/device_vector.h>
#include <thrust/unique.h>
#include <thrust/sort.h>

inline __host__ __device__ float dot(const float4& a, const float4& b)
{
    return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w;
}

struct identity_4d
{
  __host__ __device__
  bool operator()(const float4& a, const float4& b) const
  {
    // based on the norm function you provided in the discussion
    return dot(a,b) < (0.1f*0.1f);
  }
};

struct z_order_4d
{
  __host__ __device__
  bool operator()(const float4& p, const float4& q) const
  {
    // you need to implement the z-order algorithm here
    // ...
  }
};

int main()
{
  const int N = 100;
  thrust::device_vector<float4> data(N);
  // fill the data
  // ...

  thrust::sort(data.begin(),data.end(), z_order_4d());

  thrust::unique(data.begin(),data.end(), identity_4d());

}

Answer 2

我建议您遵循@m.s 提出的方法。在那里发布的答案中。正如我在评论中所说，元素排序是一种非常有用的机制，有助于降低此类问题的复杂性。

然而，所提出的问题询问是否可以在不排序的情况下对相似的元素进行分组。对于像 GPU 这样的固有并行处理器，我花了一些时间思考如何在没有排序的情况下完成它。

如果我们既有大量的对象，又有大量独特的对象类型，那么我认为可以为问题带来一定程度的并行性，但是我在这里概述的方法仍然会有残暴的、分散的内存访问模式。对于只有少量不同或唯一对象类型的情况，我在这里讨论的算法几乎没有什么可称道的。这只是一种可能的方法。可能还有其他更好的方法：

起点是为每个元素开发一组 "linked lists"，指示左侧的匹配邻居和右侧的匹配邻居。这是通过我的 search_functor 和 thrust::for_each 在整个数据集上完成的。这一步是相当并行的，对于大数据集也有合理的内存访问效率，但它确实需要从头到尾遍历整个数据集的最坏情况（副作用，我会称之为，无法使用排序；我们必须将每个元素与其他元素进行比较，直到找到匹配为止）。两个链表的生成让我们避免了all-to-all比较
一旦我们从第 1 步构建了列表（右邻居和左邻居），使用 thrust::count.[= 计算唯一对象的数量就很容易了。 20=]
然后我们使用thrust::copy_if流压缩获得每个唯一元素的起始索引（即数据集中每种类型唯一元素的最左边索引）。
下一步是计算每个唯一元素的实例数。这一步是做列表遍历，每个元素列表一个线程。如果我有少量独特的元素，这将无法有效地利用 GPU。此外，列表遍历将导致糟糕的访问模式。
在我们计算了每种类型对象的数量之后，我们可以通过thrust::exclusive_scan在输出列表中为每种对象类型构建一个起始索引序列每种类型的对象。
最后，我们可以将每个输入元素复制到输出列表中的适当位置。由于我们还没有办法对元素进行分组或排序，我们必须再次求助于列表遍历。再一次，如果唯一对象类型的数量很少，这将是 GPU 的低效使用，并且还会有糟糕的内存访问模式。

这是一个完整的示例，使用您的示例字符数据集。为了帮助阐明我们打算对没有固有顺序的对象进行分组的想法，我创建了一个有点随意的对象定义 (my_obj)，它定义了 == 比较运算符，但没有定义 < 或 >.

$ cat t707.cu
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/for_each.h>
#include <thrust/transform.h>
#include <thrust/transform_scan.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>
#include <thrust/count.h>
#include <iostream>


template <typename T>
class my_obj
{
  T element;
  int index;
  public:
  __host__ __device__ my_obj() : element(0), index(0) {};
  __host__ __device__ my_obj(T a) : element(a), index(0) {};
  __host__ __device__ my_obj(T a, int idx) : element(a), index(idx) {};
  __host__ __device__
  T get() {
    return element;}
  __host__ __device__
  void set(T a) {
    element = a;}
  __host__ __device__
  int get_idx() {
    return index;}
  __host__ __device__
  void set_idx(int idx) {
    index = idx;}
  __host__ __device__
  bool operator ==(my_obj &e2)
  {
    return (e2.get() == this->get());
  }
};

template <typename T>
struct search_functor
{
    my_obj<T> *data;
    int end;
    int *rn;
    int *ln;
    search_functor(my_obj<T> *_a, int *_rn, int *_ln, int len) : data(_a), rn(_rn), ln(_ln), end(len) {};
    __host__ __device__
    void operator()(int idx){
      for (int i = idx+1; i < end; i++)
        if (data[idx] == data[i]) {
          ln[i] = idx;
          rn[idx] = i;
          return;}
      return;
      }
};

template <typename T>
struct copy_functor
{
    my_obj<T> *data;
    my_obj<T> *result;
    int *rn;
    copy_functor(my_obj<T> *_in, my_obj<T> *_out, int *_rn) : data(_in), result(_out), rn(_rn) {};
    __host__ __device__
    void operator()(const thrust::tuple<int, int> &t1) const {
      int idx1 = thrust::get<0>(t1);
      int idx2 = thrust::get<1>(t1);
      result[idx1] = data[idx2];
      int i = rn[idx2];
      int j = 1;
      while (i != -1){
        result[idx1+(j++)] = data[i];
        i = rn[i];}
      return;
      }
};

struct count_functor
{
    int *rn;
    int *ot;
    count_functor(int *_rn, int *_ot) : rn(_rn), ot(_ot) {};
    __host__ __device__
    int operator()(int idx1, int idx2){
      ot[idx1] = idx2;
      int i = rn[idx1];
      int count = 1;
      while (i != -1) {
        ot[i] = idx2;
        count++;
        i = rn[i];}
      return count;
      }
};

using namespace thrust::placeholders;

int main(){

  // data setup
  char data[] = { 'a' ,  'e' ,  't' ,  'b' ,  'c' ,  'a' ,  'c' , 'e' ,  't' ,  'a' };
  int sz = sizeof(data)/sizeof(char);
  for (int i = 0; i < sz; i++) std::cout << data[i] << ",";
  std::cout << std::endl;
  thrust::host_vector<my_obj<char> > h_data(sz);
  for (int i = 0; i < sz; i++) { h_data[i].set(data[i]); h_data[i].set_idx(i); }
  thrust::device_vector<my_obj<char> > d_data = h_data;


  // create left and right neighbor indices
  thrust::device_vector<int> ln(d_data.size(), -1);
  thrust::device_vector<int> rn(d_data.size(), -1);
  thrust::for_each(thrust::counting_iterator<int>(0), thrust::counting_iterator<int>(0) + sz, search_functor<char>(thrust::raw_pointer_cast(d_data.data()), thrust::raw_pointer_cast(rn.data()), thrust::raw_pointer_cast(ln.data()), d_data.size()));
  // determine number of unique objects
  int uni_objs = thrust::count(ln.begin(), ln.end(), -1);
  // determine the number of instances of each unique object
    // get object starting indices
  thrust::device_vector<int> uni_obj_idxs(uni_objs);
  thrust::copy_if(thrust::counting_iterator<int>(0), thrust::counting_iterator<int>(0)+d_data.size(), ln.begin(), uni_obj_idxs.begin(), (_1 == -1));
    // count each object list
  thrust::device_vector<int> num_objs(uni_objs);
  thrust::device_vector<int> obj_type(d_data.size());
  thrust::transform(uni_obj_idxs.begin(), uni_obj_idxs.end(), thrust::counting_iterator<int>(0),  num_objs.begin(), count_functor(thrust::raw_pointer_cast(rn.data()), thrust::raw_pointer_cast(obj_type.data())));

  // at this point, we have built object lists that have allowed us to identify a unique, orderable "type" for each object
  // the sensible thing to do would be to employ a sort_by_key on obj_type and an index sequence at this point
  // and use the reordered index sequence to reorder the original objects, thus grouping them
  // however...  without sorting...

  // build output vector indices
  thrust::device_vector<int> copy_start(num_objs.size());
  thrust::exclusive_scan(num_objs.begin(), num_objs.end(), copy_start.begin());
  // copy (by object type) input to output
  thrust::device_vector<my_obj<char> > d_result(d_data.size());
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(copy_start.begin(), uni_obj_idxs.begin())), thrust::make_zip_iterator(thrust::make_tuple(copy_start.end(), uni_obj_idxs.end())), copy_functor<char>(thrust::raw_pointer_cast(d_data.data()), thrust::raw_pointer_cast(d_result.data()), thrust::raw_pointer_cast(rn.data())));

  // display results
  std::cout << "Grouped: " << std::endl;
  for (int i = 0; i < d_data.size(); i++){
    my_obj<char> temp = d_result[i];
    std::cout << temp.get() << ",";}
  std::cout << std::endl;
  for (int i = 0; i < d_data.size(); i++){
    my_obj<char> temp = d_result[i];
    std::cout << temp.get_idx() << ",";}
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t707 t707.cu
$ ./t707
a,e,t,b,c,a,c,e,t,a,
Grouped:
a,a,a,e,e,t,t,b,c,c,
0,5,9,1,7,2,8,3,4,6,
$

如何在不排序的情况下使用推力将相等的元素放在一起

How to bring equal elements together using thrust without sort

cuda

thrust