Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

Merge pull request #2391 from ilya-lavrenov:tapi_match_template

Browse Source
pull/2436/merge
Andrey Pavlenko 11 years ago committed by OpenCV Buildbot
parent c00f0faf24 653b99c9bd
commit c9d8025fab
3 changed files with 378 additions and 303 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 488
      modules/imgproc/src/opencl/match_template.cl
  2. 190
      modules/imgproc/src/templmatch.cpp
  3. 3
      modules/imgproc/test/ocl/test_match_template.cpp

488
modules/imgproc/src/opencl/match_template.cl
Unescape View File

@ -32,259 +32,315 @@
#define DATA_SIZE ((int)sizeof(type)) #define DATA_SIZE ((int)sizeof(type))
#define ELEM_TYPE elem_type #define ELEM_TYPE elem_type
#define ELEM_SIZE ((int)sizeof(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 SQSUMS_PTR(ox, oy) mad24(y + oy, src_sqsums_step, mad24(x + ox, cn, src_sqsums_offset))
#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx*CN + img_sums_offset + ox*CN) #define SUMS_PTR(ox, oy) mad24(y + oy, src_sums_step, mad24(x + ox, cn, src_sums_offset))
inline float normAcc(float num, float denum) inline float normAcc(float num, float denum)
{ {
if(fabs(num) < denum) if (fabs(num) < denum)
{
return num / denum; return num / denum;
} if (fabs(num) < denum * 1.125f)
if(fabs(num) < denum * 1.125f)
{
return num > 0 ? 1 : -1; return num > 0 ? 1 : -1;
}
return 0; return 0;
} }
inline float normAcc_SQDIFF(float num, float denum) inline float normAcc_SQDIFF(float num, float denum)
{ {
if(fabs(num) < denum) if (fabs(num) < denum)
{
return num / denum; return num / denum;
} if (fabs(num) < denum * 1.125f)
if(fabs(num) < denum * 1.125f)
{
return num > 0 ? 1 : -1; return num > 0 ? 1 : -1;
}
return 1; return 1;
} }
//////////////////////////////////////////CCORR///////////////////////////////////////////////////////////////////////// #define noconvert
#if cn == 1
#define convertToDT(value) (float)(value)
#elif cn == 2
#define convertToDT(value) (float)(value.x + value.y)
#elif cn == 4
#define convertToDT(value) (float)(value.x + value.y + value.z + value.w)
#else
#error "cn should be 1, 2 or 4"
#endif
__kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step,int img_offset, #ifdef CALC_SUM
__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) __kernel void calcSum(__global const uchar * srcptr, int src_step, int src_offset,
int cols, int total, __global float * dst)
{ {
int gidx = get_global_id(0); int lid = get_local_id(0), id = 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)); __local WT localmem[WGS2_ALIGNED];
WT accumulator = (WT)(0), tmp;
if(gidx < res_cols && gidy < res_rows) for ( ; id < total; id += WGS)
{ {
for(i = 0; i < tpl_rows; i ++) int src_index = mad24(id / cols, src_step, mad24(id % cols, (int)sizeof(T), src_offset));
{ __global const T * src = (__global const T *)(srcptr + src_index);
__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)); tmp = convertToWT(src[0]);
#if wdepth == 4
for(j = 0; j < tpl_cols; j ++) accumulator = mad24(tmp, tmp, accumulator);
#else
accumulator = mad(tmp, tmp, accumulator);
#endif
}
#pragma unroll if (lid < WGS2_ALIGNED)
for (int c = 0; c < CN; c++) localmem[lid] = accumulator;
barrier(CLK_LOCAL_MEM_FENCE);
sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]); if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
localmem[lid - WGS2_ALIGNED] += accumulator;
barrier(CLK_LOCAL_MEM_FENCE);
for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
{
if (lid < lsize)
{
int lid2 = lsize + lid;
localmem[lid] += localmem[lid2];
} }
__global float * result = (__global float *)(res+res_idx); barrier(CLK_LOCAL_MEM_FENCE);
*result = sum;
} }
if (lid == 0)
dst[0] = convertToDT(localmem[0]);
} }
__kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, #elif defined CCORR
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, float tpl_sqsum) __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_step, int src_offset,
__global const uchar * templateptr, int template_step, int template_offset, int template_rows, int template_cols,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{ {
int gidx = get_global_id(0); int x = get_global_id(0);
int gidy = get_global_id(1); int y = get_global_id(1);
img_sqsums_step /= sizeof(float); if (x < dst_cols && y < dst_rows)
img_sqsums_offset /= sizeof(float); {
WT sum = (WT)(0);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); __global const T * src = (__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset)));
__global const T * template = (__global const T *)(templateptr + template_offset);
if(gidx < res_cols && gidy < res_rows) for (int i = 0; i < template_rows; ++i)
{ {
__global float * sqsum = (__global float*)(img_sqsums); for (int j = 0; j < template_cols; ++j)
float image_sqsum_ = (float)( #if wdepth == 4
(sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - sum = mad24(convertToWT(src[j]), convertToWT(template[j]), sum);
(sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); #else
sum = mad(convertToWT(src[j]), convertToWT(template[j]), sum);
#endif
src = (__global const T *)((__global const uchar *)src + src_step);
template = (__global const T *)((__global const uchar *)template + template_step);
}
__global float * result = (__global float *)(res+res_idx); int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
*result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum)); *(__global float *)(dst + dst_idx) = convertToDT(sum);
} }
} }
////////////////////////////////////////////SQDIFF//////////////////////////////////////////////////////////////////////// #elif defined CCORR_NORMED
__kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step,int img_offset, __kernel void matchTemplate_CCORR_NORMED(__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
__global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
__global uchar * res,int res_step,int res_offset,int res_rows,int res_cols) int template_rows, int template_cols, __global const float * template_sqsum)
{ {
int gidx = get_global_id(0); int x = get_global_id(0);
int gidy = get_global_id(1); int y = 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) if (x < dst_cols && y < dst_rows)
{ {
for(i = 0; i < tpl_rows; i ++) __global const float * sqsum = (__global const float *)(src_sqsums);
{
__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 ++) src_sqsums_step /= sizeof(float);
src_sqsums_offset /= sizeof(float);
float image_sqsum_ = (float)(sqsum[SQSUMS_PTR(template_cols, template_rows)] - sqsum[SQSUMS_PTR(template_cols, 0)] -
sqsum[SQSUMS_PTR(0, template_rows)] + sqsum[SQSUMS_PTR(0, 0)]);
#pragma unroll int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
for (int c = 0; c < CN; c++) __global float * dstult = (__global float *)(dst + dst_idx);
{ *dstult = normAcc(*dstult, sqrt(image_sqsum_ * template_sqsum[0]));
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, #elif defined SQDIFF
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int tpl_rows, int tpl_cols, float 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)); __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * srcptr, int src_step, int src_offset,
__global const uchar * templateptr, int template_step, int template_offset, int template_rows, int template_cols,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
int x = get_global_id(0);
int y = get_global_id(1);
if(gidx < res_cols && gidy < res_rows) if (x < dst_cols && y < dst_rows)
{ {
__global float * sqsum = (__global float*)(img_sqsums); __global const T * src = (__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset)));
float image_sqsum_ = (float)( __global const T * template = (__global const T *)(templateptr + template_offset);
(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); WT sum = (WT)(0), value;
*result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum)); for (int i = 0; i < template_rows; ++i)
{
for (int j = 0; j < template_cols; ++j)
{
value = convertToWT(src[j]) - convertToWT(template[j]);
#if wdepth == 4
sum = mad24(value, value, sum);
#else
sum = mad(value, value, sum);
#endif
}
src = (__global const T *)((__global const uchar *)src + src_step);
template = (__global const T *)((__global const uchar *)template + template_step);
}
int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
*(__global float *)(dst + dst_idx) = convertToDT(sum);
} }
} }
////////////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// #elif defined SQDIFF_NORMED
__kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __kernel void matchTemplate_SQDIFF_NORMED(__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int tpl_rows, int tpl_cols, float tpl_sum) int template_rows, int template_cols, __global const float * template_sqsum)
{ {
int gidx = get_global_id(0); int x = get_global_id(0);
int gidy = get_global_id(1); int y = 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)); if (x < dst_cols && y < dst_rows)
float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); src_sqsums_step /= sizeof(float);
src_sqsums_offset /= sizeof(float);
image_sum_ += (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)])- __global const float * sqsum = (__global const float *)(src_sqsums);
(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])) * tpl_sum; float image_sqsum_ = (float)(
(sqsum[SQSUMS_PTR(template_cols, template_rows)] - sqsum[SQSUMS_PTR(template_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, template_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
float template_sqsum_value = template_sqsum[0];
__global float * result = (__global float *)(res+res_idx); int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
*result -= image_sum_; __global float * dstult = (__global float *)(dst + dst_idx);
*dstult = normAcc_SQDIFF(image_sqsum_ - 2.0f * dstult[0] + template_sqsum_value, sqrt(image_sqsum_ * template_sqsum_value));
} }
} }
__kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, #elif defined CCOEFF
__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; #if cn == 1
img_sums_offset /= ELEM_SIZE;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); __kernel void matchTemplate_Prepared_CCOEFF(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
float image_sum_ = 0; __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int template_rows, int template_cols, float template_sum)
{
int x = get_global_id(0);
int y = get_global_id(1);
if(gidx < res_cols && gidy < res_rows) if (x < dst_cols && y < dst_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_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)])); src_sums_step /= ELEM_SIZE;
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])); src_sums_offset /= ELEM_SIZE;
float image_sum_ = (float)((sum[SUMS_PTR(template_cols, template_rows)] - sum[SUMS_PTR(template_cols, 0)])-
(sum[SUMS_PTR(0, template_rows)] - sum[SUMS_PTR(0, 0)])) * template_sum;
__global float * result = (__global float *)(res+res_idx); int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
__global float * dstult = (__global float *)(dst + dst_idx);
*result -= image_sum_; *dstult -= image_sum_;
} }
} }
__kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, #elif cn == 2
__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) __kernel void matchTemplate_Prepared_CCOEFF(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int template_rows, int template_cols, float template_sum_0, float template_sum_1)
{ {
int gidx = get_global_id(0); int x = get_global_id(0);
int gidy = get_global_id(1); int y = get_global_id(1);
if (x < dst_cols && y < dst_rows)
{
src_sums_step /= ELEM_SIZE;
src_sums_offset /= ELEM_SIZE;
img_sums_step /= ELEM_SIZE; __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
img_sums_offset /= ELEM_SIZE;
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float image_sum_ = template_sum_0 * (float)((sum[SUMS_PTR(template_cols, template_rows)] - sum[SUMS_PTR(template_cols, 0)]) -(sum[SUMS_PTR(0, template_rows)] - sum[SUMS_PTR(0, 0)]));
float image_sum_ = 0; image_sum_ += template_sum_1 * (float)((sum[SUMS_PTR(template_cols, template_rows)+1] - sum[SUMS_PTR(template_cols, 0)+1])-(sum[SUMS_PTR(0, template_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
if(gidx < res_cols && gidy < res_rows)
int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
__global float * dstult = (__global float *)(dst+dst_idx);
*dstult -= image_sum_;
}
}
#elif cn == 4
__kernel void matchTemplate_Prepared_CCOEFF(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int template_rows, int template_cols, float template_sum_0, float template_sum_1, float template_sum_2, float template_sum_3)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < dst_cols && y < dst_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); src_sums_step /= ELEM_SIZE;
src_sums_offset /= ELEM_SIZE;
int c_r = SUMS_PTR(tpl_cols, tpl_rows); __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
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])); int c_r = SUMS_PTR(template_cols, template_rows);
image_sum_ += tpl_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1])); int c_o = SUMS_PTR(template_cols, 0);
image_sum_ += tpl_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2])); int o_r = SUMS_PTR(0,template_rows);
image_sum_ += tpl_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3])); int oo = SUMS_PTR(0, 0);
__global float * result = (__global float *)(res+res_idx); float image_sum_ = template_sum_0 * (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[oo]));
image_sum_ += template_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1]));
image_sum_ += template_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2]));
image_sum_ += template_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3]));
*result -= image_sum_; int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
__global float * dstult = (__global float *)(dst+dst_idx);
*dstult -= image_sum_;
} }
} }
__kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, #else
__global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, #error "cn should be 1, 2 or 4"
__global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, #endif
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; #elif defined CCOEFF_NORMED
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 cn == 1
__kernel void matchTemplate_CCOEFF_NORMED(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int t_rows, int t_cols, float weight, float template_sum, float template_sqsum)
{
int x = get_global_id(0);
int y = get_global_id(1);
if(gidx < res_cols && gidy < res_rows) if (x < dst_cols && y < dst_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); src_sums_offset /= ELEM_SIZE;
__global float * sqsum = (__global float*)(img_sqsums); src_sums_step /= ELEM_SIZE;
src_sqsums_step /= sizeof(float);
src_sqsums_offset /= sizeof(float);
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
__global float * sqsum = (__global float*)(src_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)])); (sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)]));
@ -292,35 +348,35 @@ __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)]) - 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)])); (sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
__global float * result = (__global float *)(res+res_idx); int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
__global float * dstult = (__global float *)(dst+dst_idx);
*result = normAcc((*result) - image_sum_ * tpl_sum, *dstult = normAcc((*dstult) - image_sum_ * template_sum,
sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_))); sqrt(template_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, #elif cn == 2
__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)); __kernel void matchTemplate_CCOEFF_NORMED(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int t_rows, int t_cols, float weight, float template_sum_0, float template_sum_1, float template_sqsum)
{
int x = get_global_id(0);
int y = get_global_id(1);
float sum_[2]; float sum_[2];
float sqsum_[2]; float sqsum_[2];
if(gidx < res_cols && gidy < res_rows) if (x < dst_cols && y < dst_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); src_sums_offset /= ELEM_SIZE;
__global float * sqsum = (__global float*)(img_sqsums); src_sums_step /= ELEM_SIZE;
src_sqsums_step /= sizeof(float);
src_sqsums_offset /= sizeof(float);
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
__global float * sqsum = (__global float*)(src_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_[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_[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]));
@ -328,40 +384,41 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i
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_[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])); 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 num = sum_[0]*template_sum_0 + sum_[1]*template_sum_1;
float denum = sqrt( tpl_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] + float denum = sqrt( template_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] +
sqsum_[1] - weight * sum_[1]* sum_[1])); sqsum_[1] - weight * sum_[1]* sum_[1]));
__global float * result = (__global float *)(res+res_idx); int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
*result = normAcc((*result) - num, denum); __global float * dstult = (__global float *)(dst+dst_idx);
*dstult = normAcc((*dstult) - num, denum);
} }
} }
__kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, #elif cn == 4
__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)); __kernel void matchTemplate_CCOEFF_NORMED(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int t_rows, int t_cols, float weight,
float template_sum_0, float template_sum_1, float template_sum_2, float template_sum_3,
float template_sqsum)
{
int x = get_global_id(0);
int y = get_global_id(1);
float sum_[4]; float sum_[4];
float sqsum_[4]; float sqsum_[4];
if(gidx < res_cols && gidy < res_rows) if (x < dst_cols && y < dst_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); src_sums_offset /= ELEM_SIZE;
__global float * sqsum = (__global float*)(img_sqsums); src_sums_step /= ELEM_SIZE;
src_sqsums_step /= sizeof(float);
src_sqsums_offset /= sizeof(float);
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
__global float * sqsum = (__global float*)(src_sqsums);
int c_r = SUMS_PTR(t_cols, t_rows); int c_r = SUMS_PTR(t_cols, t_rows);
int c_o = SUMS_PTR(t_cols, 0); int c_o = SUMS_PTR(t_cols, 0);
@ -383,15 +440,22 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i
sqsum_[2] = (float)((sqsum[c_r+2] - sqsum[c_o+2])-(sqsum[o_r+2] - sqsum[o_o+2])); 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])); 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 num = sum_[0]*template_sum_0 + sum_[1]*template_sum_1 + sum_[2]*template_sum_2 + sum_[3]*template_sum_3;
float denum = sqrt( tpl_sqsum * ( float denum = sqrt( template_sqsum * (
sqsum_[0] - weight * sum_[0]* sum_[0] + sqsum_[0] - weight * sum_[0]* sum_[0] +
sqsum_[1] - weight * sum_[1]* sum_[1] + sqsum_[1] - weight * sum_[1]* sum_[1] +
sqsum_[2] - weight * sum_[2]* sum_[2] + sqsum_[2] - weight * sum_[2]* sum_[2] +
sqsum_[3] - weight * sum_[3]* sum_[3] )); sqsum_[3] - weight * sum_[3]* sum_[3] ));
__global float * result = (__global float *)(res+res_idx); int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
*result = normAcc((*result) - num, denum); __global float * dstult = (__global float *)(dst+dst_idx);
*dstult = normAcc((*dstult) - num, denum);
} }
} }
#else
#error "cn should be 1, 2 or 4"
#endif
#endif

190
modules/imgproc/src/templmatch.cpp
Unescape View File

@ -49,80 +49,95 @@ namespace cv
#ifdef HAVE_OPENCL #ifdef HAVE_OPENCL
static bool useNaive(int method, int depth, const Size & size) /////////////////////////////////////////////////// CCORR //////////////////////////////////////////////////////////////
enum
{ {
#ifdef HAVE_CLAMDFFT SUM_1 = 0, SUM_2 = 1
if (method == TM_SQDIFF && depth == CV_32F) };
return true;
else if(method == TM_CCORR || (method == TM_SQDIFF && depth == CV_8U)) static bool sumTemplate(InputArray _src, UMat & result)
return size.height < 18 && size.width < 18; {
else int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
int wdepth = std::max(CV_32S, depth), wtype = CV_MAKE_TYPE(wdepth, cn);
size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
int wgs2_aligned = 1;
while (wgs2_aligned < (int)wgs)
wgs2_aligned <<= 1;
wgs2_aligned >>= 1;
char cvt[40];
ocl::Kernel k("calcSum", ocl::imgproc::match_template_oclsrc,
format("-D CALC_SUM -D T=%s -D WT=%s -D cn=%d -D convertToWT=%s -D WGS=%d -D WGS2_ALIGNED=%d -D wdepth=%d",
ocl::typeToStr(type), ocl::typeToStr(wtype), cn,
ocl::convertTypeStr(depth, wdepth, cn, cvt),
(int)wgs, wgs2_aligned, wdepth));
if (k.empty())
return false; return false;
#else
(void)(method);
(void)(depth);
(void)(size);
return true;
#endif
}
/////////////////////////////////////////////////// CCORR ////////////////////////////////////////////////////////////// UMat src = _src.getUMat();
result.create(1, 1, CV_32FC1);
ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
resarg = ocl::KernelArg::PtrWriteOnly(result);
k.args(srcarg, src.cols, (int)src.total(), resarg);
size_t globalsize = wgs;
return k.run(1, &globalsize, &wgs, false);
}
static bool matchTemplateNaive_CCORR(InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplateNaive_CCORR(InputArray _image, InputArray _templ, OutputArray _result)
{ {
int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
int wdepth = std::max(depth, CV_32S), wtype = CV_MAKE_TYPE(wdepth, cn);
char cvt[40];
ocl::Kernel k("matchTemplate_Naive_CCORR", ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_Naive_CCORR", ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn)); format("-D CCORR -D T=%s -D WT=%s -D convertToWT=%s -D cn=%d -D wdepth=%d", ocl::typeToStr(type), ocl::typeToStr(wtype),
ocl::convertTypeStr(depth, wdepth, cn, cvt), cn, wdepth));
if (k.empty()) if (k.empty())
return false; return false;
UMat image = _image.getUMat(), templ = _templ.getUMat(); UMat image = _image.getUMat(), templ = _templ.getUMat();
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32FC1);
UMat result = _result.getUMat(); UMat result = _result.getUMat();
k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ),
ocl::KernelArg::WriteOnly(result));
size_t globalsize[2] = { result.cols, result.rows }; size_t globalsize[2] = { result.cols, result.rows };
return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), return k.run(2, globalsize, NULL, false);
ocl::KernelArg::WriteOnly(result)).run(2, globalsize, NULL, false);
} }
static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
{ {
matchTemplate(_image, _templ, _result, CV_TM_CCORR); matchTemplate(_image, _templ, _result, CV_TM_CCORR);
int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = _image.type(), cn = CV_MAT_CN(type);
ocl::Kernel k("matchTemplate_CCORR_NORMED", ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_CCORR_NORMED", ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), format("-D CCORR_NORMED -D T=%s -D cn=%d", ocl::typeToStr(type), cn));
ocl::typeToStr(depth), cn));
if (k.empty()) if (k.empty())
return false; return false;
UMat image = _image.getUMat(), templ = _templ.getUMat(); UMat image = _image.getUMat(), templ = _templ.getUMat();
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32FC1);
UMat result = _result.getUMat(); UMat result = _result.getUMat();
UMat image_sums, image_sqsums; UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
UMat temp; UMat templ_sqsum;
multiply(templ, templ, temp, 1, CV_32F); if (!sumTemplate(templ, templ_sqsum))
Scalar s = sum(temp); return false;
float templ_sqsum = 0;
for (int i = 0; i < cn; ++i)
templ_sqsum += static_cast<float>(s[i]);
size_t globalsize[2] = { result.cols, result.rows }; k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result),
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, ocl::KernelArg::PtrReadOnly(templ_sqsum));
templ.rows, templ.cols, templ_sqsum).run(2, globalsize, NULL, false);
}
static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArray _result) size_t globalsize[2] = { result.cols, result.rows };
{ return k.run(2, globalsize, NULL, false);
if (useNaive(TM_CCORR, _image.depth(), _templ.size()) )
return matchTemplateNaive_CCORR(_image, _templ, _result);
else
return false;
} }
////////////////////////////////////// SQDIFF ////////////////////////////////////////////////////////////// ////////////////////////////////////// SQDIFF //////////////////////////////////////////////////////////////
@ -130,10 +145,12 @@ static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArra
static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result)
{ {
int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
int wdepth = std::max(depth, CV_32S), wtype = CV_MAKE_TYPE(wdepth, cn);
char cvt[40];
ocl::Kernel k("matchTemplate_Naive_SQDIFF", ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_Naive_SQDIFF", ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), format("-D SQDIFF -D T=%s -D WT=%s -D convertToWT=%s -D cn=%d -D wdepth=%d", ocl::typeToStr(type),
ocl::typeToStr(depth), cn)); ocl::typeToStr(wtype), ocl::convertTypeStr(depth, wdepth, cn, cvt), cn, wdepth));
if (k.empty()) if (k.empty())
return false; return false;
@ -141,20 +158,21 @@ static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, Outp
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
UMat result = _result.getUMat(); UMat result = _result.getUMat();
k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ),
ocl::KernelArg::WriteOnly(result));
size_t globalsize[2] = { result.cols, result.rows }; size_t globalsize[2] = { result.cols, result.rows };
return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), return k.run(2, globalsize, NULL, false);
ocl::KernelArg::WriteOnly(result)).run(2, globalsize, NULL, false);
} }
static bool matchTemplate_SQDIFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplate_SQDIFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
{ {
matchTemplate(_image, _templ, _result, CV_TM_CCORR); matchTemplate(_image, _templ, _result, CV_TM_CCORR);
int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = _image.type(), cn = CV_MAT_CN(type);
ocl::Kernel k("matchTemplate_SQDIFF_NORMED", ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_SQDIFF_NORMED", ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", format("-D SQDIFF_NORMED -D T=%s -D cn=%d", ocl::typeToStr(type), cn));
ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty()) if (k.empty())
return false; return false;
@ -165,24 +183,15 @@ static bool matchTemplate_SQDIFF_NORMED(InputArray _image, InputArray _templ, Ou
UMat image_sums, image_sqsums; UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
UMat temp; UMat templ_sqsum;
multiply(templ, templ, temp, 1, CV_32F); if (!sumTemplate(_templ, templ_sqsum))
Scalar s = sum(temp); return false;
float templ_sqsum = 0;
for (int i = 0; i < cn; ++i)
templ_sqsum += (float)s[i];
size_t globalsize[2] = { result.cols, result.rows }; k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result),
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, ocl::KernelArg::PtrReadOnly(templ_sqsum));
templ.rows, templ.cols, templ_sqsum).run(2, globalsize, NULL, false);
}
static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result) size_t globalsize[2] = { result.cols, result.rows };
{ return k.run(2, globalsize, NULL, false);
if (useNaive(TM_SQDIFF, _image.depth(), _templ.size()))
return matchTemplateNaive_SQDIFF(_image, _templ, _result);
else
return false;
} }
///////////////////////////////////// CCOEFF ///////////////////////////////////////////////////////////////// ///////////////////////////////////// CCOEFF /////////////////////////////////////////////////////////////////
@ -194,15 +203,15 @@ static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArr
UMat image_sums, temp; UMat image_sums, temp;
integral(_image, temp); integral(_image, temp);
if(temp.depth() == CV_64F) if (temp.depth() == CV_64F)
temp.convertTo(image_sums, CV_32F); temp.convertTo(image_sums, CV_32F);
else else
image_sums = temp; image_sums = temp;
int type = image_sums.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = image_sums.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
ocl::Kernel k(cv::format("matchTemplate_Prepared_CCOEFF_C%d", cn).c_str(), ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_Prepared_CCOEFF", ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn)); format("-D CCOEFF -D T=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty()) if (k.empty())
return false; return false;
@ -211,25 +220,28 @@ static bool matchTemplate_CCOEFF(InputArray _image, InputArray _templ, OutputArr
_result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F); _result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F);
UMat result = _result.getUMat(); UMat result = _result.getUMat();
size_t globalsize[2] = { result.cols, result.rows };
if (cn == 1) if (cn == 1)
{ {
float templ_sum = static_cast<float>(sum(_templ)[0]) / tsize.area(); float templ_sum = static_cast<float>(sum(_templ)[0]) / tsize.area();
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result),
templ.rows, templ.cols, templ_sum).run(2, globalsize, NULL, false); k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result),
templ.rows, templ.cols, templ_sum);
} }
else else
{ {
Vec4f templ_sum = Vec4f::all(0); Vec4f templ_sum = Vec4f::all(0);
templ_sum = sum(templ) / tsize.area(); templ_sum = sum(templ) / tsize.area();
if (cn == 2)
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(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::ReadWrite(result), templ.rows, templ.cols, if (cn == 2)
templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3]).run(2, globalsize, NULL, false); k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols,
templ_sum[0], templ_sum[1]);
else
k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols,
templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3]);
} }
size_t globalsize[2] = { result.cols, result.rows };
return k.run(2, globalsize, NULL, false);
} }
static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
@ -241,8 +253,8 @@ static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, Ou
int type = image_sums.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = image_sums.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
ocl::Kernel k(format("matchTemplate_CCOEFF_NORMED_C%d", cn).c_str(), ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_CCOEFF_NORMED", ocl::imgproc::match_template_oclsrc,
format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn)); format("-D CCOEFF_NORMED -D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty()) if (k.empty())
return false; return false;
@ -251,7 +263,6 @@ static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, Ou
_result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F); _result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F);
UMat result = _result.getUMat(); UMat result = _result.getUMat();
size_t globalsize[2] = { result.cols, result.rows };
float scale = 1.f / tsize.area(); float scale = 1.f / tsize.area();
if (cn == 1) if (cn == 1)
@ -270,9 +281,8 @@ static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, Ou
return true; return true;
} }
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum) ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum);
.run(2,globalsize,NULL,false);
} }
else else
{ {
@ -295,15 +305,17 @@ static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, Ou
} }
if (cn == 2) if (cn == 2)
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale, ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale,
templ_sum[0], templ_sum[1], templ_sqsum_sum).run(2, globalsize, NULL, false); templ_sum[0], templ_sum[1], templ_sqsum_sum);
else
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale, ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale,
templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3], templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3], templ_sqsum_sum);
templ_sqsum_sum).run(2, globalsize, NULL, false);
} }
size_t globalsize[2] = { result.cols, result.rows };
return k.run(2, globalsize, NULL, false);
} }
/////////////////////////////////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////////////////////////////////
@ -319,7 +331,7 @@ static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _
static const Caller callers[] = static const Caller callers[] =
{ {
matchTemplate_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplate_CCORR, matchTemplateNaive_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplateNaive_CCORR,
matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED
}; };
const Caller caller = callers[method]; const Caller caller = callers[method];

3
modules/imgproc/test/ocl/test_match_template.cpp
Unescape View File

@ -53,8 +53,7 @@ namespace ocl {
///////////////////////////////////////////// matchTemplate ////////////////////////////////////////////////////////// ///////////////////////////////////////////// matchTemplate //////////////////////////////////////////////////////////
CV_ENUM(MatchTemplType, CV_TM_SQDIFF, CV_TM_SQDIFF_NORMED, CV_TM_CCORR, CV_ENUM(MatchTemplType, CV_TM_CCORR, CV_TM_CCORR_NORMED, CV_TM_SQDIFF, CV_TM_SQDIFF_NORMED, CV_TM_CCOEFF, CV_TM_CCOEFF_NORMED)
CV_TM_CCORR_NORMED, CV_TM_CCOEFF, CV_TM_CCOEFF_NORMED)
PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, MatchTemplType, bool) PARAM_TEST_CASE(MatchTemplate, MatDepth, Channels, MatchTemplType, bool)
{ {

Write Preview
Loading…
Cancel
Save