generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Optimize OpenCL version of morfology and box filters for small filter kernels

Browse Source
This commit is contained in:
vbystricky
2014-07-28 15:48:02 +04:00
parent 7da0e37ba8
commit 878dec65c8
4 changed files with 400 additions and 64 deletions
Download Patch File Download Diff File Expand all files Collapse all files

176
modules/imgproc/src/morph.cpp
View File

@@ -1339,20 +1339,188 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst,
#ifdef HAVE_OPENCL
#define ROUNDUP(sz, n) ((sz) + (n) - 1 - (((sz) + (n) - 1) % (n)))
static bool ocl_morphSmall( InputArray _src, OutputArray _dst, InputArray _kernel, Point anchor, int borderType,
int op, int actual_op = -1, InputArray _extraMat = noArray())
{
const ocl::Device & dev = ocl::Device::getDefault();
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), esz = CV_ELEM_SIZE(type);
bool doubleSupport = dev.doubleFPConfig() > 0;
if (cn > 4 || (!doubleSupport && depth == CV_64F) ||
_src.offset() % esz != 0 || _src.step() % esz != 0)
return false;
Size ksize = _kernel.size();
if (anchor.x < 0)
anchor.x = ksize.width / 2;
if (anchor.y < 0)
anchor.y = ksize.height / 2;
Size size = _src.size(), wholeSize;
bool isolated = (borderType & BORDER_ISOLATED) != 0;
borderType &= ~BORDER_ISOLATED;
int wdepth = depth, wtype = type;
if (depth == CV_8U)
{
wdepth = CV_32S;
wtype = CV_MAKETYPE(wdepth, cn);
}
char cvt[2][40];
bool haveExtraMat = !_extraMat.empty();
CV_Assert(actual_op <= 3 || haveExtraMat);
const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE",
"BORDER_REFLECT", 0, "BORDER_REFLECT_101" };
size_t globalsize[2] = { size.width, size.height };
UMat src = _src.getUMat();
if (!isolated)
{
Point ofs;
src.locateROI(wholeSize, ofs);
}
int h = isolated ? size.height : wholeSize.height;
int w = isolated ? size.width : wholeSize.width;
if (w < ksize.width || h < ksize.height)
return false;
// Figure out what vector size to use for loading the pixels.
int pxLoadNumPixels = cn != 1 || size.width % 4 ? 1 : 4;
int pxLoadVecSize = cn * pxLoadNumPixels;
// Figure out how many pixels per work item to compute in X and Y
// directions. Too many and we run out of registers.
int pxPerWorkItemX = 1, pxPerWorkItemY = 1;
if (cn <= 2 && ksize.width <= 4 && ksize.height <= 4)
{
pxPerWorkItemX = size.width % 8 ? size.width % 4 ? size.width % 2 ? 1 : 2 : 4 : 8;
pxPerWorkItemY = size.height % 2 ? 1 : 2;
}
else if (cn < 4 || (ksize.width <= 4 && ksize.height <= 4))
{
pxPerWorkItemX = size.width % 2 ? 1 : 2;
pxPerWorkItemY = size.height % 2 ? 1 : 2;
}
globalsize[0] = size.width / pxPerWorkItemX;
globalsize[1] = size.height / pxPerWorkItemY;
// Need some padding in the private array for pixels
int privDataWidth = ROUNDUP(pxPerWorkItemX + ksize.width - 1, pxLoadNumPixels);
// Make the global size a nice round number so the runtime can pick
// from reasonable choices for the workgroup size
const int wgRound = 256;
globalsize[0] = ROUNDUP(globalsize[0], wgRound);
if (actual_op < 0)
actual_op = op;
// build processing
String processing;
Mat kernel8u;
_kernel.getMat().convertTo(kernel8u, CV_8U);
for (int y = 0; y < kernel8u.rows; ++y)
for (int x = 0; x < kernel8u.cols; ++x)
if (kernel8u.at<uchar>(y, x) != 0)
processing += format("PROCESS(%d,%d)", y, x);
static const char * const op2str[] = { "OP_ERODE", "OP_DILATE", NULL, NULL, "OP_GRADIENT", "OP_TOPHAT", "OP_BLACKHAT" };
String opts = format("-D cn=%d "
"-D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d "
"-D PX_LOAD_VEC_SIZE=%d -D PX_LOAD_NUM_PX=%d -D DEPTH_%d "
"-D PX_PER_WI_X=%d -D PX_PER_WI_Y=%d -D PRIV_DATA_WIDTH=%d -D %s -D %s "
"-D PX_LOAD_X_ITERATIONS=%d -D PX_LOAD_Y_ITERATIONS=%d "
"-D srcT=%s -D srcT1=%s -D dstT=srcT -D dstT1=srcT1 -D WT=%s -D WT1=%s "
"-D convertToWT=%s -D convertToDstT=%s -D PROCESS_ELEM_=%s -D %s%s",
cn, anchor.x, anchor.y, ksize.width, ksize.height,
pxLoadVecSize, pxLoadNumPixels, depth,
pxPerWorkItemX, pxPerWorkItemY, privDataWidth, borderMap[borderType],
isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED",
privDataWidth / pxLoadNumPixels, pxPerWorkItemY + ksize.height - 1,
ocl::typeToStr(type), ocl::typeToStr(depth),
haveExtraMat ? ocl::typeToStr(wtype):"srcT",//to prevent overflow - WT
haveExtraMat ? ocl::typeToStr(wdepth):"srcT1",//to prevent overflow - WT1
haveExtraMat ? ocl::convertTypeStr(depth, wdepth, cn, cvt[0]) : "noconvert",//to prevent overflow - src to WT
haveExtraMat ? ocl::convertTypeStr(wdepth, depth, cn, cvt[1]) : "noconvert",//to prevent overflow - WT to dst
processing.c_str(), op2str[op],
actual_op == op ? "" : cv::format(" -D %s", op2str[actual_op]).c_str());
ocl::Kernel kernel("filterSmall", cv::ocl::imgproc::filterSmall_oclsrc, opts);
if (kernel.empty())
return false;
_dst.create(size, type);
UMat dst = _dst.getUMat();
UMat source;
if(src.u != dst.u)
source = src;
else
{
Point ofs;
int cols = src.cols, rows = src.rows;
src.locateROI(wholeSize, ofs);
src.adjustROI(ofs.y, wholeSize.height - rows - ofs.y, ofs.x, wholeSize.width - cols - ofs.x);
src.copyTo(source);
src.adjustROI(-ofs.y, -wholeSize.height + rows + ofs.y, -ofs.x, -wholeSize.width + cols + ofs.x);
source.adjustROI(-ofs.y, -wholeSize.height + rows + ofs.y, -ofs.x, -wholeSize.width + cols + ofs.x);
source.locateROI(wholeSize, ofs);
}
UMat extraMat = _extraMat.getUMat();
int idxArg = kernel.set(0, ocl::KernelArg::PtrReadOnly(source));
idxArg = kernel.set(idxArg, (int)source.step);
int srcOffsetX = (int)((source.offset % source.step) / source.elemSize());
int srcOffsetY = (int)(source.offset / source.step);
int srcEndX = isolated ? srcOffsetX + size.width : wholeSize.width;
int srcEndY = isolated ? srcOffsetY + size.height : wholeSize.height;
idxArg = kernel.set(idxArg, srcOffsetX);
idxArg = kernel.set(idxArg, srcOffsetY);
idxArg = kernel.set(idxArg, srcEndX);
idxArg = kernel.set(idxArg, srcEndY);
idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst));
if (haveExtraMat)
{
idxArg = kernel.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(extraMat));
}
return kernel.run(2, globalsize, NULL, false);
}
static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel,
Point anchor, int iterations, int op, int borderType,
const Scalar &, int actual_op = -1, InputArray _extraMat = noArray())
{
const ocl::Device & dev = ocl::Device::getDefault();
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
bool doubleSupport = dev.doubleFPConfig() > 0;
int type = _src.type(), depth = CV_MAT_DEPTH(type),
cn = CV_MAT_CN(type), esz = CV_ELEM_SIZE(type);
Mat kernel = _kernel.getMat();
Size ksize = kernel.data ? kernel.size() : Size(3, 3), ssize = _src.size();
// try to use OpenCL kernel adopted for small morph kernel
if (dev.isIntel() && !(dev.type() & ocl::Device::TYPE_CPU) &&
((ksize.width < 5 && ksize.height < 5 && esz <= 4) ||
(ksize.width == 5 && ksize.height == 5 && cn == 1)) &&
(iterations == 1))
{
if (ocl_morphSmall(_src, _dst, _kernel, anchor, borderType, op, actual_op, _extraMat))
return true;
}
bool doubleSupport = dev.doubleFPConfig() > 0;
if ((depth == CV_64F && !doubleSupport) || borderType != BORDER_CONSTANT)
return false;
Mat kernel = _kernel.getMat();
bool haveExtraMat = !_extraMat.empty();
Size ksize = kernel.data ? kernel.size() : Size(3, 3), ssize = _src.size();
CV_Assert(actual_op <= 3 || haveExtraMat);
if (iterations == 0 || kernel.rows*kernel.cols == 1)

214
modules/imgproc/src/opencl/boxFilterSmall.cl → modules/imgproc/src/opencl/filterSmall.cl
View File

@@ -153,35 +153,10 @@ inline bool isBorder(const struct RectCoords bounds, int2 coord, int numPixels)
}
#endif
inline WT getBorderPixel(const struct RectCoords bounds, int2 coord,
__global const uchar * srcptr, int srcstep)
{
#ifdef BORDER_CONSTANT
return (WT)(0);
#else
int selected_col = coord.x;
int selected_row = coord.y;
EXTRAPOLATE(selected_col, selected_row,
bounds.x1, bounds.y1,
bounds.x2, bounds.y2);
__global const uchar* ptr = srcptr + mad24(selected_row, srcstep, selected_col * SRCSIZE);
return convertToWT(loadpix(ptr));
#endif
}
inline WT readSrcPixelSingle(int2 pos, __global const uchar * srcptr,
int srcstep, const struct RectCoords srcCoords)
{
if (!isBorder(srcCoords, pos, 1))
{
__global const uchar * ptr = srcptr + mad24(pos.y, srcstep, pos.x * SRCSIZE);
return convertToWT(loadpix(ptr));
}
else
return getBorderPixel(srcCoords, pos, srcptr, srcstep);
}
#define float1 float
#define uchar1 uchar
#define int1 int
#define uint1 unit
#define __CAT(x, y) x##y
#define CAT(x, y) __CAT(x, y)
@@ -191,7 +166,7 @@ inline WT readSrcPixelSingle(int2 pos, __global const uchar * srcptr,
#define PX_LOAD_FLOAT_VEC_TYPE CAT(WT1, PX_LOAD_VEC_SIZE)
#define PX_LOAD_FLOAT_VEC_CONV CAT(convert_, PX_LOAD_FLOAT_VEC_TYPE)
#define PX_LOAD CAT(vload, PX_LOAD_VEC_SIZE)
#define float1 float
inline PX_LOAD_FLOAT_VEC_TYPE readSrcPixelGroup(int2 pos, __global const uchar * srcptr,
int srcstep, const struct RectCoords srcCoords)
@@ -218,12 +193,150 @@ inline PX_LOAD_FLOAT_VEC_TYPE readSrcPixelGroup(int2 pos, __global const uchar *
#define LOOP(N, VAR, STMT) CAT(LOOP, N)((VAR), (STMT))
__kernel void boxFilterSmall(__global const uchar * srcptr, int src_step, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY,
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols
#ifdef NORMALIZE
, float alpha
#ifdef OP_BOX_FILTER
#define PROCESS_ELEM \
WT total_sum = (WT)(0); \
int sy = 0; \
LOOP(KERNEL_SIZE_Y, sy, \
{ \
int sx = 0; \
LOOP(KERNEL_SIZE_X, sx, \
{ \
total_sum += privateData[py + sy][px + sx]; \
}); \
})
#elif defined OP_FILTER2D
#define DIG(a) a,
__constant WT1 kernelData[] = { COEFF };
#define PROCESS_ELEM \
WT total_sum = 0; \
int sy = 0; \
int kernelIndex = 0; \
LOOP(KERNEL_SIZE_Y, sy, \
{ \
int sx = 0; \
LOOP(KERNEL_SIZE_X, sx, \
{ \
total_sum = fma(kernelData[kernelIndex++], privateData[py + sy][px + sx], total_sum); \
}); \
})
#elif defined OP_ERODE || defined OP_DILATE
#ifdef DEPTH_0
#define MIN_VAL 0
#define MAX_VAL UCHAR_MAX
#elif defined DEPTH_1
#define MIN_VAL SCHAR_MIN
#define MAX_VAL SCHAR_MAX
#elif defined DEPTH_2
#define MIN_VAL 0
#define MAX_VAL USHRT_MAX
#elif defined DEPTH_3
#define MIN_VAL SHRT_MIN
#define MAX_VAL SHRT_MAX
#elif defined DEPTH_4
#define MIN_VAL INT_MIN
#define MAX_VAL INT_MAX
#elif defined DEPTH_5
#define MIN_VAL (-FLT_MAX)
#define MAX_VAL FLT_MAX
#elif defined DEPTH_6
#define MIN_VAL (-DBL_MAX)
#define MAX_VAL DBL_MAX
#endif
)
#ifdef OP_ERODE
#define VAL (WT)MAX_VAL
#elif defined OP_DILATE
#define VAL (WT)MIN_VAL
#else
#error "Unknown operation"
#endif
#define convert_float1 convert_float
#define convert_uchar1 convert_uchar
#define convert_int1 convert_int
#define convert_uint1 convert_uint
#ifdef OP_ERODE
#if defined INTEL_DEVICE && defined DEPTH_0
// workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older)
#define WA_CONVERT_1 CAT(convert_uint, cn)
#define WA_CONVERT_2 CAT(convert_, srcT)
#define MORPH_OP(A, B) WA_CONVERT_2(min(WA_CONVERT_1(A), WA_CONVERT_1(B)))
#else
#define MORPH_OP(A, B) min((A), (B))
#endif
#endif
#ifdef OP_DILATE
#define MORPH_OP(A, B) max((A), (B))
#endif
#define PROCESS(_y, _x) \
total_sum = convertToWT(MORPH_OP(convertToWT(total_sum), convertToWT(privateData[py + _y][px + _x])));
#define PROCESS_ELEM \
WT total_sum = convertToWT(VAL); \
PROCESS_ELEM_
#else
#error "No processing is specified"
#endif
#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT
#define EXTRA_PARAMS , __global const uchar * matptr, int mat_step, int mat_offset
#else
#define EXTRA_PARAMS
#endif
inline WT getBorderPixel(const struct RectCoords bounds, int2 coord,
__global const uchar * srcptr, int srcstep)
{
#ifdef BORDER_CONSTANT
#ifdef OP_ERODE
return (WT)(MAX_VAL);
#elif defined OP_DILATE
return (WT)(MIN_VAL);
#else
return (WT)(0);
#endif
#else
int selected_col = coord.x;
int selected_row = coord.y;
EXTRAPOLATE(selected_col, selected_row,
bounds.x1, bounds.y1,
bounds.x2, bounds.y2);
__global const uchar* ptr = srcptr + mad24(selected_row, srcstep, selected_col * SRCSIZE);
return convertToWT(loadpix(ptr));
#endif
}
inline WT readSrcPixelSingle(int2 pos, __global const uchar * srcptr,
int srcstep, const struct RectCoords srcCoords)
{
if (!isBorder(srcCoords, pos, 1))
{
__global const uchar * ptr = srcptr + mad24(pos.y, srcstep, pos.x * SRCSIZE);
return convertToWT(loadpix(ptr));
}
else
return getBorderPixel(srcCoords, pos, srcptr, srcstep);
}
__kernel void filterSmall(__global const uchar * srcptr, int src_step, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY,
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols
#ifdef NORMALIZE
, float alpha
#endif
EXTRA_PARAMS )
{
// for non-isolated border: offsetX, offsetY, wholeX, wholeY
const struct RectCoords srcCoords = { srcOffsetX, srcOffsetY, srcEndX, srcEndY };
@@ -282,24 +395,27 @@ __kernel void boxFilterSmall(__global const uchar * srcptr, int src_step, int sr
LOOP(PX_PER_WI_X, px,
{
int x = startX + px;
int sy = 0;
int kernelIndex = 0;
WT total_sum = (WT)(0);
LOOP(KERNEL_SIZE_Y, sy,
{
int sx = 0;
LOOP(KERNEL_SIZE_X, sx,
{
total_sum += privateData[py + sy][px + sx];
});
});
__global dstT * dstPtr = (__global dstT *)(dstptr + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset)));
PROCESS_ELEM;
int dst_index = mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset));
__global dstT * dstPtr = (__global dstT *)(dstptr + dst_index);
#ifdef NORMALIZE
total_sum *= (WT)(alpha);
#endif
#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT
//for this type of operations SRCSIZE == DSTSIZE
int mat_index = mad24(y, mat_step, mad24(x, SRCSIZE, mat_offset));
WT value = convertToWT(loadpix(matptr + mat_index));
#ifdef OP_GRADIENT
storepix(convertToDstT(convertToWT(total_sum) - convertToWT(value)), dstPtr );
#elif defined OP_TOPHAT
storepix(convertToDstT(convertToWT(value) - convertToWT(total_sum)), dstPtr );
#elif defined OP_BLACKHAT
storepix(convertToDstT(convertToWT(total_sum) - convertToWT(value)), dstPtr );
#endif
#else // erode or dilate, or open-close
storepix(convertToDstT(total_sum), dstPtr);
#endif
});
});
}

4
modules/imgproc/src/smooth.cpp
View File

@@ -720,7 +720,7 @@ static bool ocl_boxFilter( InputArray _src, OutputArray _dst, int ddepth,
"-D PX_PER_WI_X=%d -D PX_PER_WI_Y=%d -D PRIV_DATA_WIDTH=%d -D %s -D %s "
"-D PX_LOAD_X_ITERATIONS=%d -D PX_LOAD_Y_ITERATIONS=%d "
"-D srcT=%s -D srcT1=%s -D dstT=%s -D dstT1=%s -D WT=%s -D WT1=%s "
"-D convertToWT=%s -D convertToDstT=%s%s%s",
"-D convertToWT=%s -D convertToDstT=%s%s%s -D OP_BOX_FILTER",
cn, anchor.x, anchor.y, ksize.width, ksize.height,
pxLoadVecSize, pxLoadNumPixels,
pxPerWorkItemX, pxPerWorkItemY, privDataWidth, borderMap[borderType],
@@ -734,7 +734,7 @@ static bool ocl_boxFilter( InputArray _src, OutputArray _dst, int ddepth,
if (!kernel.create("boxFilterSmall", cv::ocl::imgproc::boxFilterSmall_oclsrc, build_options))
if (!kernel.create("filterSmall", cv::ocl::imgproc::filterSmall_oclsrc, build_options))
return false;
}
else