c++ - CUB 选择是否有返回的索引

Question

我最近在使用 Thrust 时遇到性能问题图书馆。这些来自在大型嵌套循环结构的基础上分配内存的推力。这显然是不需要的，理想情况下使用预分配的全局内存块执行。我想通过以下三种方式之一删除或改进有问题的代码:

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

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

我一直听说 CUB 很棒，所以我认为这是在愤怒中使用它的理想机会。我想知道的是:

如何实现 CUB select_if返回索引？

这可以用 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 的评论，仿函数应该采用取消引用 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));

Answer 1

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

$ 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，cub v1.2.2，CUDA 5.5