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.

Answer 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.

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

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

2. 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