Сжатие потока с помощью Thrust; лучшие практики и самый быстрый способ?

Question

Я заинтересован в портировании некоторого существующего кода для использования Thrust, чтобы посмотреть, смогу ли я сравнительно легко ускорить его на GPU.

Чего я хочу добиться, так это операции сжатия потока, в которой будут храниться только ненулевые элементы. У меня это в основном работает, в соответствии с примером кода ниже. Часть, в которой я не уверен, как справиться, имеет дело со всем дополнительным пространством заполнения, которое находится в d_res и, следовательно, в h_res, после сжатия.

В примере просто используется последовательность 0-99 со всеми четными записями, установленными на ноль. Это всего лишь пример, и настоящей проблемой будет общий разреженный массив.

Этот ответ мне очень помог, хотя когда дело доходит до считывания данных, размер, как известно, постоянен: как быстро сжать разреженный массив с помощью CUDA C?

Я подозреваю, что могу обойти это, подсчитав количество нулей в d_src, а затем выделив только d_res для этого размера или выполнив подсчет после сжатия, и скопировав только столько элементов. Это действительно правильный способ сделать это?

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

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>

//Predicate functor
struct is_not_zero
{
    __host__ __device__
        bool operator()(const int x)
    {
        return (x != 0);
    }
};

using namespace std;

int main(void)
{
    size_t N = 100;

    //Host Vector
    thrust::host_vector<int> h_src(N);

    //Fill with some zero and some nonzero data, as an example
    for (int i = 0; i < N; i++){
        if (i % 2 == 0){
            h_src[i] = 0;
        }
        else{
            h_src[i] = i;
        }
    }

    //Print out source data
    cout << "Source:" << endl;

    for (int i = 0; i < N; i++){
        cout << h_src[i] << " ";
    }
    cout << endl;

    //copies to device
    thrust::device_vector<int> d_src = h_src;

    //Result vector
    thrust::device_vector<int> d_res(d_src.size());

    //Copy non-zero elements from d_src to d_res
    thrust::copy_if(d_src.begin(), d_src.end(), d_res.begin(), is_not_zero());

    //Copy back to host
    thrust::host_vector<int> h_res(d_res.begin(), d_res.end());
    //thrust::host_vector<int> h_res = d_res; //Or just this?

    //Show results
    cout << "h_res size is " << h_res.size() << endl;
    cout << "Result after remove:" << endl;

    for (int i = 0; i < h_res.size(); i++){
        cout << h_res[i] << " ";
    }
    cout << endl;

    return 0;
}

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

Точно так же скорость всегда интересна. Читая некоторые из различных руководств по тяге, кажется, что здесь небольшие изменения, и могут быть большие спасатели скорости или растраты. Итак, пожалуйста, дайте мне знать, если есть разумный способ ускорить это.

Answer 1

То, что вы, казалось, упустили из виду, это то, что copy_if возвращает итератор, который указывает на конец скопированных данных из операции сжатия потока. Так что все, что требуется, это:

//copies to device
thrust::device_vector<int> d_src = h_src;

//Result vector
thrust::device_vector<int> d_res(d_src.size());

//Copy non-zero elements from d_src to d_res
auto result_end = thrust::copy_if(d_src.begin(), d_src.end(), d_res.begin(), is_not_zero());

//Copy back to host
thrust::host_vector<int> h_res(d_res.begin(), result_end);

Делать это размеры h_res хранить только ненулевые значения и только копирует ненулевые значения из выходных данных сжатия потока. Никаких дополнительных вычислений не требуется.