CUB выберите, если с возвращенными индексами

Я недавно столкнулся с проблемами производительности при использовании 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 blockint 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));

Будет ли это попытаться сделать какое-либо распределение памяти под капотом?

РЕДАКТИРОВАТЬ:

Таким образом, исходя из комментария Роберта Кровеллы, функтор должен взять продукт разыменования 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));