在 Thrust 比较运算符中加速 device 函数

Question

我是运行数组上的 Thrust 并行二进制搜索类型例程：

// array and array2 are raw pointers to device memory
thrust::device_ptr<int> array_ptr(array);

// Search for first position where 0 could be inserted in array
// without violating the ordering
thrust::device_vector<int>::iterator iter;
iter = thrust::lower_bound(array_ptr, array_ptr+length, 0, cmp(array2));

一个自定义函数对象cmp定义了一个自定义比较运算符：

struct cmp
{
    cmp(int *array2){ this->array2 = array2; }

    __device__ bool operator()(const int& x, const int& y)
    {
        return device_function(array2,x) <= device_function(array2,y);
    }

    int *array2;
};

比较依赖于对在设备上编译的函数的调用：

__device__ int device_function( const int* array2, const int value ){
    int quantity = 0;

    for (int i = 0; i < 50000; ++i){
        if ( array2[i] > value ){ quantity += array2[i]; }
    }

    return quantity;
}

我的问题是：device_function 中的总和减少在设备上执行了什么（如果有的话）并行执行？如果函数是这样串行执行的，我怎样才能引入并行来加速函数评估？

Answer 1

My question is: what (if any) parallel execution is done on the device for the sum-reduction in device_function?

None。 C/C++ __device__ 函数中的普通代码（无论是在 CUDA 还是 Thrust 中）从单个 CUDA 线程的上下文中顺序执行。

If the function executes serially as such, how can I introduce parallelism to speed up the function evaluation?

一种可能的方法是使用 Thrust v1.8（可从 github 或 CUDA 7 RC 获得）并在传递给 thrust::lower_bound.

是一个在传递给另一个推力函数的自定义仿函数中使用 thrust::sort 的有效示例。

使用此方法进行并行化需要在支持 CUDA 动态并行化的设备上进行编译和执行。并且不能保证整体加速，就像任何 CUDA 动态并行代码一样。这种级别的并行性是否会提供任何好处将取决于许多因素，例如先前级别的并行性是否已经最大限度地利用了设备。

例如，您的 device_function 中包含的函数似乎可以通过对 thrust::transform_reduce 的单个调用来替换。然后，您的 cmp 函数可以重写为如下内容（在浏览器中编码，未经测试）：

struct cmp
{
    cmp(int *array2){ this->array2 = array2; }

    __device__ bool operator()(const int& x, const int& y)
    {
        return (thrust::transform_reduce(thrust::device, array2,array2+50000, my_greater_op(x), 0, thrust::plus<int>()) <= thrust::transform_reduce(thrust::device, array2,array2+50000, my_greater_op(y), 0, thrust::plus<int>()));
    }

    int *array2;

};

并且您必须提供适当的 my_greater_op 仿函数：

struct my_greater_op
{
  int val;
  my_greater_op(int _val) {val = _val;}
  __host__ __device__ int operator(const int& x)
  {
     return (x>val)?x:0;
  }
};

在 Thrust 比较运算符中加速 device 函数

Accelerating device function in Thrust comparison operator

parallel-processing

cuda

gpgpu

thrust

在 Thrust 比较运算符中加速 __device__ 函数

Accelerating __device__ function in Thrust comparison operator

parallel-processing

cuda

gpgpu

thrust

在 Thrust 比较运算符中加速 device 函数

Accelerating device function in Thrust comparison operator