在 __host__ __device__ 仿函数中创建 Thrust::device_vectors
creating Thrust::device_vectors in a __host__ __device__ functor
我目前正在尝试并行化当前在主函数中按顺序运行的推力 cuda 代码(因此不利用 GPU 的强大功能)。我本质上已经采用了功能代码并将其放入 thrust::for_each 可以使用 cuda 流调用的仿函数中。但是,如果我使用
定义仿函数
__host__ __device__
VS2013 抛出各种警告说我正在尝试从设备启动主机功能。这些错误发生在我使用
定义向量的地方
thrust::device_vector vect (size_vector);
以及一些 thrust::transform 功能。它特别引用了 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__
函数,其中包含单独的主机和设备定义代码。也有可能我在研究这个主题时遗漏了一些东西。任何建议将不胜感激。
These errors are occurring in places were I am defining a vector
正如编译器清楚地告诉您的那样,问题在于 thrust::vector
中定义的构造函数和所有运算符当前是 仅主机 函数。试图在 __device__
函数中使用它们是非法的。
除了不尝试在设备代码中实例化向量外,没有其他解决方案。
Thrust 为其所有算法提供主机和设备路径,但算法只能从主机启动。
在编译时,Thrust 会查看迭代器的类型以确定要构建的路径。如果它构建了一个设备路径,那么同样的限制适用于常规 CUDA 代码,其中之一是设备代码不能调用主机上的函数。
因此,像 thrust::sort()
这样的语句会启动一个算法,并且只能存在于主机代码中。在编译时,检查传递给 sort()
的迭代器,并且使用 Thrust 模板构建处理您的特定类型的 sort()
的主机或设备版本。如果构建了一个设备版本并且它采用了一个仿函数,那么它也必须可以构建该仿函数的设备版本,这意味着该仿函数不能包含启动新算法的 Thrust 语句。
在运行时,设备版本的声明如 thrust::sort()
将启动一个或多个 CUDA 内核,因此您可能想要研究的是 Thrust 将不同算法组合到同一个内核中的能力,Thrust调用内核融合。有几种方法可以做到这一点,其中之一是使用转换迭代器。有关详细信息,请参阅 Thrust 文档。
我目前正在尝试并行化当前在主函数中按顺序运行的推力 cuda 代码(因此不利用 GPU 的强大功能)。我本质上已经采用了功能代码并将其放入 thrust::for_each 可以使用 cuda 流调用的仿函数中。但是,如果我使用
定义仿函数__host__ __device__
VS2013 抛出各种警告说我正在尝试从设备启动主机功能。这些错误发生在我使用
定义向量的地方thrust::device_vector vect (size_vector);
以及一些 thrust::transform 功能。它特别引用了 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__
函数,其中包含单独的主机和设备定义代码。也有可能我在研究这个主题时遗漏了一些东西。任何建议将不胜感激。
These errors are occurring in places were I am defining a vector
正如编译器清楚地告诉您的那样,问题在于 thrust::vector
中定义的构造函数和所有运算符当前是 仅主机 函数。试图在 __device__
函数中使用它们是非法的。
除了不尝试在设备代码中实例化向量外,没有其他解决方案。
Thrust 为其所有算法提供主机和设备路径,但算法只能从主机启动。
在编译时,Thrust 会查看迭代器的类型以确定要构建的路径。如果它构建了一个设备路径,那么同样的限制适用于常规 CUDA 代码,其中之一是设备代码不能调用主机上的函数。
因此,像 thrust::sort()
这样的语句会启动一个算法,并且只能存在于主机代码中。在编译时,检查传递给 sort()
的迭代器,并且使用 Thrust 模板构建处理您的特定类型的 sort()
的主机或设备版本。如果构建了一个设备版本并且它采用了一个仿函数,那么它也必须可以构建该仿函数的设备版本,这意味着该仿函数不能包含启动新算法的 Thrust 语句。
在运行时,设备版本的声明如 thrust::sort()
将启动一个或多个 CUDA 内核,因此您可能想要研究的是 Thrust 将不同算法组合到同一个内核中的能力,Thrust调用内核融合。有几种方法可以做到这一点,其中之一是使用转换迭代器。有关详细信息,请参阅 Thrust 文档。