结构中指向数组的 cudaFree 指针上的 CUDA 分段错误

Posted

技术标签:

【中文标题】结构中指向数组的 cudaFree 指针上的 CUDA 分段错误【英文标题】:CUDA Segmentation fault on cudaFree pointer to array in struct 【发布时间】:2021-05-22 20:44:05 【问题描述】:

我在 CUDA 设备上得到了一个结构,它包含一个指向数组的指针。计算、访问元素和一切正常,但是当我尝试成为一个好孩子并打电话时

cudaFree(my_struct->pointer_to_array)

我遇到了分段错误。 cudaFree(my_struct) 但是工作得很好。有什么我遗漏的吗?

请找到以下最小示例:

#include <stdio.h>

#include <cuda.h>
#include <cuda_runtime.h>
#include <cassert>

typedef struct 
  int n;
  float *arr;
 DummyStruct;

__global__ void check(DummyStruct *d) 
  printf("EL %f", d->arr[0]);


int main() 
  cudaError_t status;

  // create host pointer to dummy struct
  DummyStruct *dummy;
  dummy = (DummyStruct *)malloc(sizeof(DummyStruct));

  int arr_size = 32;

  dummy->n = 0;
  float *arr = (float *) malloc(sizeof(float) * arr_size);

  for (int i=0; i < 32; i++) 
    arr[i] = i;
  

  // allocate device array
  float *d_arr;
  status = cudaMalloc(&d_arr, arr_size * sizeof(float));
  assert( status == cudaSuccess );

  status = cudaMemcpy(d_arr, arr, arr_size * sizeof(float), cudaMemcpyHostToDevice);
  assert( status == cudaSuccess );

  free(arr);

  // for some reason this should happen here and not d_sp->coeff = d_coeff ...
  dummy->arr = d_arr;

  // allocate and ship struct to device
  DummyStruct* d_dummy;
  status = cudaMalloc(&d_dummy, sizeof(DummyStruct));
  assert( status == cudaSuccess );

  status = cudaMemcpy(d_dummy, dummy, sizeof(DummyStruct), cudaMemcpyHostToDevice);
  assert( status == cudaSuccess );

  // free host struct
  free(dummy);


  // check whether array access works
  check<<<1, 1>>>(d_dummy);


  // THIS causes Segmentation fault (core dumped)
  status = cudaFree(d_dummy->arr);
  assert( status == cudaSuccess );

  status = cudaFree(d_dummy);
  assert( status == cudaSuccess );

【问题讨论】:

【参考方案1】:

此声明:

status = cudaFree(d_dummy->arr);

需要在 host 代码中取消引用 device 指针(d_dummy - 使用设备分配器分配,即 cudaMalloc)。这在 CUDA 中是非法的。

既然您已经知道 (d_dummy-&gt;arr) == d_arr,那么释放嵌入指针的一种可能方法是:

status = cudaFree(d_arr);

类似的概念(取消引用主机代码中的设备指针)是此处注释的基础:

// for some reason this should happen here and not d_sp->coeff = d_coeff ...

【讨论】:

感谢您的及时回复。问题是,在我的真实代码中,我无法再访问 d_arr 我想做cudaFree (仅在上面的示例中)。是否可以引入辅助指针并在其上执行 cudaFree? 您可以将您的结构 (d_dummy) 复制回主机,然后从该主机副本中检索指针。或者在您执行dummy-&gt;arr = d_arr; 时,您可以添加另一行代码,例如float *d_helper_pointer = d_arr;,然后再执行cudaFree(d_helper_pointer);,我没有其他任何想法。我只能使用您显示的代码或您提供的描述。 确实,将结构复制回主机是有意义的。现在可以用了,谢谢!我想知道这种生活在设备上的结构是否被认为是好的风格...... 根据我的经验,使用结构作为内核参数是相当普遍的。带有嵌入式指针的结构可能有点挑战性,但您似乎已经正确导航了其中的大部分内容。如果您只打算将 1 个结构传递给内核,则通过值而不是指针传递可能更简单。如果您打算传递一个结构数组 (AoS),那么通常 SoA 方案比 AoS 更受欢迎,但我们在这里与您的问题相去甚远。 cuda SO 标签上有很多问题,讨论了这些不同的概念。

以上是关于结构中指向数组的 cudaFree 指针上的 CUDA 分段错误的主要内容,如果未能解决你的问题,请参考以下文章

如何访问 C 中指向列表的指针结构数组?

Arduino指向结构变量的指针和指向结构数组的指针打印输出

C:指向结构指针数组的指针(分配/解除分配问题)

C ++在指向结构的指针中访问类数组指针

如何初始化指向结构的指针数组?

分配指向结构数组的指针