Accessing Class Member Array Pointer on Device from Host in CUDA

Question

I've been pretty stumped on this problem for a some time now. This involves CUDA device pointers. I have an instance of a custom class sitting on my device, and it has a member variable which is a pointer to an array (which is on the device).

class MyClass {
public:
    int* array;
    // Other variables and functions, etc.
};

It needs to be a dynamically allocated array because the size of the array depends on some input at the beginning of the program. For the duration of the program I am modifying the class using kernel functions, but eventually I want to get a copy of this class on the host to output to file. However I can't seem to be able to get cudaMemCpy to work for me.

I can get a copy of the class by using this code (where dc is a pointer to the class on the device):

MyClass hc;
cudaMemcpy(&hc, dc, sizeof(dc), cudaMemcpyDeviceToHost);

But this only gets the information in the class that isn't a pointer, which makes sense since the pointer retrieved in hc would still be pointing to data on the device. So I figured that I could use this code to actually get the array.

int* h_array;
cudaMemcpy(h_array, dc->array, sizeof(dc->array), cudaMemcpyDeviceToHost);

This only returns an empty array, plus I get a cudaFree error ("Cuda error: cuda free operations: invalid argument"). I've tried a bunch of variations of this, including using hc->array, with no success. Is there any way that I can get this array without having to a write a kernel function to copy each individual entry? I'm working with CUDA 5.0.

Is the device pointer in the class (so the value of MyClass.array) allocated using the host API or using malloc/new inside a kernel on the device? — talonmies
– talonmies, Commented Jan 10, 2013 at 16:13
The pointer is allocated using the host API. After I read in the input I allocate an array on the host and another array on the device (using cudaMalloc). I store some initial values to the host array, then use cudaMemCpy to copy this information on the device array. Then I use a simple <<<1,1>>> kernel to set the value of MyClass.array to the device array. I free the host array after that since I don't need it anymore. — C.G.
– C.G., Commented Jan 10, 2013 at 16:54
dc is a pointer to device mem. You can not deference it on the host like this dc->array — kangshiyin
– kangshiyin, Commented Jan 10, 2013 at 19:11
Having this into account, I advise you to rethink the way your approaching the problem. To begin, it is simpler (but not necessarily more efficient) to work with arrays of structures/classes, than with structures/classes of arrays (in which case, the arrays would have to have a fixed size anyway). — sandeep.ganage
– sandeep.ganage, Commented Jan 11, 2013 at 8:00

kangshiyin · Accepted Answer · 2013-01-11 13:26:51Z

1

I think you use sizeof and pointers in a wrong way.

sizeof(dc) and sizeof(dc->array) in your code could be replaced by sizeof(MyClass) & ArraySize * sizeof(int).

For pointers, you have to do cudaMemcpy twice to get your array.

First get object hc, which stores the addr of your array.

cudaMemcpy(&hc, dc, sizeof(MyClass), cudaMemcpyDeviceToHost);

Then get the array itself.

cudaMemcpy(h_array, hc.array, ArraySize*sizeof(int),D2H);

Also, dc is a pointer to device mem. You can not dereference it on the host like this dc->array

edited Jan 11, 2013 at 13:26

answered Jan 10, 2013 at 19:08

kangshiyin

9,8291 gold badge19 silver badges30 bronze badges

Sign up to request clarification or add additional context in comments.

4 Comments

C.G. Over a year ago

This sort of helps! While your method did not work for my original code, it did make me think to change how the MyClass object was initialized. Originally it was set up so that it received the device pointer after I copied it to the device, but I changed it to giving it the pointer before copying it to the device. For whatever reason though I can still access the array while I have the original host copy of the class, but when I copy it back from the device the pointer doesn't work anymore (results in the same error).

kangshiyin Over a year ago

@C.G. It's strange. the value (the addr in this case) stored in array should not change after copying to and back from device mem.

kangshiyin Over a year ago

@C.G. I think sizeof(dc) and sizeof(dc->array) in your code could also be a problem. Try sizeof(MyClass) & ArraySize*sizeof(int) instead.

C.G. Over a year ago

Yep, that's what the problem was. Switching to sizeof(MyClass) from sizeof(dc) solved it. Thank you for your help!

Collectives™ on Stack Overflow

Accessing Class Member Array Pointer on Device from Host in CUDA

1 Answer 1

4 Comments

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

4 Comments

Your Answer

Sign up or log in

Post as a guest

Related