CUDA Thrust and sort_by_key

Question

我是一名有用的助手，可以为您翻译文本。

我正在寻找一种基于CUDA的排序算法，可以对元素数组A（双精度）进行排序，并返回该数组A的键数组B。 我知道Thrust库中的sort_by_key函数，但我希望我的元素数组A保持不变。 我该怎么做？

我的代码如下：

void sortCUDA(double V[], int P[], int N) {

        real_t *Vcpy = (double*) malloc(N*sizeof(double));
        memcpy(Vcpy,V,N*sizeof(double));

        thrust::sort_by_key(V, V + N, P);
        free(Vcpy);
}

我正在将推力算法与我在顺序CPU上拥有的其他算法进行比较。

N               mergesort       sortCUDA
113             0.000008        0.000010
226             0.000018        0.000016
452             0.000036        0.000020
905             0.000061        0.000034
1810            0.000135        0.000071
3621            0.000297        0.000156
7242            0.000917        0.000338
14484           0.001421        0.000853
28968           0.003069        0.001931
57937           0.006666        0.003939
115874          0.014435        0.008025
231749          0.031059        0.016718
463499          0.067407        0.039848
926999          0.148170        0.118003
1853998         0.329005        0.260837
3707996         0.731768        0.544357
7415992         1.638445        1.073755
14831984        3.668039        2.150179
115035495       39.276560       19.812200
230070990       87.750377       39.762915
460141980       200.940501      74.605219

推进性能还不错，但我认为如果使用 OMP 可能可以轻松获得更好的 CPU 时间。

我认为这是由于 memcpy 引起的。

解决方案：

void thrustSort(double V[], int P[], int N)
{
        thrust::device_vector<int> d_P(N);
        thrust::device_vector<double> d_V(V, V + N);
        thrust::sequence(d_P.begin(), d_P.end());

        thrust::sort_by_key(d_V.begin(), d_V.end(), d_P.begin());

        thrust::copy(d_P.begin(),d_P.end(),P);
}

其中V是需要排序的双精度数值

- Ignacio Molina Cuquerella

---

在排序之前复制A？此外，如果您是Thrust用户，可以考虑加入thrust google group。 - Robert Crovella

也许你应该发布一些代码并回答有关大小的问题。我希望排序操作的成本显着高于向量复制的成本。 - Robert Crovella

看起来你根本没有使用CUDA设备。Thrust有主机端算法和设备端算法。此外，你说添加向量复制使它“变得非常慢”。但我没有看到数据或证据表明你已经计时了差异。 - Robert Crovella

你需要学习更多关于thrust，也许看一下快速入门指南。向量可以存在于主机或设备上。如果您传递的向量（或数组指针）是基于主机的，Thrust将使用基于主机的算法进行排序（使GPU空闲）。如果您传递的向量或指针是基于设备的，则Thrust将使用基于设备的算法进行排序（即在GPU上）。您发布的代码给我留下了指针是基于主机的印象。 - Robert Crovella

我实际上对Thrust比你的归并排序更快感到印象深刻，即使对于大小为226的数据，特别是因为你正在添加向量复制的成本（不知道你是否在使用归并排序时也这样做 - 你没有发布那段代码）。如果您使用Thrust设备排序，将会有一个将向量复制到设备的成本。这将惩罚您的小型排序，但可能会大幅改善大型排序。此外，Thrust的开发版本应该在排序方面更快。 - Robert Crovella

Answer 1

您可以修改比较运算符，以按键而不是值进行排序。@Robert Crovella正确指出，无法从主机分配原始设备指针。修改后的算法如下：

struct cmp : public binary_function<int,int,bool>
{
  cmp(const double *ptr) : rawA(ptr) { }

  __host__ __device__ bool operator()(const int i, const int j) const 
  {return rawA[i] > rawA[j];}

   const double *rawA; // an array in global mem
}; 

void sortkeys(double *A, int n) {
  // move data to the gpu
  thrust::device_vector<double> devA(A, A + n);
  double *rawA = thrust::raw_pointer_cast(devA.data());

  thrust::device_vector<int> B(n);
  // initialize keys
  thrust::sequence(B.begin(), B.end());
  thrust::sort(B.begin(), B.end(), cmp(rawA));
  // B now contains the sorted keys
 }

这里提供了使用arrayfire的替代方案。虽然我不确定哪个更有效，因为arrayfire解决方案使用了两个额外的数组：

void sortkeys(double *A, int n) {
   af::array devA(n, A, af::afHost);
   af::array vals, indices;
   // sort and populate vals/indices arrays
   af::sort(vals, indices, devA);
   std::cout << devA << "\n" << indices << "\n";
}

- user1545642

---

我在尝试让这个工作，但是除此之外，如果键不是序列（0、1、2、...），这个会起作用吗？假设sort_by_key的一般情况不需要这些键。 - Robert Crovella

此外，我认为这行代码并没有实现你想要的功能：rawA = thrust::raw_pointer_cast(devA.data()); 我无法让它正常工作。虽然它可以编译通过，但是在该行之后如果你尝试解引用rawA，thrust会抛出异常。我已经成功地使用了另一种方法，基本上是相同的方法，但是使用了cudaMemcpyToSymbol，而不是那行代码。 - Robert Crovella

嗯，你说得对，在主机上分配原始设备指针没有太多意义...但是感谢您提供一个可行的示例。我不确定您所说的键不是序列（0、1、2、...）的含义是什么？如果有一个序列[0, 1, 2, .. n]，则可以为任何其他键序列提供一对一映射。 - user1545642

sort_by_key的一般推进并不要求键值为0、1、2。请查看此处的示例。你的方法在这些键上无法使用。如果对生成的索引进行排序，确实可以（稍后）重新排列键集，但是这是否比仅复制原始数据集并进行排序更快，我不确定。 - Robert Crovella

@asm 我认为这一行代码：rawA = thrust::raw_pointer_cast(devA.data()); 应该改成这样：double * rawA = thrust::raw_pointer_cast(devA.data()); 无论如何，你现在发布的代码对我来说无法编译（在那一行中未定义rawA），但如果我进行这个更改，它就可以工作了。 - Robert Crovella

哦，你说得对，我漏掉了声明。我在我的机器上编译了算法，它对我来说运行良好。 - user1545642

Answer 2

在@asm提供的答案基础上（我没能让它工作），这段代码对我来说似乎有效，并且只对键进行排序。然而，我认为它仅限于键按照0、1、2、3、4...的顺序排列，对应于（双精度）值。由于这是一种“索引-值”排序，可以通过进行索引复制来扩展到任意键序列的情况，也许可以。但是，我不确定生成索引序列然后重新排列原始键的过程是否比仅将原始值数据复制到新向量中更快（对于任意键的情况）。

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/sort.h>

using namespace std;

__device__  double *rawA; // an array in global mem

struct cmp : public binary_function<int, int, bool>
{
  __host__ __device__  bool operator()(const int i, const int j) const
  {return ( rawA[i] < rawA[j]);}
};

void sortkeys(double *A, int n) {
  // move data to the gpu
  thrust::device_vector<double> devA(A, A + n);
//  rawA = thrust::raw_pointer_cast(&(devA[0]));
  double *test = raw_pointer_cast(devA.data());
  cudaMemcpyToSymbol(rawA, &test, sizeof(double *));

  thrust::device_vector<int> B(n);
  // initialize keys
  thrust::sequence(B.begin(), B.end());
  thrust::sort(B.begin(), B.end(), cmp());
  // B now contains the sorted keys
  thrust::host_vector<int> hostB = B;
  for (int i=0; i<hostB.size(); i++)
    std::cout << hostB[i] << " ";
  std::cout<<std::endl;
  for (int i=0; i<hostB.size(); i++)
    std::cout << A[hostB[i]] << " ";
  std::cout<<std::endl;
 }


int main(){

  double C[] = {0.7, 0.3, 0.4, 0.2, 0.6, 1.2, -0.5, 0.5, 0.0, 10.0};
  sortkeys(C, 9);
  std::cout << std::endl;
  return 0;
}

- Robert Crovella

Answer 3

这个数组有多大？如果内存允许，从速度的角度来看，最有效的方法可能是在排序之前先复制原始数组。

- DRobinson

---

这是我第一次做的，但速度太慢了。 - Ignacio Molina Cuquerella