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

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

实现自定义推力内存分配器

用CUB代码替换推力代码(带有预先分配的临时存储)

编写自定义内核以执行我想要的操作

尽管第三个选项是我通常的首选，但我要执行的操作是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>应该是使仿函数现在成为d_out的产品:

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);
    }
};

在代码中，cub::ItemOffsetPair<int,int>应该是ojit_code的设备数组:

//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));

最佳答案

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

$ 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

关于c++ - CUB选择是否返回索引，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/22476069/