2

我最近在使用该Thrust库时遇到了性能问题。这些来自于在大型嵌套循环结构的基础上分配内存的推力。这显然是不希望的,理想的执行是使用预先分配的全局内存板。我想通过以下三种方式之一删除或改进有问题的代码:

  1. 实现自定义推力内存分配器
  2. 用 CUB 代码替换推力代码(带有预先分配的临时存储)
  3. 编写一个自定义内核来做我想做的事

虽然第三个选项是我通常的首选,但我想要执行的操作是copy_if/select_if类型操作,其中返回数据和索引。编写自定义内核可能会重新发明轮子,因此我更愿意选择其他两个选项之一。

我一直听到关于 CUB 的好消息,所以我认为这是在愤怒中使用它的理想机会。我想知道的是:

如何实现select_if带有返回索引的 CUB?

这可以用一个ArgIndexInputIterator和这样的仿函数来完成吗?

struct GreaterThan
{
    int compare;

    __host__ __device__ __forceinline__
    GreaterThan(int compare) : compare(compare) {}

    __host__ __device__ __forceinline__
    bool operator()(const cub::ArgIndexInputIterator<int> &a) const {
        return (a.value > compare);
    }
};

在代码的主体中包含以下内容:

//d_in = device int array
//d_temp_storage = some preallocated block


int threshold_value;
GreaterThan select_op(threshold_value);

cub::ArgIndexInputIterator<int> input_itr(d_in);
cub::ArgIndexInputIterator<int> output_itr(d_out); //????


CubDebugExit(DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, output_itr, d_num_selected, num_items, select_op));

这会尝试在后台进行任何内存分配吗?

编辑:

因此,根据 Robert Crovella 的评论,仿函数应该采用取消引用 a 的产品cub::ArgIndexInputIterator<int>,这应该是cub::ItemOffsetPair<int>现在的仿函数:

struct GreaterThan
{
    int compare;

    __host__ __device__ __forceinline__
    GreaterThan(int compare) : compare(compare) {}

    __host__ __device__ __forceinline__
    bool operator()(const cub::ItemOffsetPair<int,int> &a) const {
        return (a.value > compare);
    }
};

在代码中,d_out应该是一个设备数组cub::ItemOffsetPair<int,int>

//d_in = device int array
//d_temp_storage = some preallocated block

cub::ItemOffsetPair<int,int> * d_out;
//allocate d_out

int threshold_value;
GreaterThan select_op(threshold_value);

cub::ArgIndexInputIterator<int,int> input_itr(d_in);
CubDebugExit(DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, d_out, d_num_selected, num_items, select_op));
4

2 回答 2

3

我最近在使用 Thrust 库时遇到了性能问题。这些来自于在大型嵌套循环结构的基础上分配内存的推力。这显然是不希望的,理想的执行是使用预先分配的全局内存板。

Thrust 允许您自定义在算法执行期间如何分配临时内存。

请参阅custom_temporary_allocation 示例以了解如何为预分配的平板创建缓存。

于 2014-03-19T05:38:40.133 回答
3

经过一番摆弄和询问后,我能够按照您建议的工作方式获得一个简单的代码:

$ cat t348.cu
#include <cub/cub.cuh>
#include <stdio.h>
#define DSIZE 6

struct GreaterThan
{

    __host__ __device__ __forceinline__
    bool operator()(const cub::ItemOffsetPair<int, ptrdiff_t> &a) const {
        return (a.value > DSIZE/2);
    }
};

int main(){

  int num_items = DSIZE;
  int *d_in;
  cub::ItemOffsetPair<int,ptrdiff_t> * d_out;
  int *d_num_selected;
  int *d_temp_storage = NULL;
  size_t temp_storage_bytes = 0;

  cudaMalloc((void **)&d_in, num_items*sizeof(int));
  cudaMalloc((void **)&d_num_selected, sizeof(int));
  cudaMalloc((void **)&d_out, num_items*sizeof(cub::ItemOffsetPair<int,ptrdiff_t>));

  int h_in[DSIZE] = {5, 4, 3, 2, 1, 0};
  cudaMemcpy(d_in, h_in, num_items*sizeof(int), cudaMemcpyHostToDevice);

  cub::ArgIndexInputIterator<int *> input_itr(d_in);


  cub::DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, d_out, d_num_selected, num_items, GreaterThan());

  cudaMalloc(&d_temp_storage, temp_storage_bytes);

  cub::DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, d_out, d_num_selected, num_items, GreaterThan());
  int h_num_selected = 0;
  cudaMemcpy(&h_num_selected, d_num_selected, sizeof(int), cudaMemcpyDeviceToHost);
  cub::ItemOffsetPair<int, ptrdiff_t> h_out[h_num_selected];
  cudaMemcpy(h_out, d_out, h_num_selected*sizeof(cub::ItemOffsetPair<int, ptrdiff_t>), cudaMemcpyDeviceToHost);
  for (int i =0 ; i < h_num_selected; i++)
    printf("index: %d, offset: %d, value: %d\n", i, h_out[i].offset, h_out[i].value);

  return 0;
}
$ nvcc -arch=sm_20 -o t348 t348.cu
$ ./t348
index: 0, offset: 0, value: 5
index: 1, offset: 1, value: 4
$

RHEL 6.2、幼崽 v1.2.2、CUDA 5.5

于 2014-03-20T15:50:02.630 回答