mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-25 19:50:38 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Added ocl_matchTemplate( without dft)

Browse Source
This commit is contained in:
Elena Gvozdeva 2014-01-14 14:45:01 +04:00
parent ee88cc2c52
commit 86636dc265
3 changed files with 927 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

427
modules/imgproc/src/opencl/match_template.cl Normal file
View File

@ -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;
}
}

377
modules/imgproc/src/templmatch.cpp
View File

@ -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();

130
modules/imgproc/test/ocl/test_match_template.cpp Normal file
View File

@ -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