sending struct array to cuda kernel

Question

I'm working on a project and I have to sent a struct array to cuda kernel. The struct also contains an array. To test it I have written a simple program.

struct Point {
    short     x;
    short     *y;
};

my kernel code:

__global__ void addKernel(Point *a, Point *b, Point *c)
{
    int i = threadIdx.x;

    c[i].x = a[i].x + b[i].x;
    for (int j = 0; j<4; j++){
        c[i].y[j] = a[i].y[j] + a[i].y[j];
    }
}

my main code:

int main()
{
    const int arraySize = 4;
    const int arraySize2 = 4;

    short *ya, *yb, *yc;
    short *dev_ya, *dev_yb, *dev_yc;

    Point *a;
    Point *b;
    Point *c;
    Point *dev_a;
    Point *dev_b;
    Point *dev_c;

    size_t sizeInside = sizeof(short) * arraySize2;

    ya = (short *)malloc(sizeof(short) * arraySize2);
    yb = (short *)malloc(sizeof(short) * arraySize2);
    yc = (short *)malloc(sizeof(short) * arraySize2);

    ya[0] = 1; ya[1] =2; ya[2]=3; ya[3]=4;
    yb[0] = 2; yb[1] =3; yb[2]=4; yb[3]=5;

    size_t sizeGeneral = (sizeInside+sizeof(short)) * arraySize;

    a = (Point *)malloc( sizeGeneral );  
    b = (Point *)malloc( sizeGeneral );
    c = (Point *)malloc( sizeGeneral );


    a[0].x = 2;  a[0].y = ya;
    a[1].x = 2;  a[1].y = ya;
    a[2].x = 2;  a[2].y = ya;
    a[3].x = 2;  a[3].y = ya;

    b[0].x = 4;  b[0].y = yb;
    b[1].x = 4;  b[1].y = yb;
    b[2].x = 4;  b[2].y = yb;
    b[3].x = 4;  b[3].y = yb;

    cudaMalloc((void**)&dev_a, sizeGeneral);
    cudaMalloc((void**)&dev_b, sizeGeneral);
    cudaMalloc((void**)&dev_c, sizeGeneral);

    cudaMemcpy(dev_a, a, sizeGeneral, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, sizeGeneral, cudaMemcpyHostToDevice);

    addKernel<<<1, 4>>>(dev_a, dev_b, dev_c);

    cudaError_t err = cudaMemcpy(c, dev_c, sizeGeneral, cudaMemcpyDeviceToHost);   

    printf("{%d-->%d,%d,%d,%d} \n err= %d",c[0].x,c[0].y[0],c[1].y[1],c[1].y[2],c[2].y[3], err);        

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;
}

It seems cuda kernel is not working. Actually I can access structs 'x' variable but I cannot access 'y' array. What can I do to access the 'y' array? Thanks in advance.

What kind of error occurs? What does your application output? — Dariusz, Apr 12 '13 at 07:49
it has not a compile error. In kernel cannot access 'y' so, when i want to access 'c[i].y[j]', it crash. — Soner İlica, Apr 12 '13 at 07:54

score 1 · Answer 1 · answered Apr 12 '13 at 07:41

When you are sending this struct to kernel you send short and pointer to short in host memory not device. This is crucial. For simple type - as short this works, because kernel has its local copy in memory designated to accept parameters. So when you call this kernel you have moved x and y to device, but not the area pointed by y. This you have to do manually by allocating space for it and updating pointer y to point to device memory.

score 0 · Accepted Answer · answered Apr 12 '13 at 07:51

0

You are not passin the array to the device. You can either make the array a part of the struct, by defining it like this:

struct {
  short normalVal;
  short inStructArr[4];
}

Or pass the array into the device memory and update the pointer in the struct.

answered Apr 12 '13 at 07:51

Dariusz

21,561
9
74
114

Handling pointers to device memory inside of data structures that reside in device memory is somewhat complicated. The modified struct method proposed here is *much* simpler. If you insist on passing the pointer, as Adrian indicates in his answer, then you need to use a method similar to what I described [here](http://stackoverflow.com/questions/15113960/cuda-allocating-array-of-pointers-to-images-and-the-images/15115362#15115362). Nevertheless I don't recommend it due to it's complexity, if you can avoid it. – Robert Crovella Apr 12 '13 at 14:26

sending struct array to cuda kernel

2 Answers2