c++ - 使用 Thrust 进行流压缩；最佳实践和最快方法？

Question

我有兴趣移植一些现有代码以使用推力，看看是否可以相对轻松地在 GPU 上加速。

我想要完成的是流压缩操作，其中仅保留非零元素。根据下面的示例代码，我基本上可以正常工作。我不确定如何解决的部分是在压缩发生后处理 d_res 以及 h_res 中的所有额外填充空间。

该示例仅使用 0-99 序列，并将所有偶数条目设置为零。这只是一个示例，真正的问题将是一般的稀疏数组。

这里的答案对我帮助很大，尽管在读取数据时，大小已知是恒定的: How to quickly compact a sparse array with CUDA C?

我怀疑我可以通过计算 d_src 中 0 的数量，然后仅将 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;
}

另外，我是一个推力新手，所以如果上面的代码有任何明显的缺陷，不符合使用推力的推荐做法，请告诉我。

同样，速度始终是人们关注的焦点。阅读了一些不同的推力教程，似乎这里的变化很小，但可能会节省或浪费很大的速度。因此，请告诉我是否有一种聪明的方法可以加快速度。

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 仅保存非零值，并且仅从流压缩的输出中复制非零值。不需要额外的计算。