Как скопировать указатель переменной структуры с хоста на устройство в cuda

Question

У меня есть структура, содержащая некоторые переменные и некоторые переменные указателя. Я хочу скопировать эту структуру с хоста на устройство в 2 разных функциях. В первой функции мне нужно скопировать всю структуру, кроме одной переменной указателя, а затем во второй функции мне нужно скопировать этот оставшийся указатель.

Я могу скопировать всю структуру, но не могу скопировать оставшуюся переменную указателя во второй функции.

#include<iostream>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

struct MultiSGDKernelParam {
  int count;
  size_t sizes;
  float *weights;
  float *mom; 
  float lrs;
};


__global__ void Launch(MultiSGDKernelParam *param, int N, MultiSGDKernelParam *result)
{
  for(int i=0; i<N; i++)
  {
     result[i] =param[i];     
  }
}

MultiSGDKernelParam *fillStructure(float *temp, const int N)
{       
    MultiSGDKernelParam *param;
        param = (MultiSGDKernelParam*) malloc( N * sizeof(MultiSGDKernelParam));
        for( int i=0; i< N ; i++)
        {
            param[i].count = i;
            param[i].sizes =  i*2;
            param[i].lrs =  param[i].sizes - i;
            param[i].weights = &temp[i];
        }

    std::cout<<"Inside the function"<<"\n"; 
        for(int i=0; i< N; i++)
        {
                std::cout<<param[i].sizes<<" ,"<<param[i].lrs<<"\t";
        }

    std::cout<<std::endl;   
        for(int i =0 ; i<N;i++)
        {
          std::cout<<*(param[i].weights)<<"\t";

        }
        std::cout<<std::endl;
    MultiSGDKernelParam *d_param;
        cudaMalloc((void**)&d_param, N  * sizeof(MultiSGDKernelParam));
        cudaMemcpy(d_param,param,N  * sizeof(MultiSGDKernelParam),cudaMemcpyHostToDevice);

    return d_param;

}

MultiSGDKernelParam * fillFullStructure(float *tweight, float *tmom,  const int N )
{
  MultiSGDKernelParam *param = fillStructure( tweight, N );

 /* float *d_mom;

   cudaMalloc((void**)&d_mom,N*sizeof(float));
   cudaCheckErrors("cudaMalloc1 fail");
   cudaMemcpy(d_mom,tmom,N*sizeof(float), cudaMemcpyHostToDevice);
   cudaCheckErrors("cudaMemcpy1 fail");*/
   for( int i=0; i< N ; i++)
        {
          cudaMemcpy(&(param[i].mom),&(tmom[i]),sizeof(float), cudaMemcpyHostToDevice);
      cudaCheckErrors("cudaMempcpy2 fail");
        }

    std::cout<<"Momentum Values copied"<<"\n";
   /*cudaMemcpy(&(param->mom),tmom,N*sizeof(float), cudaMemcpyHostToDevice);
   cudaCheckErrors("cudaMempcpy1fail");*/
   return param;
}



int main()
{
    static const  int N =5;
    float tempweight [N], tempmom[N] ;
    for(int i=0; i< N; i++)
    {
            tempweight[i] = i*3 +1;
        tempmom[i] = i+3;
    }

    MultiSGDKernelParam *result;
    MultiSGDKernelParam *param = fillFullStructure( tempweight,tempmom, N ); 
     const unsigned blocks = 1;
         const unsigned threadsPerBlock = 4;
    cudaMalloc(&result, N  * sizeof(MultiSGDKernelParam));
    Launch<<<blocks,threadsPerBlock>>>(param, N, result);
        cudaDeviceSynchronize();
    MultiSGDKernelParam *paramresult;
    paramresult = (MultiSGDKernelParam*) malloc( N * sizeof(MultiSGDKernelParam));
    cudaMemcpy(paramresult,result, N * sizeof(MultiSGDKernelParam),cudaMemcpyDeviceToHost);
    std::cout<<"Inside Main"<<"\n";
    for(int i=0; i< N; i++)
        {
           std::cout<<paramresult[i].sizes<<" ,"<<paramresult[i].lrs<<"\t";
        }
    std::cout<<std::endl;
    for(int i =0 ; i<N;i++)
    { 
          std::cout<<*(paramresult[i].weights)<<"\t";
          std::cout<<*(paramresult[i].mom)<<"\t";
    }
         std::cout<<std::endl;

    return 0;
}

Выходные данные выдаются как

Inside the function    
0 ,0    2 ,1    4 ,2    6 ,3    8 ,4    
1   4   7   10  13  
Momentum Values copied
Inside Main
0 ,0    2 ,1    4 ,2    6 ,3    8 ,4    
Segmentation fault (core dumped)

Iкод скомпилирован, но выдает ошибку сегментации при печати значений. Если копирование выполнено успешно Если нет, то в чем проблема.

Answer 1

• Я не рекомендую писать ядра CUDA следующим образом:

  __global__ void Launch(MultiSGDKernelParam *param, int N, MultiSGDKernelParam *result)
  {
    for(int i=0; i<N; i++)
    {
       result[i] =param[i];     
    }
  }
  

  Даже если это только для демонстрации, вы должны сделать одну из двух вещей: либо написать ядро ​​какчто (без специализации для потоков CUDA) и запуск только 1 блока из 1 потока (тогда очевидно, что это только для демонстрации), или же использовать правильное индексирование потоков CUDA (например, int i = threadIdx.x+blockDim.x*blockIdx.x;) и избавиться от цикла for,и запустите свой блок с несколькими потоками. В нынешнем виде вы не сделали ни того, ни другого. У вас есть обычный цикл for без специализации, работающий в нескольких потоках. Конечно, это не предмет вашего вопроса, возможно, но такое поведение у вас сейчас означает, что потоки будут наступать друг на друга, когда они пытаются записать в result[i]. Даже если весь остальной код верен, это может затмить понимание того, правильно ли все работает. Мы исправим это, переключив конфигурацию запуска на <<<1,1>>>

• Это:

  param[i].weights = &temp[i];
  

  не может быть правильным. Вы устанавливаете указатель внутри структуры, чтобы он указывал на то, что находится в памяти хоста . (Элемент temp здесь указывает на ваш хост-массив tempweight.) Такой указатель никак нельзя использовать в коде устройства. Это фундаментальный принцип CUDA. Когда вы копируете эту структуру на устройство, числовое значение этого указателя никак не изменится, то есть оно все еще будет указывать на память хоста. Если вы собираетесь использовать этот указатель в любой точке кода устройства, вам нужно научиться работать с операцией глубокого копирования CUDA. И этот ответ проходит этот шаг за шагом. Как это происходит, вы на самом деле не пытаетесь разыменовать этот указатель в коде устройства - вы просто копируете структуры из одного места в другое. Поэтому нам не нужно углубляться в это, чтобы код, который вы показали, работал правильно.

• Проксимальная причина ошибки seg в том, что вы не инициализировали mom член структуры в любом месте вашего кода, но вы пытаетесь разыменовать его здесь:

  std::cout<<*(paramresult[i].mom)<<"\t";
  

  В C или C ++, если вы пытаетесь разыменовать указатель, который вы не инициализировали, вероятно, произойдут плохие вещи. Мы можем исправить это, закомментировав эту строку кода. Мы также можем «исправить» его, скопировав только числовое значение указателя из элемента структуры weights в элемент структуры mom в коде устройства. Однако мы не можем использовать эти указатели непосредственно в коде устройства, поскольку они являются указателями хоста, как указано выше.

В следующем коде указаны первый и третий элементы, указанные выше. Кажется, он работает правильно для меня.

$ cat t1529.cu
#include<iostream>
#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

struct MultiSGDKernelParam {
  int count;
  size_t sizes;
  float *weights;
  float *mom;
  float lrs;
};


__global__ void Launch(MultiSGDKernelParam *param, int N, MultiSGDKernelParam *result)
{
  for(int i=0; i<N; i++)
  {
     result[i] =param[i];
  }
}

MultiSGDKernelParam *fillStructure(float *temp, const int N)
{
    MultiSGDKernelParam *param;
        param = (MultiSGDKernelParam*) malloc( N * sizeof(MultiSGDKernelParam));
        for( int i=0; i< N ; i++)
        {
            param[i].count = i;
            param[i].sizes =  i*2;
            param[i].lrs =  param[i].sizes - i;
            param[i].weights = &temp[i];
        }

    std::cout<<"Inside the function"<<"\n";
        for(int i=0; i< N; i++)
        {
                std::cout<<param[i].sizes<<" ,"<<param[i].lrs<<"\t";
        }

    std::cout<<std::endl;
        for(int i =0 ; i<N;i++)
        {
          std::cout<<*(param[i].weights)<<"\t";

        }
        std::cout<<std::endl;
    MultiSGDKernelParam *d_param;
        cudaMalloc((void**)&d_param, N  * sizeof(MultiSGDKernelParam));
        cudaMemcpy(d_param,param,N  * sizeof(MultiSGDKernelParam),cudaMemcpyHostToDevice);

    return d_param;

}

MultiSGDKernelParam * fillFullStructure(float *tweight, float *tmom,  const int N )
{
  MultiSGDKernelParam *param = fillStructure( tweight, N );

 /* float *d_mom;

   cudaMalloc((void**)&d_mom,N*sizeof(float));
   cudaCheckErrors("cudaMalloc1 fail");
   cudaMemcpy(d_mom,tmom,N*sizeof(float), cudaMemcpyHostToDevice);
   cudaCheckErrors("cudaMemcpy1 fail");*/
   for( int i=0; i< N ; i++)
        {
          cudaMemcpy(&(param[i].mom),&(tmom[i]),sizeof(float), cudaMemcpyHostToDevice);
      cudaCheckErrors("cudaMempcpy2 fail");
        }

    std::cout<<"Momentum Values copied"<<"\n";
   /*cudaMemcpy(&(param->mom),tmom,N*sizeof(float), cudaMemcpyHostToDevice);
   cudaCheckErrors("cudaMempcpy1fail");*/
   return param;
}



int main()
{
    static const  int N =5;
    float tempweight [N], tempmom[N] ;
    for(int i=0; i< N; i++)
    {
            tempweight[i] = i*3 +1;
        tempmom[i] = i+3;
    }

    MultiSGDKernelParam *result;
    MultiSGDKernelParam *param = fillFullStructure( tempweight,tempmom, N );
    const unsigned blocks = 1;
    const unsigned threadsPerBlock = 1;
    cudaMalloc(&result, N  * sizeof(MultiSGDKernelParam));
    Launch<<<blocks,threadsPerBlock>>>(param, N, result);
    cudaDeviceSynchronize();
    MultiSGDKernelParam *paramresult;
    paramresult = (MultiSGDKernelParam*) malloc( N * sizeof(MultiSGDKernelParam));
    cudaMemcpy(paramresult,result, N * sizeof(MultiSGDKernelParam),cudaMemcpyDeviceToHost);
    std::cout<<"Inside Main"<<"\n";
    for(int i=0; i< N; i++)
        {
           std::cout<<paramresult[i].sizes<<" ,"<<paramresult[i].lrs<<"\t";
        }
    std::cout<<std::endl;
    for(int i =0 ; i<N;i++)
    {
          std::cout<<*(paramresult[i].weights)<<"\t";
        //  std::cout<<*(paramresult[i].mom)<<"\t";
    }
         std::cout<<std::endl;

    return 0;
}

$ nvcc -o t1529 t1529.cu
$ cuda-memcheck ./t1529
========= CUDA-MEMCHECK
Inside the function
0 ,0    2 ,1    4 ,2    6 ,3    8 ,4
1       4       7       10      13
Momentum Values copied
Inside Main
0 ,0    2 ,1    4 ,2    6 ,3    8 ,4
1       4       7       10      13
========= ERROR SUMMARY: 0 errors
$

Если вы действительно хотите использовать структурные элементы (указатели) weights и mom в коде устройства, вам нужно начать пытаться понять глубокийОперация копирования в CUDA. Я уже дал вам ссылку, которая объясняет процесс, шаг за шагом, с работающим примером. Прямо сейчас вы не указали в своем коде, что вы реализовали что-либо из этого, и написание кода для вас выходит за рамки того, что я собираюсь здесь ответить, поскольку вы не пытались это сделать.