   unsigned int j = __umul24(blockIdx_y, blockDim.y) + threadIdx.y;
   unsigned int k = __umul24(blockIdx_z, blockDim.z) + threadIdx.z;
 
-  if (i >= vol_dim.x || j >= vol_dim.y || k >= vol_dim.z)
+  if (i >= c_vol_size.x || j >= c_vol_size.y || k >= c_vol_size.z)
     {
     return;
     }
 
   // Index row major into the volume
-  long int vol_idx = i + (j + k*vol_dim.y)*(vol_dim.x);
+  long int vol_idx = i + (j + k*c_vol_size.y)*(c_vol_size.x);
 
   float3 ip;
   float  voxel_data = 0;

#include <cuda.h>

// T E X T U R E S ////////////////////////////////////////////////////////
texture<float, 3, cudaReadModeElementType> tex_img;
//texture<float, 1, cudaReadModeElementType> tex_matrix;
///////////////////////////////////////////////////////////////////////////

__constant__ float c_matrices[1024 * 12]; //Can process stacks of at most 1024 projections
__constant__ int3 c_projSize;

//_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_
// K E R N E L S -_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_

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

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