Skip to content

O
opencv
Commit ac3f06bc authored by Andrey Pavlenko's avatar Andrey Pavlenko Committed by OpenCV Buildbot
Browse files
Options

Merge pull request #2149 from ElenaGvozdeva:ocl_matchTemplate

parents 22c804d9 07a88d48
Hide whitespace changes
Inline Side-by-side
Showing with 932 additions and 8 deletions
modules/imgproc/perf/opencl/perf_matchTemplate.cpp 0 → 100644
View file @ ac3f06bc
#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<Size, Size, MethodType> ImgSize_TmplSize_Method_t;
typedef TestBaseWithParam<ImgSize_TmplSize_Method_t> ImgSize_TmplSize_Method;
OCL_PERF_TEST_P(ImgSize_TmplSize_Method, MatchTemplate,
::testing::Combine(
testing::Values(szSmall128, cv::Size(320, 240),
cv::Size(640, 480), cv::Size(800, 600),
cv::Size(1024, 768), cv::Size(1280, 1024)),
testing::Values(cv::Size(12, 12), cv::Size(28, 9),
cv::Size(8, 30), cv::Size(16, 16)),
MethodType::all()
)
)
{
Size imgSz = get<0>(GetParam());
Size tmplSz = get<1>(GetParam());
int method = get<2>(GetParam());
UMat img(imgSz, CV_8UC1);
UMat tmpl(tmplSz, CV_8UC1);
UMat result(imgSz - tmplSz + Size(1,1), CV_32F);
declare
.in(img, WARMUP_RNG)
.in(tmpl, WARMUP_RNG)
.out(result)
.time(30);
OCL_TEST_CYCLE() matchTemplate(img, tmpl, result, method);
bool isNormed =
method == TM_CCORR_NORMED ||
method == TM_SQDIFF_NORMED ||
method == TM_CCOEFF_NORMED;
double eps = isNormed ? 3e-2
: 255 * 255 * tmpl.total() * 1e-4;
if (isNormed)
SANITY_CHECK(result,eps,ERROR_RELATIVE);
else
SANITY_CHECK(result, eps);
}
}
}
#endif // HAVE_OPENCL
\ No newline at end of file
modules/imgproc/src/opencl/match_template.cl 0 → 100644
View file @ ac3f06bc
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors as is and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
#define DATA_SIZE ((int)sizeof(type))
#define ELEM_TYPE elem_type
#define ELEM_SIZE ((int)sizeof(elem_type))
#define CN cn
#define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, gidx*CN + img_sqsums_offset + ox*CN)
#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx*CN + img_sums_offset + ox*CN)
inline float normAcc(float num, float denum)
{
if(fabs(num) < denum)
{
return num / denum;
}
if(fabs(num) < denum * 1.125f)
{
return num > 0 ? 1 : -1;
}
return 0;
}
inline float normAcc_SQDIFF(float num, float denum)
{
if(fabs(num) < denum)
{
return num / denum;
}
if(fabs(num) < denum * 1.125f)
{
return num > 0 ? 1 : -1;
}
return 1;
}
//////////////////////////////////////////CCORR/////////////////////////////////////////////////////////////////////////
__kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step,int img_offset,
__global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols,
__global uchar * res,int res_step,int res_offset,int res_rows,int res_cols)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
int i,j;
float sum = 0;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows)
{
for(i = 0; i < tpl_rows; i ++)
{
__global const ELEM_TYPE * img_ptr = (__global const ELEM_TYPE *)(img + mad24(gidy + i, img_step, gidx*DATA_SIZE + img_offset));
__global const ELEM_TYPE * tpl_ptr = (__global const ELEM_TYPE *)(tpl + mad24(i, tpl_step, tpl_offset));
for(j = 0; j < tpl_cols; j ++)
#pragma unroll
for (int c = 0; c < CN; c++)
sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]);
}
__global float * result = (__global float *)(res+res_idx);
*result = sum;
}
}
__kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, ulong tpl_sqsum)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows)
{
__global float * sqsum = (__global float*)(img_sqsums);
float image_sqsum_ = (float)(
(sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
__global float * result = (__global float *)(res+res_idx);
*result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum));
}
}
////////////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////////////////
__kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step,int img_offset,
__global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols,
__global uchar * res,int res_step,int res_offset,int res_rows,int res_cols)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
int i,j;
float delta;
float sum = 0;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows)
{
for(i = 0; i < tpl_rows; i ++)
{
__global const ELEM_TYPE * img_ptr = (__global const ELEM_TYPE *)(img + mad24(gidy + i, img_step, gidx*DATA_SIZE + img_offset));
__global const ELEM_TYPE * tpl_ptr = (__global const ELEM_TYPE *)(tpl + mad24(i, tpl_step, tpl_offset));
for(j = 0; j < tpl_cols; j ++)
#pragma unroll
for (int c = 0; c < CN; c++)
{
delta = (float)(img_ptr[j*CN+c] - tpl_ptr[j*CN+c]);
sum += delta*delta;
}
}
__global float * result = (__global float *)(res+res_idx);
*result = sum;
}
}
__kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, ulong tpl_sqsum)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows)
{
__global float * sqsum = (__global float*)(img_sqsums);
float image_sqsum_ = (float)(
(sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
__global float * result = (__global float *)(res+res_idx);
*result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum));
}
}
////////////////////////////////////////////CCOEFF/////////////////////////////////////////////////////////////////
__kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, float tpl_sum)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sums_step /= ELEM_SIZE;
img_sums_offset /= ELEM_SIZE;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows)
{
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
image_sum_ += (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)])-
(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])) * tpl_sum;
__global float * result = (__global float *)(res+res_idx);
*result -= image_sum_;
}
}
__kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sums_step /= ELEM_SIZE;
img_sums_offset /= ELEM_SIZE;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows)
{
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
image_sum_ += tpl_sum_0 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)]) -(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)]));
image_sum_ += tpl_sum_1 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+1] - sum[SUMS_PTR(tpl_cols, 0)+1])-(sum[SUMS_PTR(0, tpl_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
__global float * result = (__global float *)(res+res_idx);
*result -= image_sum_;
}
}
__kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sums_step /= ELEM_SIZE;
img_sums_offset /= ELEM_SIZE;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows)
{
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
int c_r = SUMS_PTR(tpl_cols, tpl_rows);
int c_o = SUMS_PTR(tpl_cols, 0);
int o_r = SUMS_PTR(0,tpl_rows);
int oo = SUMS_PTR(0, 0);
image_sum_ += tpl_sum_0 * (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[oo]));
image_sum_ += tpl_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1]));
image_sum_ += tpl_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2]));
image_sum_ += tpl_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3]));
__global float * result = (__global float *)(res+res_idx);
*result -= image_sum_;
}
}
__kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
__global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int t_rows, int t_cols, float weight, float tpl_sum, float tpl_sqsum)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sums_offset /= ELEM_SIZE;
img_sums_step /= ELEM_SIZE;
img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows)
{
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
__global float * sqsum = (__global float*)(img_sqsums);
float image_sum_ = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)]) -
(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)]));
float image_sqsum_ = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
__global float * result = (__global float *)(res+res_idx);
*result = normAcc((*result) - image_sum_ * tpl_sum,
sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_)));
}
}
__kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
__global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int t_rows, int t_cols, float weight, float tpl_sum_0, float tpl_sum_1, float tpl_sqsum)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sums_offset /= ELEM_SIZE;
img_sums_step /= ELEM_SIZE;
img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float sum_[2];
float sqsum_[2];
if(gidx < res_cols && gidy < res_rows)
{
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
__global float * sqsum = (__global float*)(img_sqsums);
sum_[0] = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)])-(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)]));
sum_[1] = (float)((sum[SUMS_PTR(t_cols, t_rows)+1] - sum[SUMS_PTR(t_cols, 0)+1])-(sum[SUMS_PTR(0, t_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
sqsum_[0] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)])-(sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
sqsum_[1] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)+1] - sqsum[SQSUMS_PTR(t_cols, 0)+1])-(sqsum[SQSUMS_PTR(0, t_rows)+1] - sqsum[SQSUMS_PTR(0, 0)+1]));
float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1;
float denum = sqrt( tpl_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] +
sqsum_[1] - weight * sum_[1]* sum_[1]));
__global float * result = (__global float *)(res+res_idx);
*result = normAcc((*result) - num, denum);
}
}
__kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
__global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int t_rows, int t_cols, float weight,
float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3,
float tpl_sqsum)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
img_sums_offset /= ELEM_SIZE;
img_sums_step /= ELEM_SIZE;
img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float sum_[4];
float sqsum_[4];
if(gidx < res_cols && gidy < res_rows)
{
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
__global float * sqsum = (__global float*)(img_sqsums);
int c_r = SUMS_PTR(t_cols, t_rows);
int c_o = SUMS_PTR(t_cols, 0);
int o_r = SUMS_PTR(0, t_rows);
int o_o = SUMS_PTR(0, 0);
sum_[0] = (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[o_o ]));
sum_[1] = (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[o_o +1]));
sum_[2] = (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[o_o +2]));
sum_[3] = (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[o_o +3]));
c_r = SQSUMS_PTR(t_cols, t_rows);
c_o = SQSUMS_PTR(t_cols, 0);
o_r = SQSUMS_PTR(0, t_rows);
o_o = SQSUMS_PTR(0, 0);
sqsum_[0] = (float)((sqsum[c_r] - sqsum[c_o]) -(sqsum[o_r] - sqsum[o_o]));
sqsum_[1] = (float)((sqsum[c_r+1] - sqsum[c_o+1])-(sqsum[o_r+1] - sqsum[o_o+1]));
sqsum_[2] = (float)((sqsum[c_r+2] - sqsum[c_o+2])-(sqsum[o_r+2] - sqsum[o_o+2]));
sqsum_[3] = (float)((sqsum[c_r+3] - sqsum[c_o+3])-(sqsum[o_r+3] - sqsum[o_o+3]));
float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1 + sum_[2]*tpl_sum_2 + sum_[3]*tpl_sum_3;
float denum = sqrt( tpl_sqsum * (
sqsum_[0] - weight * sum_[0]* sum_[0] +
sqsum_[1] - weight * sum_[1]* sum_[1] +
sqsum_[2] - weight * sum_[2]* sum_[2] +
sqsum_[3] - weight * sum_[3]* sum_[3] ));
__global float * result = (__global float *)(res+res_idx);
*result = normAcc((*result) - num, denum);
}
}
\ No newline at end of file
modules/imgproc/src/templmatch.cpp
View file @ ac3f06bc
...@@ -40,6 +40,337 @@ ...@@ -40,6 +40,337 @@
//M*/ //M*/
#include "precomp.hpp" #include "precomp.hpp"
#include "opencl_kernels.hpp"
//////////////////////////////////////////////////matchTemplate//////////////////////////////////////////////////////////
namespace cv
{
static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArray _result);
static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result);
static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result);
static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result);
static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArray _result);
static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result);
static bool matchTemplateNaive_CCORR (InputArray _image, InputArray _templ, OutputArray _result, int cn);
static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result, int cn);
static bool useNaive(int method, int depth, Size size)
{
#ifdef HAVE_CLAMDFFT
if (method == TM_SQDIFF && depth == CV_32F)
return true;
else if(method == TM_CCORR || (method == TM_SQDIFF && depth == CV_8U))
return size.height < 18 && size.width < 18;
else
return false;
#else
#define UNUSED(x) (void)(x);
UNUSED(method) UNUSED(depth) UNUSED(size)
#undef UNUSED
return true;
#endif
}
///////////////////////////////////////////////////CCORR//////////////////////////////////////////////////////////////
static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArray _result)
{
if (useNaive(TM_CCORR, _image.depth(), _templ.size()) )
return matchTemplateNaive_CCORR(_image, _templ, _result, _image.channels());
else
return false;
}
static bool matchTemplateNaive_CCORR (InputArray _image, InputArray _templ, OutputArray _result, int cn)
{
int type = _image.type();
int depth = CV_MAT_DEPTH(type);
const char * kernelName = "matchTemplate_Naive_CCORR";
ocl::Kernel k (kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
UMat image = _image.getUMat();
UMat templ = _templ.getUMat(), result;
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false);
}
static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
{
matchTemplate(_image, _templ, _result, CV_TM_CCORR);
int type = _image.type();
int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
const char * kernelName = "matchTemplate_CCORR_NORMED";
ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
UMat image = _image.getUMat();
UMat templ = _templ.getUMat(), result;
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
result = _result.getUMat();
UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
UMat templ_resh, temp;
templ.reshape(1).convertTo(templ_resh, CV_32F);
multiply(templ_resh, templ_resh, temp);
unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0];
size_t globalsize[2] = {result.cols, result.rows};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false);
}
//////////////////////////////////////SQDIFF//////////////////////////////////////////////////////////////
static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result)
{
if (useNaive(TM_SQDIFF, _image.depth(), _templ.size()))
{
return matchTemplateNaive_SQDIFF(_image, _templ, _result, _image.channels());;
}
else
return false;
}
static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result, int cn)
{
int type = _image.type();
int depth = CV_MAT_DEPTH(type);
const char * kernelName = "matchTemplate_Naive_SQDIFF";
ocl::Kernel k (kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
UMat image = _image.getUMat();
UMat templ = _templ.getUMat(), result;
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false);
}
static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result)
{
matchTemplate(_image, _templ, _result, CV_TM_CCORR);
int type = _image.type();
int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
const char * kernelName = "matchTemplate_SQDIFF_NORMED";
ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
UMat image = _image.getUMat();
UMat templ = _templ.getUMat(), result;
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
result = _result.getUMat();
UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
UMat temp, templ_resh;
templ.reshape(1).convertTo(templ_resh, CV_32F);
multiply(templ_resh, templ_resh, temp);
unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0];
size_t globalsize[2] = {result.cols, result.rows};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false);
}
/////////////////////////////////////CCOEFF/////////////////////////////////////////////////////////////////
static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArray _result)
{
matchTemplate(_image, _templ, _result, CV_TM_CCORR);
UMat image_sums;
integral(_image, image_sums);
int type = image_sums.type();
int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
const char * kernelName;
if (cn==1)
kernelName = "matchTemplate_Prepared_CCOEFF_C1";
else if (cn==2)
kernelName = "matchTemplate_Prepared_CCOEFF_C2";
else
kernelName = "matchTemplate_Prepared_CCOEFF_C4";
ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc, format("-D type=%s -D elem_type=%s -D cn=%d",ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
UMat templ = _templ.getUMat(), result;
Size size = _image.size();
_result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F);
result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows};
if (cn==1)
{
float templ_sum = (float)sum(_templ)[0]/ _templ.size().area();
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,NULL,false);
}
else
{
Vec4f templ_sum = Vec4f::all(0);
templ_sum = sum(templ)/ templ.size().area();
if (cn==2)
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols,
templ_sum[0],templ_sum[1]).run(2,globalsize,NULL,false);
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols,
templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,NULL,false);
}
}
static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
{
UMat imagef, templf;
_image.getUMat().convertTo(imagef, CV_32F);
_templ.getUMat().convertTo(templf, CV_32F);
matchTemplate(imagef, templf, _result, CV_TM_CCORR);
const char * kernelName;
UMat temp, image_sums, image_sqsums;
integral(_image,image_sums, image_sqsums, CV_32F, CV_32F);
int type = image_sums.type();
int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
if (cn== 1)
kernelName = "matchTemplate_CCOEFF_NORMED_C1";
else if (cn==2)
kernelName = "matchTemplate_CCOEFF_NORMED_C2";
else
kernelName = "matchTemplate_CCOEFF_NORMED_C4";
ocl::Kernel k(kernelName, ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
UMat image = _image.getUMat();
UMat templ = _templ.getUMat(), result;
int image_rows = _image.size().height, image_cols = _image.size().width;
_result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F);
result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows};
float scale = 1.f / templ.size().area();
if (cn==1)
{
float templ_sum = (float)sum(templ)[0];
multiply(templf, templf, temp);
float templ_sqsum = (float)sum(temp)[0];
templ_sqsum -= scale * templ_sum * templ_sum;
templ_sum *= scale;
if (templ_sqsum < DBL_EPSILON)
{
result = Scalar::all(1);
return true;
}
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums),ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum)
.run(2,globalsize,NULL,false);
}
else
{
Vec4f templ_sum = Vec4f::all(0);
Vec4f templ_sqsum = Vec4f::all(0);
templ_sum = sum(templ);
multiply(templf, templf, temp);
templ_sqsum = sum(temp);
float templ_sqsum_sum = 0;
for(int i = 0; i < cn; i ++)
{
templ_sqsum_sum += templ_sqsum[i] - scale * templ_sum[i] * templ_sum[i];
}
templ_sum *= scale;
if (templ_sqsum_sum < DBL_EPSILON)
{
result = Scalar::all(1);
return true;
}
if (cn==2)
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale,
templ_sum[0],templ_sum[1], templ_sqsum_sum)
.run(2,globalsize,NULL,false);
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale,
templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3], templ_sqsum_sum)
.run(2,globalsize,NULL,false);
}
}
///////////////////////////////////////////////////////////////////////////////////////////////////////////
static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method)
{
int cn = CV_MAT_CN(_img.type());
if (cn == 3 || cn > 4)
return false;
typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result);
const Caller callers[] =
{
matchTemplate_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplate_CCORR,
matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED
};
Caller caller = callers[method];
return caller(_img, _templ, _result);
}
}
namespace cv namespace cv
{ {
...@@ -226,15 +557,28 @@ void crossCorr( const Mat& img, const Mat& _templ, Mat& corr, ...@@ -226,15 +557,28 @@ void crossCorr( const Mat& img, const Mat& _templ, Mat& corr,
} }
} }
} }
} }
/*****************************************************************************************/ ////////////////////////////////////////////////////////////////////////////////////////////////////////
void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method ) void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method )
{ {
CV_Assert( CV_TM_SQDIFF <= method && method <= CV_TM_CCOEFF_NORMED ); CV_Assert( CV_TM_SQDIFF <= method && method <= CV_TM_CCOEFF_NORMED );
CV_Assert( (_img.depth() == CV_8U || _img.depth() == CV_32F) && _img.type() == _templ.type() );
CV_Assert(_img.dims() <= 2);
bool swapNotNeed = (_img.size().height >= _templ.size().height && _img.size().width >= _templ.size().width);
if (!swapNotNeed)
{
CV_Assert(_img.size().height <= _templ.size().height && _img.size().width <= _templ.size().width);
}
bool use_opencl = ocl::useOpenCL() && _result.isUMat();
if ( use_opencl && (swapNotNeed ? ocl_matchTemplate(_img,_templ,_result,method) : ocl_matchTemplate(_templ,_img,_result,method)))
return;
int numType = method == CV_TM_CCORR || method == CV_TM_CCORR_NORMED ? 0 : int numType = method == CV_TM_CCORR || method == CV_TM_CCORR_NORMED ? 0 :
method == CV_TM_CCOEFF || method == CV_TM_CCOEFF_NORMED ? 1 : 2; method == CV_TM_CCOEFF || method == CV_TM_CCOEFF_NORMED ? 1 : 2;
bool isNormed = method == CV_TM_CCORR_NORMED || bool isNormed = method == CV_TM_CCORR_NORMED ||
...@@ -242,14 +586,9 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, ...@@ -242,14 +586,9 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result,
method == CV_TM_CCOEFF_NORMED; method == CV_TM_CCOEFF_NORMED;
Mat img = _img.getMat(), templ = _templ.getMat(); Mat img = _img.getMat(), templ = _templ.getMat();
if( img.rows < templ.rows || img.cols < templ.cols ) if(!swapNotNeed )
std::swap(img, templ); 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); Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1);
_result.create(corrSize, CV_32F); _result.create(corrSize, CV_32F);
Mat result = _result.getMat(); Mat result = _result.getMat();
......
modules/imgproc/test/ocl/test_match_template.cpp 0 → 100644
View file @ ac3f06bc
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors as is and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "test_precomp.hpp"
#include "opencv2/ts/ocl_test.hpp"
#include "iostream"
#include "fstream"
#ifdef HAVE_OPENCL
namespace cvtest {
namespace ocl {
/////////////////////////////////////////////matchTemplate//////////////////////////////////////////////////////////
PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, int, bool)
{
int type;
int depth;
int method;
bool use_roi;
TEST_DECLARE_INPUT_PARAMETER(image)
TEST_DECLARE_INPUT_PARAMETER(templ)
TEST_DECLARE_OUTPUT_PARAMETER(result)
virtual void SetUp()
{
type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
depth = GET_PARAM(0);
method = GET_PARAM(2);
use_roi = GET_PARAM(3);
}
virtual void generateTestData()
{
Size image_roiSize = randomSize(2, 100);
Size templ_roiSize = Size(randomInt(1, image_roiSize.width), randomInt(1, image_roiSize.height));
Size result_roiSize = Size(image_roiSize.width - templ_roiSize.width + 1,
image_roiSize.height - templ_roiSize.height + 1);
const double upValue = 256;
Border imageBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(image, image_roi, image_roiSize, imageBorder, type, -upValue, upValue);
Border templBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(templ, templ_roi, templ_roiSize, templBorder, type, -upValue, upValue);
Border resultBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(result, result_roi, result_roiSize, resultBorder, CV_32F, -upValue, upValue);
UMAT_UPLOAD_INPUT_PARAMETER(image)
UMAT_UPLOAD_INPUT_PARAMETER(templ)
UMAT_UPLOAD_OUTPUT_PARAMETER(result)
}
void Near(double threshold = 0.0)
{
OCL_EXPECT_MATS_NEAR(result,threshold);
}
};
OCL_TEST_P(MatchTemplate, Mat)
{
for (int j = 0; j < test_loop_times; j++)
{
generateTestData();
OCL_OFF(cv::matchTemplate(image_roi,templ_roi,result_roi, method));
OCL_ON(cv::matchTemplate(uimage_roi,utempl_roi,uresult_roi, method));
if (method == 0)
Near(10.0f);
else
Near(method % 2 == 1 ? 0.001f : 1.0f);
}
}
OCL_INSTANTIATE_TEST_CASE_P(ImageProc, MatchTemplate, Combine(
Values(CV_8U, CV_32F),
Values(1, 2, 4),
Values(0,1,2,3,4,5),
Bool())
);
} } // namespace cvtest::ocl
#endif
\ No newline at end of file
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment