// 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,
//M*/
#include "precomp.hpp"
+#include "opencl_kernels.hpp"
+
#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
static IppStatus sts = ippInit();
#endif
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 )
}
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 );
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 );
}
}
}
--- /dev/null
+// 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
+ }
+}