在CUDA中按键排序(小)数组

Question

我正在尝试编写一个函数,它接受一组未分类的键/值对,例如

<7, 4>
<2, 8>
<3, 1>
<2, 2>
<1, 5>
<7, 1>
<3, 8>
<7, 2>

并按键对它们进行排序,同时使用相同的键减少对的值:

<1, 5>
<2, 10>
<3, 9>
<7, 7>

目前,我正在使用__device__类似下面的一个函数,它本质上是一个bitonic排序,它将组合相同键的值并将旧数据设置为一个无限大的值(仅99用于现在),以便后续的bitonic排序将筛选它们到底部并且数组被int *删除的值切割.

__device__ void interBitonicSortReduce(int2 *sdata, int tid, int recordNum, int *removed) {
  int n = MIN(DEFAULT_DIMBLOCK, recordNum);
  for (int k = 2; k <= n; k *= 2) {
    for (int j = k / 2; j > 0; j /= 2) {
      int ixj = tid ^ j;
      if (ixj > tid) {
        if (sdata[tid].x == sdata[ixj].x && sdata[tid].x < 99) {
          atomicAdd(&sdata[tid].y, sdata[ixj].y);
          sdata[ixj].x = 99; 
          sdata[ixj].y = 99; 
          atomicAdd(removed, 1); 
        }   
        if ((tid & k) == 0 && sdata[tid].x > sdata[ixj].x)
          swapData2(sdata[tid], sdata[ixj]);
        if ((tid & k) != 0 && sdata[tid].x < sdata[ixj].x)
          swapData2(sdata[tid], sdata[ixj]);
        __syncthreads();
      }   
    }   
  }
}

这适用于小型数据集但是具有较大的集合(尽管仍然在单个块的大小内),单个调用就不会这样做.

尝试将排序和减少结合在同一个函数中是明智的吗？显然,该函数需要不止一次被调用,但有可能确切地确定需要调用多少次来根据其大小耗尽所有数据？

或者我应该用这样的东西单独预先形成减少量:

__device__ int interReduce(int2 *sdata, int tid) {
  int index = tid;
  while (sdata[index].x == sdata[tid].x) {
    index--;
    if (index < 0)
      break;
  }
  if (index+1 != tid) {
    atomicAdd(&sdata[index+1].y, sdata[tid].y);
    sdata[tid].x = 99;
    sdata[tid].y = 99;
    return 1;
  }
  return 0;
}

我正在尝试提出最有效的解决方案,但我对CUDA和并行算法的经验是有限的.

Answer 1

您可以使用推力来做到这一点。

使用thrust::sort_by_key， 然后使用thrust::reduce_by_key

这是一个例子：

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
#include <thrust/sequence.h>

#define N 12
typedef thrust::device_vector<int>::iterator dintiter;
int main(){

  thrust::device_vector<int> keys(N);
  thrust::device_vector<int> values(N);
  thrust::device_vector<int> new_keys(N);
  thrust::device_vector<int> new_values(N);
  thrust::sequence(keys.begin(), keys.end());
  thrust::sequence(values.begin(), values.end());

  keys[3] = 1;
  keys[9] = 1;
  keys[8] = 2;
  keys[7] = 4;

  thrust::sort_by_key(keys.begin(), keys.end(), values.begin());
  thrust::pair<dintiter, dintiter> new_end;
  new_end = thrust::reduce_by_key(keys.begin(), keys.end(), values.begin(), new_keys.begin(), new_values.begin());

  std::cout << "results  values:" << std::endl;
  thrust::copy(new_values.begin(), new_end.second, std::ostream_iterator<int>( std::cout, " "));
  std::cout << std::endl << "results keys:" << std::endl;
  thrust::copy(new_keys.begin(), new_end.first, std::ostream_iterator<int>( std::cout, " "));
  std::cout << std::endl;

  return 0;
}