mirror of
https://github.com/opencv/opencv.git
synced 2025-01-18 14:13:15 +08:00
added support of ksize >= 5 to cv::Laplacian
This commit is contained in:
Ilya Lavrenov
parent
d6b54ff9d1
commit
e1efed1914
@ -11,6 +11,7 @@
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2000, Intel Corporation, all rights reserved.
|
||||
// Copyright (C) 2014, Itseez, Inc, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
@ -40,6 +41,8 @@
|
||||
//M*/
|
||||
|
||||
#include "precomp.hpp"
|
||||
#include "opencl_kernels.hpp"
|
||||
|
||||
#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
|
||||
static IppStatus sts = ippInit();
|
||||
#endif
|
||||
@ -495,6 +498,58 @@ void cv::Scharr( InputArray _src, OutputArray _dst, int ddepth, int dx, int dy,
|
||||
sepFilter2D( _src, _dst, ddepth, kx, ky, Point(-1, -1), delta, borderType );
|
||||
}
|
||||
|
||||
#ifdef HAVE_OPENCL
|
||||
|
||||
namespace cv {
|
||||
|
||||
static bool ocl_Laplacian5(InputArray _src, OutputArray _dst,
|
||||
const Mat & kd, const Mat & ks, double scale, double delta,
|
||||
int borderType, int depth, int ddepth)
|
||||
{
|
||||
int iscale = cvRound(scale), idelta = cvRound(delta);
|
||||
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0,
|
||||
floatCoeff = std::fabs(delta - idelta) > DBL_EPSILON || std::fabs(scale - iscale) > DBL_EPSILON;
|
||||
int cn = _src.channels(), wdepth = std::max(depth, floatCoeff ? CV_32F : CV_32S), kercn = 1;
|
||||
|
||||
if (!doubleSupport && wdepth == CV_64F)
|
||||
return false;
|
||||
|
||||
char cvt[2][40];
|
||||
ocl::Kernel k("sumConvert", ocl::imgproc::laplacian5_oclsrc,
|
||||
format("-D srcT=%s -D WT=%s -D dstT=%s -D coeffT=%s -D wdepth=%d "
|
||||
"-D convertToWT=%s -D convertToDT=%s%s",
|
||||
ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)),
|
||||
ocl::typeToStr(CV_MAKE_TYPE(wdepth, kercn)),
|
||||
ocl::typeToStr(CV_MAKE_TYPE(ddepth, kercn)),
|
||||
ocl::typeToStr(wdepth), wdepth,
|
||||
ocl::convertTypeStr(depth, wdepth, kercn, cvt[0]),
|
||||
ocl::convertTypeStr(wdepth, ddepth, kercn, cvt[1]),
|
||||
doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
|
||||
if (k.empty())
|
||||
return false;
|
||||
|
||||
UMat d2x, d2y;
|
||||
sepFilter2D(_src, d2x, depth, kd, ks, Point(-1, -1), 0, borderType);
|
||||
sepFilter2D(_src, d2y, depth, ks, kd, Point(-1, -1), 0, borderType);
|
||||
|
||||
UMat dst = _dst.getUMat();
|
||||
|
||||
ocl::KernelArg d2xarg = ocl::KernelArg::ReadOnlyNoSize(d2x),
|
||||
d2yarg = ocl::KernelArg::ReadOnlyNoSize(d2y),
|
||||
dstarg = ocl::KernelArg::WriteOnly(dst, cn, kercn);
|
||||
|
||||
if (wdepth >= CV_32F)
|
||||
k.args(d2xarg, d2yarg, dstarg, (float)scale, (float)delta);
|
||||
else
|
||||
k.args(d2xarg, d2yarg, dstarg, iscale, idelta);
|
||||
|
||||
size_t globalsize[] = { dst.cols * cn / kercn, dst.rows };
|
||||
return k.run(2, globalsize, NULL, false);
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
void cv::Laplacian( InputArray _src, OutputArray _dst, int ddepth, int ksize,
|
||||
double scale, double delta, int borderType )
|
||||
@ -531,27 +586,28 @@ void cv::Laplacian( InputArray _src, OutputArray _dst, int ddepth, int ksize,
|
||||
}
|
||||
else
|
||||
{
|
||||
Mat src = _src.getMat(), dst = _dst.getMat();
|
||||
const size_t STRIPE_SIZE = 1 << 14;
|
||||
|
||||
int depth = src.depth();
|
||||
int ktype = std::max(CV_32F, std::max(ddepth, depth));
|
||||
int wdepth = depth == CV_8U && ksize <= 5 ? CV_16S : depth <= CV_32F ? CV_32F : CV_64F;
|
||||
int wtype = CV_MAKETYPE(wdepth, src.channels());
|
||||
int ktype = std::max(CV_32F, std::max(ddepth, sdepth));
|
||||
int wdepth = sdepth == CV_8U && ksize <= 5 ? CV_16S : sdepth <= CV_32F ? CV_32F : CV_64F;
|
||||
int wtype = CV_MAKETYPE(wdepth, cn);
|
||||
Mat kd, ks;
|
||||
getSobelKernels( kd, ks, 2, 0, ksize, false, ktype );
|
||||
int dtype = CV_MAKETYPE(ddepth, src.channels());
|
||||
|
||||
int dy0 = std::min(std::max((int)(STRIPE_SIZE/(getElemSize(src.type())*src.cols)), 1), src.rows);
|
||||
Ptr<FilterEngine> fx = createSeparableLinearFilter(src.type(),
|
||||
CV_OCL_RUN(_dst.isUMat(),
|
||||
ocl_Laplacian5(_src, _dst, kd, ks, scale,
|
||||
delta, borderType, wdepth, ddepth))
|
||||
|
||||
const size_t STRIPE_SIZE = 1 << 14;
|
||||
Ptr<FilterEngine> fx = createSeparableLinearFilter(stype,
|
||||
wtype, kd, ks, Point(-1,-1), 0, borderType, borderType, Scalar() );
|
||||
Ptr<FilterEngine> fy = createSeparableLinearFilter(src.type(),
|
||||
Ptr<FilterEngine> fy = createSeparableLinearFilter(stype,
|
||||
wtype, ks, kd, Point(-1,-1), 0, borderType, borderType, Scalar() );
|
||||
|
||||
Mat src = _src.getMat(), dst = _dst.getMat();
|
||||
int y = fx->start(src), dsty = 0, dy = 0;
|
||||
fy->start(src);
|
||||
const uchar* sptr = src.data + y*src.step;
|
||||
|
||||
int dy0 = std::min(std::max((int)(STRIPE_SIZE/(CV_ELEM_SIZE(stype)*src.cols)), 1), src.rows);
|
||||
Mat d2x( dy0 + kd.rows - 1, src.cols, wtype );
|
||||
Mat d2y( dy0 + kd.rows - 1, src.cols, wtype );
|
||||
|
||||
@ -564,7 +620,7 @@ void cv::Laplacian( InputArray _src, OutputArray _dst, int ddepth, int ksize,
|
||||
Mat dstripe = dst.rowRange(dsty, dsty + dy);
|
||||
d2x.rows = d2y.rows = dy; // modify the headers, which should work
|
||||
d2x += d2y;
|
||||
d2x.convertTo( dstripe, dtype, scale, delta );
|
||||
d2x.convertTo( dstripe, ddepth, scale, delta );
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
34
modules/imgproc/src/opencl/laplacian5.cl
Normal file
34
modules/imgproc/src/opencl/laplacian5.cl
Normal file
@ -0,0 +1,34 @@
|
||||
// This file is part of OpenCV project.
|
||||
// It is subject to the license terms in the LICENSE file found in the top-level directory
|
||||
// of this distribution and at http://opencv.org/license.html.
|
||||
|
||||
// Copyright (C) 2014, Itseez, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
|
||||
#define noconvert
|
||||
|
||||
__kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1_offset,
|
||||
__global const uchar * src2ptr, int src2_step, int src2_offset,
|
||||
__global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
|
||||
coeffT scale, coeffT delta)
|
||||
{
|
||||
int x = get_global_id(0);
|
||||
int y = get_global_id(1);
|
||||
|
||||
if (y < dst_rows && x < dst_cols)
|
||||
{
|
||||
int src1_index = mad24(y, src1_step, mad24(x, (int)sizeof(srcT), src1_offset));
|
||||
int src2_index = mad24(y, src2_step, mad24(x, (int)sizeof(srcT), src2_offset));
|
||||
int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset));
|
||||
|
||||
__global const srcT * src1 = (__global const srcT *)(src1ptr + src1_index);
|
||||
__global const srcT * src2 = (__global const srcT *)(src2ptr + src2_index);
|
||||
__global dstT * dst = (__global dstT *)(dstptr + dst_index);
|
||||
|
||||
#if wdepth <= 4
|
||||
dst[0] = convertToDT( mad24((WT)(scale), convertToWT(src1[0]) + convertToWT(src2[0]), (WT)(delta)) );
|
||||
#else
|
||||
dst[0] = convertToDT( mad((WT)(scale), convertToWT(src1[0]) + convertToWT(src2[0]), (WT)(delta)) );
|
||||
#endif
|
||||
}
|
||||
}
|
||||
@ -316,7 +316,7 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine(
|
||||
|
||||
OCL_INSTANTIATE_TEST_CASE_P(Filter, LaplacianTest, Combine(
|
||||
FILTER_TYPES,
|
||||
Values(1, 3), // kernel size
|
||||
Values(1, 3, 5), // kernel size
|
||||
Values(Size(0, 0)), // not used
|
||||
FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
|
||||
Values(1.0, 0.2, 3.0), // kernel scale
|
||||
|
||||
Loading…
Reference in New Issue
Block a user