cudaMallo c массив в структуре создает «недопустимый доступ к памяти» - PullRequest
Новая
       3

cudaMallo c массив в структуре создает «недопустимый доступ к памяти»

0 голосов
/ 23 января 2020

ниже - упрощенная версия проблемы, которую я пытаюсь решить. Оба фрагмента кода компилируются, но # 2 создает «незаконный доступ к памяти». По сути, если массив инкапсулирован в структуру, передача указателя на эту структуру в cudaMallo c создает всевозможные проблемы - по крайней мере, так, как я это делаю. Я почти уверен, что это связано с тем, что адрес dum в приведенном ниже коде находится на хосте и поэтому не доступен внутри ядра. Проблема в том, что я не знаю, как создать версию устройства dum ... Например, использование cudaMalloc( (void**)&dum , sizeof(dummy) * 1 ) вместо синтаксиса new dummy ниже не решает проблему. Я думаю, что я путаюсь с двойным указателем, используемым cudaMalloc.

Конечно, в этом примере может показаться глупым инкапсулировать массив double в структуру, в фактический код, который мне действительно нужно сделать это хотя. 

struct dummy
{
  double *arr;
};



void allocate( dummy *dum , int n )
{
  cudaMalloc( (double**)&(dum->arr) , sizeof(double) * n );
}



__global__ void test( double val , dummy *dum , int n )
{
  printf( "test\n" );
  for( int ii = 0 ; ii < n ; ii++ )
    dum->arr[ii] = val;
}


__global__ void test2( double val , double *arr , int n )
{
  printf( "test\n" );
  for( int ii = 0 ; ii < n ; ii++ )
    arr[ii] = val;
}


int main()
{

  int n = 10;

  dummy *dum = new dummy;


  /* CODE 1: the piece of code below works */
  double *p;
  gpu_err_chk( cudaMalloc( &p , sizeof(double) * n ) );
  test2<<< 1 , 1 >>>( 123.0 , p , n );
  gpu_err_chk( cudaDeviceSynchronize() );


  /* CODE 2: the piece of code below does not... */
  allocate( dum , n );
  test<<< 1 , 1 >>>( 123.0 , dum , n );
  gpu_err_chk( cudaDeviceSynchronize() );

  return 1;

}

1 Ответ

0 голосов
/ 23 января 2020

Прочитав несколько примеров в предыдущих постах Роберта, я смог переписать код так, чтобы он работал: 

struct dummy
{
  double *arr;
};



__global__ void test( dummy *dum , int n )
{
  printf( "test\n" );
  for( int ii = 0 ; ii < n ; ii++ )
    printf( "dum->arr[%d] = %f\n" , ii , dum->arr[ii] );

}



int main()
{

  int n = 10;

  dummy *dum_d , *dum_h;

  srand( time(0) );

  dum_h  = new dummy;
  dum_h->arr = new double[n];
  for( int ii = 0 ; ii < n ; ii++  ){
    dum_h->arr[ii]  = double( rand() ) / RAND_MAX;
    printf( "reference data %d = %f\n" , ii , dum_h->arr[ii] );
  }

  cudaMalloc( &dum_d , sizeof(dummy) * 1 );
  cudaMemcpy( dum_d , dum_h , sizeof(dummy) * 1 , cudaMemcpyHostToDevice );

  double *tmp;
  cudaMalloc( &tmp , sizeof(double) * n );
  cudaMemcpy( &( dum_d->arr ) , &tmp , sizeof(double*) , cudaMemcpyHostToDevice );  // copy the pointer (host) to the device structre to a device pointer               
  cudaMemcpy( tmp , dum_h->arr , sizeof(double) * n , cudaMemcpyHostToDevice );

  delete [] dum_h->arr;
  delete dum_h;

  test<<< 1 , 1 >>>( dum_d , n );
  gpu_err_chk( cudaDeviceSynchronize() );

  cudaFree( tmp );
  cudaFree( dum_d );

  return 1;

}

Однако я все еще не понимаю, почему это работает. У кого-нибудь есть визуальная схема того, что происходит? Я теряюсь с разными указателями ...

Более того, есть одна вещь, которая действительно поражает меня: я могу освободить tmp прямо перед запуском ядра, и код все еще работает, то есть: 

  cudaFree( tmp );

  test<<< 1 , 1 >>>( dum_d , n );
  gpu_err_chk( cudaDeviceSynchronize() );

Как это так? На мой взгляд (явно неправильно), массив устройств, содержащий случайные значения, пропал ...

Еще одна путаница в том, что я не могу освободить dum_d->arr напрямую (freeCuda(dum_d->arr)), это бросает ошибка сегментации.

...