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

improved cv::filter2D

Browse Source
This commit is contained in:
Ilya Lavrenov 2014-03-24 23:07:00 +04:00
parent d63a8ba018
commit f7d6d3cff5
3 changed files with 22 additions and 20 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

25
modules/core/src/system.cpp
View File

@ -414,24 +414,23 @@ const String& getBuildInformation()
String format( const char* fmt, ... )
{
char buf[1024];
AutoBuffer<char, 1024> buf;
va_list va;
va_start(va, fmt);
int len = vsnprintf(buf, sizeof(buf), fmt, va);
va_end(va);
if (len >= (int)sizeof(buf))
for ( ; ; )
{
String s(len, '\0');
va_list va;
va_start(va, fmt);
len = vsnprintf((char*)s.c_str(), len + 1, fmt, va);
(void)len;
int bsize = static_cast<int>(buf.size()),
len = vsnprintf((char *)buf, bsize, fmt, va);
va_end(va);
return s;
}
return String(buf, len);
if (len < 0 || len >= bsize)
{
buf.resize(std::max(bsize << 1, len + 1));
continue;
}
return String((char *)buf, len);
}
}
String tempfile( const char* suffix )

12
modules/imgproc/src/filter.cpp
View File

@ -42,7 +42,6 @@
#include "precomp.hpp"
#include "opencl_kernels.hpp"
#include <sstream>
/****************************************************************************************\
Base Image Filter
@ -3197,6 +3196,8 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
size_t tryWorkItems = maxWorkItemSizes[0];
char cvt[2][40];
String kerStr = ocl::kernelToStr(kernelMatDataFloat, CV_32F);
for ( ; ; )
{
size_t BLOCK_SIZE = tryWorkItems;
@ -3226,14 +3227,14 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
String opts = format("-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D cn=%d "
"-D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d "
"-D KERNEL_SIZE_Y2_ALIGNED=%d -D %s -D %s -D %s%s "
"-D KERNEL_SIZE_Y2_ALIGNED=%d -D %s -D %s -D %s%s%s "
"-D srcT=%s -D srcT1=%s -D dstT=%s -D dstT1=%s -D WT=%s -D WT1=%s "
"-D convertToWT=%s -D convertToDstT=%s",
(int)BLOCK_SIZE, (int)BLOCK_SIZE_Y, cn, anchor.x, anchor.y,
ksize.width, ksize.height, kernel_size_y2_aligned, borderMap[borderType],
extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED",
doubleSupport ? " -D DOUBLE_SUPPORT" : "",
doubleSupport ? " -D DOUBLE_SUPPORT" : "", kerStr.c_str(),
ocl::typeToStr(type), ocl::typeToStr(sdepth), ocl::typeToStr(dtype),
ocl::typeToStr(ddepth), ocl::typeToStr(wtype), ocl::typeToStr(wdepth),
ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
@ -3255,7 +3256,7 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
}
_dst.create(sz, dtype);
UMat dst = _dst.getUMat(), kernalDataUMat(kernelMatDataFloat, true);
UMat dst = _dst.getUMat();
int srcOffsetX = (int)((src.offset % src.step) / src.elemSize());
int srcOffsetY = (int)(src.offset / src.step);
@ -3263,8 +3264,7 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
int srcEndY = (isolated ? (srcOffsetY + sz.height) : wholeSize.height);
k.args(ocl::KernelArg::PtrReadOnly(src), (int)src.step, srcOffsetX, srcOffsetY,
srcEndX, srcEndY, ocl::KernelArg::WriteOnly(dst),
ocl::KernelArg::PtrReadOnly(kernalDataUMat), (float)delta);
srcEndX, srcEndY, ocl::KernelArg::WriteOnly(dst), (float)delta);
return k.run(2, globalsize, localsize, false);
}

5
modules/imgproc/src/opencl/filter2D.cl
View File

@ -200,8 +200,11 @@ inline WT readSrcPixel(int2 pos, __global const uchar * srcptr, int src_step, co
}
}
#define DIG(a) a,
__constant WT1 kernelData[] = { COEFF };
__kernel void filter2D(__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, __constant WT1 * kernelData, float delta)
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols, float delta)
{
const struct RectCoords srcCoords = { srcOffsetX, srcOffsetY, srcEndX, srcEndY }; // for non-isolated border: offsetX, offsetY, wholeX, wholeY