c++ - 我可以在 CUDA 设备上为包含 float 数组的对象分配内存吗？

Question

我正在研究并行求解具有不同初始条件的相同常微分方程。我已经用 OpenMP 解决了这个问题，现在我想在 GPU 上实现类似的代码。具体来说，我想在类构造函数中为 float 在设备上分配内存，然后在析构函数中释放它。它对我不起作用，因为我的可执行文件“被信号 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() << "s\n";

    return 0;
}

//end of file

Answer 1

区分主机端和设备端内存

正如其他答案所说:

• 您使用 cudaMalloc() 分配的 GPU(全局)内存无法通过在 CPU 上运行的代码访问；和
• 您在普通 C++ 中分配的系统内存(又名主机内存)(使用 std::vector，使用 std::make_unique，使用 new 等)在 GPU 上运行的代码无法访问。

因此，您需要同时分配主机端和设备端内存。有关同时使用设备端和主机端内存的简单示例，请参阅 CUDA vectorAdd sample program .

_{(实际上，您还可以进行一种特殊类型的分配，可以从设备和主机访问它；这是 Unified Memory。但是现在让我们忽略它，因为我们正在处理基础知识。)}

不要活在名词的王国里

Specifically, I want to allocate memory on device for floats in class constructor and then deallocate it in destructor.

我不确定您是否真的想这样做。您似乎采用了一种更像 Java 的方法，在这种方法中，您所做的一切都是以名词为中心的，即类用于一切:您不求解方程，您有一个“方程求解器”。你不“做 X”，你有一个“XDoer”类等。为什么不只是有一个(模板化的)函数来解决 ODE 系统，返回解决方案？您是否以任何其他方式使用您的“求解器”？

(这一点的灵感来自 Steve Yegge 的博文，Execution in the Kingdom of Nouns。)

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

在编写良好的现代 C++ 中，我们尝试 avoid direct, manual allocation of memory (顺便说一句，这是指向 C++ 核心编程指南的链接)。现在，您确实可以使用析构函数释放内存，所以还不错，但我真的会考虑使用 std::unique_ptr在主机上和设备上的等效项(例如来 self 的现代 C++ CUDA API 包装器 cuda-api-wrappers 库的 cuda::memory::unique_ptr)；或面向 GPU 的容器类，如 thrust的设备 vector 。

检查错误

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

所以，要么你编写错误检查，就像我在上面链接的 vectorAdd() 示例中那样，要么你得到一些库来展示更多类似于标准库的行为。 cuda-api-wrappers 和 thrust 都会这样做——在不同的抽象层次上；其他库/框架也是如此。