CUB使用返回的索引选择它

CUB select if with returned indexes

本文关键字：索引 选择 返回 CUB      更新时间：2023-10-16

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

  1. 实现自定义推力内存分配器
  2. 用CUB代码替换推力代码(预分配临时存储)
  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));

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

$ 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: %dn", 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

我最近在使用Thrust库时遇到了性能问题。这些来自于在大型嵌套循环结构的底部分配内存。这是显然不需要，理想的执行使用预先分配的全局内存。

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

查看custom_temporary_allocation示例，了解如何为预分配的slab创建缓存。

相关文章：