Commit bab5700a authored by Vadim Pisarevsky's avatar Vadim Pisarevsky

Merge pull request #3013 from ElenaGvozdeva:ocl_matchTemplate

parents d611684d 90ac88cb
...@@ -90,11 +90,8 @@ __kernel void calcSum(__global const uchar * srcptr, int src_step, int src_offse ...@@ -90,11 +90,8 @@ __kernel void calcSum(__global const uchar * srcptr, int src_step, int src_offse
T src = loadpix(srcptr + src_index); T src = loadpix(srcptr + src_index);
tmp = convertToWT(src); tmp = convertToWT(src);
#if wdepth == 4
accumulator = mad24(tmp, tmp, accumulator);
#else
accumulator = mad(tmp, tmp, accumulator); accumulator = mad(tmp, tmp, accumulator);
#endif
} }
if (lid < WGS2_ALIGNED) if (lid < WGS2_ALIGNED)
...@@ -165,11 +162,9 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s ...@@ -165,11 +162,9 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s
{ {
T temp = (T)(template[j]); T temp = (T)(template[j]);
T src = *(__global const T*)(srcptr + ind + j*(int)sizeof(T1)); T src = *(__global const T*)(srcptr + ind + j*(int)sizeof(T1));
#if wdepth == 4
sum = mad24(convertToWT(src), convertToWT(temp), sum);
#else
sum = mad(convertToWT(src), convertToWT(temp), sum); sum = mad(convertToWT(src), convertToWT(temp), sum);
#endif
} }
ind += src_step; ind += src_step;
template = (__global const T1 *)((__global const uchar *)template + template_step); template = (__global const T1 *)((__global const uchar *)template + template_step);
...@@ -195,12 +190,7 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s ...@@ -195,12 +190,7 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s
#pragma unroll #pragma unroll
for (int cx=0, x = x0; cx < PIX_PER_WI_X && x < dst_cols; ++cx, ++x) for (int cx=0, x = x0; cx < PIX_PER_WI_X && x < dst_cols; ++cx, ++x)
{ {
#if wdepth == 4
sum[cx] = mad24(convertToWT1(src[j+cx]), convertToWT1(template[j]), sum[cx]);
#else
sum[cx] = mad(convertToWT1(src[j+cx]), convertToWT1(template[j]), sum[cx]); sum[cx] = mad(convertToWT1(src[j+cx]), convertToWT1(template[j]), sum[cx]);
#endif
} }
} }
...@@ -237,11 +227,8 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s ...@@ -237,11 +227,8 @@ __kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_s
{ {
T src = loadpix(srcptr + mad24(y+i, src_step, mad24(x+j, TSIZE, src_offset))); T src = loadpix(srcptr + mad24(y+i, src_step, mad24(x+j, TSIZE, src_offset)));
T template = loadpix(templateptr + mad24(i, template_step, mad24(j, TSIZE, template_offset))); T template = loadpix(templateptr + mad24(i, template_step, mad24(j, TSIZE, template_offset)));
#if wdepth == 4
sum = mad24(convertToWT(src), convertToWT(template), sum);
#else
sum = mad(convertToWT(src), convertToWT(template), sum); sum = mad(convertToWT(src), convertToWT(template), sum);
#endif
} }
} }
...@@ -296,11 +283,8 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * srcptr, int src_ ...@@ -296,11 +283,8 @@ __kernel void matchTemplate_Naive_SQDIFF(__global const uchar * srcptr, int src_
T template = loadpix(templateptr + mad24(i, template_step, mad24(j, TSIZE, template_offset))); T template = loadpix(templateptr + mad24(i, template_step, mad24(j, TSIZE, template_offset)));
value = convertToWT(src) - convertToWT(template); value = convertToWT(src) - convertToWT(template);
#if wdepth == 4
sum = mad24(value, value, sum);
#else
sum = mad(value, value, sum); sum = mad(value, value, sum);
#endif
} }
} }
......
...@@ -79,7 +79,7 @@ static bool extractFirstChannel_32F(InputArray _image, OutputArray _result, int ...@@ -79,7 +79,7 @@ static bool extractFirstChannel_32F(InputArray _image, OutputArray _result, int
static bool sumTemplate(InputArray _src, UMat & result) static bool sumTemplate(InputArray _src, UMat & result)
{ {
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); 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); int wdepth = CV_32F, wtype = CV_MAKE_TYPE(wdepth, cn);
size_t wgs = ocl::Device::getDefault().maxWorkGroupSize(); size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
int wgs2_aligned = 1; int wgs2_aligned = 1;
...@@ -89,10 +89,10 @@ static bool sumTemplate(InputArray _src, UMat & result) ...@@ -89,10 +89,10 @@ static bool sumTemplate(InputArray _src, UMat & result)
char cvt[40]; char cvt[40];
ocl::Kernel k("calcSum", ocl::imgproc::match_template_oclsrc, ocl::Kernel k("calcSum", ocl::imgproc::match_template_oclsrc,
format("-D CALC_SUM -D T=%s -D T1=%s -D WT=%s -D cn=%d -D convertToWT=%s -D WGS=%d -D WGS2_ALIGNED=%d -D wdepth=%d", format("-D CALC_SUM -D T=%s -D T1=%s -D WT=%s -D cn=%d -D convertToWT=%s -D WGS=%d -D WGS2_ALIGNED=%d",
ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(wtype), cn, ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(wtype), cn,
ocl::convertTypeStr(depth, wdepth, cn, cvt), ocl::convertTypeStr(depth, wdepth, cn, cvt),
(int)wgs, wgs2_aligned, wdepth)); (int)wgs, wgs2_aligned));
if (k.empty()) if (k.empty())
return false; return false;
...@@ -110,12 +110,8 @@ static bool sumTemplate(InputArray _src, UMat & result) ...@@ -110,12 +110,8 @@ static bool sumTemplate(InputArray _src, UMat & result)
static bool useNaive(Size size) static bool useNaive(Size size)
{ {
if (!ocl::Device::getDefault().isIntel())
return true;
int dft_size = 18; int dft_size = 18;
return size.height < dft_size && size.width < dft_size; return size.height < dft_size && size.width < dft_size;
} }
struct ConvolveBuf struct ConvolveBuf
...@@ -129,7 +125,6 @@ struct ConvolveBuf ...@@ -129,7 +125,6 @@ struct ConvolveBuf
UMat image_block, templ_block, result_data; UMat image_block, templ_block, result_data;
void create(Size image_size, Size templ_size); void create(Size image_size, Size templ_size);
static Size estimateBlockSize(Size result_size);
}; };
void ConvolveBuf::create(Size image_size, Size templ_size) void ConvolveBuf::create(Size image_size, Size templ_size)
...@@ -137,19 +132,26 @@ void ConvolveBuf::create(Size image_size, Size templ_size) ...@@ -137,19 +132,26 @@ void ConvolveBuf::create(Size image_size, Size templ_size)
result_size = Size(image_size.width - templ_size.width + 1, result_size = Size(image_size.width - templ_size.width + 1,
image_size.height - templ_size.height + 1); image_size.height - templ_size.height + 1);
block_size = user_block_size; const double blockScale = 4.5;
if (user_block_size.width == 0 || user_block_size.height == 0) const int minBlockSize = 256;
block_size = estimateBlockSize(result_size);
dft_size.width = 1 << int(ceil(std::log(block_size.width + templ_size.width - 1.) / std::log(2.))); block_size.width = cvRound(result_size.width*blockScale);
dft_size.height = 1 << int(ceil(std::log(block_size.height + templ_size.height - 1.) / std::log(2.))); block_size.width = std::max( block_size.width, minBlockSize - templ_size.width + 1 );
block_size.width = std::min( block_size.width, result_size.width );
block_size.height = cvRound(templ_size.height*blockScale);
block_size.height = std::max( block_size.height, minBlockSize - templ_size.height + 1 );
block_size.height = std::min( block_size.height, result_size.height );
dft_size.width = getOptimalDFTSize(block_size.width + templ_size.width - 1); dft_size.width = std::max(getOptimalDFTSize(block_size.width + templ_size.width - 1), 2);
dft_size.height = getOptimalDFTSize(block_size.height + templ_size.height - 1); dft_size.height = getOptimalDFTSize(block_size.height + templ_size.height - 1);
if( dft_size.width <= 0 || dft_size.height <= 0 )
CV_Error( CV_StsOutOfRange, "the input arrays are too big" );
// To avoid wasting time doing small DFTs // recompute block size
dft_size.width = std::max(dft_size.width, 512); block_size.width = dft_size.width - templ_size.width + 1;
dft_size.height = std::max(dft_size.height, 512); block_size.width = std::min( block_size.width, result_size.width);
block_size.height = dft_size.height - templ_size.height + 1;
block_size.height = std::min( block_size.height, result_size.height );
image_block.create(dft_size, CV_32F); image_block.create(dft_size, CV_32F);
templ_block.create(dft_size, CV_32F); templ_block.create(dft_size, CV_32F);
...@@ -164,15 +166,6 @@ void ConvolveBuf::create(Size image_size, Size templ_size) ...@@ -164,15 +166,6 @@ void ConvolveBuf::create(Size image_size, Size templ_size)
block_size.height = std::min(dft_size.height - templ_size.height + 1, result_size.height); block_size.height = std::min(dft_size.height - templ_size.height + 1, result_size.height);
} }
Size ConvolveBuf::estimateBlockSize(Size result_size)
{
int width = (result_size.width + 2) / 3;
int height = (result_size.height + 2) / 3;
width = std::min(width, result_size.width);
height = std::min(height, result_size.height);
return Size(width, height);
}
static bool convolve_dft(InputArray _image, InputArray _templ, OutputArray _result) static bool convolve_dft(InputArray _image, InputArray _templ, OutputArray _result)
{ {
ConvolveBuf buf; ConvolveBuf buf;
...@@ -202,7 +195,7 @@ static bool convolve_dft(InputArray _image, InputArray _templ, OutputArray _resu ...@@ -202,7 +195,7 @@ static bool convolve_dft(InputArray _image, InputArray _templ, OutputArray _resu
copyMakeBorder(templ_roi, templ_block, 0, templ_block.rows - templ_roi.rows, 0, copyMakeBorder(templ_roi, templ_block, 0, templ_block.rows - templ_roi.rows, 0,
templ_block.cols - templ_roi.cols, BORDER_ISOLATED); templ_block.cols - templ_roi.cols, BORDER_ISOLATED);
dft(templ_block, templ_spect, 0); dft(templ_block, templ_spect, 0, templ.rows);
// Process all blocks of the result matrix // Process all blocks of the result matrix
for (int y = 0; y < result.rows; y += block_size.height) for (int y = 0; y < result.rows; y += block_size.height)
...@@ -281,8 +274,8 @@ static bool matchTemplateNaive_CCORR(InputArray _image, InputArray _templ, Outpu ...@@ -281,8 +274,8 @@ static bool matchTemplateNaive_CCORR(InputArray _image, InputArray _templ, Outpu
const char* convertToWT = ocl::convertTypeStr(depth, wdepth, rated_cn, cvt1); const char* convertToWT = ocl::convertTypeStr(depth, wdepth, rated_cn, cvt1);
ocl::Kernel k("matchTemplate_Naive_CCORR", ocl::imgproc::match_template_oclsrc, ocl::Kernel k("matchTemplate_Naive_CCORR", ocl::imgproc::match_template_oclsrc,
format("-D CCORR -D T=%s -D T1=%s -D WT=%s -D WT1=%s -D convertToWT=%s -D convertToWT1=%s -D cn=%d -D wdepth=%d -D PIX_PER_WI_X=%d", ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(wtype1), ocl::typeToStr(wtype), format("-D CCORR -D T=%s -D T1=%s -D WT=%s -D WT1=%s -D convertToWT=%s -D convertToWT1=%s -D cn=%d -D PIX_PER_WI_X=%d", ocl::typeToStr(type), ocl::typeToStr(depth), ocl::typeToStr(wtype1), ocl::typeToStr(wtype),
convertToWT, convertToWT1, cn, wdepth, pxPerWIx)); convertToWT, convertToWT1, cn, pxPerWIx));
if (k.empty()) if (k.empty())
return false; return false;
...@@ -358,8 +351,8 @@ static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, Outp ...@@ -358,8 +351,8 @@ static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, Outp
char cvt[40]; 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 SQDIFF -D T=%s -D T1=%s -D WT=%s -D convertToWT=%s -D cn=%d -D wdepth=%d", ocl::typeToStr(type), ocl::typeToStr(depth), format("-D SQDIFF -D T=%s -D T1=%s -D WT=%s -D convertToWT=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth),
ocl::typeToStr(wtype), ocl::convertTypeStr(depth, wdepth, cn, cvt), cn, wdepth)); ocl::typeToStr(wtype), ocl::convertTypeStr(depth, wdepth, cn, cvt), cn));
if (k.empty()) if (k.empty())
return false; return false;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment