 #include "vvcuda.h"
 #include "vvcudapar.h"
 
+const int MAX_SLICES = 512;
+
 __constant__ int   c_vox[5];
-__constant__ float c_start[4];
-__constant__ float c_inc[4];
+__constant__ float2 c_start[MAX_SLICES];
+static float2 h_start[MAX_SLICES];
 
 texture<uchar4, 1, cudaReadModeNormalizedFloat> tex_tf;
 

    for (int slice=firstSlice; slice!=lastSlice; slice += sliceStep)
    {
        // compute upper left image corner
        const float w1 = 1.f/cur[3];
        const int iPosX = int(cur[0]*w1+0.5f);
        const int iPosY = int(cur[1]*w1+0.5f);
        cur[0] += c_inc[0];
        cur[1] += c_inc[1];
        cur[3] += c_inc[3];

        if(line < iPosY)
            continue;

   findSlicePosition(firstSlice, &start, NULL);
   findSlicePosition(firstSlice+sliceStep, &inc, NULL);
   inc.sub(&start);
   vvVector4 cur = start;
   for(int slice=firstSlice; slice != lastSlice; slice += sliceStep)
   {
       h_start[slice].x = cur.e[0] / cur.e[3];
       h_start[slice].y = cur.e[1] / cur.e[3];
       cur.add(&inc);
   }

   bool ok = true;
   vvCuda::checkError(&ok, cudaMemcpyToSymbol(c_start, h_start, sizeof(h_start)), "cudaMemcpy start");

   // prepare intermediate image
   if (warpMode==CUDATEXTURE)
