CUDA Segmentation fault on cudaFree pointer to array in struct

Viewed 170

I got a struct living on the CUDA device which contains a pointer to an array. Computations, accessing the elements and everything works fine but when I try to be a good boy and call

cudaFree(my_struct->pointer_to_array)

I get a segmentation fault. cudaFree(my_struct) however works perfectly fine. Is there something I am missing?

Please find the following minimal example:

#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 Answers

This statement:

status = cudaFree(d_dummy->arr);

requires dereferencing of a device pointer (d_dummy - which was allocated with a device allocator i.e. cudaMalloc) in host code. That is illegal in CUDA.

Since you already know that (d_dummy->arr) == d_arr, one possible approach to freeing the embedded pointer would be:

status = cudaFree(d_arr);

A similar concept (dereferencing a device pointer in host code) underlies the comment here:

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