cudaMalloc在结构内创建一个数组会创建“非法内存访问”

Question

下面是我要解决的问题的简化版本。这两个代码段都可以编译，但是＃2抛出“非法内存访问”。基本上，如果将数组封装在结构中，则将指向该结构的指针传递给cudaMalloc会产生各种问题-至少是我这样做的方法。我很确定这是由于以下代码中的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;

}

Answer 1

回应Robert Crovella的说明：如初始线程中所述，在GPU上分配dum也会引发非法的内存访问错误。这是我的意思：

cudaMalloc( &dum , sizeof(dummy) * 1 );
cudaMalloc( &dum->arr , sizeof(double) * n );

test<<< 1 , 1 >>>( 123.0 , dum , n );
gpu_err_chk( cudaDeviceSynchronize() );

cudaFree( dum->arr );
cudaFree( dum );

因此，该问题似乎（至少不仅如此）并非源于dum是初始代码中的主机指针的事实。不知道我在做什么错...

UPDATE：我在前面的文章中仔细研究了Robert的一些示例，以下代码现已生效：

dum_h  = new dummy;
dum_h->arr = new double[n];

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               

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

尽管我仍然不完全了解此解决方案：1）创建临时数组tmp只是为了创建一些可用于访问设备上的dum_d的有效设备地址，对吗？

2]此解决方案需要创建3倍于仅主机应用程序所需的内存量：主机（dummy）上有1个完全分配的结构dum_h，设备上有1个[[0] C0]）和设备dum_d上的1个double数组。这似乎太疯狂了……尤其是在我的实际应用中，几乎完全使用了设备内存。

3）释放此内存似乎有点复杂...我尝试了明显的方法：

tmp

但是这会引发分段错误。