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

Question

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

我想为推力::排序创建我自己的比较器，但它的速度大大减慢！ 我通过从 function.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 将产生与 x 等效的结果y。

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

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