From 86636dc265a0d2c5733293a433d15c0ad73e4b10 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Tue, 14 Jan 2014 14:45:01 +0400 Subject: [PATCH 01/13] Added ocl_matchTemplate( without dft) --- modules/imgproc/src/opencl/match_template.cl | 427 ++++++++++++++++++ modules/imgproc/src/templmatch.cpp | 377 +++++++++++++++- .../imgproc/test/ocl/test_match_template.cpp | 130 ++++++ 3 files changed, 927 insertions(+), 7 deletions(-) create mode 100644 modules/imgproc/src/opencl/match_template.cl create mode 100644 modules/imgproc/test/ocl/test_match_template.cpp diff --git a/modules/imgproc/src/opencl/match_template.cl b/modules/imgproc/src/opencl/match_template.cl new file mode 100644 index 000000000..b24ec2c48 --- /dev/null +++ b/modules/imgproc/src/opencl/match_template.cl @@ -0,0 +1,427 @@ +// 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_TYPE type +#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 + img_sqsums_offset + ox) * CN) +#define SQSUMS(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; + + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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 ++) + + 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); + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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; + + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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 ++) + + 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); + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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; + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + 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; + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + 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; + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + 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])); + image_sum_ += tpl_sum_2 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+2] - sum[SUMS_PTR(tpl_cols, 0)+2])-(sum[SUMS_PTR(0, tpl_rows)+2] - sum[SUMS_PTR(0, 0)+2])); + image_sum_ += tpl_sum_3 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+3] - sum[SUMS_PTR(tpl_cols, 0)+3])-(sum[SUMS_PTR(0, tpl_rows)+3] - sum[SUMS_PTR(0, 0)+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 float * 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); + res_step /= sizeof(*res); + res_offset /= sizeof(*res); + + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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 float * 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); + res_step /= sizeof(*res); + res_offset /= sizeof(*res); + + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(0, 0)])); + sqsum_[1] = (float)((sqsum[SQSUMS(t_cols, t_rows)+1] - sqsum[SQSUMS(t_cols, 0)+1])-(sqsum[SQSUMS(0, t_rows)+1] - sqsum[SQSUMS(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 float * 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); + res_step /= sizeof(*res); + res_offset /= sizeof(*res); + + + int res_idx = mad24(gidy, res_step, res_offset + gidx); + + 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); + + 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])); + sum_[2] = (float)((sum[SUMS_PTR(t_cols, t_rows)+2] - sum[SUMS_PTR(t_cols, 0)+2])-(sum[SUMS_PTR(0, t_rows)+2] - sum[SUMS_PTR(0, 0)+2])); + sum_[3] = (float)((sum[SUMS_PTR(t_cols, t_rows)+3] - sum[SUMS_PTR(t_cols, 0)+3])-(sum[SUMS_PTR(0, t_rows)+3] - sum[SUMS_PTR(0, 0)+3])); + + sqsum_[0] = (float)((sqsum[SQSUMS(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(0, 0)])); + sqsum_[1] = (float)((sqsum[SQSUMS(t_cols, t_rows)+1] - sqsum[SQSUMS(t_cols, 0)+1])-(sqsum[SQSUMS(0, t_rows)+1] - sqsum[SQSUMS(0, 0)+1])); + sqsum_[2] = (float)((sqsum[SQSUMS(t_cols, t_rows)+2] - sqsum[SQSUMS(t_cols, 0)+2])-(sqsum[SQSUMS(0, t_rows)+2] - sqsum[SQSUMS(0, 0)+2])); + sqsum_[3] = (float)((sqsum[SQSUMS(t_cols, t_rows)+3] - sqsum[SQSUMS(t_cols, 0)+3])-(sqsum[SQSUMS(0, t_rows)+3] - sqsum[SQSUMS(0, 0)+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); + } +} + +//////////////////////////////////////////// extractFirstChannel///////////////////////////// +__kernel void extractFirstChannel( const __global float4* img, int img_step, int img_offset, + __global float* res, int res_step, int res_offset, int rows, int cols) +{ + img_step /= sizeof(float4); + img_offset /= sizeof(float4); + res_step /= sizeof(float); + res_offset /= sizeof(float); + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + if(gidx < cols && gidy < rows) + { + __global const float4 * image = (__global const float4 *)(img) + mad24(gidy, img_step, img_offset + gidx); + __global float * result = (__global float *)(res)+ mad24(gidy, res_step, res_offset + gidx); + *result = image[0].x; + } +} \ No newline at end of file diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index bfe7ce600..3e31b9c76 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -40,6 +40,365 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" + +//////////////////////////////////////////////////matchTemplate////////////////////////////////////////////////////////// +namespace cv +{ + struct MatchTemplateBuf + { + Size user_block_size; + UMat imagef, templf; + UMat image_sums; + UMat image_sqsums; + }; + + static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + + static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + + static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + + static bool matchTemplateNaive_CCORR (const UMat &image, const UMat &templ, UMat &result, int cn); + static bool matchTemplateNaive_SQDIFF(const UMat &image, const UMat &templ, UMat &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 extractFirstChannel_32F(const UMat &image, UMat &result) + { + const char * kernelName = "extractFirstChannel"; + int type = image.type(); + int depth = CV_MAT_DEPTH(type); + int cn = CV_MAT_CN(type); + + 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; + + size_t globalsize[2] = {result.cols, result.rows}; + size_t localsize[2] = {16, 16}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); + } + + static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + { + if (useNaive(TM_CCORR, image.depth(), templ.size()) ) + return matchTemplateNaive_CCORR(image, templ, result, image.channels()); + else + return false; + } + + static bool matchTemplateNaive_CCORR (const UMat &image, const UMat &templ, UMat &result, int cn) + { + int type = image.type(); + int depth = CV_MAT_DEPTH(type); + + CV_Assert(result.channels() == 1); + + 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; + + size_t globalsize[2] = {result.cols, result.rows}; + size_t localsize[2] = {16, 16}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); + } + + static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + { + if (!matchTemplate_CCORR(image, templ, result, buf)) + return false; + + 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 temp; + integral(image.reshape(1), buf.image_sums, temp); + + if(temp.depth() == CV_64F) + temp.convertTo(buf.image_sqsums, CV_32F); + else + buf.image_sqsums = temp; + + UMat 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}; + size_t localsize[2] = {16, 16}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); + } + +//////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// + + static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + { + if (useNaive(TM_SQDIFF, image.depth(), templ.size())) + { + return matchTemplateNaive_SQDIFF(image, templ, result, image.channels());; + } + else + return false; + } + + static bool matchTemplateNaive_SQDIFF(const UMat &image, const UMat &templ, UMat &result, int cn) + { + int type = image.type(); + int depth = CV_MAT_DEPTH(type); + + CV_Assert(result.channels() == 1); + + 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; + + size_t globalsize[2] = {result.cols, result.rows}; + size_t localsize[2] = {16, 16}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); + } + + static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + { + if (!matchTemplate_CCORR(image, templ, result, buf)) + return false; + + 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 temp; + integral(image.reshape(1), buf.image_sums, temp); + + if(temp.depth() == CV_64F) + temp.convertTo(buf.image_sqsums, CV_32F); + else + buf.image_sqsums = temp; + + UMat 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}; + size_t localsize[2] = {16, 16}; + + return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); + } + +/////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// + + static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + { + if (!matchTemplate_CCORR(image, templ, result, buf)) + return false; + + integral(image, buf.image_sums); + + int type = buf.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; + + size_t globalsize[2] = {result.cols, result.rows}; + size_t localsize[2] = {16, 16}; + + if (cn==1) + { + float templ_sum = (float)sum(templ)[0]/ templ.size().area(); + return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,localsize,true); + } + else + { + Vec4f templ_sum = Vec4f::all(0); + templ_sum = sum(templ)/ templ.size().area(); + if (cn==2) + return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, + templ_sum[0],templ_sum[1]).run(2,globalsize,localsize,true); + + return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.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,localsize,true); + } + } + + static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + { + image.convertTo(buf.imagef, CV_32F); + templ.convertTo(buf.templf, CV_32F); + + if(!matchTemplate_CCORR(buf.imagef, buf.templf, result, buf)) + return false; + + const char * kernelName; + + UMat temp; + integral(image, buf.image_sums, temp); + + int type = buf.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; + + if(temp.depth() == CV_64F) + temp.convertTo(buf.image_sqsums, CV_32F); + else + buf.image_sqsums = temp; + + size_t globalsize[2] = {result.cols, result.rows}; + size_t localsize[2] = {16, 16}; + + float scale = 1.f / templ.size().area(); + + if (cn==1) + { + float templ_sum = (float)sum(templ)[0]; + + multiply(buf.templf, buf.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(buf.image_sums),ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), + ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum) + .run(2,globalsize,localsize,true); + } + else + { + Vec4f templ_sum = Vec4f::all(0); + Vec4f templ_sqsum = Vec4f::all(0); + + templ_sum = sum(templ); + + multiply(buf.templf, buf.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(buf.image_sums), ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), + ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, + templ_sum[0],templ_sum[1], templ_sqsum_sum) + .run(2,globalsize,localsize,true); + + return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::ReadOnlyNoSize(buf.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,localsize,true); + } + + } + +/////////////////////////////////////////////////////////////////////////////////////////////////////////// + + static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method) + { + int type = _img.type(); + int cn = CV_MAT_CN(type); + + CV_Assert( cn == _templ.channels() && cn!=3 && cn<=4); + + typedef bool (*Caller)(const UMat &, const UMat &, UMat &, MatchTemplateBuf &); + + const Caller callers[] = + { + matchTemplate_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplate_CCORR, + matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED + }; + + Caller caller; + if (!(caller = callers[method])) + return false; + + MatchTemplateBuf buf; + + UMat image = _img.getUMat(); + UMat templ = _templ.getUMat(), result; + _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); + result = _result.getUMat(); + return caller(image, templ, result, buf); + } +} namespace cv { @@ -226,15 +585,24 @@ 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.size().height >= _templ.size().height && _img.size().width >= _templ.size().width); + + CV_Assert(_img.dims() <= 2); + + bool use_opencl = ocl::useOpenCL() && _result.isUMat(); + if ( use_opencl && ocl_matchTemplate(_img,_templ,_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 || @@ -245,11 +613,6 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, if( img.rows < templ.rows || img.cols < templ.cols ) 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..2163ac36c --- /dev/null +++ b/modules/imgproc/test/ocl/test_match_template.cpp @@ -0,0 +1,130 @@ +/*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, 20); + 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; + const double max_val = 100; + + Border imageBorder = randomBorder(0, use_roi ? max_val : 0); + randomSubMat(image, image_roi, image_roiSize, imageBorder, type, -upValue, upValue); + + Border templBorder = randomBorder(0, use_roi ? max_val : 0); + randomSubMat(templ, templ_roi, templ_roiSize, templBorder, type, -upValue, upValue); + + Border resultBorder = randomBorder(0, use_roi ? max_val : 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) + { + EXPECT_MAT_NEAR(result, uresult, threshold); + EXPECT_MAT_NEAR(result_roi, uresult_roi, 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 From ba5d480f80bee48e238a7ff42b6097cbd0bc4ccf Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Tue, 14 Jan 2014 16:31:04 +0400 Subject: [PATCH 02/13] Fixed warnings --- modules/imgproc/src/opencl/match_template.cl | 20 --- modules/imgproc/src/templmatch.cpp | 123 ++++++++----------- 2 files changed, 48 insertions(+), 95 deletions(-) diff --git a/modules/imgproc/src/opencl/match_template.cl b/modules/imgproc/src/opencl/match_template.cl index b24ec2c48..26a9cbc1e 100644 --- a/modules/imgproc/src/opencl/match_template.cl +++ b/modules/imgproc/src/opencl/match_template.cl @@ -404,24 +404,4 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i __global float * result = (__global float *)(res)+res_idx; *result = normAcc((*result) - num, denum); } -} - -//////////////////////////////////////////// extractFirstChannel///////////////////////////// -__kernel void extractFirstChannel( const __global float4* img, int img_step, int img_offset, - __global float* res, int res_step, int res_offset, int rows, int cols) -{ - img_step /= sizeof(float4); - img_offset /= sizeof(float4); - res_step /= sizeof(float); - res_offset /= sizeof(float); - - int gidx = get_global_id(0); - int gidy = get_global_id(1); - - if(gidx < cols && gidy < rows) - { - __global const float4 * image = (__global const float4 *)(img) + mad24(gidy, img_step, img_offset + gidx); - __global float * result = (__global float *)(res)+ mad24(gidy, res_step, res_offset + gidx); - *result = image[0].x; - } } \ No newline at end of file diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index 3e31b9c76..ea6bddc8d 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -45,22 +45,14 @@ //////////////////////////////////////////////////matchTemplate////////////////////////////////////////////////////////// namespace cv { - struct MatchTemplateBuf - { - Size user_block_size; - UMat imagef, templf; - UMat image_sums; - UMat image_sqsums; - }; + static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result); + static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result); - static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); - static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result); + static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result); - static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); - static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); - - static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); - static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf); + static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result); + static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result); static bool matchTemplateNaive_CCORR (const UMat &image, const UMat &templ, UMat &result, int cn); static bool matchTemplateNaive_SQDIFF(const UMat &image, const UMat &templ, UMat &result, int cn); @@ -84,24 +76,7 @@ namespace cv ///////////////////////////////////////////////////CCORR////////////////////////////////////////////////////////////// - static bool extractFirstChannel_32F(const UMat &image, UMat &result) - { - const char * kernelName = "extractFirstChannel"; - int type = image.type(); - int depth = CV_MAT_DEPTH(type); - int cn = CV_MAT_CN(type); - - 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; - - size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); - } - - static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result) { if (useNaive(TM_CCORR, image.depth(), templ.size()) ) return matchTemplateNaive_CCORR(image, templ, result, image.channels()); @@ -128,9 +103,9 @@ namespace cv return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); } - static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result) { - if (!matchTemplate_CCORR(image, templ, result, buf)) + if (!matchTemplate_CCORR(image, templ, result)) return false; int type = image.type(); @@ -142,13 +117,13 @@ namespace cv if (k.empty()) return false; - UMat temp; - integral(image.reshape(1), buf.image_sums, temp); + UMat temp, image_sums, image_sqsums; + integral(image.reshape(1), image_sums, temp); if(temp.depth() == CV_64F) - temp.convertTo(buf.image_sqsums, CV_32F); + temp.convertTo(image_sqsums, CV_32F); else - buf.image_sqsums = temp; + image_sqsums = temp; UMat templ_resh; templ.reshape(1).convertTo(templ_resh, CV_32F); @@ -159,12 +134,12 @@ namespace cv size_t globalsize[2] = {result.cols, result.rows}; size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); } //////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// - static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result) { if (useNaive(TM_SQDIFF, image.depth(), templ.size())) { @@ -193,9 +168,9 @@ namespace cv return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); } - static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result) { - if (!matchTemplate_CCORR(image, templ, result, buf)) + if (!matchTemplate_CCORR(image, templ, result)) return false; int type = image.type(); @@ -207,13 +182,13 @@ namespace cv if (k.empty()) return false; - UMat temp; - integral(image.reshape(1), buf.image_sums, temp); + UMat temp, image_sums, image_sqsums; + integral(image.reshape(1), image_sums, temp); if(temp.depth() == CV_64F) - temp.convertTo(buf.image_sqsums, CV_32F); + temp.convertTo(image_sqsums, CV_32F); else - buf.image_sqsums = temp; + image_sqsums = temp; UMat templ_resh; templ.reshape(1).convertTo(templ_resh, CV_32F); @@ -224,19 +199,20 @@ namespace cv size_t globalsize[2] = {result.cols, result.rows}; size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); } /////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// - static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result) { - if (!matchTemplate_CCORR(image, templ, result, buf)) + if (!matchTemplate_CCORR(image, templ, result)) return false; - integral(image, buf.image_sums); + UMat image_sums; + integral(image, image_sums); - int type = buf.image_sums.type(); + int type = image_sums.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); const char * kernelName; @@ -258,35 +234,36 @@ namespace cv if (cn==1) { float templ_sum = (float)sum(templ)[0]/ templ.size().area(); - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,localsize,true); } else { Vec4f templ_sum = Vec4f::all(0); templ_sum = sum(templ)/ templ.size().area(); if (cn==2) - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, + 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,localsize,true); - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, + 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,localsize,true); } } - static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result, MatchTemplateBuf &buf) + static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result) { - image.convertTo(buf.imagef, CV_32F); - templ.convertTo(buf.templf, CV_32F); + UMat imagef, templf; + image.convertTo(imagef, CV_32F); + templ.convertTo(templf, CV_32F); - if(!matchTemplate_CCORR(buf.imagef, buf.templf, result, buf)) + if(!matchTemplate_CCORR(imagef, templf, result)) return false; const char * kernelName; - UMat temp; - integral(image, buf.image_sums, temp); + UMat temp, image_sums, image_sqsums; + integral(image,image_sums, temp); - int type = buf.image_sums.type(); + int type = image_sums.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); if (cn== 1) @@ -302,9 +279,9 @@ namespace cv return false; if(temp.depth() == CV_64F) - temp.convertTo(buf.image_sqsums, CV_32F); + temp.convertTo(image_sqsums, CV_32F); else - buf.image_sqsums = temp; + image_sqsums = temp; size_t globalsize[2] = {result.cols, result.rows}; size_t localsize[2] = {16, 16}; @@ -315,7 +292,7 @@ namespace cv { float templ_sum = (float)sum(templ)[0]; - multiply(buf.templf, buf.templf, temp); + multiply(templf, templf, temp); float templ_sqsum = (float)sum(temp)[0]; templ_sqsum -= scale * templ_sum * templ_sum; @@ -327,7 +304,7 @@ namespace cv return true; } - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums),ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), + 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,localsize,true); } @@ -338,7 +315,7 @@ namespace cv templ_sum = sum(templ); - multiply(buf.templf, buf.templf, temp); + multiply(templf, templf, temp); templ_sqsum = sum(temp); float templ_sqsum_sum = 0; @@ -356,12 +333,12 @@ namespace cv } if (cn==2) - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), + 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,localsize,true); - return k.args(ocl::KernelArg::ReadOnlyNoSize(buf.image_sums), ocl::KernelArg::ReadOnlyNoSize(buf.image_sqsums), + 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,localsize,true); @@ -378,7 +355,7 @@ namespace cv CV_Assert( cn == _templ.channels() && cn!=3 && cn<=4); - typedef bool (*Caller)(const UMat &, const UMat &, UMat &, MatchTemplateBuf &); + typedef bool (*Caller)(const UMat &, const UMat &, UMat &); const Caller callers[] = { @@ -386,17 +363,13 @@ namespace cv matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED }; - Caller caller; - if (!(caller = callers[method])) - return false; - - MatchTemplateBuf buf; + Caller caller = callers[method]; UMat image = _img.getUMat(); UMat templ = _templ.getUMat(), result; _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - return caller(image, templ, result, buf); + return caller(image, templ, result); } } From 9de70c9a00e60f0ae333e23de939a2fb4c7300e9 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Tue, 14 Jan 2014 17:34:05 +0400 Subject: [PATCH 03/13] Fixed trailing whitespace --- modules/imgproc/src/opencl/match_template.cl | 4 ++-- modules/imgproc/src/templmatch.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/modules/imgproc/src/opencl/match_template.cl b/modules/imgproc/src/opencl/match_template.cl index 26a9cbc1e..89a8d618f 100644 --- a/modules/imgproc/src/opencl/match_template.cl +++ b/modules/imgproc/src/opencl/match_template.cl @@ -93,7 +93,7 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step 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; @@ -300,7 +300,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i __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)]) - + 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)]) - diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index ea6bddc8d..1abb8fb2f 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -273,7 +273,7 @@ namespace cv else kernelName = "matchTemplate_CCOEFF_NORMED_C4"; - ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, + 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; From da4d33ec2699575311474ccb42d4fa456d1fae56 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Wed, 15 Jan 2014 10:24:14 +0400 Subject: [PATCH 04/13] Fixed test_match_template --- modules/imgproc/test/ocl/test_match_template.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/modules/imgproc/test/ocl/test_match_template.cpp b/modules/imgproc/test/ocl/test_match_template.cpp index 2163ac36c..48a8027b2 100644 --- a/modules/imgproc/test/ocl/test_match_template.cpp +++ b/modules/imgproc/test/ocl/test_match_template.cpp @@ -75,20 +75,19 @@ PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, int, bool) virtual void generateTestData() { Size image_roiSize = randomSize(2, 20); - Size templ_roiSize = Size (randomInt(1,image_roiSize.width), randomInt(1,image_roiSize.height)); + 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; - const double max_val = 100; - Border imageBorder = randomBorder(0, use_roi ? max_val : 0); + 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_val : 0); + 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_val : 0); + 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) From 47b572f99f4f818ac4cc3ac8c7ac50228aa64a0d Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Thu, 16 Jan 2014 14:26:48 +0400 Subject: [PATCH 05/13] fixed --- modules/imgproc/src/templmatch.cpp | 134 ++++++++++-------- .../imgproc/test/ocl/test_match_template.cpp | 2 +- 2 files changed, 78 insertions(+), 58 deletions(-) diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index 1abb8fb2f..ae10cd929 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -45,17 +45,17 @@ //////////////////////////////////////////////////matchTemplate////////////////////////////////////////////////////////// namespace cv { - static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result); - static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result); + 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(const UMat &image, const UMat &templ, UMat &result); - static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &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(const UMat &image, const UMat &templ, UMat &result); - static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &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 (const UMat &image, const UMat &templ, UMat &result, int cn); - static bool matchTemplateNaive_SQDIFF(const UMat &image, const UMat &templ, UMat &result, int cn); + 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) { @@ -76,39 +76,41 @@ namespace cv ///////////////////////////////////////////////////CCORR////////////////////////////////////////////////////////////// - static bool matchTemplate_CCORR(const UMat &image, const UMat &templ, UMat &result) + 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()); + if (useNaive(TM_CCORR, _image.depth(), _templ.size()) ) + return matchTemplateNaive_CCORR(_image, _templ, _result, _image.channels()); else return false; } - static bool matchTemplateNaive_CCORR (const UMat &image, const UMat &templ, UMat &result, int cn) + static bool matchTemplateNaive_CCORR (InputArray _image, InputArray _templ, OutputArray _result, int cn) { - int type = image.type(); + int type = _image.type(); int depth = CV_MAT_DEPTH(type); - CV_Assert(result.channels() == 1); - 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}; size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); } - static bool matchTemplate_CCORR_NORMED(const UMat &image, const UMat &templ, UMat &result) + static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result) { - if (!matchTemplate_CCORR(image, templ, result)) - return false; + matchTemplate(_image, _templ, _result, CV_TM_CCORR); - int type = image.type(); + int type = _image.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); const char * kernelName = "matchTemplate_CCORR_NORMED"; @@ -117,6 +119,11 @@ namespace cv 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 temp, image_sums, image_sqsums; integral(image.reshape(1), image_sums, temp); @@ -134,46 +141,48 @@ namespace cv size_t globalsize[2] = {result.cols, result.rows}; size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); } //////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// - static bool matchTemplate_SQDIFF(const UMat &image, const UMat &templ, UMat &result) + static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result) { - if (useNaive(TM_SQDIFF, image.depth(), templ.size())) + if (useNaive(TM_SQDIFF, _image.depth(), _templ.size())) { - return matchTemplateNaive_SQDIFF(image, templ, result, image.channels());; + return matchTemplateNaive_SQDIFF(_image, _templ, _result, _image.channels());; } else return false; } - static bool matchTemplateNaive_SQDIFF(const UMat &image, const UMat &templ, UMat &result, int cn) + static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result, int cn) { - int type = image.type(); + int type = _image.type(); int depth = CV_MAT_DEPTH(type); - CV_Assert(result.channels() == 1); - 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}; size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); } - static bool matchTemplate_SQDIFF_NORMED (const UMat &image, const UMat &templ, UMat &result) + static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result) { - if (!matchTemplate_CCORR(image, templ, result)) - return false; + matchTemplate(_image, _templ, _result, CV_TM_CCORR); - int type = image.type(); + int type = _image.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); const char * kernelName = "matchTemplate_SQDIFF_NORMED"; @@ -182,6 +191,11 @@ namespace cv 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 temp, image_sums, image_sqsums; integral(image.reshape(1), image_sums, temp); @@ -199,18 +213,17 @@ namespace cv size_t globalsize[2] = {result.cols, result.rows}; size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,true); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); } /////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// - static bool matchTemplate_CCOEFF(const UMat &image, const UMat &templ, UMat &result) + static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArray _result) { - if (!matchTemplate_CCORR(image, templ, result)) - return false; + matchTemplate(_image, _templ, _result, CV_TM_CCORR); UMat image_sums; - integral(image, image_sums); + integral(_image, image_sums); int type = image_sums.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); @@ -228,13 +241,18 @@ namespace cv if (k.empty()) return false; + 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}; size_t localsize[2] = {16, 16}; 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,localsize,true); + 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,localsize,false); } else { @@ -242,26 +260,26 @@ namespace cv 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,localsize,true); + templ_sum[0],templ_sum[1]).run(2,globalsize,localsize,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,localsize,true); + templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,localsize,false); } } - static bool matchTemplate_CCOEFF_NORMED(const UMat &image, const UMat &templ, UMat &result) + static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result) { UMat imagef, templf; - image.convertTo(imagef, CV_32F); - templ.convertTo(templf, CV_32F); - if(!matchTemplate_CCORR(imagef, templf, result)) - return false; + _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, temp); + integral(_image,image_sums, temp); int type = image_sums.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); @@ -278,6 +296,12 @@ namespace cv 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(); + if(temp.depth() == CV_64F) temp.convertTo(image_sqsums, CV_32F); else @@ -306,7 +330,7 @@ namespace cv 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,localsize,true); + .run(2,globalsize,localsize,false); } else { @@ -336,12 +360,12 @@ namespace cv 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,localsize,true); + .run(2,globalsize,localsize,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,localsize,true); + .run(2,globalsize,localsize,false); } } @@ -355,7 +379,7 @@ namespace cv CV_Assert( cn == _templ.channels() && cn!=3 && cn<=4); - typedef bool (*Caller)(const UMat &, const UMat &, UMat &); + typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result); const Caller callers[] = { @@ -365,11 +389,7 @@ namespace cv Caller caller = callers[method]; - UMat image = _img.getUMat(); - UMat templ = _templ.getUMat(), result; - _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); - result = _result.getUMat(); - return caller(image, templ, result); + return caller(_img, _templ, _result); } } diff --git a/modules/imgproc/test/ocl/test_match_template.cpp b/modules/imgproc/test/ocl/test_match_template.cpp index 48a8027b2..2a0402342 100644 --- a/modules/imgproc/test/ocl/test_match_template.cpp +++ b/modules/imgproc/test/ocl/test_match_template.cpp @@ -74,7 +74,7 @@ PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, int, bool) virtual void generateTestData() { - Size image_roiSize = randomSize(2, 20); + 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); From 4da1ba56b3c705bd9217ee16d810d754368ff5de Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 17 Jan 2014 15:01:22 +0400 Subject: [PATCH 06/13] fixed --- modules/imgproc/src/opencl/match_template.cl | 126 +++++++++---------- modules/imgproc/src/templmatch.cpp | 77 +++++------- 2 files changed, 87 insertions(+), 116 deletions(-) diff --git a/modules/imgproc/src/opencl/match_template.cl b/modules/imgproc/src/opencl/match_template.cl index 89a8d618f..451ec0cb2 100644 --- a/modules/imgproc/src/opencl/match_template.cl +++ b/modules/imgproc/src/opencl/match_template.cl @@ -29,15 +29,13 @@ // 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_TYPE type #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 + img_sqsums_offset + ox) * CN) -#define SQSUMS(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)) +#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) { @@ -76,10 +74,7 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step int i,j; float sum = 0; - res_step /= sizeof(float); - res_offset /= sizeof(float); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -90,12 +85,13 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step 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; + __global float * result = (__global float *)(res+res_idx); *result = sum; } } @@ -109,10 +105,8 @@ __kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, in img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -121,7 +115,7 @@ __kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, in (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; + __global float * result = (__global float *)(res+res_idx); *result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum)); } } @@ -138,10 +132,7 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step float delta; float sum = 0; - res_step /= sizeof(float); - res_offset /= sizeof(float); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -152,13 +143,14 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step 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; + __global float * result = (__global float *)(res+res_idx); *result = sum; } } @@ -172,10 +164,8 @@ __kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, i img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -184,7 +174,7 @@ __kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, i (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; + __global float * result = (__global float *)(res+res_idx); *result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum)); } @@ -201,10 +191,8 @@ __kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, img_sums_step /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE; - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float image_sum_ = 0; if(gidx < res_cols && gidy < res_rows) @@ -214,7 +202,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * 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; + __global float * result = (__global float *)(res+res_idx); *result -= image_sum_; } } @@ -228,10 +216,8 @@ __kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, img_sums_step /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE; - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float image_sum_ = 0; if(gidx < res_cols && gidy < res_rows) @@ -241,37 +227,40 @@ __kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * 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; + __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) + __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; - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + 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])); - image_sum_ += tpl_sum_2 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+2] - sum[SUMS_PTR(tpl_cols, 0)+2])-(sum[SUMS_PTR(0, tpl_rows)+2] - sum[SUMS_PTR(0, 0)+2])); - image_sum_ += tpl_sum_3 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+3] - sum[SUMS_PTR(tpl_cols, 0)+3])-(sum[SUMS_PTR(0, tpl_rows)+3] - sum[SUMS_PTR(0, 0)+3])); + 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); - __global float * result = (__global float *)(res)+res_idx; + 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_; } @@ -279,7 +268,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, __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 float * res, int res_step, int res_offset, int res_rows, int res_cols, + __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); @@ -289,11 +278,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i img_sums_step /= ELEM_SIZE; img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(*res); - res_offset /= sizeof(*res); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -306,7 +292,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i 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; + __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_))); @@ -315,7 +301,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i __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 float * res, int res_step, int res_offset, int res_rows, int res_cols, + __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); @@ -325,11 +311,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i img_sums_step /= ELEM_SIZE; img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(*res); - res_offset /= sizeof(*res); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float sum_[2]; float sqsum_[2]; @@ -342,22 +325,22 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i 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(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(0, 0)])); - sqsum_[1] = (float)((sqsum[SQSUMS(t_cols, t_rows)+1] - sqsum[SQSUMS(t_cols, 0)+1])-(sqsum[SQSUMS(0, t_rows)+1] - sqsum[SQSUMS(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; + __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 float * res, int res_step, int res_offset, int res_rows, int res_cols, + __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) @@ -369,11 +352,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i img_sums_step /= ELEM_SIZE; img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(*res); - res_offset /= sizeof(*res); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float sum_[4]; float sqsum_[4]; @@ -383,15 +363,25 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i __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])); - sum_[2] = (float)((sum[SUMS_PTR(t_cols, t_rows)+2] - sum[SUMS_PTR(t_cols, 0)+2])-(sum[SUMS_PTR(0, t_rows)+2] - sum[SUMS_PTR(0, 0)+2])); - sum_[3] = (float)((sum[SUMS_PTR(t_cols, t_rows)+3] - sum[SUMS_PTR(t_cols, 0)+3])-(sum[SUMS_PTR(0, t_rows)+3] - sum[SUMS_PTR(0, 0)+3])); + 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); - sqsum_[0] = (float)((sqsum[SQSUMS(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(0, 0)])); - sqsum_[1] = (float)((sqsum[SQSUMS(t_cols, t_rows)+1] - sqsum[SQSUMS(t_cols, 0)+1])-(sqsum[SQSUMS(0, t_rows)+1] - sqsum[SQSUMS(0, 0)+1])); - sqsum_[2] = (float)((sqsum[SQSUMS(t_cols, t_rows)+2] - sqsum[SQSUMS(t_cols, 0)+2])-(sqsum[SQSUMS(0, t_rows)+2] - sqsum[SQSUMS(0, 0)+2])); - sqsum_[3] = (float)((sqsum[SQSUMS(t_cols, t_rows)+3] - sqsum[SQSUMS(t_cols, 0)+3])-(sqsum[SQSUMS(0, t_rows)+3] - sqsum[SQSUMS(0, 0)+3])); + 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; @@ -401,7 +391,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i sqsum_[2] - weight * sum_[2]* sum_[2] + sqsum_[3] - weight * sum_[3]* sum_[3] )); - __global float * result = (__global float *)(res)+res_idx; + __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 ae10cd929..bb3423bb2 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -101,9 +101,8 @@ namespace cv result = _result.getUMat(); size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); + 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) @@ -124,24 +123,18 @@ namespace cv _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - UMat temp, image_sums, image_sqsums; - integral(image.reshape(1), image_sums, temp); + UMat image_sums, image_sqsums; + integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); - if(temp.depth() == CV_64F) - temp.convertTo(image_sqsums, CV_32F); - else - image_sqsums = temp; - - UMat templ_resh; + 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}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false); } //////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// @@ -173,9 +166,8 @@ namespace cv result = _result.getUMat(); size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); + 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) @@ -196,24 +188,18 @@ namespace cv _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - UMat temp, image_sums, image_sqsums; - integral(image.reshape(1), image_sums, temp); + UMat image_sums, image_sqsums; + integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); - if(temp.depth() == CV_64F) - temp.convertTo(image_sqsums, CV_32F); - else - image_sqsums = temp; - - UMat templ_resh; + 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}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false); } /////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// @@ -242,17 +228,16 @@ namespace cv return false; 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); + 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}; - size_t localsize[2] = {16, 16}; 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,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,NULL,false); } else { @@ -260,10 +245,10 @@ namespace cv 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,localsize,false); + 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,localsize,false); + templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,NULL,false); } } @@ -279,7 +264,7 @@ namespace cv const char * kernelName; UMat temp, image_sums, image_sqsums; - integral(_image,image_sums, temp); + 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); @@ -302,13 +287,7 @@ namespace cv _result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - if(temp.depth() == CV_64F) - temp.convertTo(image_sqsums, CV_32F); - else - image_sqsums = temp; - size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; float scale = 1.f / templ.size().area(); @@ -330,7 +309,7 @@ namespace cv 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,localsize,false); + .run(2,globalsize,NULL,false); } else { @@ -360,12 +339,12 @@ namespace cv 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,localsize,false); + .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,localsize,false); + .run(2,globalsize,NULL,false); } } @@ -374,10 +353,10 @@ namespace cv static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method) { - int type = _img.type(); - int cn = CV_MAT_CN(type); + int cn = CV_MAT_CN(_img.type()); - CV_Assert( cn == _templ.channels() && cn!=3 && cn<=4); + if (cn == 3 || cn > 4) + return false; typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result); @@ -588,13 +567,15 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, CV_Assert( (_img.depth() == CV_8U || _img.depth() == CV_32F) && _img.type() == _templ.type() ); - CV_Assert(_img.size().height >= _templ.size().height && _img.size().width >= _templ.size().width); - 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 && ocl_matchTemplate(_img,_templ,_result,method)) - return; + 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; @@ -603,7 +584,7 @@ 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); Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1); From 7e2bdb590f781de8c2fb6c2109641e81fae4146b Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 17 Jan 2014 16:26:11 +0400 Subject: [PATCH 07/13] fixed --- modules/imgproc/src/templmatch.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index bb3423bb2..aaaa83345 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -571,7 +571,9 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, 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))) From 2c0765b4d13260a53103b2693a13d997efd15784 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 17 Jan 2014 18:37:57 +0400 Subject: [PATCH 08/13] fixed trailing whitespace --- modules/imgproc/src/templmatch.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index aaaa83345..9953c6851 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -568,7 +568,7 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, 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) { From b4652e2cb32a3824ec70cd0c5f7ca6df09c99685 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Mon, 20 Jan 2014 15:48:16 +0400 Subject: [PATCH 09/13] Added perf test for ocl_matchTemplate --- .../perf/opencl/perf_matchTemplate.cpp | 54 +++++++++++++++++++ .../imgproc/test/ocl/test_match_template.cpp | 3 +- 2 files changed, 55 insertions(+), 2 deletions(-) create mode 100644 modules/imgproc/perf/opencl/perf_matchTemplate.cpp diff --git a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp new file mode 100644 index 000000000..b790f36cc --- /dev/null +++ b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp @@ -0,0 +1,54 @@ +#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()); + + Mat img(imgSz, CV_8UC1); + Mat tmpl(tmplSz, CV_8UC1); + Mat 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 ? 1e-6 + : 255 * 255 * tmpl.total() * 1e-6; + + SANITY_CHECK(result, eps); + } +} +} + +#endif // HAVE_OPENCL \ No newline at end of file diff --git a/modules/imgproc/test/ocl/test_match_template.cpp b/modules/imgproc/test/ocl/test_match_template.cpp index 2a0402342..7656a1446 100644 --- a/modules/imgproc/test/ocl/test_match_template.cpp +++ b/modules/imgproc/test/ocl/test_match_template.cpp @@ -97,8 +97,7 @@ PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, int, bool) void Near(double threshold = 0.0) { - EXPECT_MAT_NEAR(result, uresult, threshold); - EXPECT_MAT_NEAR(result_roi, uresult_roi, threshold); + OCL_EXPECT_MATS_NEAR(result,threshold); } }; From 8c5e19c27058e9ec9fb0b3272d6a4175222d10a8 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Tue, 21 Jan 2014 16:57:01 +0400 Subject: [PATCH 10/13] fixed pref test --- modules/imgproc/perf/opencl/perf_matchTemplate.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp index b790f36cc..b5ffecc0e 100644 --- a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp +++ b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp @@ -27,9 +27,9 @@ namespace ocl { Size tmplSz = get<1>(GetParam()); int method = get<2>(GetParam()); - Mat img(imgSz, CV_8UC1); - Mat tmpl(tmplSz, CV_8UC1); - Mat result(imgSz - tmplSz + Size(1,1), CV_32F); + UMat img(imgSz, CV_8UC1); + UMat tmpl(tmplSz, CV_8UC1); + UMat result(imgSz - tmplSz + Size(1,1), CV_32F); declare .in(img, WARMUP_RNG) From 7bd8f9294d9d28895456b13f31d116f13c7651b7 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Wed, 22 Jan 2014 13:34:40 +0400 Subject: [PATCH 11/13] fixed --- modules/imgproc/perf/opencl/perf_matchTemplate.cpp | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp index b5ffecc0e..b28645cd9 100644 --- a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp +++ b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp @@ -12,7 +12,7 @@ namespace ocl { typedef std::tr1::tuple ImgSize_TmplSize_Method_t; typedef TestBaseWithParam ImgSize_TmplSize_Method; - OCL_PERF_TEST_P(ImgSize_TmplSize_Method, matchTemplate, + 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), @@ -43,10 +43,13 @@ namespace ocl { method == TM_CCORR_NORMED || method == TM_SQDIFF_NORMED || method == TM_CCOEFF_NORMED; - double eps = isNormed ? 1e-6 - : 255 * 255 * tmpl.total() * 1e-6; + double eps = isNormed ? 1e-3 + : 255 * 255 * tmpl.total() * 1e-3; - SANITY_CHECK(result, eps); + if (isNormed) + SANITY_CHECK(result,eps,ERROR_RELATIVE); + else + SANITY_CHECK(result, eps); } } } From 4d86804f1120e655031a372d14f754ab756f45b5 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Thu, 23 Jan 2014 15:31:32 +0400 Subject: [PATCH 12/13] Fixed merge conflicts --- modules/imgproc/test/ocl/test_match_template.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/imgproc/test/ocl/test_match_template.cpp b/modules/imgproc/test/ocl/test_match_template.cpp index 7656a1446..782509cc9 100644 --- a/modules/imgproc/test/ocl/test_match_template.cpp +++ b/modules/imgproc/test/ocl/test_match_template.cpp @@ -75,7 +75,7 @@ PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, int, bool) virtual void generateTestData() { Size image_roiSize = randomSize(2, 100); - Size templ_roiSize = Size(randomInt(1,image_roiSize.width), randomInt(1,image_roiSize.height)); + 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); From 07a88d48e9bdb63bf5b3ddf677296bfcd16248e4 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 24 Jan 2014 11:41:23 +0400 Subject: [PATCH 13/13] fixed --- modules/imgproc/perf/opencl/perf_matchTemplate.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp index b28645cd9..948774743 100644 --- a/modules/imgproc/perf/opencl/perf_matchTemplate.cpp +++ b/modules/imgproc/perf/opencl/perf_matchTemplate.cpp @@ -43,8 +43,8 @@ namespace ocl { method == TM_CCORR_NORMED || method == TM_SQDIFF_NORMED || method == TM_CCOEFF_NORMED; - double eps = isNormed ? 1e-3 - : 255 * 255 * tmpl.total() * 1e-3; + double eps = isNormed ? 3e-2 + : 255 * 255 * tmpl.total() * 1e-4; if (isNormed) SANITY_CHECK(result,eps,ERROR_RELATIVE);