 const dim3 Patch(16, 16);
 
 #define PATCHES
+#define UNROLL
 //#define NOSHMEM
 //#define SHMCLASS
 //#define NOOP

#ifdef NOSHMEM
    for (int ix=threadIdx.x+from.x; ix<to.x; ix+=blockDim.x)
    {
        Pixel d;
        initPixel(&d);
        float sf = -1.f;

        // composite slices for this image line
        for (int slice=firstSlice; slice!=lastSlice; slice += sliceStep)
        {
            if(earlyRayTerm && isOpaque(d))
                break;

            // compute upper left image corner

}


template<typename Pixel, int principal, bool preInt>
struct Ray
{
};

template<typename Pixel, int principal>
struct Ray<Pixel, principal, false>
{
    Pixel d;

    __device__ Ray()
    {
        initPixel(&d);
    }

    __device__ void accumulate(int x, int y, int z)
    {
        const float v = volume(x, y, z, principal);
        const float4 c = classify<float4>(v);

        // blend
        const float w = d.w*c.w;
        d.x += w*c.x;
        d.y += w*c.y;
        d.z += w*c.z;
        d.w -= w;
    }
};

template<typename Pixel, int principal>
struct Ray<Pixel, principal, true>
{
    Pixel d;
    float sf;

    __device__ Ray()
        : sf(-1.f)
    {
        initPixel(&d);
    }

    __device__ void accumulate(int x, int y, int z)
    {
        const float sb = volume(x, y, z, principal);
        if(sf >= 0.f)
        {
            const float4 c = tex2D(tex_preint, sf, sb);
            // blend
            const float w = d.w*c.w;
            d.x += w*c.x;
