 #endif
 
    // transfer function is stored as a texture
-   vvCuda::checkError(&ok, cudaMalloc(&d_tf, 4096*4), "cudaMalloc tf");
-   vvCuda::checkError(&ok, cudaBindTexture(NULL, tex_tf, d_tf, 4096), "bind tf tex");
+   vvCuda::checkError(&ok, cudaMalloc(&d_tf, Base::getLUTSize()*sizeof(LutEntry)), "cudaMalloc tf");
+   vvCuda::checkError(&ok, cudaBindTexture(NULL, tex_tf, d_tf, Base::getLUTSize()*sizeof(LutEntry)), "bind tf tex");
 
    // pre-integration table
-   cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar4>();
+   cudaChannelFormatDesc desc = cudaCreateChannelDesc<LutEntry>();
    vvCuda::checkError(&ok, cudaMallocArray(&d_preint, &desc, Base::PRE_INT_TABLE_SIZE, Base::PRE_INT_TABLE_SIZE), "cudaMalloc preint");
    tex_preint.normalized = true;
    tex_preint.filterMode = Base::bilinLookup ? cudaFilterModeLinear : cudaFilterModePoint;

#define SHMLOAD
#endif

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

#ifdef VOLTEX3D
texture<uchar4, 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, Base::getLUTSize()*sizeof(LutEntry), 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*sizeof(LutEntry), cudaMemcpyHostToDevice), "cudaMemcpy preint");
    }
}

