 
                 h_tcStart[slice].x = 1.f + (h_start[slice].x - sx + 0.5f)*h_tcStep[slice].x;
                 h_tcStart[slice].y = 1.f + (h_start[slice].y - sy + 0.5f)*h_tcStep[slice].y;
-
-                h_tc3[slice] = 1.f-(slice+0.5f)*1.f/Base::vd->vox[Base::principal];
                 break;
             case 1:
                 h_tcStep[slice].x = -1.f/(ex-sx);

// device code (CUDA)
//----------------------------------------------------------------------------

__device__ int2 coord(int2 from)
{
#if 0
    return make_int2(threadIdx.x + blockDim.x*blockIdx.x + from.x,
            threadIdx.y + blockDim.y*blockIdx.y + from.y);
#else
    const int xr = threadIdx.x & 3;
    const int x4 = threadIdx.x >> 2;
    const int yr = threadIdx.y & 3;
    const int y4 = threadIdx.y >> 2;
    return make_int2(xr + 4 * yr + blockDim.x*blockIdx.x + from.x,
            x4 + 4 * y4 + blockDim.y*blockIdx.y + from.y);
#endif
}

__global__ void clearImage(uchar4 * __restrict__ img, int width, int height,

    tc_inc.x *= scale;
    tc_inc.y *= scale;

    float3 orig = make_float3(tc.x, tc.y, z);
    float3 dir = make_float3(tc_inc.x, tc_inc.y, -c_zStep);
    float tnear, tfar;
    bool isect = intersectTexCube(orig, dir, &tnear, &tfar);

    if(!isect)
    {
        // copy pixel to intermediate image
        if(p.x >= from.x && p.x < to.x && p.y >= from.y && p.y < to.y)
        {
            uchar4 *dest = img + p.y*width+p.x;
            setpixel(dest, ray.d);
        }
        return;
    }

    int slstart = 0;
    int slstop = nslice;

#if 1
    int2 c[4];
    c[0] = coordMin(from);
    c[1] = coordMax(from);
    c[2] = coordMin(from); c[2].x = c[1].x;
    c[3] = coordMax(from); c[3].x = c[0].x;

    slstart = nslice;
    slstop = 0;
    int unistop = nslice;
    int unistart = 0;
    int nisect = 0;
    for(int i=0; i<4; ++i)
    {
        float2 tc1 = texcoord(c[i], 0);
        float2 tc1_inc = texcoord(c[i], 1);
        tc1_inc.x -= tc1.x;
        tc1_inc.y -= tc1.y;
        tc1_inc.x *= scale;
        tc1_inc.y *= scale;
        float3 orig1 = make_float3(tc1.x, tc1.y, z);
        float3 dir1 = make_float3(tc1_inc.x, tc1_inc.y, -c_zStep);
        float tnear1, tfar1;
        bool isect1 = intersectTexCube(orig1, dir1, &tnear1, &tfar1);
        if(isect1)
        {
            ++nisect;
            slstart = fminf(tnear1, slstart);
            unistart = fmaxf(tnear1, unistart);

            slstop = fmaxf(tfar1+1.f, slstop);
            unistop = fminf(tfar1, unistop);
        }
    }

    if(nisect==4)
    {
        unistart += 1;
        unistop -= 1;
        if(slstop < unistop)
            unistop = slstop;
        if(unistop > unistart)
        {
            unistop = (unistop-unistart)/unroll*unroll+unistart;
        }
        else
        {
            unistart=nslice;
            unistop=nslice;
        }
    }
    else
    {
        unistart=nslice;
        unistop=nslice;
    }

    z -= slstart* c_zStep;
    tc.x += slstart*tc_inc.x;
    tc.y += slstart*tc_inc.y;
#else
    float3 orig = make_float3(tc.x, tc.y, z);
    float3 dir = make_float3(tc_inc.x, tc_inc.y, -c_zStep);
    float tnear, tfar;
    bool isect = intersectTexCube(orig, dir, &tnear, &tfar);

    if(!isect)
        return;

    slstart=tnear;
    z -= slstart* c_zStep;
    tc.x += slstart*tc_inc.x;
    tc.y += slstart*tc_inc.y;
    slstop=tfar;
#endif

    for(int sl=slstart; sl<unistart; ++sl)
    {
        if(earlyRayTerm && isOpaque(ray.d))
            break;;

        if(tc.x >= 0.f && tc.x < 1.f
                && tc.y >= 0.f && tc.y < 1.f)
            ray.accumulate(tc.x, tc.y, z);

        tc.x += tc_inc.x;
        tc.y += tc_inc.y;
#if VOLTEX3D != 3
