Cuda array of pointers allocation

Accelerated Computing CUDA CUDA Programming and Performance
1

How can I allocate this structure on Cuda?

struct scene{
Object * object;
…
}
struct Object{
Triangle * triangles;
…
}

struct Triangle{
…
}

2 
$ cat t2175.cu
#include <cstdio>

struct Triangle{
  int h,w;
};

struct Object{
  Triangle *triangles;
};

struct scene{
  Object *object;
};

__global__ void k(scene *s){

  printf("%d\n", s[0].object[0].triangles[0].h);
}

const int num_scene=2;
const int num_obj=3;
const int num_tri=4;

int main(){

  Triangle *d_tri, *h_tri = new Triangle[num_tri];
  cudaMalloc(&d_tri, num_tri*sizeof(Triangle));
  h_tri[0].h=8675309;
  cudaMemcpy(d_tri, h_tri, num_tri*sizeof(Triangle), cudaMemcpyHostToDevice);
  Object *d_obj, *h_obj = new Object[num_obj];
  cudaMalloc(&d_obj, num_obj*sizeof(Object));
  h_obj[0].triangles = d_tri;
  cudaMemcpy(d_obj, h_obj, num_obj*sizeof(Object), cudaMemcpyHostToDevice);
  scene *d_s, *h_s = new scene[num_scene];
  cudaMalloc(&d_s, num_scene*sizeof(scene));
  h_s[0].object = d_obj;
  cudaMemcpy(d_s, h_s, num_scene*sizeof(scene), cudaMemcpyHostToDevice);
  k<<<1,1>>>(d_s);
  cudaDeviceSynchronize();
}



$ nvcc -o t2175 t2175.cu
$ compute-sanitizer ./t2175
========= COMPUTE-SANITIZER
8675309
========= ERROR SUMMARY: 0 errors
$

Note that in the above example I have only fully allocated for the first scene (allocation is provided for 2 scenes, but only the first has its object pointer properly set). Likewise in the first scene, I have only fully allocated for the first object (3 objects are allocated for, but only the first has its triangles pointer properly set). Likewise in the first object of the first scene, I have allocated for 4 triangles, but only the first has any values (h) set.

This survey may also be of interest.

3

Thanks a lot Robert, it is exactaly what i need!