如何强制仿函数看到整个 thrust::vector 以便可以进行排序？

Question

我是 CUDA 的新手，在使用函子时遇到了一些麻烦。我正在尝试将 thrust::vector 中的 thrust::vector 输入到函子中。目前我可以输入一个向量并对每个元素和 return 使用 thrust::for_each 修改后的向量做一些事情，但是如果我想在一个函子中对一个向量进行排序，我需要能够输入一次整个向量，这样函子就可以作为一个整体作用于它。有没有办法做到这一点？

下面的代码可以编译，但不会 return 对向量进行排序。

#include <iostream>
#include <fstream>
#include <sstream>
#include <string>
#include <vector>
#include <iterator>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/reduce.h>
#include <thrust/transform_reduce.h>
#include <thrust/transform.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
#include <thrust/system/cuda/execution_policy.h>
#include <thrust/tuple.h>
#include <thrust/count.h>
#include <thrust/sequence.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/for_each.h>
#include <ctime>
#include <cstdio>
#include <cassert> 

using namespace std;
template<typename T>
struct sort_vector
{
    __host__ __device__ thrust::device_vector<float> operator()    (thrust::tuple<thrust::device_vector<float>, thrust::device_vector<float>>     x)
    {
    thrust::device_vector<float> y = thrust::get<0>(x);
    thrust::sort(y.begin(), y.end());
    return thrust::get<1>(x) = y;
    }
};

int main() {
    thrust::device_vector<float> d_fraction(5);
    d_fraction[0] = 1;
    d_fraction[1] = 5;
    d_fraction[2] = 3;
    d_fraction[3] = 2;
    d_fraction[4] = 4;

    cout << "original" << endl;
    int f = 0;
    while (f < 5){
        cout << d_fraction[f] << endl;
        f++;
    }

    cudaStream_t s1;
    cudaStreamCreate(&s1);
    thrust::device_vector<float> result1(5);

    thrust::for_each(thrust::cuda::par.on(s1), 
    thrust::make_zip_iterator(thrust::make_tuple(d_fraction.begin(), result1.begin())),
    thrust::make_zip_iterator(thrust::make_tuple(d_fraction.end(), result1.end())), sort_vector<thrust::device_vector<float>>());

    cudaStreamSynchronize(s1);

    cout << "sorted" << endl;
    int d = 0;
    while (d < 5){
        cout << Sresult2[d] << endl;
        d++;
    }

    cudaStreamDestroy(s1);
    return 0;
}

然而，当我尝试使用诸如

的参考时

 _host__ __device__ thrust::device_vector<float> operator()    (thrust::tuple<thrust::device_vector<float> &, thrust::device_vector<float> &>     x)

代码不再编译。

我是否需要为向量转换一个引用指针，以便仿函数可以看到整个向量？或者问题是否可能是我按值传递向量，而我不知道有一种不同的方式将向量传递给仿函数？

Answer 1

仿函数通常在单个线程的上下文中运行。如果使用 CUDA 后端，我们谈论的是 CUDA 线程。

对向量进行排序的典型方法是直接在向量上使用 thrust::sort。在最普通的用法中，根本不需要函子定义。

如果您想对向量 "within a functor" 进行排序，则需要将指向该向量的指针传递给函子，然后让函子处理它。

推力设备代码（在仿函数上下文中执行的代码）通常不能直接处理 thrust::device_vector 等结构。目前也不可能构建设备向量的设备向量。

因此，我已将您的代码修改为可行且排序的代码 "within a functor"。我选择将要分类的向量连接成一个向量。我们把这个vector的地址传给排序函子，然后每个线程计算出自己要排序的范围，传给thrust::sort在线程中进行顺序排序：

$ cat t1211.cu
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
#include <thrust/for_each.h>
#include <thrust/sequence.h>
#include <cstdlib>

const int num_segs = 3;  // number of segments to sort
const int num_vals = 5;  // number of values in each segment
const int range = 100;   // range of values

using namespace std;

template <typename T>
struct sort_vector
{
    T *data;
    sort_vector(T *_data) : data(_data) {};
    __host__ __device__ void operator()(int idx)
    {
    thrust::sort(thrust::seq, data+idx*num_vals, data+((idx+1)*num_vals));
    }
};

int main() {
    thrust::device_vector<float> d_data(num_segs*num_vals);
    for (int i = 0; i < num_segs*num_vals; i++)
      d_data[i] = rand()%range;

    cout << "original" << endl;
    int f = 0;
    while (f < num_segs*num_vals){
        cout << d_data[f] << endl;
        f++;
    }
    thrust::device_vector<int> d_idxs(num_segs);
    thrust::sequence(d_idxs.begin(), d_idxs.end());
    cudaStream_t s1;
    cudaStreamCreate(&s1);
    //thrust::device_vector<float> result1(5);

    thrust::for_each(thrust::cuda::par.on(s1),
    d_idxs.begin(),
    d_idxs.end(), sort_vector<float>(thrust::raw_pointer_cast(d_data.data())));

    cudaStreamSynchronize(s1);

    cout << "sorted" << endl;
    int d = 0;
    while (d < num_segs*num_vals){
        cout << d_data[d] << endl;
        d++;
    }

    cudaStreamDestroy(s1);
    return 0;
}
$ nvcc -o t1211 t1211.cu
$ ./t1211
original
83
86
77
15
93
35
86
92
49
21
62
27
90
59
63
sorted
15
77
83
86
93
21
35
49
86
92
27
59
62
63
90
$

在这种情况下，正如 thrust::seq 所表明的那样，每个线程中正在完成的工作是按顺序完成的。（总的来说，这里的线程是并行运行的，但它们并不合作——每个线程都在处理一个独立的问题）。

这不是唯一可能的解决方案。您可能对 question/answer 感兴趣，其中包含各种其他相关想法。

明确地说，我认为您在这里讨论的是 "vectorized"（或分段）排序。这不是最快的方法，但我正在尝试向您展示一些可行的概念，作为对您所展示内容的直接扩展，以回答您的问题（"How to force a functor to see an entire thrust::vector so that sorting is possible?"）讨论了一种更快的矢量化排序方法在上面链接 question/answer.

如何强制仿函数看到整个 thrust::vector 以便可以进行排序？

How to force a functor to see an entire thrust::vector so that sorting is possible?

c++

cuda

functor

thrust