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

fixed

Browse Source
pull/2149/head
Elena Gvozdeva 11 years ago
parent 47b572f99f
commit 4da1ba56b3
2 changed files with 87 additions and 116 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. 122
      modules/imgproc/src/opencl/match_template.cl
  2. 75
      modules/imgproc/src/templmatch.cpp

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

@ -29,15 +29,13 @@
// or tort (including negligence or otherwise) arising in any way out of // 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. // 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 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 CN cn
#define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, (gidx + img_sqsums_offset + ox) * CN) #define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, gidx*CN + img_sqsums_offset + ox*CN)
#define SQSUMS(ox, oy) mad24(gidy + oy, img_sqsums_step, (gidx*CN + img_sqsums_offset + ox*CN)) #define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx*CN + img_sums_offset + ox*CN)
#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, (gidx*CN + img_sums_offset + ox*CN))
inline float normAcc(float num, float denum) inline float normAcc(float num, float denum)
{ {
@ -76,10 +74,7 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step
int i,j; int i,j;
float sum = 0; float sum = 0;
res_step /= sizeof(float); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx);
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
{ {
@ -90,12 +85,13 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step
for(j = 0; j < tpl_cols; j ++) for(j = 0; j < tpl_cols; j ++)
#pragma unroll
for (int c = 0; c < CN; c++) for (int c = 0; c < CN; c++)
sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]); sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]);
} }
__global float * result = (__global float *)(res)+res_idx; __global float * result = (__global float *)(res+res_idx);
*result = sum; *result = sum;
} }
} }
@ -109,10 +105,8 @@ __kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, in
img_sqsums_step /= sizeof(float); img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float); img_sqsums_offset /= sizeof(float);
res_step /= sizeof(float);
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
{ {
@ -121,7 +115,7 @@ __kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, in
(sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
__global float * result = (__global float *)(res)+res_idx; __global float * result = (__global float *)(res+res_idx);
*result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum)); *result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum));
} }
} }
@ -138,10 +132,7 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step
float delta; float delta;
float sum = 0; float sum = 0;
res_step /= sizeof(float); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx);
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
{ {
@ -152,13 +143,14 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step
for(j = 0; j < tpl_cols; j ++) for(j = 0; j < tpl_cols; j ++)
#pragma unroll
for (int c = 0; c < CN; c++) for (int c = 0; c < CN; c++)
{ {
delta = (float)(img_ptr[j*CN+c] - tpl_ptr[j*CN+c]); delta = (float)(img_ptr[j*CN+c] - tpl_ptr[j*CN+c]);
sum += delta*delta; sum += delta*delta;
} }
} }
__global float * result = (__global float *)(res)+res_idx; __global float * result = (__global float *)(res+res_idx);
*result = sum; *result = sum;
} }
} }
@ -172,10 +164,8 @@ __kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, i
img_sqsums_step /= sizeof(float); img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float); img_sqsums_offset /= sizeof(float);
res_step /= sizeof(float);
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
{ {
@ -184,7 +174,7 @@ __kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, i
(sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
__global float * result = (__global float *)(res)+res_idx; __global float * result = (__global float *)(res+res_idx);
*result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum)); *result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum));
} }
@ -201,10 +191,8 @@ __kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums,
img_sums_step /= ELEM_SIZE; img_sums_step /= ELEM_SIZE;
img_sums_offset /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE;
res_step /= sizeof(float);
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float image_sum_ = 0; float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
@ -214,7 +202,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums,
image_sum_ += (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)])- 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; (sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])) * tpl_sum;
__global float * result = (__global float *)(res)+res_idx; __global float * result = (__global float *)(res+res_idx);
*result -= image_sum_; *result -= image_sum_;
} }
} }
@ -228,10 +216,8 @@ __kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums,
img_sums_step /= ELEM_SIZE; img_sums_step /= ELEM_SIZE;
img_sums_offset /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE;
res_step /= sizeof(float);
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float image_sum_ = 0; float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
@ -241,7 +227,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums,
image_sum_ += tpl_sum_0 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)]) -(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])); image_sum_ += tpl_sum_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_1 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+1] - sum[SUMS_PTR(tpl_cols, 0)+1])-(sum[SUMS_PTR(0, tpl_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
__global float * result = (__global float *)(res)+res_idx; __global float * result = (__global float *)(res+res_idx);
*result -= image_sum_; *result -= image_sum_;
} }
@ -256,22 +242,25 @@ __kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums,
img_sums_step /= ELEM_SIZE; img_sums_step /= ELEM_SIZE;
img_sums_offset /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE;
res_step /= sizeof(float);
res_offset /= sizeof(float);
int res_idx = mad24(gidy, res_step, res_offset + gidx); int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
float image_sum_ = 0; float image_sum_ = 0;
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
{ {
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); __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)])); int c_r = SUMS_PTR(tpl_cols, tpl_rows);
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])); int c_o = SUMS_PTR(tpl_cols, 0);
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])); int o_r = SUMS_PTR(0,tpl_rows);
image_sum_ += tpl_sum_3 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+3] - sum[SUMS_PTR(tpl_cols, 0)+3])-(sum[SUMS_PTR(0, tpl_rows)+3] - sum[SUMS_PTR(0, 0)+3])); int oo = SUMS_PTR(0, 0);
__global float * result = (__global float *)(res)+res_idx; image_sum_ += tpl_sum_0 * (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[oo]));
image_sum_ += tpl_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1]));
image_sum_ += tpl_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2]));
image_sum_ += tpl_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3]));
__global float * result = (__global float *)(res+res_idx);
*result -= image_sum_; *result -= image_sum_;
} }
@ -279,7 +268,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums,
__kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __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 const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global float * res, int res_step, int res_offset, int res_rows, int res_cols, __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int t_rows, int t_cols, float weight, float tpl_sum, float tpl_sqsum) int t_rows, int t_cols, float weight, float tpl_sum, float tpl_sqsum)
{ {
int gidx = get_global_id(0); int gidx = get_global_id(0);
@ -289,11 +278,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i
img_sums_step /= ELEM_SIZE; img_sums_step /= ELEM_SIZE;
img_sqsums_step /= sizeof(float); img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float); img_sqsums_offset /= sizeof(float);
res_step /= sizeof(*res);
res_offset /= sizeof(*res);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
int res_idx = mad24(gidy, res_step, res_offset + gidx);
if(gidx < res_cols && gidy < res_rows) if(gidx < res_cols && gidy < res_rows)
{ {
@ -306,7 +292,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i
float image_sqsum_ = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)]) - 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; __global float * result = (__global float *)(res+res_idx);
*result = normAcc((*result) - image_sum_ * tpl_sum, *result = normAcc((*result) - image_sum_ * tpl_sum,
sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_))); sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_)));
@ -315,7 +301,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i
__kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __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 const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global float * res, int res_step, int res_offset, int res_rows, int res_cols, __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int t_rows, int t_cols, float weight, float tpl_sum_0, float tpl_sum_1, float tpl_sqsum) int 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 gidx = get_global_id(0);
@ -325,11 +311,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i
img_sums_step /= ELEM_SIZE; img_sums_step /= ELEM_SIZE;
img_sqsums_step /= sizeof(float); img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float); img_sqsums_offset /= sizeof(float);
res_step /= sizeof(*res);
res_offset /= sizeof(*res);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
int res_idx = mad24(gidy, res_step, res_offset + gidx);
float sum_[2]; float sum_[2];
float sqsum_[2]; float sqsum_[2];
@ -342,22 +325,22 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i
sum_[0] = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)])-(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)])); sum_[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]));
sqsum_[0] = (float)((sqsum[SQSUMS(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(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(t_cols, t_rows)+1] - sqsum[SQSUMS(t_cols, 0)+1])-(sqsum[SQSUMS(0, t_rows)+1] - sqsum[SQSUMS(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]*tpl_sum_0 + sum_[1]*tpl_sum_1;
float denum = sqrt( tpl_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] + float denum = sqrt( tpl_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; __global float * result = (__global float *)(res+res_idx);
*result = normAcc((*result) - num, denum); *result = normAcc((*result) - num, denum);
} }
} }
__kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __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 const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
__global float * res, int res_step, int res_offset, int res_rows, int res_cols, __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
int t_rows, int t_cols, float weight, 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_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3,
float tpl_sqsum) float tpl_sqsum)
@ -369,11 +352,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i
img_sums_step /= ELEM_SIZE; img_sums_step /= ELEM_SIZE;
img_sqsums_step /= sizeof(float); img_sqsums_step /= sizeof(float);
img_sqsums_offset /= sizeof(float); img_sqsums_offset /= sizeof(float);
res_step /= sizeof(*res);
res_offset /= sizeof(*res);
int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
int res_idx = mad24(gidy, res_step, res_offset + gidx);
float sum_[4]; float sum_[4];
float sqsum_[4]; float sqsum_[4];
@ -383,15 +363,25 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i
__global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
__global float * sqsum = (__global float*)(img_sqsums); __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)])); int c_r = SUMS_PTR(t_cols, t_rows);
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])); int c_o = SUMS_PTR(t_cols, 0);
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])); int o_r = SUMS_PTR(0, t_rows);
sum_[3] = (float)((sum[SUMS_PTR(t_cols, t_rows)+3] - sum[SUMS_PTR(t_cols, 0)+3])-(sum[SUMS_PTR(0, t_rows)+3] - sum[SUMS_PTR(0, 0)+3])); int 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[SQSUMS(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(0, 0)])); sqsum_[0] = (float)((sqsum[c_r] - sqsum[c_o]) -(sqsum[o_r] - sqsum[o_o]));
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_[1] = (float)((sqsum[c_r+1] - sqsum[c_o+1])-(sqsum[o_r+1] - sqsum[o_o+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_[2] = (float)((sqsum[c_r+2] - sqsum[c_o+2])-(sqsum[o_r+2] - sqsum[o_o+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])); 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]*tpl_sum_0 + sum_[1]*tpl_sum_1 + sum_[2]*tpl_sum_2 + sum_[3]*tpl_sum_3;
@ -401,7 +391,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i
sqsum_[2] - weight * sum_[2]* sum_[2] + sqsum_[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; __global float * result = (__global float *)(res+res_idx);
*result = normAcc((*result) - num, denum); *result = normAcc((*result) - num, denum);
} }
} }

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

@ -101,9 +101,8 @@ namespace cv
result = _result.getUMat(); result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows}; size_t globalsize[2] = {result.cols, result.rows};
size_t localsize[2] = {16, 16};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false);
} }
static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
@ -124,24 +123,18 @@ namespace cv
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
result = _result.getUMat(); result = _result.getUMat();
UMat temp, image_sums, image_sqsums; UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, temp); integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
if(temp.depth() == CV_64F)
temp.convertTo(image_sqsums, CV_32F);
else
image_sqsums = temp;
UMat templ_resh; UMat templ_resh, temp;
templ.reshape(1).convertTo(templ_resh, CV_32F); templ.reshape(1).convertTo(templ_resh, CV_32F);
multiply(templ_resh, templ_resh, temp); multiply(templ_resh, templ_resh, temp);
unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0]; unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0];
size_t globalsize[2] = {result.cols, result.rows}; size_t globalsize[2] = {result.cols, result.rows};
size_t localsize[2] = {16, 16};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false);
} }
//////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// //////////////////////////////////////SQDIFF//////////////////////////////////////////////////////////////
@ -173,9 +166,8 @@ namespace cv
result = _result.getUMat(); result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows}; size_t globalsize[2] = {result.cols, result.rows};
size_t localsize[2] = {16, 16};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false);
} }
static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result) static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result)
@ -196,24 +188,18 @@ namespace cv
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
result = _result.getUMat(); result = _result.getUMat();
UMat temp, image_sums, image_sqsums; UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, temp); integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
if(temp.depth() == CV_64F) UMat temp, templ_resh;
temp.convertTo(image_sqsums, CV_32F);
else
image_sqsums = temp;
UMat templ_resh;
templ.reshape(1).convertTo(templ_resh, CV_32F); templ.reshape(1).convertTo(templ_resh, CV_32F);
multiply(templ_resh, templ_resh, temp); multiply(templ_resh, templ_resh, temp);
unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0]; unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0];
size_t globalsize[2] = {result.cols, result.rows}; size_t globalsize[2] = {result.cols, result.rows};
size_t localsize[2] = {16, 16};
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false);
} }
/////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// /////////////////////////////////////CCOEFF/////////////////////////////////////////////////////////////////
@ -242,17 +228,16 @@ namespace cv
return false; return false;
UMat templ = _templ.getUMat(), result; UMat templ = _templ.getUMat(), result;
int image_rows = _image.size().height, image_cols = _image.size().width; Size size = _image.size();
_result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F); _result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F);
result = _result.getUMat(); result = _result.getUMat();
size_t globalsize[2] = {result.cols, result.rows}; size_t globalsize[2] = {result.cols, result.rows};
size_t localsize[2] = {16, 16};
if (cn==1) if (cn==1)
{ {
float templ_sum = (float)sum(_templ)[0]/ _templ.size().area(); float templ_sum = (float)sum(_templ)[0]/ _templ.size().area();
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,localsize,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,NULL,false);
} }
else else
{ {
@ -260,10 +245,10 @@ namespace cv
templ_sum = sum(templ)/ templ.size().area(); templ_sum = sum(templ)/ templ.size().area();
if (cn==2) if (cn==2)
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols,
templ_sum[0],templ_sum[1]).run(2,globalsize,localsize,false); templ_sum[0],templ_sum[1]).run(2,globalsize,NULL,false);
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols,
templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,localsize,false); templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,NULL,false);
} }
} }
@ -279,7 +264,7 @@ namespace cv
const char * kernelName; const char * kernelName;
UMat temp, image_sums, image_sqsums; UMat temp, image_sums, image_sqsums;
integral(_image,image_sums, temp); integral(_image,image_sums, image_sqsums, CV_32F, CV_32F);
int type = image_sums.type(); int type = image_sums.type();
int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
@ -302,13 +287,7 @@ namespace cv
_result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F); _result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F);
result = _result.getUMat(); result = _result.getUMat();
if(temp.depth() == CV_64F)
temp.convertTo(image_sqsums, CV_32F);
else
image_sqsums = temp;
size_t globalsize[2] = {result.cols, result.rows}; size_t globalsize[2] = {result.cols, result.rows};
size_t localsize[2] = {16, 16};
float scale = 1.f / templ.size().area(); float scale = 1.f / templ.size().area();
@ -330,7 +309,7 @@ namespace cv
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums),ocl::KernelArg::ReadOnlyNoSize(image_sqsums), 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) ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum)
.run(2,globalsize,localsize,false); .run(2,globalsize,NULL,false);
} }
else else
{ {
@ -360,12 +339,12 @@ namespace cv
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale,
templ_sum[0],templ_sum[1], templ_sqsum_sum) templ_sum[0],templ_sum[1], templ_sqsum_sum)
.run(2,globalsize,localsize,false); .run(2,globalsize,NULL,false);
return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale,
templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3], templ_sqsum_sum) templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3], templ_sqsum_sum)
.run(2,globalsize,localsize,false); .run(2,globalsize,NULL,false);
} }
} }
@ -374,10 +353,10 @@ namespace cv
static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method) static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method)
{ {
int type = _img.type(); int cn = CV_MAT_CN(_img.type());
int cn = CV_MAT_CN(type);
CV_Assert( cn == _templ.channels() && cn!=3 && cn<=4); if (cn == 3 || cn > 4)
return false;
typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result); typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result);
@ -588,12 +567,14 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result,
CV_Assert( (_img.depth() == CV_8U || _img.depth() == CV_32F) && _img.type() == _templ.type() ); CV_Assert( (_img.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); 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(); bool use_opencl = ocl::useOpenCL() && _result.isUMat();
if ( use_opencl && ocl_matchTemplate(_img,_templ,_result,method)) if ( use_opencl && (swapNotNeed ? ocl_matchTemplate(_img,_templ,_result,method) : ocl_matchTemplate(_templ,_img,_result,method)))
return; 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 :
@ -603,7 +584,7 @@ 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);
Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1); Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1);

Write Preview
Loading…
Cancel
Save