快速CUDA推力定制比较运算符

Question

我正在评估CUDA并且目前使用Thrust库对数字进行排序.

我想为thrust :: sort创建我自己的比较器,但它会大幅减速!我只是从functional.h复制代码,创建了自己较少的实现.然而,它似乎以其他方式编译并且工作非常缓慢.

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

我正在使用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),而比较函数是内置的小于操作--Trust识别出thrust::less<int>会产生等效的结果x < y.

在案例2中,Thrust对用户提供的内容一无所知less<int>,并且必须使用基于具有不同渐近复杂度的比较排序的更保守的算法,即使事实上你less<int>的等价于thrust::less<int>.

通常,用户定义的比较运算符不能与更严格,更快速的排序一起使用,这些排序操纵数据的二进制表示,例如基数排序.在这些情况下,Thrust回归到更普遍但更慢的类型.