       const Scalar * __restrict__ voxels,
 #endif
       int firstSlice, int lastSlice,
-      int from, int to);
+      int2 from, int2 to);
 
 //----------------------------------------------------------------------------
 // device code (CUDA)
 //----------------------------------------------------------------------------
 
-__device__ int2 coord(int from)
-{
-    return make_int2(threadIdx.x + blockDim.x*blockIdx.x,
-            threadIdx.y + blockDim.y*blockIdx.y + from);
+__device__ int2 coord(int2 from)
+{
+    return make_int2(threadIdx.x + blockDim.x*blockIdx.x + from.x,
+            threadIdx.y + blockDim.y*blockIdx.y + from.y);
 }
 
 __global__ void clearImage(uchar4 * __restrict__ img, int width, int height,
-      int from, int to)
-{
-    const int2 p = coord(from);
-    if (p.y >= to)
+      int fromY, int toY)
+{
+    const int2 p = coord(make_int2(0, fromY));
+    if (p.y >= toY)
         return;
     if (p.x >= width)
         return;


   Base::earlyRayTermination = 0;

   if (from == -1)
       from = 0;
   if (to == -1)
       to = Base::intImg->height;

   // compute data for determining upper left image corner of each slice and copy it to device
   vvVector4 start, end;

      const Scalar * __restrict__ voxels,
#endif
      int firstSlice, int lastSlice,
      int2 from, int2 to)
{
#ifndef VOLTEX3D
#ifdef PITCHED
