return k.run(2, globalsize, localsize, false);
}
-static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, int anchor)
+static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, double delta, int anchor)
{
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
if (dst.depth() == CV_64F && !doubleSupport)
if (k.empty())
return false;
- k.args(ocl::KernelArg::ReadOnly(buf), ocl::KernelArg::WriteOnly(dst));
+ k.args(ocl::KernelArg::ReadOnly(buf), ocl::KernelArg::WriteOnly(dst),
+ static_cast<float>(delta));
return k.run(2, globalsize, localsize, false);
}
static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
Mat row_kernel, Mat col_kernel,
- int borderType, int ddepth)
+ double delta, int borderType, int ddepth)
{
Size size = _src.size(), wholeSize;
Point origin;
src.locateROI(wholeSize, origin);
k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y,
- wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst));
+ wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst),
+ static_cast<float>(delta));
return k.run(2, gt2, lt2, false);
}
const ocl::Device & d = ocl::Device::getDefault();
Size imgSize = _src.size();
- if (abs(delta)> FLT_MIN)
- return false;
-
int type = _src.type(), sdepth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
if (cn > 4)
return false;
imgSize.height > optimizedSepFilterLocalSize + (kernelY.cols >> 1) &&
(!(borderType & BORDER_ISOLATED) || _src.offset() == 0) && anchor == Point(-1, -1) &&
(d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())),
- ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, borderType, ddepth), true)
+ ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, delta,
+ borderType & ~BORDER_ISOLATED, ddepth), true)
if (anchor.x < 0)
anchor.x = kernelX.cols >> 1;
_dst.create(srcSize, CV_MAKETYPE(ddepth, cn));
UMat dst = _dst.getUMat();
- return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y);
+ return ocl_sepColFilter2D(buf, dst, kernelY, delta, anchor.y);
}
#endif
__constant float mat_kernel[] = { COEFF };
__kernel void col_filter(__global const uchar * src, int src_step, int src_offset, int src_whole_rows, int src_whole_cols,
- __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, float delta)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < dst_cols && y < dst_rows)
{
start_addr = mad24(y, dst_step, mad24(DSTSIZE, x, dst_offset));
- storepix(convertToDstT(sum), dst + start_addr);
+ storepix(convertToDstT(sum + (srcT)(delta)), dst + start_addr);
}
}
__constant float mat_kernelY[] = { KERNEL_MATRIX_Y };
__kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width,
- __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+ __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, float delta)
{
// RADIUSX, RADIUSY are filter dimensions
// BLK_X, BLK_Y are local wrogroup sizes
for (i=0; i<=2*RADIUSX; i++)
sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum);
- //store result into destination image
- storepix(convertToDstT(sum), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset)));
+ // store result into destination image
+ storepix(convertToDstT(sum + (WT)(delta)), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset)));
}
int borderType;
bool useRoi;
Mat kernelX, kernelY;
+ double delta;
TEST_DECLARE_INPUT_PARAMETER(src);
TEST_DECLARE_OUTPUT_PARAMETER(dst);
randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
anchor.x = anchor.y = -1;
+ delta = randomDouble(-100, 100);
UMAT_UPLOAD_INPUT_PARAMETER(src);
UMAT_UPLOAD_OUTPUT_PARAMETER(dst);
{
random_roi();
- OCL_OFF(cv::sepFilter2D(src_roi, dst_roi, -1, kernelX, kernelY, anchor, 0.0, borderType));
- OCL_ON(cv::sepFilter2D(usrc_roi, udst_roi, -1, kernelX, kernelY, anchor, 0.0, borderType));
+ OCL_OFF(cv::sepFilter2D(src_roi, dst_roi, -1, kernelX, kernelY, anchor, delta, borderType));
+ OCL_ON(cv::sepFilter2D(usrc_roi, udst_roi, -1, kernelX, kernelY, anchor, delta, borderType));
Near(1.0);
}