CUDA/Thrust double pointer problem (vector of pointers)

Hey all, I am using CUDA and the Thrust library. I am running into a problem when I try to access a double pointer on the CUDA kernel loaded with a thrust::device_vector of type Object* (vector of pointers) from the host. When compiled with ‘nvcc -o thrust main.cpp’ i receive the warning ‘Warning: Cannot tell what pointer points to, assuming global memory space’ and a launch error upon attempting to run the program.

I have read the Nvidia forums and the solution seems to be ‘Don’t use double pointers in a CUDA kernel’. I am not looking to collapse the double pointer into a 1D pointer before sending to the kernel…Has anyone found a solution to this problem? The required code is below, thanks in advance!




Sphere * parseSphere(int i)


  Sphere * s = new Sphere();

  s->a = 1+i;

  s->b = 2+i;

  s->c = 3+i;

  return s;


int main( int argc, char** argv ) {

int i;

  thrust::host_vector<Sphere *> spheres_h;

  thrust::host_vector<Sphere> spheres_resh(NUM_OBJECTS);

//initialize spheres_h


    Sphere * sphere = parseSphere(i);



//initialize spheres_resh


    spheres_resh[i].a = 1;

    spheres_resh[i].b = 1;

    spheres_resh[i].c = 1;


thrust::device_vector<Sphere *> spheres_dv = spheres_h;

  thrust::device_vector<Sphere> spheres_resv = spheres_resh;

  Sphere ** spheres_d = thrust::raw_pointer_cast(&spheres_dv[0]);

  Sphere * spheres_res = thrust::raw_pointer_cast(&spheres_resv[0]);




bool result = true;


    result &= (spheres_resh[i].a == i+1);

    result &= (spheres_resh[i].b == i+2);

    result &= (spheres_resh[i].c == i+3);




    cout << "Data GOOD!" << endl;


    cout << "Data BAD!" << endl;


return 0;




__global__ void deviceBegin(Sphere ** spheres_d, Sphere * spheres_res, float num_objects)


  int index = threadIdx.x + blockIdx.x*blockDim.x;

spheres_res[index].a = (*(spheres_d+index))->a; //causes warning/launch error

  spheres_res[index].b = (*(spheres_d+index))->b; 

  spheres_res[index].c = (*(spheres_d+index))->c; 


void kernelBegin(Sphere ** spheres_d, Sphere * spheres_res, float num_objects)


int threads = 512;//per block

 int grids = ((num_objects)/threads)+1;//blocks per grid

deviceBegin<<<grids,threads>>>(spheres_d, spheres_res, num_objects);