values passed into kernel getting messy

Question

I'm trying to initialize some structs on the device, but when they come back, theyre all messy. I know that something is wrong with how I'm using my pointers, but I cant seem to figure it out.

heres my memory allocation on the host

body* devBodies;
body** devBodyList;
float* devRandoms;
cudaMalloc((void**)&devBodies, n * sizeof(body));
cudaMalloc((void**)&devBodyList, n * sizeof(body*));
cudaMalloc((void**)&devRandoms, n * 3 * sizeof(float));

heres my function call, where devRandoms is just a list of floats.

CUDAInitBodies<<<n,1>>>(devBodyList, devBodies, devRandoms);

and heres my method:

__global__ void CUDAInitBodies(body** devBodyList, body* devBody, float* rand)
{
    int j;
    int tid = blockIdx.x;
    *(devBodyList[tid]) = devBody[tid];
    devBodyList[tid]->id = tid;
    devBodyList[tid]->m = 10;
    for(j = 0; j < 3; j++) 
    {
        devBodyList[tid]->a[j] = 0;
        devBodyList[tid]->v[j] = 0;
        devBodyList[tid]->pos[j] = rand[(tid * 3) + j];
    }   

}

when I then copy the data back to the host:

body* bodies = (body*)malloc(n * sizeof(body));
cudaMemcpy(bodies, devBodies, n * sizeof(body), cudaMemcpyDeviceToHost);

when I print out the variables of my "body", I get this:

====BODY[-581043205]====
    M = -42522218781525353518415985938704367616.000000
    V = [-311601248975690312470654313562112.000000, 17269896096570671104.000000, 307939529506715418513587721849804947456.000000]
    X = -19247336126697914498972549714433540096.000000
    Y = 17731266573644159438123340575306416128.000000
    Z = -544771403677696.000000

I've trying different ways of doing this for quite a while, but nothing seems to be doing the trick

Answer 1

The reason you are getting garbage output is that you are dereferencing an uninitialized pointer at this line of code in your kernel:

*(devBodyList[tid]) = devBody[tid];

On the host, you allocated storage for an array of pointers here:

cudaMalloc((void**)&devBodyList, n * sizeof(body*));

That only creates storage for the array of pointers. It does not set any of the pointers in the array to point to anything valid.

Each member of the array is a pointer to a body struct, using the first member as an example:

devBodyList[0]

But that pointer does not point to anything (valid) unless you initialize it, with a statement like (in device code):

devBodyList[0] = &(devBody[0]);

Now I can use devBodyList[0] as a pointer to the body structure allocated at devBody[0]. Note that devBodyList[0] is stored on the device, so to initialize it I must either:

1. Initialize it in device code before using it (such as the above line of code).
2. Initialize it on the host, but setting up a valid pointer and then using an operation like cudaMemcpy to copy that initialized value to the device storage that was allocted.

Unless you account for the above, when you go to use the pointer, it will contain a garbage value, and dereferencing it will produce invalid accesses. You can see evidence of these invalid accesses by running your code with cuda-memcheck. You'll get a message like "invalid global write..."

Here's a fully worked code, based on your pieces, that does not produce garbage results. (I'm not saying it's particularly sensible, since using devBodyList[x] to refer to devBody[x] seems unnecessary to me.) But it is legal code and will not produce any kernel errors:

#include <stdio.h>

struct body {
int id;
int m;
int a[3];
int v[3];
float pos[3];
};

__global__ void CUDAInitBodies(body** devBodyList, body* devBody, float* rand)
{
    int j;
    int tid = blockIdx.x;
    devBodyList[tid] = &(devBody[tid]);
    *(devBodyList[tid]) = devBody[tid];
    devBodyList[tid]->id = tid;
    devBodyList[tid]->m = 10;
    for(j = 0; j < 3; j++)
    {
        devBodyList[tid]->a[j] = 0;
        devBodyList[tid]->v[j] = 0;
        devBodyList[tid]->pos[j] = rand[(tid * 3) + j];
    }

}

int main(){
  int n = 1;
  body *devBodies;
  body **devBodyList;
  float *devRandoms;

  cudaMalloc((void**)&devBodies, n * sizeof(body));
  cudaMalloc((void**)&devBodyList, n * sizeof(body*));
  cudaMalloc((void**)&devRandoms, n * 3 * sizeof(float));

  CUDAInitBodies<<<n,1>>>(devBodyList, devBodies, devRandoms);

  body* bodies = (body*)malloc(n * sizeof(body));
  cudaMemcpy(bodies, devBodies, n * sizeof(body), cudaMemcpyDeviceToHost);

  printf("Body %d\n", bodies[0].id);
  printf("M : %d\n", bodies[0].m);
  printf("V : %d\n", bodies[0].v[0]);
  return 0;
}