 #define SHMLOAD
 #endif
 
+typedef uchar4 LutEntry;
+
 #ifdef SHMCLASS
 #define SHMLOAD
-texture<uchar4, 1, cudaReadModeElementType> tex_tf;
-#else
-texture<uchar4, 1, cudaReadModeNormalizedFloat> tex_tf;
+texture<LutEntry, 1, cudaReadModeElementType> tex_tf;
+#else
+texture<LutEntry, 1, cudaReadModeNormalizedFloat> tex_tf;
 #endif
 
 #ifdef VOLTEX3D
-texture<uchar4, 2, cudaReadModeNormalizedFloat> tex_preint;
+texture<LutEntry, 2, cudaReadModeNormalizedFloat> tex_preint;
 #ifdef FLOATDATA
 texture<float, 3, cudaReadModeElementType> tex_raw;
 #else

        Base::makeLookupTextureCorrect(dist);   // use this line for slow but more correct pre-integration LUT
    }

    vvCuda::checkError(NULL, cudaMemcpy(d_tf, Base::rgbaConv, sizeof(Base::rgbaConv), cudaMemcpyHostToDevice), "cudaMemcpy tf");
    if(Base::preIntegration)
    {
        vvCuda::checkError(NULL, cudaMemcpyToArray(d_preint, 0, 0, &Base::preIntTable[0][0][0],
                    Base::PRE_INT_TABLE_SIZE*Base::PRE_INT_TABLE_SIZE*4, cudaMemcpyHostToDevice), "cudaMemcpy preint");
    }
}


{
    uchar4 c = *dst;
    const float w = src.w/255.f;
    *dst = make_uchar4(src.x + c.x*w,
            src.y + c.y*w,
            src.z + c.z*w,
