From 4da1ba56b3c705bd9217ee16d810d754368ff5de Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 17 Jan 2014 15:01:22 +0400 Subject: [PATCH] fixed --- modules/imgproc/src/opencl/match_template.cl | 126 ++++++++++++--------------- modules/imgproc/src/templmatch.cpp | 77 ++++++---------- 2 files changed, 87 insertions(+), 116 deletions(-) diff --git a/modules/imgproc/src/opencl/match_template.cl b/modules/imgproc/src/opencl/match_template.cl index 89a8d61..451ec0c 100644 --- a/modules/imgproc/src/opencl/match_template.cl +++ b/modules/imgproc/src/opencl/match_template.cl @@ -29,15 +29,13 @@ // or tort (including negligence or otherwise) arising in any way out of // the use of this software, even if advised of the possibility of such damage. -#define DATA_TYPE type #define DATA_SIZE ((int)sizeof(type)) #define ELEM_TYPE elem_type #define ELEM_SIZE ((int)sizeof(elem_type)) #define CN cn -#define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, (gidx + img_sqsums_offset + ox) * CN) -#define SQSUMS(ox, oy) mad24(gidy + oy, img_sqsums_step, (gidx*CN + img_sqsums_offset + ox*CN)) -#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, (gidx*CN + img_sums_offset + ox*CN)) +#define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, gidx*CN + img_sqsums_offset + ox*CN) +#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx*CN + img_sums_offset + ox*CN) inline float normAcc(float num, float denum) { @@ -76,10 +74,7 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step int i,j; float sum = 0; - res_step /= sizeof(float); - res_offset /= sizeof(float); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -90,12 +85,13 @@ __kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step for(j = 0; j < tpl_cols; j ++) +#pragma unroll for (int c = 0; c < CN; c++) sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]); } - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = sum; } } @@ -109,10 +105,8 @@ __kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, in img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -121,7 +115,7 @@ __kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, in (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum)); } } @@ -138,10 +132,7 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step float delta; float sum = 0; - res_step /= sizeof(float); - res_offset /= sizeof(float); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -152,13 +143,14 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step for(j = 0; j < tpl_cols; j ++) +#pragma unroll for (int c = 0; c < CN; c++) { delta = (float)(img_ptr[j*CN+c] - tpl_ptr[j*CN+c]); sum += delta*delta; } } - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = sum; } } @@ -172,10 +164,8 @@ __kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, i img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -184,7 +174,7 @@ __kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, i (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) - (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)])); - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum)); } @@ -201,10 +191,8 @@ __kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, img_sums_step /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE; - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float image_sum_ = 0; if(gidx < res_cols && gidy < res_rows) @@ -214,7 +202,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, image_sum_ += (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)])- (sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])) * tpl_sum; - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result -= image_sum_; } } @@ -228,10 +216,8 @@ __kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, img_sums_step /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE; - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float image_sum_ = 0; if(gidx < res_cols && gidy < res_rows) @@ -241,37 +227,40 @@ __kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, image_sum_ += tpl_sum_0 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)]) -(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])); image_sum_ += tpl_sum_1 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+1] - sum[SUMS_PTR(tpl_cols, 0)+1])-(sum[SUMS_PTR(0, tpl_rows)+1] - sum[SUMS_PTR(0, 0)+1])); - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result -= image_sum_; } } __kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, - __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, - int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3) + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, + int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3) { int gidx = get_global_id(0); int gidy = get_global_id(1); img_sums_step /= ELEM_SIZE; img_sums_offset /= ELEM_SIZE; - res_step /= sizeof(float); - res_offset /= sizeof(float); - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float image_sum_ = 0; if(gidx < res_cols && gidy < res_rows) { __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); - image_sum_ += tpl_sum_0 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)]) -(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])); - image_sum_ += tpl_sum_1 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+1] - sum[SUMS_PTR(tpl_cols, 0)+1])-(sum[SUMS_PTR(0, tpl_rows)+1] - sum[SUMS_PTR(0, 0)+1])); - image_sum_ += tpl_sum_2 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+2] - sum[SUMS_PTR(tpl_cols, 0)+2])-(sum[SUMS_PTR(0, tpl_rows)+2] - sum[SUMS_PTR(0, 0)+2])); - image_sum_ += tpl_sum_3 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+3] - sum[SUMS_PTR(tpl_cols, 0)+3])-(sum[SUMS_PTR(0, tpl_rows)+3] - sum[SUMS_PTR(0, 0)+3])); + int c_r = SUMS_PTR(tpl_cols, tpl_rows); + int c_o = SUMS_PTR(tpl_cols, 0); + int o_r = SUMS_PTR(0,tpl_rows); + int oo = SUMS_PTR(0, 0); + + 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; + __global float * result = (__global float *)(res+res_idx); *result -= image_sum_; } @@ -279,7 +268,7 @@ __kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, - __global float * res, int res_step, int res_offset, int res_rows, int res_cols, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, int t_rows, int t_cols, float weight, float tpl_sum, float tpl_sqsum) { int gidx = get_global_id(0); @@ -289,11 +278,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i img_sums_step /= ELEM_SIZE; img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(*res); - res_offset /= sizeof(*res); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); if(gidx < res_cols && gidy < res_rows) { @@ -306,7 +292,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i float image_sqsum_ = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)]) - (sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)])); - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = normAcc((*result) - image_sum_ * tpl_sum, sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_))); @@ -315,7 +301,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, i __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, - __global float * res, int res_step, int res_offset, int res_rows, int res_cols, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, int t_rows, int t_cols, float weight, float tpl_sum_0, float tpl_sum_1, float tpl_sqsum) { int gidx = get_global_id(0); @@ -325,11 +311,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i img_sums_step /= ELEM_SIZE; img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(*res); - res_offset /= sizeof(*res); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float sum_[2]; float sqsum_[2]; @@ -342,22 +325,22 @@ __kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, i sum_[0] = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)])-(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)])); sum_[1] = (float)((sum[SUMS_PTR(t_cols, t_rows)+1] - sum[SUMS_PTR(t_cols, 0)+1])-(sum[SUMS_PTR(0, t_rows)+1] - sum[SUMS_PTR(0, 0)+1])); - sqsum_[0] = (float)((sqsum[SQSUMS(t_cols, t_rows)] - sqsum[SQSUMS(t_cols, 0)])-(sqsum[SQSUMS(0, t_rows)] - sqsum[SQSUMS(0, 0)])); - sqsum_[1] = (float)((sqsum[SQSUMS(t_cols, t_rows)+1] - sqsum[SQSUMS(t_cols, 0)+1])-(sqsum[SQSUMS(0, t_rows)+1] - sqsum[SQSUMS(0, 0)+1])); + sqsum_[0] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)])-(sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)])); + sqsum_[1] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)+1] - sqsum[SQSUMS_PTR(t_cols, 0)+1])-(sqsum[SQSUMS_PTR(0, t_rows)+1] - sqsum[SQSUMS_PTR(0, 0)+1])); float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1; float denum = sqrt( tpl_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] + sqsum_[1] - weight * sum_[1]* sum_[1])); - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = normAcc((*result) - num, denum); } } __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset, __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset, - __global float * res, int res_step, int res_offset, int res_rows, int res_cols, + __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols, int t_rows, int t_cols, float weight, float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3, float tpl_sqsum) @@ -369,11 +352,8 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i img_sums_step /= ELEM_SIZE; img_sqsums_step /= sizeof(float); img_sqsums_offset /= sizeof(float); - res_step /= sizeof(*res); - res_offset /= sizeof(*res); - - int res_idx = mad24(gidy, res_step, res_offset + gidx); + int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float)); float sum_[4]; float sqsum_[4]; @@ -383,15 +363,25 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums); __global float * sqsum = (__global float*)(img_sqsums); - sum_[0] = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)])-(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)])); - sum_[1] = (float)((sum[SUMS_PTR(t_cols, t_rows)+1] - sum[SUMS_PTR(t_cols, 0)+1])-(sum[SUMS_PTR(0, t_rows)+1] - sum[SUMS_PTR(0, 0)+1])); - sum_[2] = (float)((sum[SUMS_PTR(t_cols, t_rows)+2] - sum[SUMS_PTR(t_cols, 0)+2])-(sum[SUMS_PTR(0, t_rows)+2] - sum[SUMS_PTR(0, 0)+2])); - sum_[3] = (float)((sum[SUMS_PTR(t_cols, t_rows)+3] - sum[SUMS_PTR(t_cols, 0)+3])-(sum[SUMS_PTR(0, t_rows)+3] - sum[SUMS_PTR(0, 0)+3])); + int c_r = SUMS_PTR(t_cols, t_rows); + int c_o = SUMS_PTR(t_cols, 0); + int o_r = SUMS_PTR(0, t_rows); + int o_o = SUMS_PTR(0, 0); + + 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_[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])); + sqsum_[0] = (float)((sqsum[c_r] - sqsum[c_o]) -(sqsum[o_r] - sqsum[o_o])); + sqsum_[1] = (float)((sqsum[c_r+1] - sqsum[c_o+1])-(sqsum[o_r+1] - sqsum[o_o+1])); + sqsum_[2] = (float)((sqsum[c_r+2] - sqsum[c_o+2])-(sqsum[o_r+2] - sqsum[o_o+2])); + sqsum_[3] = (float)((sqsum[c_r+3] - sqsum[c_o+3])-(sqsum[o_r+3] - sqsum[o_o+3])); float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1 + sum_[2]*tpl_sum_2 + sum_[3]*tpl_sum_3; @@ -401,7 +391,7 @@ __kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, i sqsum_[2] - weight * sum_[2]* sum_[2] + sqsum_[3] - weight * sum_[3]* sum_[3] )); - __global float * result = (__global float *)(res)+res_idx; + __global float * result = (__global float *)(res+res_idx); *result = normAcc((*result) - num, denum); } } \ No newline at end of file diff --git a/modules/imgproc/src/templmatch.cpp b/modules/imgproc/src/templmatch.cpp index ae10cd9..bb3423b 100644 --- a/modules/imgproc/src/templmatch.cpp +++ b/modules/imgproc/src/templmatch.cpp @@ -101,9 +101,8 @@ namespace cv result = _result.getUMat(); size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false); } static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result) @@ -124,24 +123,18 @@ namespace cv _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - UMat temp, image_sums, image_sqsums; - integral(image.reshape(1), image_sums, temp); - - if(temp.depth() == CV_64F) - temp.convertTo(image_sqsums, CV_32F); - else - image_sqsums = temp; + UMat image_sums, image_sqsums; + integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); - UMat templ_resh; + UMat templ_resh, temp; templ.reshape(1).convertTo(templ_resh, CV_32F); multiply(templ_resh, templ_resh, temp); unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0]; size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false); } //////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////// @@ -173,9 +166,8 @@ namespace cv result = _result.getUMat(); size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::WriteOnly(result)).run(2,globalsize,NULL,false); } static bool matchTemplate_SQDIFF_NORMED (InputArray _image, InputArray _templ, OutputArray _result) @@ -196,24 +188,18 @@ namespace cv _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - UMat temp, image_sums, image_sqsums; - integral(image.reshape(1), image_sums, temp); + UMat image_sums, image_sqsums; + integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F); - if(temp.depth() == CV_64F) - temp.convertTo(image_sqsums, CV_32F); - else - image_sqsums = temp; - - UMat templ_resh; + UMat temp, templ_resh; templ.reshape(1).convertTo(templ_resh, CV_32F); multiply(templ_resh, templ_resh, temp); unsigned long long templ_sqsum = (unsigned long long)sum(temp)[0]; size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sqsum).run(2,globalsize,NULL,false); } /////////////////////////////////////CCOEFF///////////////////////////////////////////////////////////////// @@ -242,17 +228,16 @@ namespace cv return false; UMat templ = _templ.getUMat(), result; - int image_rows = _image.size().height, image_cols = _image.size().width; - _result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F); + Size size = _image.size(); + _result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F); result = _result.getUMat(); size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; if (cn==1) { float templ_sum = (float)sum(_templ)[0]/ _templ.size().area(); - return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,localsize,false); + return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, templ_sum).run(2,globalsize,NULL,false); } else { @@ -260,10 +245,10 @@ namespace cv templ_sum = sum(templ)/ templ.size().area(); if (cn==2) return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, - templ_sum[0],templ_sum[1]).run(2,globalsize,localsize,false); + templ_sum[0],templ_sum[1]).run(2,globalsize,NULL,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, - templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,localsize,false); + templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3]).run(2,globalsize,NULL,false); } } @@ -279,7 +264,7 @@ namespace cv const char * kernelName; UMat temp, image_sums, image_sqsums; - integral(_image,image_sums, temp); + integral(_image,image_sums, image_sqsums, CV_32F, CV_32F); int type = image_sums.type(); int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); @@ -302,13 +287,7 @@ namespace cv _result.create(image_rows - templ.rows + 1, image_cols - templ.cols + 1, CV_32F); result = _result.getUMat(); - if(temp.depth() == CV_64F) - temp.convertTo(image_sqsums, CV_32F); - else - image_sqsums = temp; - size_t globalsize[2] = {result.cols, result.rows}; - size_t localsize[2] = {16, 16}; float scale = 1.f / templ.size().area(); @@ -330,7 +309,7 @@ namespace cv return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums),ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum) - .run(2,globalsize,localsize,false); + .run(2,globalsize,NULL,false); } else { @@ -360,12 +339,12 @@ namespace cv return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum[0],templ_sum[1], templ_sqsum_sum) - .run(2,globalsize,localsize,false); + .run(2,globalsize,NULL,false); return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::WriteOnly(result), templ.rows, templ.cols, scale, templ_sum[0],templ_sum[1],templ_sum[2],templ_sum[3], templ_sqsum_sum) - .run(2,globalsize,localsize,false); + .run(2,globalsize,NULL,false); } } @@ -374,10 +353,10 @@ namespace cv static bool ocl_matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, int method) { - int type = _img.type(); - int cn = CV_MAT_CN(type); + int cn = CV_MAT_CN(_img.type()); - CV_Assert( cn == _templ.channels() && cn!=3 && cn<=4); + if (cn == 3 || cn > 4) + return false; typedef bool (*Caller)(InputArray _img, InputArray _templ, OutputArray _result); @@ -588,13 +567,15 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, CV_Assert( (_img.depth() == CV_8U || _img.depth() == CV_32F) && _img.type() == _templ.type() ); - CV_Assert(_img.size().height >= _templ.size().height && _img.size().width >= _templ.size().width); - CV_Assert(_img.dims() <= 2); + + bool swapNotNeed = (_img.size().height >= _templ.size().height && _img.size().width >= _templ.size().width); + if (!swapNotNeed) + CV_Assert(_img.size().height <= _templ.size().height && _img.size().width <= _templ.size().width); bool use_opencl = ocl::useOpenCL() && _result.isUMat(); - if ( use_opencl && ocl_matchTemplate(_img,_templ,_result,method)) - return; + if ( use_opencl && (swapNotNeed ? ocl_matchTemplate(_img,_templ,_result,method) : ocl_matchTemplate(_templ,_img,_result,method))) + return; int numType = method == CV_TM_CCORR || method == CV_TM_CCORR_NORMED ? 0 : method == CV_TM_CCOEFF || method == CV_TM_CCOEFF_NORMED ? 1 : 2; @@ -603,7 +584,7 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result, method == CV_TM_CCOEFF_NORMED; Mat img = _img.getMat(), templ = _templ.getMat(); - if( img.rows < templ.rows || img.cols < templ.cols ) + if(!swapNotNeed ) std::swap(img, templ); Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1); -- 2.7.4