 {
     cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar4>();
     cudaBindTexture2D(0, resize8UC4_tex, src.data, desc, src.cols, src.rows, src.step);
+    resize8UC4_tex.filterMode = cudaFilterModeLinear;
 
     dim3 threads(32, 8);
     dim3 grid(div_up(dst.cols, threads.x), div_up(dst.rows, threads.y));

//-------------------------------------------------------------------
// Resize

texture<uchar4, 2> resize8UC4_tex;
texture<unsigned char, 2> resize8UC1_tex;


extern "C" __global__ void resize_8UC4_kernel(float sx, float sy, DevMem2D dst)

    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < dst.cols && y < dst.rows)
        ((unsigned char*)dst.ptr(y))[x] = tex2D(resize8UC1_tex, x * sx, y * sy) * 255;
}


