     // Note: cbi->img AND cbi->matrix are passed via texture memory
     //-------------------------------------
     kernel_fdk_optim <<< dimGrid, dimBlock >>> ( dev_vol,
-                                                 make_int2(img_dim[0], img_dim[1]),
                                                  make_int3(vol_dim[0], vol_dim[1], vol_dim[2]) );
     }
   else

CUDA_reconstruct_conebeam_init(
  int img_dim[2],
  int vol_dim[3],
  float *&dev_vol,             // Holds voxels on device
  float *&dev_img,         // Holds image pixels on device
  float *&dev_matrix           // Holds matrix on device
  )
{
  // Size of volume Malloc

//_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_

__global__
void kernel_fdk(float *dev_vol, int3 vol_dim, unsigned int Blocks_Y, float invBlocks_Y)
{
  // CUDA 2.0 does not allow for a 3D grid, which severely
  // limits the manipulation of large 3D arrays of data.  The
