added first version of gpu::matchTemplate, currently it works only with 8UC1…

added first version of gpu::matchTemplate, currently it works only with 8UC1 images and SQDIFF method

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -656,6 +656,10 @@ namespace cv
         CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType=BORDER_REFLECT101);
 
 
+        //! computes the proximity map for the raster template and the image where the template is searched for
+        CV_EXPORTS void matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method);
+
+
         //////////////////////////////// Filter Engine ////////////////////////////////
 
         /*!

modules/gpu/src/border_interpolate.cpp

@@ -40,7 +40,7 @@
 //
 //M*/
 
-#include "internal_common.hpp"
+#include "internal_shared.hpp"
 #include "border_interpolate.hpp"
 #include "opencv2/gpu/gpu.hpp"
 

modules/gpu/src/cuda/border_interpolate.hpp

@@ -43,7 +43,7 @@
 #ifndef __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
 #define __OPENCV_GPU_BORDER_INTERPOLATE_HPP__
 
-#include "../internal_common.hpp"
+#include "../internal_shared.hpp"
 
 namespace cv { namespace gpu {
 

modules/gpu/src/cuda/match_template.cu

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "cuda_shared.hpp"
+
+using namespace cv::gpu;
+
+namespace cv { namespace gpu { namespace imgproc {
+
+texture<unsigned char, 2> imageTex_8U;
+texture<unsigned char, 2> templTex_8U;
+
+
+__global__ void matchTemplateKernel_8U_SqDiff(int w, int h, DevMem2Df result)
+{
+    int x = blockDim.x * blockIdx.x + threadIdx.x;
+    int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+    if (x < result.cols && y < result.rows)
+    {
+        float sum = 0.f;
+        float delta;
+
+        for (int i = 0; i < h; ++i)
+        {
+            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;
+            }
+        }
+
+        result.ptr(y)[x] = sum;
+    }
+}
+
+
+void matchTemplateCaller_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, image.data, desc, image.cols, image.rows, image.step);
+    cudaBindTexture2D(0, templTex_8U, templ.data, desc, templ.cols, templ.rows, templ.step);
+    imageTex_8U.filterMode = cudaFilterModePoint;
+    templTex_8U.filterMode = cudaFilterModePoint;
+
+    matchTemplateKernel_8U_SqDiff<<<grid, threads>>>(templ.cols, templ.rows, result);
+    cudaSafeCall(cudaThreadSynchronize());
+    cudaSafeCall(cudaUnbindTexture(imageTex_8U));
+    cudaSafeCall(cudaUnbindTexture(templTex_8U));
+}
+
+
+}}}

modules/gpu/src/internal_common.hpp → modules/gpu/src/internal_shared.hpp

@@ -40,8 +40,8 @@
 //
 //M*/
 
-#ifndef __OPENCV_GPU_INTERNAL_COMMON_HPP__
-#define __OPENCV_GPU_INTERNAL_COMMON_HPP__
+#ifndef __OPENCV_GPU_INTERNAL_SHARED_HPP__
+#define __OPENCV_GPU_INTERNAL_SHARED_HPP__
 
 namespace cv { namespace gpu {
 

modules/gpu/src/match_template.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+#if !defined (HAVE_CUDA)
+
+void cv::gpu::matchTemplate(const GpuMat&, const GpuMat&, GpuMat&, int) { throw_nogpu(); }
+
+#else
+
+namespace cv { namespace gpu { namespace imgproc {
+    
+    void matchTemplateCaller_8U_SqDiff(const DevMem2D, const DevMem2D, DevMem2Df);
+
+}}}
+
+void cv::gpu::matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method)
+{
+    CV_Assert(image.type() == CV_8U);
+    CV_Assert(method == CV_TM_SQDIFF);
+
+    CV_Assert(image.type() == templ.type());
+    CV_Assert(image.cols >= templ.cols && image.rows >= templ.rows);
+
+    result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
+    imgproc::matchTemplateCaller_8U_SqDiff(image, templ, result);
+}
+
+#endif

tests/gpu/src/gputest_main.cpp

@@ -56,7 +56,10 @@ const char* blacklist[] =
 int main( int argc, char** argv )
 
 {
-    return test_system.run( argc, argv, blacklist );
+    argc = 3;
+    const char* manargs[] = { "gputest", "-tn", "GPU-MatchTemplateTest"};
+    
+    return test_system.run( argc, (char**)manargs, blacklist );
 }
 
 /* End of file. */

tests/gpu/src/match_template.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other GpuMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or bpied warranties, including, but not limited to, the bpied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "gputest.hpp"
+
+using namespace cv;
+using namespace std;
+
+struct CV_GpuMatchTemplateTest: CvTest 
+{
+    CV_GpuMatchTemplateTest(): CvTest("GPU-MatchTemplateTest", "matchTemplate") {}
+
+    void run(int)
+    {
+        try
+        {
+            Mat image, templ;
+            Mat dst_gold;
+            gpu::GpuMat dst;
+            int n, m, h, w;
+
+            for (int i = 0; i < 4; ++i)
+            {
+                n = 1 + rand() % 100;
+                m = 1 + rand() % 100;
+                do h = 1 + rand() % 20; while (h > n);
+                do w = 1 + rand() % 20; while (w > m);
+                gen(image, n, m, CV_8U);
+                gen(templ, h, w, CV_8U);
+
+                match_template_naive(image, templ, dst_gold);
+                gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_SQDIFF);
+                if (!check8U(dst_gold, Mat(dst))) return;
+            }
+        }
+        catch (const Exception& e)
+        {
+            if (!check_and_treat_gpu_exception(e, ts)) throw;
+            return;
+        }
+    }
+
+
+    void gen(Mat& a, int rows, int cols, int type)
+    {
+        RNG rng;
+        a.create(rows, cols, type);
+        if (type == CV_8U)
+            rng.fill(a, RNG::UNIFORM, Scalar(0), Scalar(10));
+    }
+
+    // Naive version for unsigned char
+    // Time complexity is O(a.size().area() * b.size().area()).
+    void match_template_naive(const Mat& a, const Mat& b, Mat& c)
+    {
+        c.create(a.rows - b.rows + 1, a.cols - b.cols + 1, CV_32F);         
+        for (int i = 0; i < c.rows; ++i)
+        {
+            for (int j = 0; j < c.cols; ++j)
+            {
+                float delta;
+                float sum = 0.f;
+                for (int y = 0; y < b.rows; ++y)
+                {
+                    const unsigned char* arow = a.ptr(i + y);
+                    const unsigned char* brow = b.ptr(y);
+                    for (int x = 0; x < b.cols; ++x)
+                    {
+                        delta = (float)(arow[j + x] - brow[x]);
+                        sum += delta * delta;
+                    }
+                }
+                c.at<float>(i, j) = sum;
+            }
+        }
+    }
+
+
+    bool check8U(const Mat& a, const Mat& b)
+    {
+        if (a.size() != b.size())
+        {
+            ts->printf(CvTS::CONSOLE, "bad size");
+            ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
+            return false;
+        }
+
+        for (int i = 0; i < a.rows; ++i)
+        {
+            for (int j = 0; j < a.cols; ++j)
+            {
+                float v1 = a.at<float>(i, j);
+                float v2 = b.at<float>(i, j);
+                if (fabs(v1 - v2) > 1e-3f)
+                {
+                    ts->printf(CvTS::CONSOLE, "(gold)%f != %f, pos: (%d, %d) size: (%d, %d)\n", 
+                               v1, v2, j, i, a.cols, a.rows);
+                    ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
+                    return false;
+                }
+            }
+        }
+
+        return true;
+    }
+} match_template_test;