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

        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");
    }
}


#define SHMLOAD
#endif

typedef uchar4 LutEntry;

#ifdef SHMCLASS
#define SHMLOAD
texture<LutEntry, 1, cudaReadModeElementType> tex_tf;
#else
texture<LutEntry, 1, cudaReadModeNormalizedFloat> tex_tf;
#endif

#ifdef VOLTEX3D
texture<LutEntry, 2, cudaReadModeNormalizedFloat> tex_preint;
#ifdef FLOATDATA
texture<float, 3, cudaReadModeElementType> tex_raw;
#else
