快速 CUDA 推力自定义比较运算符

Question

我正在评估 CUDA，目前正在使用 Thrust 库对数字进行排序。

我想为推力::排序创建自己的比较器，但它的速度大大减慢！ 我通过从 functional.h 复制代码创建了自己的 less 实现。 但是，它似乎以其他方式编译并且运行速度非常慢。

1. 默认比较器：thrust::less() - 94ms
2. 我自己的比较器：less() - 906ms

我使用的是 Visual Studio 2010。我应该怎么做才能获得与选项 1 相同的性能？

完整代码：

#include <stdio.h>

#include <cuda.h>

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

int myRand()
{
        static int counter = 0;
        if ( counter++ % 10000 == 0 )
                srand(time(NULL)+counter);
        return (rand()<<16) | rand();
}

template<typename T>
struct less : public thrust::binary_function<T,T,bool>
{
  __host__ __device__ bool operator()(const T &lhs, const T &rhs) const {
     return lhs < rhs;
  }
}; 

int main()
{
    thrust::host_vector<int> h_vec(10 * 1000 * 1000);
    thrust::generate(h_vec.begin(), h_vec.end(), myRand);

    thrust::device_vector<int> d_vec = h_vec;

    int clc = clock();
    thrust::sort(d_vec.begin(), d_vec.end(), less<int>());
    printf("%dms\n", (clock()-clc) * 1000 / CLOCKS_PER_SEC);

    return 0;
}

Answer 1

您观察到性能差异的原因是 Thrust 使用不同的算法实现排序，具体取决于提供给 thrust::sort 的参数。

在情况 1. 中，Thrust 可以证明可以使用基数排序在线性时间内实现排序。这是因为要排序的数据类型是内置的数值类型（int），比较函数是内置的小于运算——Thrust 识别出thrust::less<int> 会产生与@等价的结果987654325@.

在情况 2 中，Thrust 对您的用户提供的 less<int> 一无所知，并且必须使用基于具有不同渐近复杂度的比较排序的更保守算法，即使实际上您的 less<int> 相当于thrust::less<int>.

一般来说，用户定义的比较运算符不能与更严格、更快的排序一起使用，这些排序操作数据的二进制表示，例如基数排序。在这些情况下，Thrust 会使用更通用但更慢的排序方式。