     for (int ix=threadIdx.x; ix<width; ix+=blockDim.x)
     {
         uchar4 *dest = img + line*width+ix;
-        uchar4 c = *dest;
-#ifdef FLOATIMG
-        *dest = make_uchar4(imgLine[ix].x * 255.f + c.x*imgLine[ix].w,
-                imgLine[ix].y * 255.f + c.y*imgLine[ix].w,
-                imgLine[ix].z * 255.f + c.z*imgLine[ix].w,
-                (1.f-imgLine[ix].w) * (255-c.w) + c.w);
-#else
-        const float w= imgLine[ix].w/255.f;
-        imgLine[ix].w = 255 - imgLine[ix].w;
-        *dest = make_uchar4(imgLine[ix].x + c.x*w,
-                imgLine[ix].y + c.y*w,
-                imgLine[ix].z + c.z*w,
-                (1.f-w) * (255-c.w) + c.w);
-#endif
+        blend(dest, imgLine[ix]);
     }
 }
 #endif

        }
#endif
    }
    }

#ifndef NODISPLAY
    // copy line to intermediate image
    for (int ix=threadIdx.x; ix<width; ix+=blockDim.x)
    {
        uchar4 *dest = img + line*width+ix;
        uchar4 c = *dest;
#ifdef FLOATIMG
        *dest = make_uchar4(imgLine[ix].x * 255.f + c.x*imgLine[ix].w,
                imgLine[ix].y * 255.f + c.y*imgLine[ix].w,
                imgLine[ix].z * 255.f + c.z*imgLine[ix].w,
                (1.f-imgLine[ix].w) * (255-c.w) + c.w);
#else
        const float w= imgLine[ix].w/255.f;
        imgLine[ix].w = 255 - imgLine[ix].w;
        *dest = make_uchar4(imgLine[ix].x + c.x*w,
                imgLine[ix].y + c.y*w,
                imgLine[ix].z + c.z*w,
                (1.f-w) * (255-c.w) + c.w);
#endif
    }
#endif
}


#ifdef VOLTEX3D
template<typename Scalar, int BPV, int sliceStep, int principal>
__global__ void compositeSlicesBilinear(
      uchar4 * __restrict__ img, int width, int height,
#ifdef PITCHED

{
#ifdef VOLTEX3D
    if(sliceInterpol)
        return compositeSlicesBilinear<Scalar, 1, Pixel, sliceStep, principal>;
    else
#endif
        return compositeSlicesNearest<Scalar, 1, Pixel, sliceStep, principal>;
}

template<int principal>
