Passing an struct including a pointer to another struct, to kernel in CUDA

Question

I have two structs as

struct collapsed {
    char **seq;
    int num;
};


struct data {
    collapsed *x;
    int num;
    int numblocks;
    int *blocksizes;
    float *regmult;
    float *learnmult;
};

I am passing it to my kernel as;

__global__ void KERNEL(data* X,...){
    ...
    collapsed x = X->x[0]; // GIVES CUDA_EXPECTION_1:Lane Illegal Address
}

data X;
//init X
data *X_dev;
cudaMalloc((data **) & X_dev, sizeof(data));
cudaMemcpy(X_dev, &X, sizeof(data), cudaMemcpyHostToDevice);
KERNEL<<<...>>>(X_dev,...);

This code gives CUDA_EXPECTION_1:Lane Illegal Address in the kernel code. What is wrong or what is the right way to do it ? Any idea?

Answer 1

You're dereferencing a host pointer on the device. X is a valid device pointer.

But when you copied the X struct to the device, you copied x along with it, which contains a host pointer. When you dereference that pointer:

collapsed x = X->x[0];
                 ^ this is dereferencing the x pointer

the device code throws an error.

More detail is given here as well as instructions on how to fix it.