我可以在CUDA设备上为包含浮点数数组的对象分配内存吗

Can I allocate memory on CUDA device for objects containing arrays of float numbers?

本文关键字:数组 对象 分配 内存 浮点数 包含 CUDA 我可以      更新时间:2023-10-16

我正在研究具有不同初始条件的相同常微分方程的并行求解。我已经用OpenMP解决了这个问题,现在我想在GPU上实现类似的代码。具体来说,我想在类构造函数中为浮点分配设备上的内存,然后在析构函数中解除分配。这对我来说不起作用,因为我的可执行文件"被信号SIGSEGV(地址边界错误)终止"。CUDA中是否可以使用类、构造函数和析构函数?

顺便说一句,我是CUDA的新手,在C++方面也不是很有经验。

我附上了代码,以防我对我的问题描述不好。

#include <cmath>
#include <iostream>
#include <fstream>
#include <iomanip>
#include <random>
#include <string>
#include <chrono>
#include <ctime>
using namespace std;
template<class ode_sys>
class solver: public ode_sys 
{
public:
int *nn;
float *t,*tt,*dt,*x,*xx,*m0,*m1,*m2,*m3;
using ode_sys::rhs_sys;
__host__ solver(int n): ode_sys(n)
{ //here I try to allocate memory. It works malloc() and doesn't with cudaMalloc() 
size_t size=sizeof(float)*n;
cudaMalloc((void**)&nn,sizeof(int));
*nn=n;
cudaMalloc((void**)&t,sizeof(float));
cudaMalloc((void**)&tt,sizeof(float));
cudaMalloc((void**)&dt,sizeof(float));
cudaMalloc((void**)&x,size);
cudaMalloc((void**)&xx,size);
cudaMalloc((void**)&m0,size);
cudaMalloc((void**)&m1,size);
cudaMalloc((void**)&m2,size);
cudaMalloc((void**)&m3,size);
}
__host__ ~solver()
{
cudaFree(nn);
cudaFree(t);
cudaFree(tt);
cudaFree(dt);
cudaFree(x);
cudaFree(xx);
cudaFree(m0);
cudaFree(m1);
cudaFree(m2);
cudaFree(m3);
}
__host__ __device__ void rk4()
{//this part is not important now. 
}
};
class ode 
{
private:
int *nn;
public:
float *eps,*d;
__host__ ode(int n)
{
cudaMalloc((void**)&nn,sizeof(int));
*nn=n;
cudaMalloc((void**)&eps,sizeof(float));
size_t size=sizeof(float)*n;
cudaMalloc((void**)&d,size);
}
__host__ ~ode()
{
cudaFree(nn);
cudaFree(eps);
cudaFree(d);
}
__host__ __device__ float f(float x_,float y_,float z_,float d_)
{
return d_+*eps*(sinf(x_)+sinf(z_)-2*sinf(y_));
}
__host__ __device__ void rhs_sys(float *t,float *dt,float *x,float *dx)
{
}
};
//const float pi=3.14159265358979f;
__global__ void solver_kernel(int m,int n,solver<ode> *sys_d)
{
int index = threadIdx.x;
int stride = blockDim.x;
//actually ode numerical evaluation should be here
for (int l=index;l<m;l+=stride)
{//this is just to check that i can run kernel
printf("%d Hello n", l);
}
}
int main ()
{
auto start = std::chrono::system_clock::now();
std::time_t start_time = std::chrono::system_clock::to_time_t(start);
cout << "started computation at " << std::ctime(&start_time);
int m=128,n=4,l;// i want to run 128 threads, n is dimension of ode
size_t size=sizeof(solver<ode>(n));
solver<ode> *sys_d;   //an array of objects
cudaMalloc(&sys_d,size*m);    //nvprof shows that this array is allocated
for (l=0;l<m;l++)
{
new (sys_d+l) solver<ode>(n);   //it doesn't work as it meant to
}
solver_kernel<<<1,m>>>(m,n,sys_d);   
for (l=0;l<m;l++)
{
(sys_d+l)->~solver<ode>();    //it doesn't work as it meant to
}
cudaFree(sys_d);    //it works
auto end = std::chrono::system_clock::now();
std::chrono::duration<double> elapsed_seconds = end-start;
std::time_t end_time = std::chrono::system_clock::to_time_t(end);
std::cout << "finished computation at " << std::ctime(&end_time) << "elapsed time: " << elapsed_seconds.count() << "sn";
return 0;
}
//end of file

区分主机端和设备端内存

正如其他答案所述:

  • cudaMalloc()分配的GPU(全局)内存无法通过CPU上运行的代码访问;以及
  • 您在纯C++(使用std::vectorstd::make_uniquenew等)中分配的系统内存(也称为主机内存Y)是GPU上运行的代码无法访问的

因此,您需要同时分配主机端和设备端内存。有关同时使用设备端和主机端内存的简单示例,请参阅CUDAvectorAdd示例程序。

(实际上,您也可以进行一种特殊类型的分配,可以从设备和主机访问;这是统一内存。但现在我们忽略它,因为我们正在处理基本问题。)

不要生活在名词的王国里

具体来说,我想在类构造函数中为浮点分配设备上的内存,然后在析构函数中解除分配。

我不确定你是否真的想这么做。你似乎采取了一种更Java风格的方法,在这种方法中,你所做的一切都是以名词为中心的,即类用于一切:你不求解方程,你有一个"方程求解器"。你没有"做X",你有一个"XDoer"类等等。为什么不只是有一个(模板化的)函数来解决ODE系统,并返回解决方案呢?您是否以其他方式使用"解算器"?

(这一点的灵感来自Steve Yegge的博客文章《努恩斯王国的处决》。)

尽量避免自己分配和取消分配

在编写良好的现代C++中,我们尽量避免直接手动分配内存(顺便说一句,这是C++核心编程指南的链接)。现在,确实可以使用析构函数来释放内存,所以也没那么糟糕,但我真的会考虑在主机上使用std::unique_ptr,在设备上使用等效的东西(比如我的Modern-C++CUDAneneneba API包装cuda-api-wrappers库中的cuda::memory::unique_ptr);或者类似于CCD_ 9的设备向量的面向GPU的容器类。

检查错误

在调用CUDA API函数后,您确实必须检查错误。在启动内核之后,这是双重必要的。当您调用C++标准库代码时,它会在出现错误时抛出异常;CUDA的运行时API是类似C的,并且不知道异常。它只会失败并设置一些需要检查的错误变量。

因此,您要么编写错误检查,就像我在上面链接的vectorAdd()示例中一样,要么获得一些库来展示更标准的类似库的行为。CCD_ 11和CCD_;其他库/框架也是如此。

您需要一个位于主机端的阵列和一个位于设备端的阵列。

初始化主机阵列,然后使用cudaMemcpy将其复制到设备阵列。必须再次在主机端进行销毁。

另一种选择是从设备初始化数组,您需要将__device__放在构造函数前面,然后只使用malloc

您不能在主机代码中取消引用指向设备内存的指针:

__host__ ode(int n)
{
cudaMalloc((void**)&nn,sizeof(int));
*nn=n; // !!! ERROR
cudaMalloc((void**)&eps,sizeof(float));
size_t size=sizeof(float)*n;
cudaMalloc((void**)&d,size);
}

您将不得不使用cudaMemcpy复制这些值。(或者使用__global__函数的参数。)