Потоковое сжатие и преобразование на основе индекса в CUDA

Question

У меня есть массив float на моем устройстве, и я хотел бы выполнить операцию сжатия stram (как представлено здесь: http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html), а затем применить преобразование на основе на значение и адрес или оригинальный элемент.

Например, у меня есть массив со значениями {10,-1, -10, 2}, и я хочу вернуть все элементы с абсолютным значением больше 5 и применить преобразование, принимающее значение и его адрес в массив. Результатом здесь будет {transform(10,0),transform(-10,2)}.

Я пытаюсь использовать Thrust с этим, но это код, который будет часто запускаться на больших массивах, поэтому в идеале он не будет использовать буферы и многократный обход массива.

Можно ли сделать то, что я хотел бы сделать, без выделения вторичного массива и выполнения нескольких обходов? Если да, то существует ли такой код? Или, по крайней мере, у кого-нибудь есть указатели, на какие функции тяги или любую другую библиотеку я мог бы составить, чтобы достичь своей цели?

Answer 1

Да, это возможно в режиме тяги с помощью одного вызова алгоритма тяги (я полагаю, это то, что ты имеешь в виду под "без... нескольких обходов") и без "выделения вторичного массива".

Один из подходов заключается в передаче массива данных плюс массив индекса /"адреса" (через thrust::counting_iterator, что позволяет избежать выделения) thrust::transform_iterator который создает вашу операцию "преобразования" (в сочетании с соответствующим функтором).

Затем вы передадите вышеупомянутый итератор преобразования в соответствующий алгоритм сжатия потока тяги, чтобы выбрать нужные значения.

Вот возможный подход:

$ cat t1044.cu
#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>
#include <math.h>

#include <iostream>

__host__ __device__ int my_transform(int data, int idx){
  return (data - idx);  //put whatever transform you want here
}

struct my_transform_func : public thrust::unary_function<thrust::tuple<int, int>, int>
{

  __host__ __device__
  int operator()(thrust::tuple<int, int> &t){
    return my_transform(thrust::get<0>(t), thrust::get<1>(t));
    }
};

struct my_test_func
{
  __host__ __device__
  bool operator()(int data){
    return (abs(data) > 5);
    }
};



int main(){

  int data[] = {10,-1,-10,2};
  int dsize = sizeof(data)/sizeof(int);

  thrust::device_vector<int> d_data(data, data+dsize);
  thrust::device_vector<int> d_result(dsize);
  int rsize = thrust::copy_if(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(d_data.begin(), thrust::counting_iterator<int>(0))), my_transform_func()), thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(d_data.end(), thrust::counting_iterator<int>(dsize))), my_transform_func()), d_data.begin(), d_result.begin(), my_test_func()) - d_result.begin();
  thrust::copy_n(d_result.begin(), rsize, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t1044 t1044.cu
$ ./t1044
10,-12,
$

Некоторые возможные критики этого подхода:

1. Казалось бы, загрузка d_data элементы дважды (один раз для операции преобразования, один раз для трафарета). Однако возможно, что оптимизирующий компилятор CUDA распознает избыточную нагрузку в коде потока, который в конечном итоге генерируется, и оптимизирует его.

2. Может показаться, что мы выполняем операцию преобразования для каждого элемента данных, независимо от того, намереваемся ли мы сохранить его или нет в результате. Еще раз, возможно, что тяга copy_if Реализация может фактически отложить операцию загрузки данных до тех пор, пока не будет принято решение по трафарету. Если бы это было так, возможно, что преобразование выполняется только по мере необходимости. Даже если это делается всегда, это может быть незначительной проблемой, так как многие операции тяги, как правило, связаны с загрузкой / хранением или пропускной способностью памяти, а не с вычислениями. Однако интересным альтернативным подходом может быть использование здесь созданной @ms адаптации, которая создает преобразование, применяемое к шагу итератора вывода, что, по-видимому, ограничивает операцию преобразования только для тех элементов данных, которые фактически сохраняются в результате. Хотя я и этого не проверял.

3. Как упомянуто в комментарии ниже, этот подход распределяет временное хранилище (это делается в рамках copy_if операция) и, конечно, я явно выделяю O(n) памяти для результата. Я подозреваю, что распределение тяги (один cudaMalloc), вероятно, также для хранения O(n). Хотя может быть возможно выполнить все запрошенные задачи (параллельная сумма префикса, сжатие потока, преобразование данных) без какого-либо дополнительного хранения данных (поэтому, возможно, запрос относится к операции на месте), я думаю, что создание Алгоритм таким образом, вероятно, будет иметь значительные негативные последствия для производительности, если он вообще выполним (мне не ясно, что параллельная сумма префикса может быть реализована без какого-либо дополнительного хранения, не говоря уже о связывании этого с уплотнением потока, т.е. данными движение параллельно). Так как Thrust освобождает все такое временное хранилище, которое он использует, не может быть большой проблемы с хранилищем, связанной с частым использованием этого метода. Единственная оставшаяся проблема (я полагаю) - это производительность. Если производительность вызывает беспокойство, то временные накладные расходы, связанные с временными распределениями, должны быть в основном устранены путем соединения вышеуказанного алгоритма с настраиваемым распределителем нагрузки (также см. Здесь), который выделит максимально необходимый буфер памяти один раз, а затем повторно использует этот буфер каждый раз, когда используется вышеупомянутый алгоритм.