From 50167f6c26d350ea6788d7c83ab24b34d231007d Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Mon, 6 Dec 2010 14:19:41 +0000 Subject: [PATCH] added first version of gpu::matchTemplate, currently it works only with 8UC1 images and SQDIFF method --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 + modules/gpu/src/border_interpolate.cpp | 2 +- modules/gpu/src/cuda/border_interpolate.hpp | 2 +- modules/gpu/src/cuda/match_template.cu | 97 ++++++++++++++ .../{internal_common.hpp => internal_shared.hpp} | 4 +- modules/gpu/src/match_template.cpp | 69 ++++++++++ tests/gpu/src/gputest_main.cpp | 5 +- tests/gpu/src/match_template.cpp | 145 +++++++++++++++++++++ 8 files changed, 323 insertions(+), 5 deletions(-) create mode 100644 modules/gpu/src/cuda/match_template.cu rename modules/gpu/src/{internal_common.hpp => internal_shared.hpp} (94%) create mode 100644 modules/gpu/src/match_template.cpp create mode 100644 tests/gpu/src/match_template.cpp diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 53a1a74..6655af1 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/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 //////////////////////////////// /*! diff --git a/modules/gpu/src/border_interpolate.cpp b/modules/gpu/src/border_interpolate.cpp index dbd8394..a162a2e 100644 --- a/modules/gpu/src/border_interpolate.cpp +++ b/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" diff --git a/modules/gpu/src/cuda/border_interpolate.hpp b/modules/gpu/src/cuda/border_interpolate.hpp index 3dad03e..d15e11d 100644 --- a/modules/gpu/src/cuda/border_interpolate.hpp +++ b/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 { diff --git a/modules/gpu/src/cuda/match_template.cu b/modules/gpu/src/cuda/match_template.cu new file mode 100644 index 0000000..ea797f0 --- /dev/null +++ b/modules/gpu/src/cuda/match_template.cu @@ -0,0 +1,97 @@ +/*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 imageTex_8U; +texture 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(); + 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<<>>(templ.cols, templ.rows, result); + cudaSafeCall(cudaThreadSynchronize()); + cudaSafeCall(cudaUnbindTexture(imageTex_8U)); + cudaSafeCall(cudaUnbindTexture(templTex_8U)); +} + + +}}} diff --git a/modules/gpu/src/internal_common.hpp b/modules/gpu/src/internal_shared.hpp similarity index 94% rename from modules/gpu/src/internal_common.hpp rename to modules/gpu/src/internal_shared.hpp index c9af363..c3a5882 100644 --- a/modules/gpu/src/internal_common.hpp +++ b/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 { diff --git a/modules/gpu/src/match_template.cpp b/modules/gpu/src/match_template.cpp new file mode 100644 index 0000000..db182bd --- /dev/null +++ b/modules/gpu/src/match_template.cpp @@ -0,0 +1,69 @@ +/*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 diff --git a/tests/gpu/src/gputest_main.cpp b/tests/gpu/src/gputest_main.cpp index a634ef0..917ce9a 100644 --- a/tests/gpu/src/gputest_main.cpp +++ b/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. */ diff --git a/tests/gpu/src/match_template.cpp b/tests/gpu/src/match_template.cpp new file mode 100644 index 0000000..c80cb29 --- /dev/null +++ b/tests/gpu/src/match_template.cpp @@ -0,0 +1,145 @@ +/*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(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(i, j); + float v2 = b.at(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; -- 2.7.4