CV_EXPORTS void cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType=BORDER_REFLECT101);\r
\r
\r
+ //! computes the proximity map for the raster template and the image where the template is searched for\r
+ CV_EXPORTS void matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method);\r
+\r
+\r
//////////////////////////////// Filter Engine ////////////////////////////////\r
\r
/*!\r
//\r
//M*/\r
\r
-#include "internal_common.hpp"\r
+#include "internal_shared.hpp"\r
#include "border_interpolate.hpp"\r
#include "opencv2/gpu/gpu.hpp"\r
\r
#ifndef __OPENCV_GPU_BORDER_INTERPOLATE_HPP__\r
#define __OPENCV_GPU_BORDER_INTERPOLATE_HPP__\r
\r
-#include "../internal_common.hpp"\r
+#include "../internal_shared.hpp"\r
\r
namespace cv { namespace gpu {\r
\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "cuda_shared.hpp"\r
+\r
+using namespace cv::gpu;\r
+\r
+namespace cv { namespace gpu { namespace imgproc {\r
+\r
+texture<unsigned char, 2> imageTex_8U;\r
+texture<unsigned char, 2> templTex_8U;\r
+\r
+\r
+__global__ void matchTemplateKernel_8U_SqDiff(int w, int h, DevMem2Df result)\r
+{\r
+ int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+ if (x < result.cols && y < result.rows)\r
+ {\r
+ float sum = 0.f;\r
+ float delta;\r
+\r
+ for (int i = 0; i < h; ++i)\r
+ {\r
+ for (int j = 0; j < w; ++j)\r
+ {\r
+ delta = (float)tex2D(imageTex_8U, x + j, y + i) - \r
+ (float)tex2D(templTex_8U, j, i);\r
+ sum += delta * delta;\r
+ }\r
+ }\r
+\r
+ result.ptr(y)[x] = sum;\r
+ }\r
+}\r
+\r
+\r
+void matchTemplateCaller_8U_SqDiff(const DevMem2D image, const DevMem2D templ, DevMem2Df result)\r
+{\r
+ dim3 threads(32, 8);\r
+ dim3 grid(divUp(image.cols - templ.cols + 1, threads.x), \r
+ divUp(image.rows - templ.rows + 1, threads.y));\r
+\r
+ cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>();\r
+ cudaBindTexture2D(0, imageTex_8U, image.data, desc, image.cols, image.rows, image.step);\r
+ cudaBindTexture2D(0, templTex_8U, templ.data, desc, templ.cols, templ.rows, templ.step);\r
+ imageTex_8U.filterMode = cudaFilterModePoint;\r
+ templTex_8U.filterMode = cudaFilterModePoint;\r
+\r
+ matchTemplateKernel_8U_SqDiff<<<grid, threads>>>(templ.cols, templ.rows, result);\r
+ cudaSafeCall(cudaThreadSynchronize());\r
+ cudaSafeCall(cudaUnbindTexture(imageTex_8U));\r
+ cudaSafeCall(cudaUnbindTexture(templTex_8U));\r
+}\r
+\r
+\r
+}}}\r
+++ /dev/null
-/*M///////////////////////////////////////////////////////////////////////////////////////\r
-//\r
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
-//\r
-// By downloading, copying, installing or using the software you agree to this license.\r
-// If you do not agree to this license, do not download, install,\r
-// copy or use the software.\r
-//\r
-//\r
-// License Agreement\r
-// For Open Source Computer Vision Library\r
-//\r
-// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
-// Third party copyrights are property of their respective owners.\r
-//\r
-// Redistribution and use in source and binary forms, with or without modification,\r
-// are permitted provided that the following conditions are met:\r
-//\r
-// * Redistribution's of source code must retain the above copyright notice,\r
-// this list of conditions and the following disclaimer.\r
-//\r
-// * Redistribution's in binary form must reproduce the above copyright notice,\r
-// this list of conditions and the following disclaimer in the documentation\r
-// and/or other materials provided with the distribution.\r
-//\r
-// * The name of the copyright holders may not be used to endorse or promote products\r
-// derived from this software without specific prior written permission.\r
-//\r
-// This software is provided by the copyright holders and contributors "as is" and\r
-// any express or implied warranties, including, but not limited to, the implied\r
-// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
-// In no event shall the Intel Corporation or contributors be liable for any direct,\r
-// indirect, incidental, special, exemplary, or consequential damages\r
-// (including, but not limited to, procurement of substitute goods or services;\r
-// loss of use, data, or profits; or business interruption) however caused\r
-// and on any theory of liability, whether in contract, strict liability,\r
-// or tort (including negligence or otherwise) arising in any way out of\r
-// the use of this software, even if advised of the possibility of such damage.\r
-//\r
-//M*/\r
-\r
-#ifndef __OPENCV_GPU_INTERNAL_COMMON_HPP__\r
-#define __OPENCV_GPU_INTERNAL_COMMON_HPP__\r
-\r
-namespace cv { namespace gpu {\r
-\r
- // Internal GPU anlagues of CPU border extrapolation types\r
- enum \r
- {\r
- BORDER_REFLECT101 = 0,\r
- BORDER_REPLICATE\r
- };\r
-\r
-}}\r
-\r
-#endif
\ No newline at end of file
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#ifndef __OPENCV_GPU_INTERNAL_SHARED_HPP__\r
+#define __OPENCV_GPU_INTERNAL_SHARED_HPP__\r
+\r
+namespace cv { namespace gpu {\r
+\r
+ // Internal GPU anlagues of CPU border extrapolation types\r
+ enum \r
+ {\r
+ BORDER_REFLECT101 = 0,\r
+ BORDER_REPLICATE\r
+ };\r
+\r
+}}\r
+\r
+#endif
\ No newline at end of file
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "precomp.hpp"\r
+\r
+#if !defined (HAVE_CUDA)\r
+\r
+void cv::gpu::matchTemplate(const GpuMat&, const GpuMat&, GpuMat&, int) { throw_nogpu(); }\r
+\r
+#else\r
+\r
+namespace cv { namespace gpu { namespace imgproc {\r
+ \r
+ void matchTemplateCaller_8U_SqDiff(const DevMem2D, const DevMem2D, DevMem2Df);\r
+\r
+}}}\r
+\r
+void cv::gpu::matchTemplate(const GpuMat& image, const GpuMat& templ, GpuMat& result, int method)\r
+{\r
+ CV_Assert(image.type() == CV_8U);\r
+ CV_Assert(method == CV_TM_SQDIFF);\r
+\r
+ CV_Assert(image.type() == templ.type());\r
+ CV_Assert(image.cols >= templ.cols && image.rows >= templ.rows);\r
+\r
+ result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);\r
+ imgproc::matchTemplateCaller_8U_SqDiff(image, templ, result);\r
+}\r
+\r
+#endif\r
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. */
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other GpuMaterials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or bpied warranties, including, but not limited to, the bpied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "gputest.hpp"\r
+\r
+using namespace cv;\r
+using namespace std;\r
+\r
+struct CV_GpuMatchTemplateTest: CvTest \r
+{\r
+ CV_GpuMatchTemplateTest(): CvTest("GPU-MatchTemplateTest", "matchTemplate") {}\r
+\r
+ void run(int)\r
+ {\r
+ try\r
+ {\r
+ Mat image, templ;\r
+ Mat dst_gold;\r
+ gpu::GpuMat dst;\r
+ int n, m, h, w;\r
+\r
+ for (int i = 0; i < 4; ++i)\r
+ {\r
+ n = 1 + rand() % 100;\r
+ m = 1 + rand() % 100;\r
+ do h = 1 + rand() % 20; while (h > n);\r
+ do w = 1 + rand() % 20; while (w > m);\r
+ gen(image, n, m, CV_8U);\r
+ gen(templ, h, w, CV_8U);\r
+\r
+ match_template_naive(image, templ, dst_gold);\r
+ gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_SQDIFF);\r
+ if (!check8U(dst_gold, Mat(dst))) return;\r
+ }\r
+ }\r
+ catch (const Exception& e)\r
+ {\r
+ if (!check_and_treat_gpu_exception(e, ts)) throw;\r
+ return;\r
+ }\r
+ }\r
+\r
+\r
+ void gen(Mat& a, int rows, int cols, int type)\r
+ {\r
+ RNG rng;\r
+ a.create(rows, cols, type);\r
+ if (type == CV_8U)\r
+ rng.fill(a, RNG::UNIFORM, Scalar(0), Scalar(10));\r
+ }\r
+\r
+ // Naive version for unsigned char\r
+ // Time complexity is O(a.size().area() * b.size().area()).\r
+ void match_template_naive(const Mat& a, const Mat& b, Mat& c)\r
+ {\r
+ c.create(a.rows - b.rows + 1, a.cols - b.cols + 1, CV_32F); \r
+ for (int i = 0; i < c.rows; ++i)\r
+ {\r
+ for (int j = 0; j < c.cols; ++j)\r
+ {\r
+ float delta;\r
+ float sum = 0.f;\r
+ for (int y = 0; y < b.rows; ++y)\r
+ {\r
+ const unsigned char* arow = a.ptr(i + y);\r
+ const unsigned char* brow = b.ptr(y);\r
+ for (int x = 0; x < b.cols; ++x)\r
+ {\r
+ delta = (float)(arow[j + x] - brow[x]);\r
+ sum += delta * delta;\r
+ }\r
+ }\r
+ c.at<float>(i, j) = sum;\r
+ }\r
+ }\r
+ }\r
+\r
+\r
+ bool check8U(const Mat& a, const Mat& b)\r
+ {\r
+ if (a.size() != b.size())\r
+ {\r
+ ts->printf(CvTS::CONSOLE, "bad size");\r
+ ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);\r
+ return false;\r
+ }\r
+\r
+ for (int i = 0; i < a.rows; ++i)\r
+ {\r
+ for (int j = 0; j < a.cols; ++j)\r
+ {\r
+ float v1 = a.at<float>(i, j);\r
+ float v2 = b.at<float>(i, j);\r
+ if (fabs(v1 - v2) > 1e-3f)\r
+ {\r
+ ts->printf(CvTS::CONSOLE, "(gold)%f != %f, pos: (%d, %d) size: (%d, %d)\n", \r
+ v1, v2, j, i, a.cols, a.rows);\r
+ ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);\r
+ return false;\r
+ }\r
+ }\r
+ }\r
+\r
+ return true;\r
+ }\r
+} match_template_test;\r
projects / platform / upstream / opencv.git / commitdiff