Copy structs from host to device (and vice versa)

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi there,

I am trying to handle structs in a CUDA code. But I have problem allocating memory, and copying them ( I am not sure which part of my code is incorrect). I get this run-time error: “copying of minl to device failed”

here is part of the code that deals with this thing:

struct threeD
{
float x, y, z;
};
constant struct threeD *maxd,*mind;
threeD *maxl,*minl;

// allocate memory
if (cudaSuccess != cudaMalloc((void **)&maxl, sizeof(threeD) * blocks)) fprintf(stderr, “could not allocate maxd\n”); CudaTest(“couldn’t allocate maxd”);
if (cudaSuccess != cudaMalloc((void **)&minl, sizeof(threeD) * blocks)) fprintf(stderr, “could not allocate mind\n”); CudaTest(“couldn’t allocate mind”);

if (cudaSuccess != cudaMemcpyToSymbol(maxd, &maxl, sizeof(threeD))) fprintf(stderr, “copying of maxl to device failed\n”); CudaTest(“maxl copy to device failed”);
if (cudaSuccess != cudaMemcpyToSymbol(mind, &minl, sizeof(threeD))) fprintf(stderr, “copying of minl to device failed\n”); CudaTest(“minl copy to device failed”);

I appreciate your help in advance.

2

There are multiple problems in the code.

First: Do you want the device variables to reside in global or in constant memory? If constant memory, they can only be statically allocated, so there is no cudaMalloc(), and the cudaMemcpyToSymbol() call is also wrong. If global memory, they can either be statically allocated, in which case the same applies as for the static allocation in constant memory. Or you can allocate the dynamically, in which case the you need to use cudaMemcpy(), not cudaMemcpyToSymbol.

Second: The code you show is a device to device copy. I assume you want host->device copy, in which case you need variables on the host whose content gets copied to the device.

Untested code for a copy to variables in constant memory:

struct threeD

{

    float x, y, z;

};

struct threeD maxh, minh;  /* host variables */

__constant__ struct threeD maxd,mind; /* device variables */

if (cudaSuccess != cudaMemcpyToSymbol("maxd", &maxh, sizeof(struct threeD), 0, cudaMemcpyHostToDevice))

        fprintf(stderr, "copying of maxd to device failed\n");

    CudaTest("maxd copy to device failed");

    if (cudaSuccess != cudaMemcpyToSymbol("mind", &minh, sizeof(struct threeD), 0, cudaMemcpyHostToDevice))

        fprintf(stderr, "copying of mind to device failed\n");

    CudaTest("mind copy to device failed");

and for a copy to dynamically allocated global memory:

struct threeD

{

    float x, y, z;

};

struct threeD maxh, minh;  /* host variables */

threeD *maxl,*minl; /* pointers to device memory */

// allocate memory

    if (cudaSuccess != cudaMalloc((void **)&maxl, sizeof(struct threeD) * blocks))

        fprintf(stderr, "could not allocate maxd\n");

    CudaTest("couldn't allocate maxd");

    if (cudaSuccess != cudaMalloc((void **)&minl, sizeof(struct threeD) * blocks)) 

        fprintf(stderr, "could not allocate mind\n");

    CudaTest("couldn't allocate mind");

if (cudaSuccess != cudaMemcpy(maxl, &maxh, sizeof(struct threeD), cudaMemcpyHostToDevice))

        fprintf(stderr, "copying of maxd to device failed\n");

    CudaTest("maxd copy to device failed");

    if (cudaSuccess != cudaMemcpy(minl, &minh, sizeof(struct threeD), cudaMemcpyHostToDevice))

        fprintf(stderr, "copying of mind to device failed\n");

    CudaTest("mind copy to device failed");
3

Thanks for your reply.
I want the variables to reside in the constant memory, but each of the maxd, mind, maxl and minl are arrays (of type threeD). I tried both of your sample codes but they don’t work because the compiler can’t identify things like “mind[5].x”. and if I apply your code and keep them as before (*mind, *maxd, …) it still gives the same error “copying of minl to device failed” in the run-time.
I could manage to make the code perfectly work when there was no structs (one dimension) using the code I provided above. I got in trouble when I used arrays of structs instead of arrays.

I appreciate if you help me out with this struct thing.

4

In your code [font=“Courier New”]maxd[/font] and [font=“Courier New”]mind[/font] are pointers, not arrays. To create arrays in constant memory, you need to declare them with an explicit size:

__constant__ struct threeD maxd[12], mind[12];
5

Thanks,

The problem with my code was that in my case since we are giving the size of address (I guess) of maxh, this part of code:

if (cudaSuccess != cudaMemcpyToSymbol("maxd", &maxh, sizeof(struct threeD), 0, cudaMemcpyHostToDevice))

has to be changed to:

if (cudaSuccess != cudaMemcpyToSymbol("maxd", &maxh, sizeof(int), 0, cudaMemcpyHostToDevice))
6

That seems completely wrong. [font=“Courier New”]maxh[/font] isn’t a pointer in your code. And even if it were, it’s size were [font=“Courier New”]sizeof (struct threeD *)[/font].

If in doubt you could write

if (cudaSuccess != cudaMemcpyToSymbol("maxd", &maxh, sizeof(maxh), 0, cudaMemcpyHostToDevice))

which is what I always do just to make mistakes less likely.