Могу ли я выделить память на устройстве CUDA для объектов, содержащих массивы чисел с плавающей запятой?

Question

Я работаю над параллельным решением одинаковых обыкновенных дифференциальных уравнений с различными начальными условиями.Я решил эту проблему с 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() << "s\n";

    return 0;
}

//end of file

Answer 1

Вы не можете разыменовать указатель на память устройства в коде хоста:

__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__.)

Answer 2

Различают память на стороне хоста и на стороне устройства

В качестве другого ответа также укажите:

• Память GPU (глобальная), которую вы выделяете с помощью cudaMalloc(), недоступна по коду, работающему на CPU;и
• Системная память (она же memorY хоста), которую вы выделяете в простом C ++ (с std::vector, с std::make_unique, с new и т. д.), недоступна из-за кода, работающего наGPU.

Итак, вам нужно выделить память как на стороне хоста, так и на стороне устройства.Простой пример работы с памятью как на стороне устройства, так и на стороне хоста см. В примере программы CUDA vectorAdd .

_{(На самом деле вы также можете сделать специальный видраспределения, который является доступным как с устройства, так и с хоста, это Unified Memory . Но давайте пока проигнорируем это, поскольку мы имеем дело с основами.)}

Не живите в царстве существительных

В частности, я хочу выделить память на устройстве для поплавков в конструкторе классов, а затем освободить ее в деструкторе.

Я не уверен, что вы действительно хотите это сделать.Похоже, вы придерживаетесь подхода Java-esque, в котором все, что вы делаете, ориентировано на существительное, то есть классы используются для всего: вы не решаете уравнения, у вас есть «решатель уравнений».Вы не «делаете X», у вас есть класс «XDoer» и т. Д. Почему бы просто не иметь (шаблонную) функцию, которая решает систему ODE, возвращая решение?Используете ли вы «решатель» каким-либо другим способом?

(этот момент вдохновлен постом Стива Йегге в блоге, Казнь в царстве существительных .)

Попробуйтечтобы избежать выделения и перераспределения себя

В хорошо написанном современном C ++ мы стараемся избегать прямого, ручного выделения памяти (кстати, это ссылка на Руководство по программированию C ++ Core).).Это правда, что вы освобождаете свою память деструктором, так что это не так уж и плохо, но я бы действительно подумал об использовании std::unique_ptr на хосте и что-то эквивалентное на устройстве (например, cuda::memory::unique_ptrиз моей оболочки API Modern-C ++ CUDA cuda-api-wrappers библиотека);или GPU-ориентированный класс контейнера, такой как thrust вектор устройства.

Проверка на ошибки

Вы действительно должны проверить на ошибки послеВы вызываете функции API CUDA.И это вдвойне необходимо после запуска ядра.Когда вы вызываете код стандартной библиотеки C ++, он выдает исключение при ошибке;API CUDA во время выполнения похож на C и не знает об исключениях.Он просто потерпит неудачу и установит некоторую переменную ошибки, которую нужно проверить.

Итак, вы либо пишете проверки ошибок, как в примере vectorAdd(), с которым я ссылался выше, либо вы получаете какую-то библиотеку для демонстрации более стандартныхбиблиотечно-подобное поведение.cuda-api-wrappers и thrust оба сделают это - на разных уровнях абстракции;как и другие библиотеки / фреймворки.

Answer 3

Вам нужен массив на стороне хоста и один на стороне устройства.

Инициализируйте массив хоста, затем скопируйте его в массив устройств с помощью cudaMemcpy.Разрушение должно быть снова выполнено на стороне хоста.

Альтернативой может быть инициализация массива с устройства, вам нужно будет поставить __device__ перед вашим конструктором, а затем просто использовать malloc.