diff --git a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp new file mode 100644 index 000000000..948774743 --- /dev/null +++ b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp @@ -0,0 +1,57 @@ +#include "perf_precomp.hpp" +#include "opencv2/ts/ocl_perf.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { + +namespace ocl { + + CV_ENUM(MethodType, TM_SQDIFF, TM_SQDIFF_NORMED, TM_CCORR, TM_CCORR_NORMED, TM_CCOEFF, TM_CCOEFF_NORMED) + + typedef std::tr1::tuple ImgSize_TmplSize_Method_t; + typedef TestBaseWithParam ImgSize_TmplSize_Method; + + OCL_PERF_TEST_P(ImgSize_TmplSize_Method, MatchTemplate, + ::testing::Combine( + testing::Values(szSmall128, cv::Size(320, 240), + cv::Size(640, 480), cv::Size(800, 600), + cv::Size(1024, 768), cv::Size(1280, 1024)), + testing::Values(cv::Size(12, 12), cv::Size(28, 9), + cv::Size(8, 30), cv::Size(16, 16)), + MethodType::all() + ) + ) + { + Size imgSz = get<0>(GetParam()); + Size tmplSz = get<1>(GetParam()); + int method = get<2>(GetParam()); + + UMat img(imgSz, CV_8UC1); + UMat tmpl(tmplSz, CV_8UC1); + UMat result(imgSz - tmplSz + Size(1,1), CV_32F); + + declare + .in(img, WARMUP_RNG) + .in(tmpl, WARMUP_RNG) + .out(result) + .time(30); + + OCL_TEST_CYCLE() matchTemplate(img, tmpl, result, method); + + bool isNormed = + method == TM_CCORR_NORMED || + method == TM_SQDIFF_NORMED || + method == TM_CCOEFF_NORMED; + double eps = isNormed ? 3e-2 + : 255 * 255 * tmpl.total() * 1e-4; + + if (isNormed) + SANITY_CHECK(result,eps,ERROR_RELATIVE); + else + SANITY_CHECK(result, eps); + } +} +} + +#endif // HAVE_OPENCL \ No newline at end of file diff --git a/modules/imgproc/src/opencl/match_template.cl b/modules/imgproc/src/opencl/match_template.cl new file mode 100644 index 000000000..451ec0cb2 --- /dev/null +++ b/modules/imgproc/src/opencl/match_template.cl @@ -0,0 +1,397 @@ +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, 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. + +#define DATA_SIZE ((int)sizeof(type)) +#define ELEM_TYPE elem_type +#define ELEM_SIZE ((int)sizeof(elem_type)) +#define CN cn + +#define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, gidx*CN + img_sqsums_offset + ox*CN) +#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx*CN + img_sums_offset + ox*CN) + +inline float normAcc(float num, float denum) +{ + if(fabs(num) < denum) + { + return num / denum; + } + if(fabs(num) < denum * 1.125f) + { + return num > 0 ? 1 : -1; + } + return 0; +} + +inline float normAcc_SQDIFF(float num, float denum) +{ + if(fabs(num) < denum) + { + return num / denum; + } + if(fabs(num) < denum * 1.125f) + { + return num > 0 ? 1 : -1; + } + return 1; +} + +//////////////////////////////////////////CCORR///////////////////////////////////////////////////////////////////////// + +__kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step,int img_offset, + __global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols, + __global uchar * res,int res_step,int res_offset,int res_rows,int res_cols) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + int i,j; + float sum = 0; + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + if(gidx < res_cols && gidy < res_rows) + { + for(i = 0; i < tpl_rows; i ++) + { + __global const ELEM_TYPE * img_ptr = (__global const ELEM_TYPE *)(img + mad24(gidy + i, img_step, gidx*DATA_SIZE + img_offset)); + __global const ELEM_TYPE * tpl_ptr = (__global const ELEM_TYPE *)(tpl + mad24(i, tpl_step, tpl_offset)); + + for(j = 0; j < tpl_cols; j ++) + +#pragma unroll + for (int c = 0; c < CN; c++) + + sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]); + + } + __global float * result = (__global float *)(res+res_idx); + *result = sum; + } +} + +__kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int tpl_rows, int tpl_cols, ulong tpl_sqsum) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sqsums_step /= sizeof(float); + img_sqsums_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + if(gidx < res_cols && gidy < res_rows) + { + __global float * sqsum = (__global float*)(img_sqsums); + float image_sqsum_ = (float)( + (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - + (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); + + __global float * result = (__global float *)(res+res_idx); + *result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum)); + } +} + +////////////////////////////////////////////SQDIFF//////////////////////////////////////////////////////////////////////// + +__kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step,int img_offset, + __global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols, + __global uchar * res,int res_step,int res_offset,int res_rows,int res_cols) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + int i,j; + float delta; + float sum = 0; + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + if(gidx < res_cols && gidy < res_rows) + { + for(i = 0; i < tpl_rows; i ++) + { + __global const ELEM_TYPE * img_ptr = (__global const ELEM_TYPE *)(img + mad24(gidy + i, img_step, gidx*DATA_SIZE + img_offset)); + __global const ELEM_TYPE * tpl_ptr = (__global const ELEM_TYPE *)(tpl + mad24(i, tpl_step, tpl_offset)); + + for(j = 0; j < tpl_cols; j ++) + +#pragma unroll + for (int c = 0; c < CN; c++) + { + delta = (float)(img_ptr[j*CN+c] - tpl_ptr[j*CN+c]); + sum += delta*delta; + } + } + __global float * result = (__global float *)(res+res_idx); + *result = sum; + } +} + +__kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int tpl_rows, int tpl_cols, ulong tpl_sqsum) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sqsums_step /= sizeof(float); + img_sqsums_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + if(gidx < res_cols && gidy < res_rows) + { + __global float * sqsum = (__global float*)(img_sqsums); + float image_sqsum_ = (float)( + (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - + (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); + + __global float * result = (__global float *)(res+res_idx); + + *result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum)); + } +} + +////////////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// + +__kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int tpl_rows, int tpl_cols, float tpl_sum) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sums_step /= ELEM_SIZE; + img_sums_offset /= ELEM_SIZE; + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + float image_sum_ = 0; + + if(gidx < res_cols && gidy < res_rows) + { + __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); + + image_sum_ += (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)])- + (sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])) * tpl_sum; + + __global float * result = (__global float *)(res+res_idx); + *result -= image_sum_; + } +} + +__kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sums_step /= ELEM_SIZE; + img_sums_offset /= ELEM_SIZE; + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + float image_sum_ = 0; + + if(gidx < res_cols && gidy < res_rows) + { + __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); + + image_sum_ += tpl_sum_0 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)]) -(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])); + image_sum_ += tpl_sum_1 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+1] - sum[SUMS_PTR(tpl_cols, 0)+1])-(sum[SUMS_PTR(0, tpl_rows)+1] - sum[SUMS_PTR(0, 0)+1])); + + __global float * result = (__global float *)(res+res_idx); + + *result -= image_sum_; + } +} + +__kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sums_step /= ELEM_SIZE; + img_sums_offset /= ELEM_SIZE; + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + float image_sum_ = 0; + + if(gidx < res_cols && gidy < res_rows) + { + __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); + + int c_r = SUMS_PTR(tpl_cols, tpl_rows); + int c_o = SUMS_PTR(tpl_cols, 0); + int o_r = SUMS_PTR(0,tpl_rows); + int oo = SUMS_PTR(0, 0); + + image_sum_ += tpl_sum_0 * (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[oo])); + image_sum_ += tpl_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1])); + image_sum_ += tpl_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2])); + image_sum_ += tpl_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3])); + + __global float * result = (__global float *)(res+res_idx); + + *result -= image_sum_; + } +} + +__kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, + __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int t_rows, int t_cols, float weight, float tpl_sum, float tpl_sqsum) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sums_offset /= ELEM_SIZE; + img_sums_step /= ELEM_SIZE; + img_sqsums_step /= sizeof(float); + img_sqsums_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + if(gidx < res_cols && gidy < res_rows) + { + __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); + __global float * sqsum = (__global float*)(img_sqsums); + + float image_sum_ = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)]) - + (sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)])); + + float image_sqsum_ = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)]) - + (sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)])); + + __global float * result = (__global float *)(res+res_idx); + + *result = normAcc((*result) - image_sum_ * tpl_sum, + sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_))); + } +} + +__kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, + __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int t_rows, int t_cols, float weight, float tpl_sum_0, float tpl_sum_1, float tpl_sqsum) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sums_offset /= ELEM_SIZE; + img_sums_step /= ELEM_SIZE; + img_sqsums_step /= sizeof(float); + img_sqsums_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + float sum_[2]; + float sqsum_[2]; + + if(gidx < res_cols && gidy < res_rows) + { + __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); + __global float * sqsum = (__global float*)(img_sqsums); + + sum_[0] = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)])-(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)])); + sum_[1] = (float)((sum[SUMS_PTR(t_cols, t_rows)+1] - sum[SUMS_PTR(t_cols, 0)+1])-(sum[SUMS_PTR(0, t_rows)+1] - sum[SUMS_PTR(0, 0)+1])); + + sqsum_[0] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)])-(sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)])); + sqsum_[1] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)+1] - sqsum[SQSUMS_PTR(t_cols, 0)+1])-(sqsum[SQSUMS_PTR(0, t_rows)+1] - sqsum[SQSUMS_PTR(0, 0)+1])); + + float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1; + + float denum = sqrt( tpl_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] + + sqsum_[1] - weight * sum_[1]* sum_[1])); + + __global float * result = (__global float *)(res+res_idx); + *result = normAcc((*result) - num, denum); + } +} + +__kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, + __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int t_rows, int t_cols, float weight, + float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3, + float tpl_sqsum) +{ + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + img_sums_offset /= ELEM_SIZE; + img_sums_step /= ELEM_SIZE; + img_sqsums_step /= sizeof(float); + img_sqsums_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); + + float sum_[4]; + float sqsum_[4]; + + if(gidx < res_cols && gidy < res_rows) + { + __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); + __global float * sqsum = (__global float*)(img_sqsums); + + int c_r = SUMS_PTR(t_cols, t_rows); + int c_o = SUMS_PTR(t_cols, 0); + int o_r = SUMS_PTR(0, t_rows); + int o_o = SUMS_PTR(0, 0); + + sum_[0] = (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[o_o ])); + sum_[1] = (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[o_o +1])); + sum_[2] = (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[o_o +2])); + sum_[3] = (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[o_o +3])); + + c_r = SQSUMS_PTR(t_cols, t_rows); + c_o = SQSUMS_PTR(t_cols, 0); + o_r = SQSUMS_PTR(0, t_rows); + o_o = SQSUMS_PTR(0, 0); + + sqsum_[0] = (float)((sqsum[c_r] - sqsum[c_o]) -(sqsum[o_r] - sqsum[o_o])); + sqsum_[1] = (float)((sqsum[c_r+1] - sqsum[c_o+1])-(sqsum[o_r+1] - sqsum[o_o+1])); + sqsum_[2] = (float)((sqsum[c_r+2] - sqsum[c_o+2])-(sqsum[o_r+2] - sqsum[o_o+2])); + sqsum_[3] = (float)((sqsum[c_r+3] - sqsum[c_o+3])-(sqsum[o_r+3] - sqsum[o_o+3])); + + float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1 + sum_[2]*tpl_sum_2 + sum_[3]*tpl_sum_3; + + float denum = sqrt( tpl_sqsum * ( + sqsum_[0] - weight * sum_[0]* sum_[0] + + sqsum_[1] - weight * sum_[1]* sum_[1] + + sqsum_[2] - weight * sum_[2]* sum_[2] + + sqsum_[3] - weight * sum_[3]* sum_[3] )); + + __global float * result = (__global float *)(res+res_idx); + *result = normAcc((*result) - num, denum); + } +} \ No newline at end of file diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index bfe7ce600..9953c6851 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -40,6 +40,337 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" + +//////////////////////////////////////////////////matchTemplate////////////////////////////////////////////////////////// +namespace cv +{ + static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArray _result); + static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result); + + static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result); + static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result); + + static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArray _result); + static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result); + + static bool matchTemplateNaive_CCORR (InputArray _image, InputArray _templ, OutputArray _result, int cn); + static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result, int cn); + + static bool useNaive(int method, int depth, Size size) + { +#ifdef HAVE_CLAMDFFT + if (method == TM_SQDIFF && depth == CV_32F) + return true; + else if(method == TM_CCORR || (method == TM_SQDIFF && depth == CV_8U)) + return size.height < 18 && size.width < 18; + else + return false; +#else +#define UNUSED(x) (void)(x); + UNUSED(method) UNUSED(depth) UNUSED(size) +#undef UNUSED + return true; +#endif + } + +///////////////////////////////////////////////////CCORR////////////////////////////////////////////////////////////// + + static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArray _result) + { + if (useNaive(TM_CCORR, _image.depth(), _templ.size()) ) + return matchTemplateNaive_CCORR(_image, _templ, _result, _image.channels()); + else + return false; + } + + static bool matchTemplateNaive_CCORR (InputArray _image, InputArray _templ, OutputArray _result, int cn) + { + int type = _image.type(); + int depth = CV_MAT_DEPTH(type); + + const char * kernelName = "matchTemplate_Naive_CCORR"; + + ocl::Kernel k (kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn)); + if (k.empty()) + return false; + + UMat image = _image.getUMat(); + UMat templ = _templ.getUMat(), result; + _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); + result = _result.getUMat(); + + size_t globalsize[2] = {result.cols, result.rows}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false); + } + + static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result) + { + matchTemplate(_image, _templ, _result, CV_TM_CCORR); + + int type = _image.type(); + int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + + const char * kernelName = "matchTemplate_CCORR_NORMED"; + + ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn)); + if (k.empty()) + return false; + + UMat image = _image.getUMat(); + UMat templ = _templ.getUMat(), result; + _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); + result = _result.getUMat(); + + UMat image_sums, image_sqsums; + integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); + + UMat templ_resh, temp; + templ.reshape(1).convertTo(templ_resh, CV_32F); + + multiply(templ_resh, templ_resh, temp); + unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0]; + + size_t globalsize[2] = {result.cols, result.rows}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false); + } + +//////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// + + static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result) + { + if (useNaive(TM_SQDIFF, _image.depth(), _templ.size())) + { + return matchTemplateNaive_SQDIFF(_image, _templ, _result, _image.channels());; + } + else + return false; + } + + static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result, int cn) + { + int type = _image.type(); + int depth = CV_MAT_DEPTH(type); + + const char * kernelName = "matchTemplate_Naive_SQDIFF"; + + ocl::Kernel k (kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn)); + if (k.empty()) + return false; + + UMat image = _image.getUMat(); + UMat templ = _templ.getUMat(), result; + _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); + result = _result.getUMat(); + + size_t globalsize[2] = {result.cols, result.rows}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false); + } + + static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result) + { + matchTemplate(_image, _templ, _result, CV_TM_CCORR); + + int type = _image.type(); + int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + + const char * kernelName = "matchTemplate_SQDIFF_NORMED"; + + ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn)); + if (k.empty()) + return false; + + UMat image = _image.getUMat(); + UMat templ = _templ.getUMat(), result; + _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); + result = _result.getUMat(); + + UMat image_sums, image_sqsums; + integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); + + UMat temp, templ_resh; + templ.reshape(1).convertTo(templ_resh, CV_32F); + + multiply(templ_resh, templ_resh, temp); + unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0]; + + size_t globalsize[2] = {result.cols, result.rows}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false); + } + +/////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// + + static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArray _result) + { + matchTemplate(_image, _templ, _result, CV_TM_CCORR); + + UMat image_sums; + integral(_image, image_sums); + + int type = image_sums.type(); + int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + + const char * kernelName; + + if (cn==1) + kernelName = "matchTemplate_Prepared_CCOEFF_C1"; + else if (cn==2) + kernelName = "matchTemplate_Prepared_CCOEFF_C2"; + else + kernelName = "matchTemplate_Prepared_CCOEFF_C4"; + + ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn)); + if (k.empty()) + return false; + + UMat templ = _templ.getUMat(), result; + Size size = _image.size(); + _result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F); + result = _result.getUMat(); + + size_t globalsize[2] = {result.cols, result.rows}; + + if (cn==1) + { + float templ_sum = (float)sum(_templ)[0]/ _templ.size().area(); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,NULL,false); + } + else + { + Vec4f templ_sum = Vec4f::all(0); + templ_sum = sum(templ)/ templ.size().area(); + if (cn==2) + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, + templ_sum[0],templ_sum[1]).run(2,globalsize,NULL,false); + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, + templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,NULL,false); + } + } + + static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result) + { + UMat imagef, templf; + + _image.getUMat().convertTo(imagef, CV_32F); + _templ.getUMat().convertTo(templf, CV_32F); + + matchTemplate(imagef, templf, _result, CV_TM_CCORR); + + const char * kernelName; + + UMat temp, image_sums, image_sqsums; + integral(_image,image_sums, image_sqsums, CV_32F, CV_32F); + + int type = image_sums.type(); + int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + + if (cn== 1) + kernelName = "matchTemplate_CCOEFF_NORMED_C1"; + else if (cn==2) + kernelName = "matchTemplate_CCOEFF_NORMED_C2"; + else + kernelName = "matchTemplate_CCOEFF_NORMED_C4"; + + ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, + format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn)); + if (k.empty()) + return false; + + UMat image = _image.getUMat(); + UMat templ = _templ.getUMat(), result; + int image_rows = _image.size().height, image_cols = _image.size().width; + _result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F); + result = _result.getUMat(); + + size_t globalsize[2] = {result.cols, result.rows}; + + float scale = 1.f / templ.size().area(); + + if (cn==1) + { + float templ_sum = (float)sum(templ)[0]; + + multiply(templf, templf, temp); + float templ_sqsum = (float)sum(temp)[0]; + + templ_sqsum -= scale * templ_sum * templ_sum; + templ_sum *= scale; + + if (templ_sqsum < DBL_EPSILON) + { + result = Scalar::all(1); + return true; + } + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums),ocl::KernelArg::ReadOnlyNoSize(image_sqsums), + ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum) + .run(2,globalsize,NULL,false); + } + else + { + Vec4f templ_sum = Vec4f::all(0); + Vec4f templ_sqsum = Vec4f::all(0); + + templ_sum = sum(templ); + + multiply(templf, templf, temp); + templ_sqsum = sum(temp); + + float templ_sqsum_sum = 0; + for(int i = 0; i < cn; i ++) + { + templ_sqsum_sum += templ_sqsum[i] - scale * templ_sum[i] * templ_sum[i]; + } + + templ_sum *= scale; + + if (templ_sqsum_sum < DBL_EPSILON) + { + result = Scalar::all(1); + return true; + } + + if (cn==2) + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), + ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, + templ_sum[0],templ_sum[1], templ_sqsum_sum) + .run(2,globalsize,NULL,false); + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), + ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, + templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3], templ_sqsum_sum) + .run(2,globalsize,NULL,false); + } + + } + +/////////////////////////////////////////////////////////////////////////////////////////////////////////// + + static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method) + { + int cn = CV_MAT_CN(_img.type()); + + if (cn == 3 || cn > 4) + return false; + + typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result); + + const Caller callers[] = + { + matchTemplate_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplate_CCORR, + matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED + }; + + Caller caller = callers[method]; + + return caller(_img, _templ, _result); + } +} namespace cv { @@ -226,15 +557,28 @@ void crossCorr( const Mat& img, const Mat& _templ, Mat& corr, } } } - } -/*****************************************************************************************/ +//////////////////////////////////////////////////////////////////////////////////////////////////////// void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method ) { CV_Assert( CV_TM_SQDIFF <= method && method <= CV_TM_CCOEFF_NORMED ); + CV_Assert( (_img.depth() == CV_8U || _img.depth() == CV_32F) && _img.type() == _templ.type() ); + + CV_Assert(_img.dims() <= 2); + + bool swapNotNeed = (_img.size().height >= _templ.size().height && _img.size().width >= _templ.size().width); + if (!swapNotNeed) + { + CV_Assert(_img.size().height <= _templ.size().height && _img.size().width <= _templ.size().width); + } + + bool use_opencl = ocl::useOpenCL() && _result.isUMat(); + if ( use_opencl && (swapNotNeed ? ocl_matchTemplate(_img,_templ,_result,method) : ocl_matchTemplate(_templ,_img,_result,method))) + return; + int numType = method == CV_TM_CCORR || method == CV_TM_CCORR_NORMED ? 0 : method == CV_TM_CCOEFF || method == CV_TM_CCOEFF_NORMED ? 1 : 2; bool isNormed = method == CV_TM_CCORR_NORMED || @@ -242,14 +586,9 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, method == CV_TM_CCOEFF_NORMED; Mat img = _img.getMat(), templ = _templ.getMat(); - if( img.rows < templ.rows || img.cols < templ.cols ) + if(!swapNotNeed ) std::swap(img, templ); - CV_Assert( (img.depth() == CV_8U || img.depth() == CV_32F) && - img.type() == templ.type() ); - - CV_Assert( img.rows >= templ.rows && img.cols >= templ.cols); - Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1); _result.create(corrSize, CV_32F); Mat result = _result.getMat(); diff --git a/modules/imgproc/test/ocl/test_match_template.cpp b/modules/imgproc/test/ocl/test_match_template.cpp new file mode 100644 index 000000000..782509cc9 --- /dev/null +++ b/modules/imgproc/test/ocl/test_match_template.cpp @@ -0,0 +1,128 @@ +/*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) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, 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 "test_precomp.hpp" +#include "opencv2/ts/ocl_test.hpp" +#include "iostream" +#include "fstream" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +/////////////////////////////////////////////matchTemplate////////////////////////////////////////////////////////// + +PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, int, bool) +{ + int type; + int depth; + int method; + bool use_roi; + + TEST_DECLARE_INPUT_PARAMETER(image) + TEST_DECLARE_INPUT_PARAMETER(templ) + TEST_DECLARE_OUTPUT_PARAMETER(result) + + virtual void SetUp() + { + type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1)); + depth = GET_PARAM(0); + method = GET_PARAM(2); + use_roi = GET_PARAM(3); + } + + virtual void generateTestData() + { + Size image_roiSize = randomSize(2, 100); + Size templ_roiSize = Size(randomInt(1, image_roiSize.width), randomInt(1, image_roiSize.height)); + Size result_roiSize = Size(image_roiSize.width - templ_roiSize.width + 1, + image_roiSize.height - templ_roiSize.height + 1); + + const double upValue = 256; + + Border imageBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(image, image_roi, image_roiSize, imageBorder, type, -upValue, upValue); + + Border templBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(templ, templ_roi, templ_roiSize, templBorder, type, -upValue, upValue); + + Border resultBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(result, result_roi, result_roiSize, resultBorder, CV_32F, -upValue, upValue); + + UMAT_UPLOAD_INPUT_PARAMETER(image) + UMAT_UPLOAD_INPUT_PARAMETER(templ) + UMAT_UPLOAD_OUTPUT_PARAMETER(result) + } + + void Near(double threshold = 0.0) + { + OCL_EXPECT_MATS_NEAR(result,threshold); + } +}; + +OCL_TEST_P(MatchTemplate, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::matchTemplate(image_roi,templ_roi,result_roi, method)); + OCL_ON(cv::matchTemplate(uimage_roi,utempl_roi,uresult_roi, method)); + + if (method == 0) + Near(10.0f); + else + Near(method % 2 == 1 ? 0.001f : 1.0f); + } +} + +OCL_INSTANTIATE_TEST_CASE_P(ImageProc, MatchTemplate, Combine( + Values(CV_8U, CV_32F), + Values(1, 2, 4), + Values(0,1,2,3,4,5), + Bool()) + ); +} } // namespace cvtest::ocl + +#endif \ No newline at end of file