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 cukernel.cu’ 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!

--------------------------

        main.cpp

--------------------------

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

  for(i=0;i<NUM_OBJECTS;i++){

    Sphere * sphere = parseSphere(i);

    spheres_h.push_back(sphere);

  }

//initialize spheres_resh

  for(i=0;i<NUM_OBJECTS;i++){

    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]);

kernelBegin(spheres_d,spheres_res,NUM_OBJECTS);

thrust::copy(spheres_dv.begin(),spheres_dv.end(),spheres_h.begin());

  thrust::copy(spheres_resv.begin(),spheres_resv.end(),spheres_resh.begin());

bool result = true;

for(i=0;i<NUM_OBJECTS;i++){

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

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

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

  }

if(result)

  {

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

  }else{

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

  }

return 0;

}

--------------------------

        cukernel.cu

--------------------------

__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);

}