         {
             for (int j = 0; j < w; ++j)
             {
-                delta = tex2D(imageTex_32F, x + j, y + i) - 
-                        tex2D(templTex_32F, j, i);
+                delta = tex2D(imageTex_32F_SQDIFF, x + j, y + i) - 
+                        tex2D(templTex_32F_SQDIFF, j, i);
                 sum += delta * delta;
             }
         }

        {
            for (int j = 0; j < w; ++j)
            {
                delta = (float)tex2D(imageTex_8U, x + j, y + i) - 
                        (float)tex2D(templTex_8U, j, i);
                sum += delta * delta;
            }
        }



void matchTemplateNaive_8U_SQDIFF(const DevMem2D image, const DevMem2D templ,
                                  DevMem2Df result)
{
    dim3 threads(32, 8);
    dim3 grid(divUp(image.cols - templ.cols + 1, threads.x), 
              divUp(image.rows - templ.rows + 1, threads.y));

    cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>();
    cudaBindTexture2D(0, imageTex_8U_SQDIFF, image.data, desc, image.cols, image.rows, image.step);
    cudaBindTexture2D(0, templTex_8U_SQDIFF, templ.data, desc, templ.cols, templ.rows, templ.step);
    imageTex_8U_SQDIFF.filterMode = cudaFilterModePoint;
    templTex_8U_SQDIFF.filterMode = cudaFilterModePoint;

    matchTemplateNaiveKernel_8U_SQDIFF<<<grid, threads>>>(templ.cols, templ.rows, result);
    cudaSafeCall(cudaThreadSynchronize());
    cudaSafeCall(cudaUnbindTexture(imageTex_8U_SQDIFF));
    cudaSafeCall(cudaUnbindTexture(templTex_8U_SQDIFF));
}


texture<float, 2> imageTex_32F_SQDIFF;
texture<float, 2> templTex_32F_SQDIFF;


__global__ void matchTemplateNaiveKernel_32F_SQDIFF(int w, int h, 
