在 __host__ __device__ 函子中创建 Thrust::device_vectors

Question

我目前正在尝试并行化当前在主函数中顺序运行的推力 cuda 代码（因此不利用 GPU 的功能）。我基本上已经将功能代码放入一个仿函数中，thrust::for_each 可以使用 cuda 流调用。但是，如果我使用

定义函子

__host__ __device__ 

VS2013 抛出各种警告，说我正在尝试从设备启动主机功能。这些错误发生在我使用定义向量的地方

thrust::device_vector vect (size_vector); 

以及一些推力::转换函数。它特别引用了thrust::device_malloc_allocator 的问题。如果我将函子严格定义为 host 函子，这些错误都会消失，但是当我使用分析器时，很明显只有 0.01% 的设备被使用，这让我相信 for_each 实际上并不是在函子中启动推力代码。

编辑 下面是一些编译并显示此错误的代码

#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>
#include <ctime>
#include <vector>
#include <algorithm>
#include <memory.h>
#include <cstdio>
#include <thread>
#include <thrust/copy.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/reduce.h>


using namespace std;

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


template <typename T> 
struct sort_vector
{
    T *Ddata;
    T *vect3;
    T *answer;

    sort_vector(T *_Ddata, T *_vect3, float *a) : Ddata(_Ddata), vect3(_vect3), answer(a) {};


    __host__ __device__ void operator()(int idx)
    {
        thrust::sort(thrust::seq, Ddata + idx*num_vals, Ddata + ((idx + 1)*num_vals));
        thrust::device_ptr<float> vect3_ptr = thrust::device_pointer_cast(vect3);
        thrust::device_vector<float> vect(10, 1);
        thrust::device_vector<float> vect2(10, 3);
        thrust::transform(thrust::device, vect.begin(), vect.end(), vect2.begin(), vect3_ptr, thrust::minus<float>());
        *answer = thrust::reduce(thrust::device, Ddata + idx*num_vals, Ddata + ((idx + 1)*num_vals));

    }
};

int main() {

    thrust::device_vector<float> d_Ddata(num_segs*num_vals);
    d_Ddata[0] = 50;
    d_Ddata[1] = 9.5;
    d_Ddata[2] = 30;
    d_Ddata[3] = 8.1;
    d_Ddata[4] = 1;

    thrust::device_vector<float> d_Ddata2(num_segs*num_vals);
    d_Ddata2[0] = 50;
    d_Ddata2[1] = 20.5;
    d_Ddata2[2] = 70;
    d_Ddata2[3] = 8.1;
    d_Ddata2[4] = 1;

    thrust::device_vector<float> vect3(10, 0);
    thrust::device_vector<float> vect4(10, 0);

    cout << "original dut" << endl;
    int g = 0;
        while (g < num_segs*num_vals){
            cout << d_Ddata[g] << endl;
            g++;
        }

        thrust::device_vector<int> d_idxs(num_segs);
        thrust::sequence(d_idxs.begin(), d_idxs.end());

        thrust::device_vector<float> dv_answer(1);
        thrust::device_vector<float> dv_answer2(1);
        cudaStream_t s1, s2;
        cudaStreamCreate(&s1);
        cudaStreamCreate(&s2);

        clock_t start;
        double duration;
        start = clock();

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

        thrust::for_each(thrust::cuda::par.on(s2),
            d_idxs.begin(),
            d_idxs.end(), sort_vector<float>(thrust::raw_pointer_cast(d_Ddata2.data()), thrust::raw_pointer_cast(vect4.data()), thrust::raw_pointer_cast(dv_answer2.data())));

        cudaStreamSynchronize(s1);
        cudaStreamSynchronize(s2);

        cout << "sorted dut" << endl;
        int n = 0;
        while (n < num_segs*num_vals){
            cout << d_Ddata[n] << endl;
            n++;
        } 
        cout << "sum" << endl;
        cout << dv_answer[0] << endl;
        cout << dv_answer2[0] << endl;

        cout << "vector subtraction" << endl;
        int e = 0;
        while (e < 10){
            cout << vect3[e] << endl;
            e++;
        }

        cudaStreamDestroy(s1);
        cudaStreamDestroy(s2);

        duration = (clock() - start) / (double)CLOCKS_PER_SEC;
        cout << "time " << duration << endl;

        cin.get();
        return 0;
    }

是否有可能thrust::for_each 不能调用__host__ 函子？

某些推力呼叫是否与幕后的主机天生相关？

我能看到的唯一可能的解决方法是创建一个__host__ __device__ fucntor，其中包含单独的主机和设备定义代码。在研究这个主题时，我也可能遗漏了一些东西。任何建议将不胜感激。

Answer 1

这些错误发生在我定义向量的地方

正如编译器清楚地告诉你的那样，问题在于thrust::vector 中定义的构造函数和所有运算符目前是仅限主机 函数。尝试在 __device__function 中使用它们是非法的。

除了不尝试在设备代码中实例化向量之外，没有其他解决方案。

Answer 2

Thrust 为其所有算法提供主机和设备路径，但算法只能从主机启动。

在编译时，Thrust 查看迭代器的类型以确定要构建的路径。如果它构建设备路径，则与常规 CUDA 代码相同的限制适用，其中之一是设备代码不能调用主机上的函数。

因此，像thrust::sort() 这样的语句会启动一个算法，并且只能存在于主机代码中。在编译时，传递给sort() 的迭代器被检查并且推力模板用于构建处理您的特定类型的sort() 的主机或设备版本。如果构建了一个设备版本并且它需要一个仿函数，那么它也必须可以构建一个仿函数的设备版本，这意味着仿函数不能包含启动新算法的 Thrust 语句。

在运行时，像 thrust::sort() 这样的语句的设备版本将启动一个或多个 CUDA 内核，因此您可能想要研究的是 Thrust 将单独的算法组合到同一个内核中的能力，这被 Thrust 称为内核融合。有几种方法可以做到这一点，其中一种是使用转换迭代器。有关详细信息，请参阅 Thrust 文档。