 #include <cuda.h>
 
 // P R O T O T Y P E S ////////////////////////////////////////////////////
-__global__ void kernel_fdk(float *dev_vol, int2 img_dim, int3 vol_dim, unsigned int Blocks_Y, float invBlocks_Y);
+__global__ void kernel_fdk(float *dev_vol, int3 vol_dim, unsigned int Blocks_Y, float invBlocks_Y);
 
 ///////////////////////////////////////////////////////////////////////////
 

    // 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

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

__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
