From cb6ea8bfa1dffe39afe820082727c5b2591f3890 Mon Sep 17 00:00:00 2001
From: Alexander Alekhin <alexander.alekhin@itseez.com>
Date: Sat, 26 Oct 2013 11:37:02 +0400
Subject: [PATCH 1/6] ocl: update filter tests
---
modules/ocl/test/test_filters.cpp | 235 ++++++++++++++++--------------
modules/ocl/test/utility.hpp | 14 ++
2 files changed, 140 insertions(+), 109 deletions(-)
diff --git a/modules/ocl/test/test_filters.cpp b/modules/ocl/test/test_filters.cpp
index 86ff834d3..3cf7d37b8 100644
--- a/modules/ocl/test/test_filters.cpp
+++ b/modules/ocl/test/test_filters.cpp
@@ -59,10 +59,15 @@ using namespace cv;
PARAM_TEST_CASE(FilterTestBase, MatType,
int, // kernel size
Size, // dx, dy
- int, // border type, or iteration
+ int, // border type
+ double, // optional parameter
bool) // roi or not
{
+ bool isFP;
+
int type, borderType, ksize;
+ Size size;
+ double param;
bool useRoi;
Mat src, dst_whole, src_roi, dst_roi;
@@ -72,31 +77,53 @@ PARAM_TEST_CASE(FilterTestBase, MatType,
{
type = GET_PARAM(0);
ksize = GET_PARAM(1);
+ size = GET_PARAM(2);
borderType = GET_PARAM(3);
- useRoi = GET_PARAM(4);
+ param = GET_PARAM(4);
+ useRoi = GET_PARAM(5);
+
+ isFP = (CV_MAT_DEPTH(type) == CV_32F || CV_MAT_DEPTH(type) == CV_64F);
}
- void random_roi()
+ void random_roi(int minSize = 1)
{
- Size roiSize = randomSize(1, MAX_VALUE);
+ if (minSize == 0)
+ minSize = ksize;
+ Size roiSize = randomSize(minSize, MAX_VALUE);
Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
- randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256);
+ randomSubMat(src, src_roi, roiSize, srcBorder, type, isFP ? 0 : 5, isFP ? 1 : 256);
Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
- randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, 5, 16);
+ randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, isFP ? 0.20 : 60, isFP ? 0.25 : 70);
generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder);
generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
}
- void Near(double threshold = 0.0)
+ void Near()
+ {
+ if (isFP)
+ Near(1e-6, true);
+ else
+ Near(1, false);
+ }
+
+ void Near(double threshold, bool relative)
{
Mat roi, whole;
gdst_whole.download(whole);
gdst_roi.download(roi);
- EXPECT_MAT_NEAR(dst_whole, whole, threshold);
- EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+ if (relative)
+ {
+ EXPECT_MAT_NEAR_RELATIVE(dst_whole, whole, threshold);
+ EXPECT_MAT_NEAR_RELATIVE(dst_roi, roi, threshold);
+ }
+ else
+ {
+ EXPECT_MAT_NEAR(dst_whole, whole, threshold);
+ EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+ }
}
};
@@ -111,12 +138,12 @@ OCL_TEST_P(Blur, Mat)
for (int j = 0; j < LOOP_TIMES; j++)
{
- random_roi();
+ random_roi(0); // TODO NOTE: min value for size is kernel size (temporary bypass border issues in CPU implementation)
blur(src_roi, dst_roi, kernelSize, Point(-1, -1), borderType);
ocl::blur(gsrc_roi, gdst_roi, kernelSize, Point(-1, -1), borderType); // TODO anchor
- Near(1.0);
+ Near();
}
}
@@ -127,64 +154,51 @@ typedef FilterTestBase LaplacianTest;
OCL_TEST_P(LaplacianTest, Accuracy)
{
+ double scale = param;
+
for (int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
- // border type is used as a scale factor for the Laplacian kernel
- double scale = static_cast<double>(borderType);
+ Laplacian(src_roi, dst_roi, -1, ksize, scale); // TODO FIXIT , 0, borderType);
+ ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale); // TODO FIXIT , 0, borderType);
- Laplacian(src_roi, dst_roi, -1, ksize, scale);
- ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale);
-
- Near(1e-5);
+ Near();
}
}
/////////////////////////////////////////////////////////////////////////////////////////////////
// erode & dilate
-struct ErodeDilate :
- public FilterTestBase
-{
- int iterations;
-
- virtual void SetUp()
- {
- type = GET_PARAM(0);
- ksize = GET_PARAM(1);
- iterations = GET_PARAM(3);
- useRoi = GET_PARAM(4);
- }
-};
-
-typedef ErodeDilate Erode;
+typedef FilterTestBase Erode;
OCL_TEST_P(Erode, Mat)
{
// erode or dilate kernel
Size kernelSize(ksize, ksize);
Mat kernel;
+ int iterations = (int)param;
for (int j = 0; j < LOOP_TIMES; j++)
{
- kernel = randomMat(kernelSize, CV_8UC1, 0, 3);
-
random_roi();
- cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations);
- ocl::erode(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations); // TODO iterations, borderType
+ kernel = randomMat(kernelSize, CV_8UC1, 0, 3);
- Near(1e-5);
+ cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations);//, borderType);
+ ocl::erode(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations);//, borderType);
+
+ Near();
}
}
-typedef ErodeDilate Dilate;
+typedef FilterTestBase Dilate;
OCL_TEST_P(Dilate, Mat)
{
// erode or dilate kernel
Mat kernel;
+ int iterations = (int)param;
for (int j = 0; j < LOOP_TIMES; j++)
{
@@ -195,79 +209,56 @@ OCL_TEST_P(Dilate, Mat)
cv::dilate(src_roi, dst_roi, kernel, Point(-1, -1), iterations);
ocl::dilate(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations); // TODO iterations, borderType
- Near(1e-5);
+ Near();
}
}
/////////////////////////////////////////////////////////////////////////////////////////////////
// Sobel
-struct SobelTest :
- public FilterTestBase
-{
- int dx, dy;
-
- virtual void SetUp()
- {
- type = GET_PARAM(0);
- ksize = GET_PARAM(1);
- borderType = GET_PARAM(3);
- useRoi = GET_PARAM(4);
-
- Size d = GET_PARAM(2);
- dx = d.width, dy = d.height;
- }
-};
+typedef FilterTestBase SobelTest;
OCL_TEST_P(SobelTest, Mat)
{
+ int dx = size.width, dy = size.height;
+ double scale = param;
+
for (int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
- Sobel(src_roi, dst_roi, -1, dx, dy, ksize, /* scale */ 0.00001, /* delta */0, borderType);
- ocl::Sobel(gsrc_roi, gdst_roi, -1, dx, dy, ksize, /* scale */ 0.00001, /* delta */ 0, borderType);
+ Sobel(src_roi, dst_roi, -1, dx, dy, ksize, scale, /* delta */0, borderType);
+ ocl::Sobel(gsrc_roi, gdst_roi, -1, dx, dy, ksize, scale, /* delta */0, borderType);
- Near(1);
+ Near();
}
}
/////////////////////////////////////////////////////////////////////////////////////////////////
// Scharr
-typedef SobelTest ScharrTest;
+typedef FilterTestBase ScharrTest;
OCL_TEST_P(ScharrTest, Mat)
{
+ int dx = size.width, dy = size.height;
+ double scale = param;
+
for (int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
- Scharr(src_roi, dst_roi, -1, dx, dy, /* scale */ 1, /* delta */ 0, borderType);
- ocl::Scharr(gsrc_roi, gdst_roi, -1, dx, dy, /* scale */ 1, /* delta */ 0, borderType);
+ Scharr(src_roi, dst_roi, -1, dx, dy, scale, /* delta */ 0, borderType);
+ ocl::Scharr(gsrc_roi, gdst_roi, -1, dx, dy, scale, /* delta */ 0, borderType);
- Near(1);
+ Near();
}
}
/////////////////////////////////////////////////////////////////////////////////////////////////
// GaussianBlur
-struct GaussianBlurTest :
- public FilterTestBase
-{
- double sigma1, sigma2;
-
- virtual void SetUp()
- {
- type = GET_PARAM(0);
- ksize = GET_PARAM(1);
- borderType = GET_PARAM(3);
-
- sigma1 = rng.uniform(0.1, 1.0);
- sigma2 = rng.uniform(0.1, 1.0);
- }
-};
+typedef FilterTestBase GaussianBlurTest;
OCL_TEST_P(GaussianBlurTest, Mat)
{
@@ -275,10 +266,13 @@ OCL_TEST_P(GaussianBlurTest, Mat)
{
random_roi();
+ double sigma1 = rng.uniform(0.1, 1.0);
+ double sigma2 = rng.uniform(0.1, 1.0);
+
GaussianBlur(src_roi, dst_roi, Size(ksize, ksize), sigma1, sigma2, borderType);
ocl::GaussianBlur(gsrc_roi, gdst_roi, Size(ksize, ksize), sigma1, sigma2, borderType);
- Near(1);
+ Near();
}
}
@@ -289,19 +283,24 @@ typedef FilterTestBase Filter2D;
OCL_TEST_P(Filter2D, Mat)
{
- const Size kernelSize(ksize, ksize);
- Mat kernel;
-
for (int j = 0; j < LOOP_TIMES; j++)
{
- kernel = randomMat(kernelSize, CV_32FC1, 0.0, 1.0);
-
random_roi();
- cv::filter2D(src_roi, dst_roi, -1, kernel, Point(-1, -1), 0.0, borderType); // TODO anchor
- ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, Point(-1, -1), borderType);
+ Point anchor(-1, -1);
+ if (size.width >= 0)
+ anchor.x = size.width % ksize;
+ if (size.height >= 0)
+ anchor.y = size.height % ksize;
- Near(1);
+ const Size kernelSize(ksize, ksize);
+ Mat kernel = randomMat(kernelSize, CV_32FC1, 0, 1.0);
+ kernel *= 1.0 / (double)(ksize * ksize);
+
+ cv::filter2D(src_roi, dst_roi, -1, kernel, anchor, 0.0, borderType);
+ ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, anchor, /* TODO FIXIT 0.0,*/ borderType);
+
+ Near();
}
}
@@ -322,7 +321,7 @@ OCL_TEST_P(Bilateral, Mat)
cv::bilateralFilter(src_roi, dst_roi, ksize, sigmacolor, sigmaspace, borderType);
ocl::bilateralFilter(gsrc_roi, gdst_roi, ksize, sigmacolor, sigmaspace, borderType);
- Near(1);
+ Near();
}
}
@@ -342,7 +341,7 @@ OCL_TEST_P(AdaptiveBilateral, Mat)
adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, Point(-1, -1), borderType); // TODO anchor
ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, Point(-1, -1), borderType);
- Near(1);
+ Near();
}
}
@@ -366,80 +365,97 @@ OCL_TEST_P(MedianFilter, Mat)
//////////////////////////////////////////////////////////////////////////////////////////////////////////////
+#define FILTER_BORDER_SET_NO_ISOLATED \
+ Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_WRAP, (int)BORDER_REFLECT_101/*, \
+ (int)BORDER_CONSTANT|BORDER_ISOLATED, (int)BORDER_REPLICATE|BORDER_ISOLATED, \
+ (int)BORDER_REFLECT|BORDER_ISOLATED, (int)BORDER_WRAP|BORDER_ISOLATED, \
+ (int)BORDER_REFLECT_101|BORDER_ISOLATED*/) // WRAP and ISOLATED are not supported by cv:: version
+
+#define FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED \
+ Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, /*(int)BORDER_WRAP,*/ (int)BORDER_REFLECT_101/*, \
+ (int)BORDER_CONSTANT|BORDER_ISOLATED, (int)BORDER_REPLICATE|BORDER_ISOLATED, \
+ (int)BORDER_REFLECT|BORDER_ISOLATED, (int)BORDER_WRAP|BORDER_ISOLATED, \
+ (int)BORDER_REFLECT_101|BORDER_ISOLATED*/) // WRAP and ISOLATED are not supported by cv:: version
+
+
INSTANTIATE_TEST_CASE_P(Filter, Blur, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
Values(3, 5, 7),
Values(Size(0, 0)), // not used
- Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT_101),
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(0.0), // not used
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, LaplacianTest, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
Values(1, 3),
Values(Size(0, 0)), // not used
- Values(1, 2), // value is used as scale factor for kernel
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(1.0, 0.2, 3.0), // scalar
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
Values(3, 5, 7),
Values(Size(0, 0)), // not used
- testing::Range(1, 4),
+ Values(0), // not used
+ Values(1.0, 2.0, 3.0),
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
Values(3, 5, 7),
Values(Size(0, 0)), // not used
- testing::Range(1, 4),
+ Values(0), // not used
+ Values(1.0, 2.0, 3.0),
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
Values(3, 5),
- Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)),
- Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
- (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+ Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)), // dx, dy
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(0.0), // not used
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
- Values(0), // not used
- Values(Size(0, 1), Size(1, 0)),
- Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
- (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+ Values(1),
+ Values(Size(0, 1), Size(1, 0)), // dx, dy
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(1.0, 0.2), // scalar
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
Values(3, 5),
Values(Size(0, 0)), // not used
- Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
- (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(0.0), // not used
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, Filter2D, testing::Combine(
Values(CV_8UC1, CV_32FC1, CV_32FC4),
- Values(3, 15, 25),
- Values(Size(0, 0)), // not used
- Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
- (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+ Values(3, 15), // TODO 25: CPU implementation has some issues
+ Values(Size(-1, -1), Size(0, 0), Size(2, 1)), // anchor
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(0.0), // not used
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine(
Values(CV_8UC1, CV_8UC3),
Values(5, 9),
Values(Size(0, 0)), // not used
- Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE,
- (int)BORDER_REFLECT, (int)BORDER_WRAP, (int)BORDER_REFLECT_101),
+ FILTER_BORDER_SET_NO_ISOLATED,
+ Values(0.0), // not used
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, AdaptiveBilateral, Combine(
Values(CV_8UC1, CV_8UC3),
Values(5, 9),
Values(Size(0, 0)), // not used
- Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE,
- (int)BORDER_REFLECT, (int)BORDER_REFLECT_101),
+ FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED,
+ Values(0.0), // not used
Bool()));
INSTANTIATE_TEST_CASE_P(Filter, MedianFilter, Combine(
@@ -447,6 +463,7 @@ INSTANTIATE_TEST_CASE_P(Filter, MedianFilter, Combine(
Values(3, 5),
Values(Size(0, 0)), // not used
Values(0), // not used
+ Values(0.0), // not used
Bool()));
#endif // HAVE_OPENCL
diff --git a/modules/ocl/test/utility.hpp b/modules/ocl/test/utility.hpp
index 5ad97b08a..5eeb075c5 100644
--- a/modules/ocl/test/utility.hpp
+++ b/modules/ocl/test/utility.hpp
@@ -72,6 +72,13 @@ double checkNorm(const cv::Mat &m);
double checkNorm(const cv::Mat &m1, const cv::Mat &m2);
double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2);
+inline double checkNormRelative(const Mat &m1, const Mat &m2)
+{
+ return cv::norm(m1, m2, cv::NORM_INF) /
+ std::max((double)std::numeric_limits<float>::epsilon(),
+ (double)std::max(cv::norm(m1, cv::NORM_INF), norm(m2, cv::NORM_INF)));
+}
+
#define EXPECT_MAT_NORM(mat, eps) \
{ \
EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \
@@ -84,6 +91,13 @@ double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2);
EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \
}
+#define EXPECT_MAT_NEAR_RELATIVE(mat1, mat2, eps) \
+{ \
+ ASSERT_EQ(mat1.type(), mat2.type()); \
+ ASSERT_EQ(mat1.size(), mat2.size()); \
+ EXPECT_LE(checkNormRelative(cv::Mat(mat1), cv::Mat(mat2)), eps); \
+}
+
#define EXPECT_MAT_SIMILAR(mat1, mat2, eps) \
{ \
ASSERT_EQ(mat1.type(), mat2.type()); \
From 0bf9ece998b62f6265a789c260bae5ad146e2143 Mon Sep 17 00:00:00 2001
From: Alexander Alekhin <alexander.alekhin@itseez.com>
Date: Sat, 26 Oct 2013 23:31:51 +0400
Subject: [PATCH 2/6] ocl: rewrite boxFilter
---
modules/ocl/include/opencv2/ocl/ocl.hpp | 6 +-
modules/ocl/src/filtering.cpp | 334 +++-------
modules/ocl/src/opencl/filtering_boxFilter.cl | 628 ++++++++----------
3 files changed, 354 insertions(+), 614 deletions(-)
diff --git a/modules/ocl/include/opencv2/ocl/ocl.hpp b/modules/ocl/include/opencv2/ocl/ocl.hpp
index 5ccab64cb..05bd061ca 100644
--- a/modules/ocl/include/opencv2/ocl/ocl.hpp
+++ b/modules/ocl/include/opencv2/ocl/ocl.hpp
@@ -722,7 +722,7 @@ namespace cv
CV_EXPORTS void Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1);
//! returns 2D box filter
- // supports CV_8UC1 and CV_8UC4 source type, dst type must be the same as source type
+ // dst type must be the same as source type
CV_EXPORTS Ptr<BaseFilter_GPU> getBoxFilter_GPU(int srcType, int dstType,
const Size &ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
@@ -740,8 +740,6 @@ namespace cv
const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
//! smooths the image using the normalized box filter
- // supports data type: CV_8UC1, CV_8UC4, CV_32FC1 and CV_32FC4
- // supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP
CV_EXPORTS void boxFilter(const oclMat &src, oclMat &dst, int ddepth, Size ksize,
Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
@@ -757,8 +755,6 @@ namespace cv
const Point &anchor = Point(-1, -1), int iterations = 1);
//! a synonym for normalized box filter
- // supports data type: CV_8UC1, CV_8UC4, CV_32FC1 and CV_32FC4
- // supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101
static inline void blur(const oclMat &src, oclMat &dst, Size ksize, Point anchor = Point(-1, -1),
int borderType = BORDER_CONSTANT)
{
diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp
index d7502496f..fdddc1674 100644
--- a/modules/ocl/src/filtering.cpp
+++ b/modules/ocl/src/filtering.cpp
@@ -11,7 +11,7 @@
// For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
-// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
@@ -713,276 +713,126 @@ Ptr<FilterEngine_GPU> cv::ocl::createSeparableFilter_GPU(const Ptr<BaseRowFilter
return Ptr<FilterEngine_GPU>(new SeparableFilterEngine_GPU(rowFilter, columnFilter));
}
-/*
-**data type supported: CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4
-**support four border types: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_REFLECT_101
-*/
-
-static void GPUFilterBox_8u_C1R(const oclMat &src, oclMat &dst,
+static void GPUFilterBox(const oclMat &src, oclMat &dst,
Size &ksize, const Point anchor, const int borderType)
{
//Normalize the result by default
- float alpha = ksize.height * ksize.width;
+ float alpha = 1.0f / (ksize.height * ksize.width);
CV_Assert(src.clCxt == dst.clCxt);
CV_Assert((src.cols == dst.cols) &&
(src.rows == dst.rows));
- Context *clCxt = src.clCxt;
+ CV_Assert(src.oclchannels() == dst.oclchannels());
- string kernelName = "boxFilter_C1_D0";
+ size_t BLOCK_SIZE = src.clCxt->getDeviceInfo().maxWorkItemSizes[0];
+ size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices
+ while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * src.clCxt->getDeviceInfo().maxComputeUnits * 32 < (size_t)src.rows)
+ BLOCK_SIZE_Y *= 2;
- char btype[30];
+ CV_Assert((size_t)ksize.width <= BLOCK_SIZE);
- switch (borderType)
+ bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
+
+ vector<pair<size_t , const void *> > args;
+
+ args.push_back( make_pair( sizeof(cl_mem), (void *)&src.data));
+ cl_uint stepBytes = src.step;
+ args.push_back( make_pair( sizeof(cl_uint), (void *)&stepBytes));
+ int offsetXBytes = src.offset % src.step;
+ int offsetX = offsetXBytes / src.elemSize();
+ CV_Assert((int)(offsetX * src.elemSize()) == offsetXBytes);
+ int offsetY = src.offset / src.step;
+ int endX = (offsetX + src.cols);
+ int endY = (offsetY + src.rows);
+ cl_int rect[4] = {offsetX, offsetY, endX, endY};
+ if (!isIsolatedBorder)
{
- case 0:
- sprintf(btype, "BORDER_CONSTANT");
+ rect[2] = src.wholecols;
+ rect[3] = src.wholerows;
+ }
+ args.push_back( make_pair( sizeof(cl_int)*4, (void *)&rect[0]));
+
+ args.push_back( make_pair( sizeof(cl_mem), (void *)&dst.data));
+ cl_uint _stepBytes = dst.step;
+ args.push_back( make_pair( sizeof(cl_uint), (void *)&_stepBytes));
+ int _offsetXBytes = dst.offset % dst.step;
+ int _offsetX = _offsetXBytes / dst.elemSize();
+ CV_Assert((int)(_offsetX * dst.elemSize()) == _offsetXBytes);
+ int _offsetY = dst.offset / dst.step;
+ int _endX = (_offsetX + dst.cols);
+ int _endY = (_offsetY + dst.rows);
+ cl_int _rect[4] = {_offsetX, _offsetY, _endX, _endY};
+ args.push_back( make_pair( sizeof(cl_int)*4, (void *)&_rect[0]));
+
+ bool useDouble = src.depth() == CV_64F;
+
+ float borderValue[4] = {0, 0, 0, 0}; // DON'T move into 'if' body
+ double borderValueDouble[4] = {0, 0, 0, 0}; // DON'T move into 'if' body
+ if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT)
+ {
+ if (useDouble)
+ args.push_back( make_pair( sizeof(double) * src.oclchannels(), (void *)&borderValue[0]));
+ else
+ args.push_back( make_pair( sizeof(float) * src.oclchannels(), (void *)&borderValueDouble[0]));
+ }
+
+ double alphaDouble = alpha; // DON'T move into 'if' body
+ if (useDouble)
+ args.push_back( make_pair( sizeof(double), (void *)&alphaDouble));
+ else
+ args.push_back( make_pair( sizeof(float), (void *)&alpha));
+
+ const char* btype = NULL;
+
+ switch (borderType & ~BORDER_ISOLATED)
+ {
+ case BORDER_CONSTANT:
+ btype = "BORDER_CONSTANT";
break;
- case 1:
- sprintf(btype, "BORDER_REPLICATE");
+ case BORDER_REPLICATE:
+ btype = "BORDER_REPLICATE";
break;
- case 2:
- sprintf(btype, "BORDER_REFLECT");
+ case BORDER_REFLECT:
+ btype = "BORDER_REFLECT";
break;
- case 3:
+ case BORDER_WRAP:
CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
return;
- case 4:
- sprintf(btype, "BORDER_REFLECT_101");
+ case BORDER_REFLECT101:
+ btype = "BORDER_REFLECT_101";
break;
}
- char build_options[150];
- sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype);
+ int requiredTop = anchor.y;
+ int requiredLeft = BLOCK_SIZE; // not this: anchor.x;
+ int requiredBottom = ksize.height - 1 - anchor.y;
+ int requiredRight = BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x;
+ int h = isIsolatedBorder ? src.rows : src.wholerows;
+ int w = isIsolatedBorder ? src.cols : src.wholecols;
+ bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight;
- size_t blockSizeX = 256, blockSizeY = 1;
- size_t gSize = blockSizeX - (ksize.width - 1);
- size_t threads = (dst.offset % dst.step % 4 + dst.cols + 3) / 4;
- size_t globalSizeX = threads % gSize == 0 ? threads / gSize * blockSizeX : (threads / gSize + 1) * blockSizeX;
- size_t globalSizeY = ((dst.rows + 1) / 2) % blockSizeY == 0 ? ((dst.rows + 1) / 2) : (((dst.rows + 1) / 2) / blockSizeY + 1) * blockSizeY;
+ CV_Assert(w >= ksize.width && h >= ksize.height); // TODO Other cases are not tested well
- size_t globalThreads[3] = { globalSizeX, globalSizeY, 1 };
- size_t localThreads[3] = { blockSizeX, blockSizeY, 1 };
+ char build_options[1024];
+ sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%d -D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d -D %s -D %s -D %s",
+ (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y,
+ src.depth(), src.oclchannels(), useDouble ? 1 : 0,
+ anchor.x, anchor.y, ksize.width, ksize.height,
+ btype,
+ extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
+ isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED");
- vector<pair<size_t , const void *> > args;
- args.push_back(make_pair(sizeof(cl_mem), &src.data));
- args.push_back(make_pair(sizeof(cl_mem), &dst.data));
- args.push_back(make_pair(sizeof(cl_float), (void *)&alpha));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.step));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step));
-
- openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options);
+ size_t gt[3] = {divUp(dst.cols, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE, divUp(dst.rows, BLOCK_SIZE_Y), 1}, lt[3] = {BLOCK_SIZE, 1, 1};
+ openCLExecuteKernel(src.clCxt, &filtering_boxFilter, "boxFilter", gt, lt, args, -1, -1, build_options);
}
-static void GPUFilterBox_8u_C4R(const oclMat &src, oclMat &dst,
- Size &ksize, const Point anchor, const int borderType)
-{
- //Normalize the result by default
- float alpha = ksize.height * ksize.width;
-
- CV_Assert(src.clCxt == dst.clCxt);
- CV_Assert((src.cols == dst.cols) &&
- (src.rows == dst.rows));
- Context *clCxt = src.clCxt;
-
- string kernelName = "boxFilter_C4_D0";
-
- char btype[30];
-
- switch (borderType)
- {
- case 0:
- sprintf(btype, "BORDER_CONSTANT");
- break;
- case 1:
- sprintf(btype, "BORDER_REPLICATE");
- break;
- case 2:
- sprintf(btype, "BORDER_REFLECT");
- break;
- case 3:
- CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
- return;
- case 4:
- sprintf(btype, "BORDER_REFLECT_101");
- break;
- }
-
- char build_options[150];
- sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype);
-
- size_t blockSizeX = 256, blockSizeY = 1;
- size_t gSize = blockSizeX - ksize.width / 2 * 2;
- size_t globalSizeX = (src.cols) % gSize == 0 ? src.cols / gSize * blockSizeX : (src.cols / gSize + 1) * blockSizeX;
- size_t rows_per_thread = 2;
- size_t globalSizeY = ((src.rows + rows_per_thread - 1) / rows_per_thread) % blockSizeY == 0 ? ((src.rows + rows_per_thread - 1) / rows_per_thread) : (((src.rows + rows_per_thread - 1) / rows_per_thread) / blockSizeY + 1) * blockSizeY;
-
- size_t globalThreads[3] = { globalSizeX, globalSizeY, 1};
- size_t localThreads[3] = { blockSizeX, blockSizeY, 1};
-
- vector<pair<size_t , const void *> > args;
- args.push_back(make_pair(sizeof(cl_mem), &src.data));
- args.push_back(make_pair(sizeof(cl_mem), &dst.data));
- args.push_back(make_pair(sizeof(cl_float), (void *)&alpha));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.step));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step));
-
- openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options);
-}
-
-static void GPUFilterBox_32F_C1R(const oclMat &src, oclMat &dst,
- Size &ksize, const Point anchor, const int borderType)
-{
- //Normalize the result by default
- float alpha = ksize.height * ksize.width;
-
- CV_Assert(src.clCxt == dst.clCxt);
- CV_Assert((src.cols == dst.cols) &&
- (src.rows == dst.rows));
- Context *clCxt = src.clCxt;
-
- string kernelName = "boxFilter_C1_D5";
-
- char btype[30];
-
- switch (borderType)
- {
- case 0:
- sprintf(btype, "BORDER_CONSTANT");
- break;
- case 1:
- sprintf(btype, "BORDER_REPLICATE");
- break;
- case 2:
- sprintf(btype, "BORDER_REFLECT");
- break;
- case 3:
- CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
- return;
- case 4:
- sprintf(btype, "BORDER_REFLECT_101");
- break;
- }
-
- char build_options[150];
- sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype);
-
- size_t blockSizeX = 256, blockSizeY = 1;
- size_t gSize = blockSizeX - ksize.width / 2 * 2;
- size_t globalSizeX = (src.cols) % gSize == 0 ? src.cols / gSize * blockSizeX : (src.cols / gSize + 1) * blockSizeX;
- size_t rows_per_thread = 2;
- size_t globalSizeY = ((src.rows + rows_per_thread - 1) / rows_per_thread) % blockSizeY == 0 ? ((src.rows + rows_per_thread - 1) / rows_per_thread) : (((src.rows + rows_per_thread - 1) / rows_per_thread) / blockSizeY + 1) * blockSizeY;
-
-
- size_t globalThreads[3] = { globalSizeX, globalSizeY, 1};
- size_t localThreads[3] = { blockSizeX, blockSizeY, 1};
-
- vector<pair<size_t , const void *> > args;
- args.push_back(make_pair(sizeof(cl_mem), &src.data));
- args.push_back(make_pair(sizeof(cl_mem), &dst.data));
- args.push_back(make_pair(sizeof(cl_float), (void *)&alpha));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.step));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step));
-
- openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options);
-}
-
-static void GPUFilterBox_32F_C4R(const oclMat &src, oclMat &dst,
- Size &ksize, const Point anchor, const int borderType)
-{
- //Normalize the result by default
- float alpha = ksize.height * ksize.width;
-
- CV_Assert(src.clCxt == dst.clCxt);
- CV_Assert((src.cols == dst.cols) &&
- (src.rows == dst.rows));
- Context *clCxt = src.clCxt;
-
- string kernelName = "boxFilter_C4_D5";
-
- char btype[30];
-
- switch (borderType)
- {
- case 0:
- sprintf(btype, "BORDER_CONSTANT");
- break;
- case 1:
- sprintf(btype, "BORDER_REPLICATE");
- break;
- case 2:
- sprintf(btype, "BORDER_REFLECT");
- break;
- case 3:
- CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
- return;
- case 4:
- sprintf(btype, "BORDER_REFLECT_101");
- break;
- }
-
- char build_options[150];
- sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype);
-
- size_t blockSizeX = 256, blockSizeY = 1;
- size_t gSize = blockSizeX - ksize.width / 2 * 2;
- size_t globalSizeX = (src.cols) % gSize == 0 ? src.cols / gSize * blockSizeX : (src.cols / gSize + 1) * blockSizeX;
- size_t rows_per_thread = 2;
- size_t globalSizeY = ((src.rows + rows_per_thread - 1) / rows_per_thread) % blockSizeY == 0 ? ((src.rows + rows_per_thread - 1) / rows_per_thread) : (((src.rows + rows_per_thread - 1) / rows_per_thread) / blockSizeY + 1) * blockSizeY;
-
-
- size_t globalThreads[3] = { globalSizeX, globalSizeY, 1};
- size_t localThreads[3] = { blockSizeX, blockSizeY, 1};
-
- vector<pair<size_t , const void *> > args;
- args.push_back(make_pair(sizeof(cl_mem), &src.data));
- args.push_back(make_pair(sizeof(cl_mem), &dst.data));
- args.push_back(make_pair(sizeof(cl_float), (void *)&alpha));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.step));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step));
-
- openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options);
-}
-
-
-Ptr<BaseFilter_GPU> cv::ocl::getBoxFilter_GPU(int srcType, int dstType,
+Ptr<BaseFilter_GPU> cv::ocl::getBoxFilter_GPU(int /*srcType*/, int /*dstType*/,
const Size &ksize, Point anchor, int borderType)
{
- static const FilterBox_t FilterBox_callers[2][5] = {{0, GPUFilterBox_8u_C1R, 0, GPUFilterBox_8u_C4R, GPUFilterBox_8u_C4R},
- {0, GPUFilterBox_32F_C1R, 0, GPUFilterBox_32F_C4R, GPUFilterBox_32F_C4R}
- };
- //Remove this check if more data types need to be supported.
- CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC3 || srcType == CV_8UC4 || srcType == CV_32FC1 ||
- srcType == CV_32FC3 || srcType == CV_32FC4) && dstType == srcType);
-
normalizeAnchor(anchor, ksize);
return Ptr<BaseFilter_GPU>(new GPUBoxFilter(ksize, anchor,
- borderType, FilterBox_callers[(CV_MAT_DEPTH(srcType) == CV_32F)][CV_MAT_CN(srcType)]));
+ borderType, GPUFilterBox));
}
Ptr<FilterEngine_GPU> cv::ocl::createBoxFilter_GPU(int srcType, int dstType,
diff --git a/modules/ocl/src/opencl/filtering_boxFilter.cl b/modules/ocl/src/opencl/filtering_boxFilter.cl
index 030c13cc5..7f7fd018d 100644
--- a/modules/ocl/src/opencl/filtering_boxFilter.cl
+++ b/modules/ocl/src/opencl/filtering_boxFilter.cl
@@ -10,13 +10,9 @@
// License Agreement
// For Open Source Computer Vision Library
//
-// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
-// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
-// @Authors
-// Zhang Ying, zhangying913@gmail.com
-//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
@@ -79,400 +75,298 @@
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
#endif
-#define THREADS 256
-#define ELEM(i, l_edge, r_edge, elem1, elem2) (i) >= (l_edge) && (i) < (r_edge) ? (elem1) : (elem2)
-
-inline void update_dst_C1_D0(__global uchar *dst, __local uint* temp,
- int dst_rows, int dst_cols,
- int dst_startX, int dst_x_off,
- float alpha)
-{
- if(get_local_id(0) < anX || get_local_id(0) >= (THREADS-ksX+anX+1))
- {
- return;
- }
-
- uint4 tmp_sum = 0;
- int posX = dst_startX - dst_x_off + (get_local_id(0)-anX)*4;
- int posY = (get_group_id(1) << 1);
-
- for(int i=-anX; i<=anX; i++)
- {
- tmp_sum += vload4(get_local_id(0), temp+i);
- }
-
- if(posY < dst_rows && posX < dst_cols)
- {
- tmp_sum /= (uint4) alpha;
- if(posX >= 0 && posX < dst_cols)
- *(dst) = tmp_sum.x;
- if(posX+1 >= 0 && posX+1 < dst_cols)
- *(dst + 1) = tmp_sum.y;
- if(posX+2 >= 0 && posX+2 < dst_cols)
- *(dst + 2) = tmp_sum.z;
- if(posX+3 >= 0 && posX+3 < dst_cols)
- *(dst + 3) = tmp_sum.w;
- }
-}
-
-
-inline void update_dst_C4_D0(__global uchar4 *dst, __local uint4* temp,
- int dst_rows, int dst_cols,
- int dst_startX, int dst_x_off,
- float alpha)
-{
- if(get_local_id(0) >= (THREADS-ksX+1))
- {
- return;
- }
-
- int posX = dst_startX - dst_x_off + get_local_id(0);
- int posY = (get_group_id(1) << 1);
-
- uint4 temp_sum = 0;
- for(int i=-anX; i<=anX; i++)
- {
- temp_sum += temp[get_local_id(0) + anX + i];
- }
-
- if(posX >= 0 && posX < dst_cols && posY >= 0 && posY < dst_rows)
- *dst = convert_uchar4(convert_float4(temp_sum)/alpha);
-}
-
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////////////8uC1////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////////////////////////
-__kernel void boxFilter_C1_D0(__global const uchar * restrict src, __global uchar *dst, float alpha,
- int src_offset, int src_whole_rows, int src_whole_cols, int src_step,
- int dst_offset, int dst_rows, int dst_cols, int dst_step
- )
-{
-
- int col = get_local_id(0);
- const int gX = get_group_id(0);
- const int gY = get_group_id(1);
- int src_x_off = src_offset % src_step;
- int src_y_off = src_offset / src_step;
- int dst_x_off = dst_offset % dst_step;
- int dst_y_off = dst_offset / dst_step;
-
- int head_off = dst_x_off%4;
- int startX = ((gX * (THREADS-ksX+1)-anX) * 4) - head_off + src_x_off;
- int startY = (gY << 1) - anY + src_y_off;
- int dst_startX = (gX * (THREADS-ksX+1) * 4) - head_off + dst_x_off;
- int dst_startY = (gY << 1) + dst_y_off;
-
- uint4 data[ksY+1];
- __local uint4 temp[2][THREADS];
-
+#ifdef EXTRA_EXTRAPOLATION // border > src image size
#ifdef BORDER_CONSTANT
+// None
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ x = max(min(x, maxX - 1), minX); \
+ y = max(min(y, maxY - 1), minY); \
+ }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ if (x < minX) \
+ x -= ((x - maxX + 1) / maxX) * maxX; \
+ if (x >= maxX) \
+ x %= maxX; \
+ if (y < minY) \
+ y -= ((y - maxY + 1) / maxY) * maxY; \
+ if (y >= maxY) \
+ y %= maxY; \
+ }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \
+ { \
+ if (maxX - minX == 1) \
+ x = minX; \
+ else \
+ do \
+ { \
+ if (x < minX) \
+ x = -(x - minX) - 1 + delta; \
+ else \
+ x = maxX - 1 - (x - maxX) - delta; \
+ } \
+ while (x >= maxX || x < minX); \
+ \
+ if (maxY - minY == 1) \
+ y = minY; \
+ else \
+ do \
+ { \
+ if (y < minY) \
+ y = -(y - minY) - 1 + delta; \
+ else \
+ y = maxY - 1 - (y - maxY) - delta; \
+ } \
+ while (y >= maxY || y < minY); \
+ }
+#ifdef BORDER_REFLECT
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0)
+#elif defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1)
+#endif
+#else
+#error No extrapolation method
+#endif
+#else
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ int _row = y - minY, _col = x - minX; \
+ _row = ADDR_H(_row, 0, maxY - minY); \
+ _row = ADDR_B(_row, maxY - minY, _row); \
+ y = _row + minY; \
+ \
+ _col = ADDR_L(_col, 0, maxX - minX); \
+ _col = ADDR_R(_col, maxX - minX, _col); \
+ x = _col + minX; \
+ }
+#endif
- for(int i=0; i < ksY+1; i++)
+#if USE_DOUBLE
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#define FPTYPE double
+#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE)
+#else
+#define FPTYPE float
+#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE)
+#endif
+
+#if DATA_DEPTH == 0
+#define BASE_TYPE uchar
+#elif DATA_DEPTH == 1
+#define BASE_TYPE char
+#elif DATA_DEPTH == 2
+#define BASE_TYPE ushort
+#elif DATA_DEPTH == 3
+#define BASE_TYPE short
+#elif DATA_DEPTH == 4
+#define BASE_TYPE int
+#elif DATA_DEPTH == 5
+#define BASE_TYPE float
+#elif DATA_DEPTH == 6
+#define BASE_TYPE double
+#else
+#error data_depth
+#endif
+
+#define __CAT(x, y) x##y
+#define CAT(x, y) __CAT(x, y)
+
+#define uchar1 uchar
+#define char1 char
+#define ushort1 ushort
+#define short1 short
+#define int1 int
+#define float1 float
+#define double1 double
+
+#define convert_uchar1_sat_rte convert_uchar_sat_rte
+#define convert_char1_sat_rte convert_char_sat_rte
+#define convert_ushort1_sat_rte convert_ushort_sat_rte
+#define convert_short1_sat_rte convert_short_sat_rte
+#define convert_int1_sat_rte convert_int_sat_rte
+#define convert_float1
+#define convert_double1
+
+#if DATA_DEPTH == 5 || DATA_DEPTH == 6
+#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE)
+#else
+#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte)
+#endif
+
+#define VEC_SIZE DATA_CHAN
+
+#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE)
+#define TYPE VEC_TYPE
+
+#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE)
+
+#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE)
+
+struct RectCoords
+{
+ int x1, y1, x2, y2;
+};
+
+//#define DEBUG
+#ifdef DEBUG
+#define DEBUG_ONLY(x) x
+#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0)
+#else
+#define DEBUG_ONLY(x)
+#define ASSERT(condition)
+#endif
+
+
+inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords
+#ifdef BORDER_CONSTANT
+ , SCALAR_TYPE borderValue
+#endif
+ )
+{
+#ifdef BORDER_ISOLATED
+ if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+#else
+ if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+#endif
{
- if(startY+i >=0 && startY+i < src_whole_rows && startX+col*4 >=0 && startX+col*4+3<src_whole_cols)
+ __global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes);
+ return CONVERT_TO_FPTYPE(*ptr);
+ }
+ else
+ {
+#ifdef BORDER_CONSTANT
+ return borderValue;
+#else
+ int selected_col = pos.x;
+ int selected_row = pos.y;
+
+ EXTRAPOLATE(selected_col, selected_row,
+#ifdef BORDER_ISOLATED
+ srcCoords.x1, srcCoords.y1,
+#else
+ 0, 0,
+#endif
+ srcCoords.x2, srcCoords.y2
+ );
+
+ // debug border mapping
+ //printf("pos=%d,%d --> %d, %d\n", pos.x, pos.y, selected_col, selected_row);
+
+ pos = (int2)(selected_col, selected_row);
+ if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
{
- data[i].x = *(src+(startY+i)*src_step + startX + col * 4);
- data[i].y = *(src+(startY+i)*src_step + startX + col * 4 + 1);
- data[i].z = *(src+(startY+i)*src_step + startX + col * 4 + 2);
- data[i].w = *(src+(startY+i)*src_step + startX + col * 4 + 3);
+ __global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes);
+ return CONVERT_TO_FPTYPE(*ptr);
}
else
{
- data[i]=0;
- int con = startY+i >=0 && startY+i < src_whole_rows && startX+col*4 >=0 && startX+col*4<src_whole_cols;
- if(con)data[i].s0 = *(src+(startY+i)*src_step + startX + col*4);
- con = startY+i >=0 && startY+i < src_whole_rows && startX+col*4+1 >=0 && startX+col*4+1<src_whole_cols;
- if(con)data[i].s1 = *(src+(startY+i)*src_step + startX + col*4+1) ;
- con = startY+i >=0 && startY+i < src_whole_rows && startX+col*4+2 >=0 && startX+col*4+2<src_whole_cols;
- if(con)data[i].s2 = *(src+(startY+i)*src_step + startX + col*4+2);
- con = startY+i >=0 && startY+i < src_whole_rows && startX+col*4+3 >=0 && startX+col*4+3<src_whole_cols;
- if(con)data[i].s3 = *(src+(startY+i)*src_step + startX + col*4+3);
+ // for debug only
+ DEBUG_ONLY(printf("BUG in boxFilter kernel\n"));
+ return (FPTYPE)(0.0f);
}
- }
-
-#else
- int not_all_in_range;
- for(int i=0; i < ksY+1; i++)
- {
- not_all_in_range = (startX+col*4<0) | (startX+col*4+3>src_whole_cols-1)
- | (startY+i<0) | (startY+i>src_whole_rows-1);
- if(not_all_in_range)
- {
- int selected_row;
- int4 selected_col;
- selected_row = ADDR_H(startY+i, 0, src_whole_rows);
- selected_row = ADDR_B(startY+i, src_whole_rows, selected_row);
-
- selected_col.x = ADDR_L(startX+col*4, 0, src_whole_cols);
- selected_col.x = ADDR_R(startX+col*4, src_whole_cols, selected_col.x);
-
- selected_col.y = ADDR_L(startX+col*4+1, 0, src_whole_cols);
- selected_col.y = ADDR_R(startX+col*4+1, src_whole_cols, selected_col.y);
-
- selected_col.z = ADDR_L(startX+col*4+2, 0, src_whole_cols);
- selected_col.z = ADDR_R(startX+col*4+2, src_whole_cols, selected_col.z);
-
- selected_col.w = ADDR_L(startX+col*4+3, 0, src_whole_cols);
- selected_col.w = ADDR_R(startX+col*4+3, src_whole_cols, selected_col.w);
-
- data[i].x = *(src + selected_row * src_step + selected_col.x);
- data[i].y = *(src + selected_row * src_step + selected_col.y);
- data[i].z = *(src + selected_row * src_step + selected_col.z);
- data[i].w = *(src + selected_row * src_step + selected_col.w);
- }
- else
- {
- data[i] = convert_uint4(vload4(col,(__global uchar*)(src+(startY+i)*src_step + startX)));
- }
- }
#endif
- uint4 tmp_sum = 0;
- for(int i=1; i < ksY; i++)
- {
- tmp_sum += (data[i]);
}
-
- int index = dst_startY * dst_step + dst_startX + (col-anX)*4;
-
- temp[0][col] = tmp_sum + (data[0]);
- temp[1][col] = tmp_sum + (data[ksY]);
- barrier(CLK_LOCAL_MEM_FENCE);
- update_dst_C1_D0(dst+index, (__local uint *)(temp[0]),
- dst_rows, dst_cols, dst_startX, dst_x_off, alpha);
- update_dst_C1_D0(dst+index+dst_step, (__local uint *)(temp[1]),
- dst_rows, dst_cols, dst_startX, dst_x_off, alpha);
-
}
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////////////8uC4////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////////////////////////
-__kernel void boxFilter_C4_D0(__global const uchar4 * restrict src, __global uchar4 *dst, float alpha,
- int src_offset, int src_whole_rows, int src_whole_cols, int src_step,
- int dst_offset, int dst_rows, int dst_cols, int dst_step
- )
-{
- int col = get_local_id(0);
- const int gX = get_group_id(0);
- const int gY = get_group_id(1);
-
- int src_x_off = (src_offset % src_step) >> 2;
- int src_y_off = src_offset / src_step;
- int dst_x_off = (dst_offset % dst_step) >> 2;
- int dst_y_off = dst_offset / dst_step;
-
- int startX = gX * (THREADS-ksX+1) - anX + src_x_off;
- int startY = (gY << 1) - anY + src_y_off;
- int dst_startX = gX * (THREADS-ksX+1) + dst_x_off;
- int dst_startY = (gY << 1) + dst_y_off;
-
- uint4 data[ksY+1];
- __local uint4 temp[2][THREADS];
+// INPUT PARAMETER: BLOCK_SIZE_Y (via defines)
+__kernel
+__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1)))
+void boxFilter(__global TYPE *src, const unsigned int srcStepBytes, const int4 srcRC,
+ __global TYPE *dst, const unsigned int dstStepBytes, const int4 dstRC,
#ifdef BORDER_CONSTANT
- bool con;
- for(int i=0; i < ksY+1; i++)
- {
- con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows;
- int cur_col = clamp(startX + col, 0, src_whole_cols);
-
- data[i].x = con ? src[(startY+i)*(src_step>>2) + cur_col].x : 0;
- data[i].y = con ? src[(startY+i)*(src_step>>2) + cur_col].y : 0;
- data[i].z = con ? src[(startY+i)*(src_step>>2) + cur_col].z : 0;
- data[i].w = con ? src[(startY+i)*(src_step>>2) + cur_col].w : 0;
- }
-#else
- for(int i=0; i < ksY+1; i++)
- {
- int selected_row;
- int selected_col;
- selected_row = ADDR_H(startY+i, 0, src_whole_rows);
- selected_row = ADDR_B(startY+i, src_whole_rows, selected_row);
-
- selected_col = ADDR_L(startX+col, 0, src_whole_cols);
- selected_col = ADDR_R(startX+col, src_whole_cols, selected_col);
-
-
- data[i] = convert_uint4(src[selected_row * (src_step>>2) + selected_col]);
- }
-
+ SCALAR_TYPE borderValue,
#endif
- uint4 tmp_sum = 0;
- for(int i=1; i < ksY; i++)
- {
- tmp_sum += (data[i]);
- }
-
- int index = dst_startY * (dst_step>>2)+ dst_startX + col;
-
- temp[0][col] = tmp_sum + (data[0]);
- temp[1][col] = tmp_sum + (data[ksY]);
- barrier(CLK_LOCAL_MEM_FENCE);
- update_dst_C4_D0(dst+index, (__local uint4 *)(temp[0]),
- dst_rows, dst_cols, dst_startX, dst_x_off, alpha);
- update_dst_C4_D0(dst+index+(dst_step>>2), (__local uint4 *)(temp[1]),
- dst_rows, dst_cols, dst_startX, dst_x_off, alpha);
-
-}
-
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////////////32fC1////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////////////////////////
-__kernel void boxFilter_C1_D5(__global const float *restrict src, __global float *dst, float alpha,
- int src_offset, int src_whole_rows, int src_whole_cols, int src_step,
- int dst_offset, int dst_rows, int dst_cols, int dst_step
- )
+ FPTYPE alpha
+ )
{
- int col = get_local_id(0);
- const int gX = get_group_id(0);
- const int gY = get_group_id(1);
+ const struct RectCoords srcCoords = {srcRC.s0, srcRC.s1, srcRC.s2, srcRC.s3}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY
+ const struct RectCoords dstCoords = {dstRC.s0, dstRC.s1, dstRC.s2, dstRC.s3};
- int src_x_off = (src_offset % src_step) >> 2;
- int src_y_off = src_offset / src_step;
- int dst_x_off = (dst_offset % dst_step) >> 2;
- int dst_y_off = dst_offset / dst_step;
+ const int x = get_local_id(0) + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X;
+ const int y = get_global_id(1) * BLOCK_SIZE_Y;
- int startX = gX * (THREADS-ksX+1) - anX + src_x_off;
- int startY = (gY << 1) - anY + src_y_off;
- int dst_startX = gX * (THREADS-ksX+1) + dst_x_off;
- int dst_startY = (gY << 1) + dst_y_off;
- float data[ksY+1];
- __local float temp[2][THREADS];
+ const int local_id = get_local_id(0);
+
+ INTERMEDIATE_TYPE data[KERNEL_SIZE_Y];
+ __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE];
+
+ int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y);
+ for(int sy = 0; sy < KERNEL_SIZE_Y; sy++, srcPos.y++)
+ {
+ data[sy] = readSrcPixel(srcPos, src, srcStepBytes, srcCoords
#ifdef BORDER_CONSTANT
- bool con;
- float ss;
- for(int i=0; i < ksY+1; i++)
- {
- con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows;
-
- int cur_col = clamp(startX + col, 0, src_whole_cols);
- ss = (startY+i)<src_whole_rows&&(startY+i)>=0&&cur_col>=0&&cur_col<src_whole_cols?src[(startY+i)*(src_step>>2) + cur_col]:(float)0;
-
- data[i] = con ? ss : 0.f;
- }
-#else
- for(int i=0; i < ksY+1; i++)
- {
- int selected_row;
- int selected_col;
- selected_row = ADDR_H(startY+i, 0, src_whole_rows);
- selected_row = ADDR_B(startY+i, src_whole_rows, selected_row);
-
- selected_col = ADDR_L(startX+col, 0, src_whole_cols);
- selected_col = ADDR_R(startX+col, src_whole_cols, selected_col);
-
- data[i] = src[selected_row * (src_step>>2) + selected_col];
- }
-
+ , borderValue
#endif
- float sum0 = 0.0, sum1 = 0.0, sum2 = 0.0;
- for(int i=1; i < ksY; i++)
- {
- sum0 += (data[i]);
+ );
}
- sum1 = sum0 + (data[0]);
- sum2 = sum0 + (data[ksY]);
- temp[0][col] = sum1;
- temp[1][col] = sum2;
- barrier(CLK_LOCAL_MEM_FENCE);
- if(col < (THREADS-(ksX-1)))
- {
- col += anX;
- int posX = dst_startX - dst_x_off + col - anX;
- int posY = (gY << 1);
- float tmp_sum[2]= {0.0, 0.0};
- for(int k=0; k<2; k++)
- for(int i=-anX; i<=anX; i++)
+ INTERMEDIATE_TYPE tmp_sum = 0;
+ for(int sy = 0; sy < KERNEL_SIZE_Y; sy++)
+ {
+ tmp_sum += (data[sy]);
+ }
+
+ sumOfCols[local_id] = tmp_sum;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y);
+ __global TYPE* dstPtr = (__global TYPE*)((__global char*)dst + pos.x * sizeof(TYPE) + pos.y * dstStepBytes); // Pointer can be out of bounds!
+
+ int sy_index = 0; // current index in data[] array
+ int stepsY = min(dstCoords.y2 - pos.y, BLOCK_SIZE_Y);
+ ASSERT(stepsY > 0);
+ for (; ;)
+ {
+ ASSERT(pos.y < dstCoords.y2);
+
+ if(local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) &&
+ pos.x >= dstCoords.x1 && pos.x < dstCoords.x2)
+ {
+ ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2);
+
+ INTERMEDIATE_TYPE total_sum = 0;
+#pragma unroll
+ for (int sx = 0; sx < KERNEL_SIZE_X; sx++)
{
- tmp_sum[k] += temp[k][col+i];
+ total_sum += sumOfCols[local_id + sx - ANCHOR_X];
}
- for(int i=0; i<2; i++)
- {
- if(posX >= 0 && posX < dst_cols && (posY+i) >= 0 && (posY+i) < dst_rows)
- dst[(dst_startY+i) * (dst_step>>2)+ dst_startX + col - anX] = tmp_sum[i]/alpha;
+ *dstPtr = CONVERT_TO_TYPE(((INTERMEDIATE_TYPE)alpha) * total_sum);
}
- }
-}
-
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////////////32fC4////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////////////////////////
-__kernel void boxFilter_C4_D5(__global const float4 *restrict src, __global float4 *dst, float alpha,
- int src_offset, int src_whole_rows, int src_whole_cols, int src_step,
- int dst_offset, int dst_rows, int dst_cols, int dst_step
- )
-{
- int col = get_local_id(0);
- const int gX = get_group_id(0);
- const int gY = get_group_id(1);
-
- int src_x_off = (src_offset % src_step) >> 4;
- int src_y_off = src_offset / src_step;
- int dst_x_off = (dst_offset % dst_step) >> 4;
- int dst_y_off = dst_offset / dst_step;
-
- int startX = gX * (THREADS-ksX+1) - anX + src_x_off;
- int startY = (gY << 1) - anY + src_y_off;
- int dst_startX = gX * (THREADS-ksX+1) + dst_x_off;
- int dst_startY = (gY << 1) + dst_y_off;
- float4 data[ksY+1];
- __local float4 temp[2][THREADS];
-#ifdef BORDER_CONSTANT
- bool con;
- float4 ss;
- for(int i=0; i < ksY+1; i++)
- {
- con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows;
-
- int cur_col = clamp(startX + col, 0, src_whole_cols);
- ss = (startY+i)<src_whole_rows&&(startY+i)>=0&&cur_col>=0&&cur_col<src_whole_cols?src[(startY+i)*(src_step>>4) + cur_col]:(float4)0;
-
- data[i] = con ? ss : (float4)(0.0,0.0,0.0,0.0);
- }
+#if BLOCK_SIZE_Y == 1
+ break;
#else
- for(int i=0; i < ksY+1; i++)
- {
- int selected_row;
- int selected_col;
- selected_row = ADDR_H(startY+i, 0, src_whole_rows);
- selected_row = ADDR_B(startY+i, src_whole_rows, selected_row);
+ if (--stepsY == 0)
+ break;
- selected_col = ADDR_L(startX+col, 0, src_whole_cols);
- selected_col = ADDR_R(startX+col, src_whole_cols, selected_col);
+ barrier(CLK_LOCAL_MEM_FENCE);
- data[i] = src[selected_row * (src_step>>4) + selected_col];
- }
+ tmp_sum = sumOfCols[local_id]; // TODO FIX IT: workaround for BUG in OpenCL compiler
+ // only works with scalars: ASSERT(fabs(tmp_sum - sumOfCols[local_id]) < (INTERMEDIATE_TYPE)1e-6);
+ tmp_sum -= data[sy_index];
+ data[sy_index] = readSrcPixel(srcPos, src, srcStepBytes, srcCoords
+#ifdef BORDER_CONSTANT
+ , borderValue
#endif
- float4 sum0 = 0.0, sum1 = 0.0, sum2 = 0.0;
- for(int i=1; i < ksY; i++)
- {
- sum0 += (data[i]);
- }
- sum1 = sum0 + (data[0]);
- sum2 = sum0 + (data[ksY]);
- temp[0][col] = sum1;
- temp[1][col] = sum2;
- barrier(CLK_LOCAL_MEM_FENCE);
- if(col < (THREADS-(ksX-1)))
- {
- col += anX;
- int posX = dst_startX - dst_x_off + col - anX;
- int posY = (gY << 1);
+ );
+ srcPos.y++;
- float4 tmp_sum[2]= {(float4)(0.0,0.0,0.0,0.0), (float4)(0.0,0.0,0.0,0.0)};
- for(int k=0; k<2; k++)
- for(int i=-anX; i<=anX; i++)
- {
- tmp_sum[k] += temp[k][col+i];
- }
- for(int i=0; i<2; i++)
- {
- if(posX >= 0 && posX < dst_cols && (posY+i) >= 0 && (posY+i) < dst_rows)
- dst[(dst_startY+i) * (dst_step>>4)+ dst_startX + col - anX] = tmp_sum[i]/alpha;
- }
+ tmp_sum += data[sy_index];
+ sumOfCols[local_id] = tmp_sum;
+ sy_index = (sy_index + 1 < KERNEL_SIZE_Y) ? sy_index + 1 : 0;
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ // next line
+ DEBUG_ONLY(pos.y++);
+ dstPtr = (__global TYPE*)((__global char*)dstPtr + dstStepBytes); // Pointer can be out of bounds!
+#endif // BLOCK_SIZE_Y == 1
}
}
From 0f95f0d8b3f4a11d94d1d34326953d89a047fbd5 Mon Sep 17 00:00:00 2001
From: Alexander Alekhin <alexander.alekhin@itseez.com>
Date: Sat, 26 Oct 2013 05:06:22 +0400
Subject: [PATCH 3/6] ocl: rewrite filter2D
---
modules/ocl/include/opencv2/ocl/ocl.hpp | 12 +-
modules/ocl/src/filtering.cpp | 275 ++++++++-----
modules/ocl/src/opencl/filtering_filter2D.cl | 370 +++++++++++++++++
modules/ocl/src/opencl/filtering_laplacian.cl | 381 ------------------
modules/ocl/test/test_filters.cpp | 6 +-
5 files changed, 544 insertions(+), 500 deletions(-)
create mode 100644 modules/ocl/src/opencl/filtering_filter2D.cl
delete mode 100644 modules/ocl/src/opencl/filtering_laplacian.cl
diff --git a/modules/ocl/include/opencv2/ocl/ocl.hpp b/modules/ocl/include/opencv2/ocl/ocl.hpp
index 05bd061ca..db386952a 100644
--- a/modules/ocl/include/opencv2/ocl/ocl.hpp
+++ b/modules/ocl/include/opencv2/ocl/ocl.hpp
@@ -718,8 +718,9 @@ namespace cv
CV_EXPORTS Ptr<FilterEngine_GPU> createDerivFilter_GPU( int srcType, int dstType, int dx, int dy, int ksize, int borderType = BORDER_DEFAULT );
//! applies Laplacian operator to the image
- // supports only ksize = 1 and ksize = 3 8UC1 8UC4 32FC1 32FC4 data type
- CV_EXPORTS void Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1);
+ // supports only ksize = 1 and ksize = 3
+ CV_EXPORTS void Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1,
+ double delta=0, int borderType=BORDER_DEFAULT);
//! returns 2D box filter
// dst type must be the same as source type
@@ -731,11 +732,12 @@ namespace cv
const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
//! returns 2D filter with the specified kernel
- // supports CV_8UC1 and CV_8UC4 types
+ // supports: dst type must be the same as source type
CV_EXPORTS Ptr<BaseFilter_GPU> getLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Size &ksize,
const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
//! returns the non-separable linear filter engine
+ // supports: dst type must be the same as source type
CV_EXPORTS Ptr<FilterEngine_GPU> createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel,
const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
@@ -762,10 +764,8 @@ namespace cv
}
//! applies non-separable 2D linear filter to the image
- // Note, at the moment this function only works when anchor point is in the kernel center
- // and kernel size supported is either 3x3 or 5x5; otherwise the function will fail to output valid result
CV_EXPORTS void filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel,
- Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT);
+ Point anchor = Point(-1, -1), double delta = 0.0, int borderType = BORDER_DEFAULT);
//! applies separable 2D linear filter to the image
CV_EXPORTS void sepFilter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernelX, const Mat &kernelY,
diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp
index fdddc1674..e1255197f 100644
--- a/modules/ocl/src/filtering.cpp
+++ b/modules/ocl/src/filtering.cpp
@@ -69,37 +69,14 @@ inline void normalizeAnchor(Point &anchor, const Size &ksize)
normalizeAnchor(anchor.y, ksize.height);
}
-inline void normalizeROI(Rect &roi, const Size &ksize, const Point &anchor, const Size &src_size)
+inline void normalizeROI(Rect &roi, const Size &ksize, const Point &/*anchor*/, const Size &src_size)
{
if (roi == Rect(0, 0, -1, -1))
roi = Rect(0, 0, src_size.width, src_size.height);
CV_Assert(ksize.height > 0 && ksize.width > 0 && ((ksize.height & 1) == 1) && ((ksize.width & 1) == 1));
- CV_Assert((anchor.x == -1 && anchor.y == -1) || (anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1));
CV_Assert(roi.x >= 0 && roi.y >= 0 && roi.width <= src_size.width && roi.height <= src_size.height);
}
-
-
-inline void normalizeKernel(const Mat &kernel, oclMat &gpu_krnl, int type = CV_8U, int *nDivisor = 0, bool reverse = false)
-{
- int scale = nDivisor && (kernel.depth() == CV_32F || kernel.depth() == CV_64F) ? 256 : 1;
-
- if (nDivisor)
- *nDivisor = scale;
- Mat temp(kernel.size(), type);
- kernel.convertTo(temp, type, scale);
- Mat cont_krnl = temp.reshape(1, 1);
-
- if (reverse)
- {
- int count = cont_krnl.cols >> 1;
-
- for (int i = 0; i < count; ++i)
- std::swap(cont_krnl.at<int>(0, i), cont_krnl.at<int>(0, cont_krnl.cols - 1 - i));
- }
-
- gpu_krnl.upload(cont_krnl);
-}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -168,7 +145,7 @@ typedef void (*GPUMorfFilter_t)(const oclMat & , oclMat & , oclMat & , Size &, c
class MorphFilter_GPU : public BaseFilter_GPU
{
public:
- MorphFilter_GPU(const Size &ksize_, const Point &anchor_, const oclMat &kernel_, GPUMorfFilter_t func_) :
+ MorphFilter_GPU(const Size &ksize_, const Point &anchor_, const Mat &kernel_, GPUMorfFilter_t func_) :
BaseFilter_GPU(ksize_, anchor_, BORDER_CONSTANT), kernel(kernel_), func(func_), rectKernel(false) {}
virtual void operator()(const oclMat &src, oclMat &dst)
@@ -355,16 +332,17 @@ Ptr<BaseFilter_GPU> cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat
CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE);
CV_Assert(type == CV_8UC1 || type == CV_8UC3 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC3 || type == CV_32FC4);
- oclMat gpu_krnl;
- normalizeKernel(kernel, gpu_krnl);
normalizeAnchor(anchor, ksize);
+ Mat kernel8U;
+ kernel.convertTo(kernel8U, CV_8U);
+ Mat cont_krnl = kernel8U.reshape(1, 1);
bool noZero = true;
for(int i = 0; i < kernel.rows * kernel.cols; ++i)
if(kernel.data[i] != 1)
noZero = false;
- MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, gpu_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]);
+ MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, cont_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]);
if(noZero)
mfgpu->rectKernel = true;
@@ -524,12 +502,12 @@ void cv::ocl::morphologyEx(const oclMat &src, oclMat &dst, int op, const Mat &ke
namespace
{
-typedef void (*GPUFilter2D_t)(const oclMat & , oclMat & , const oclMat & , const Size &, const Point&, const int);
+typedef void (*GPUFilter2D_t)(const oclMat & , oclMat & , const Mat & , const Size &, const Point&, const int);
class LinearFilter_GPU : public BaseFilter_GPU
{
public:
- LinearFilter_GPU(const Size &ksize_, const Point &anchor_, const oclMat &kernel_, GPUFilter2D_t func_,
+ LinearFilter_GPU(const Size &ksize_, const Point &anchor_, const Mat &kernel_, GPUFilter2D_t func_,
int borderType_) :
BaseFilter_GPU(ksize_, anchor_, borderType_), kernel(kernel_), func(func_) {}
@@ -543,118 +521,192 @@ public:
};
}
-static void GPUFilter2D(const oclMat &src, oclMat &dst, const oclMat &mat_kernel,
+// prepare kernel: transpose and make double rows (+align). Returns size of aligned row
+// Samples:
+// a b c
+// Input: d e f
+// g h i
+// Output, last two zeros is the alignment:
+// a d g a d g 0 0
+// b e h b e h 0 0
+// c f i c f i 0 0
+template <typename T>
+static int _prepareKernelFilter2D(std::vector<T>& data, const Mat &kernel)
+{
+ Mat _kernel; kernel.convertTo(_kernel, DataDepth<T>::value);
+ int size_y_aligned = roundUp(kernel.rows * 2, 4);
+ data.clear(); data.resize(size_y_aligned * kernel.cols, 0);
+ for (int x = 0; x < kernel.cols; x++)
+ {
+ for (int y = 0; y < kernel.rows; y++)
+ {
+ data[x * size_y_aligned + y] = _kernel.at<T>(y, x);
+ data[x * size_y_aligned + y + kernel.rows] = _kernel.at<T>(y, x);
+ }
+ }
+ return size_y_aligned;
+}
+
+static void GPUFilter2D(const oclMat &src, oclMat &dst, const Mat &kernel,
const Size &ksize, const Point& anchor, const int borderType)
{
CV_Assert(src.clCxt == dst.clCxt);
CV_Assert((src.cols == dst.cols) &&
(src.rows == dst.rows));
- CV_Assert((src.oclchannels() == dst.oclchannels()));
- CV_Assert(ksize.height > 0 && ksize.width > 0 && ((ksize.height & 1) == 1) && ((ksize.width & 1) == 1));
- CV_Assert((anchor.x == -1 && anchor.y == -1) || (anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1));
- CV_Assert(ksize.width == ksize.height);
- Context *clCxt = src.clCxt;
+ CV_Assert(src.oclchannels() == dst.oclchannels());
- int filterWidth = ksize.width;
- bool ksize_3x3 = filterWidth == 3 && src.type() != CV_32FC4 && src.type() != CV_32FC3; // CV_32FC4 is not tuned up with filter2d_3x3 kernel
+ CV_Assert(kernel.cols == ksize.width && kernel.rows == ksize.height);
+ CV_Assert(kernel.channels() == 1);
- string kernelName = ksize_3x3 ? "filter2D_3x3" : "filter2D";
+ CV_Assert(anchor.x >= 0 && anchor.x < kernel.cols);
+ CV_Assert(anchor.y >= 0 && anchor.y < kernel.rows);
- size_t src_offset_x = (src.offset % src.step) / src.elemSize();
- size_t src_offset_y = src.offset / src.step;
+ bool useDouble = src.depth() == CV_64F;
- size_t dst_offset_x = (dst.offset % dst.step) / dst.elemSize();
- size_t dst_offset_y = dst.offset / dst.step;
-
- int paddingPixels = filterWidth & (-2);
-
- size_t localThreads[3] = {ksize_3x3 ? 256 : 16, ksize_3x3 ? 1 : 16, 1};
- size_t globalThreads[3] = {src.wholecols, src.wholerows, 1};
-
- int cn = src.oclchannels();
- int src_step = (int)(src.step/src.elemSize());
- int dst_step = (int)(dst.step/src.elemSize());
-
- int localWidth = localThreads[0] + paddingPixels;
- int localHeight = localThreads[1] + paddingPixels;
-
- size_t localMemSize = ksize_3x3 ? 260 * 6 * src.elemSize() : (localWidth * localHeight) * src.elemSize();
-
- int vector_lengths[4][7] = {{4, 4, 4, 4, 4, 4, 4},
- {4, 4, 1, 1, 1, 1, 1},
- {1, 1, 1, 1, 1, 1, 1},
- {4, 4, 4, 4, 1, 1, 4}
- };
- int cols = dst.cols + ((dst_offset_x) & (vector_lengths[cn - 1][src.depth()] - 1));
-
- vector< pair<size_t, const void *> > args;
- args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data));
- args.push_back(make_pair(sizeof(cl_mem), (void *)&dst.data));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src_step));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst_step));
- args.push_back(make_pair(sizeof(cl_mem), (void *)&mat_kernel.data));
- args.push_back(make_pair(localMemSize, (void *)NULL));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src_offset_x));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src_offset_y));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst_offset_x));
- args.push_back(make_pair(sizeof(cl_int), (void *)&dst_offset_y));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.cols));
- args.push_back(make_pair(sizeof(cl_int), (void *)&src.rows));
- args.push_back(make_pair(sizeof(cl_int), (void *)&cols));
- char btype[30];
- switch (borderType)
+ std::vector<float> kernelDataFloat;
+ std::vector<double> kernelDataDouble;
+ int kernel_size_y2_aligned = useDouble ?
+ _prepareKernelFilter2D<double>(kernelDataDouble, kernel)
+ : _prepareKernelFilter2D<float>(kernelDataFloat, kernel);
+ oclMat oclKernelParameter;
+ if (useDouble)
{
- case 0:
- sprintf(btype, "BORDER_CONSTANT");
+ oclKernelParameter.createEx(1, kernelDataDouble.size(), CV_64FC1, DEVICE_MEM_R_ONLY, DEVICE_MEM_DEFAULT);
+ openCLMemcpy2D(src.clCxt, oclKernelParameter.data, kernelDataDouble.size()*sizeof(double),
+ &kernelDataDouble[0], kernelDataDouble.size()*sizeof(double),
+ kernelDataDouble.size()*sizeof(double), 1, clMemcpyHostToDevice);
+ }
+ else
+ {
+ oclKernelParameter.createEx(1, kernelDataFloat.size(), CV_32FC1, DEVICE_MEM_R_ONLY, DEVICE_MEM_DEFAULT);
+ openCLMemcpy2D(src.clCxt, oclKernelParameter.data, kernelDataFloat.size()*sizeof(float),
+ &kernelDataFloat[0], kernelDataFloat.size()*sizeof(float),
+ kernelDataFloat.size()*sizeof(float), 1, clMemcpyHostToDevice);
+ }
+
+ size_t BLOCK_SIZE = src.clCxt->getDeviceInfo().maxWorkItemSizes[0];
+#if 1 // TODO Mode with several blocks requires a much more VGPRs, so this optimization is not actual for the current devices
+ size_t BLOCK_SIZE_Y = 1;
+#else
+ size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices
+ while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * src.clCxt->getDeviceInfo().maxComputeUnits * 32 < (size_t)src.rows)
+ BLOCK_SIZE_Y *= 2;
+#endif
+
+ CV_Assert((size_t)ksize.width <= BLOCK_SIZE);
+
+ bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
+
+ vector<pair<size_t , const void *> > args;
+
+ args.push_back( make_pair( sizeof(cl_mem), (void *)&src.data));
+ cl_uint stepBytes = src.step;
+ args.push_back( make_pair( sizeof(cl_uint), (void *)&stepBytes));
+ int offsetXBytes = src.offset % src.step;
+ int offsetX = offsetXBytes / src.elemSize();
+ CV_Assert((int)(offsetX * src.elemSize()) == offsetXBytes);
+ int offsetY = src.offset / src.step;
+ int endX = (offsetX + src.cols);
+ int endY = (offsetY + src.rows);
+ cl_int rect[4] = {offsetX, offsetY, endX, endY};
+ if (!isIsolatedBorder)
+ {
+ rect[2] = src.wholecols;
+ rect[3] = src.wholerows;
+ }
+ args.push_back( make_pair( sizeof(cl_int)*4, (void *)&rect[0]));
+
+ args.push_back( make_pair( sizeof(cl_mem), (void *)&dst.data));
+ cl_uint _stepBytes = dst.step;
+ args.push_back( make_pair( sizeof(cl_uint), (void *)&_stepBytes));
+ int _offsetXBytes = dst.offset % dst.step;
+ int _offsetX = _offsetXBytes / dst.elemSize();
+ CV_Assert((int)(_offsetX * dst.elemSize()) == _offsetXBytes);
+ int _offsetY = dst.offset / dst.step;
+ int _endX = (_offsetX + dst.cols);
+ int _endY = (_offsetY + dst.rows);
+ cl_int _rect[4] = {_offsetX, _offsetY, _endX, _endY};
+ args.push_back( make_pair( sizeof(cl_int)*4, (void *)&_rect[0]));
+
+ float borderValue[4] = {0, 0, 0, 0}; // DON'T move into 'if' body
+ double borderValueDouble[4] = {0, 0, 0, 0}; // DON'T move into 'if' body
+ if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT)
+ {
+ if (useDouble)
+ args.push_back( make_pair( sizeof(double) * src.oclchannels(), (void *)&borderValue[0]));
+ else
+ args.push_back( make_pair( sizeof(float) * src.oclchannels(), (void *)&borderValueDouble[0]));
+ }
+
+ args.push_back( make_pair( sizeof(cl_mem), (void *)&oclKernelParameter.data));
+
+ const char* btype = NULL;
+
+ switch (borderType & ~BORDER_ISOLATED)
+ {
+ case BORDER_CONSTANT:
+ btype = "BORDER_CONSTANT";
break;
- case 1:
- sprintf(btype, "BORDER_REPLICATE");
+ case BORDER_REPLICATE:
+ btype = "BORDER_REPLICATE";
break;
- case 2:
- sprintf(btype, "BORDER_REFLECT");
+ case BORDER_REFLECT:
+ btype = "BORDER_REFLECT";
break;
- case 3:
+ case BORDER_WRAP:
CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
return;
- case 4:
- sprintf(btype, "BORDER_REFLECT_101");
+ case BORDER_REFLECT101:
+ btype = "BORDER_REFLECT_101";
break;
}
- int type = src.depth();
- char build_options[150];
- sprintf(build_options, "-D %s -D IMG_C_%d_%d -D CN=%d -D FILTER_SIZE=%d", btype, cn, type, cn, ksize.width);
- openCLExecuteKernel(clCxt, &filtering_laplacian, kernelName, globalThreads, localThreads, args, -1, -1, build_options);
+
+ int requiredTop = anchor.y;
+ int requiredLeft = BLOCK_SIZE; // not this: anchor.x;
+ int requiredBottom = ksize.height - 1 - anchor.y;
+ int requiredRight = BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x;
+ int h = isIsolatedBorder ? src.rows : src.wholerows;
+ int w = isIsolatedBorder ? src.cols : src.wholecols;
+ bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight;
+
+ char build_options[1024];
+ sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%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",
+ (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y,
+ src.depth(), src.oclchannels(), useDouble ? 1 : 0,
+ anchor.x, anchor.y, ksize.width, ksize.height, kernel_size_y2_aligned,
+ btype,
+ extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
+ isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED");
+
+ size_t gt[3] = {divUp(dst.cols, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE, divUp(dst.rows, BLOCK_SIZE_Y), 1}, lt[3] = {BLOCK_SIZE, 1, 1};
+ openCLExecuteKernel(src.clCxt, &filtering_filter2D, "filter2D", gt, lt, args, -1, -1, build_options);
}
-Ptr<BaseFilter_GPU> cv::ocl::getLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Size &ksize,
+Ptr<BaseFilter_GPU> cv::ocl::getLinearFilter_GPU(int /*srcType*/, int /*dstType*/, const Mat &kernel, const Size &ksize,
const Point &anchor, int borderType)
{
- static const GPUFilter2D_t GPUFilter2D_callers[] = {0, GPUFilter2D, 0, GPUFilter2D, GPUFilter2D};
-
- CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC3 || srcType == CV_8UC4 || srcType == CV_32FC1 || srcType == CV_32FC3 || srcType == CV_32FC4) && dstType == srcType);
-
- oclMat gpu_krnl;
Point norm_archor = anchor;
- normalizeKernel(kernel, gpu_krnl, CV_32FC1);
normalizeAnchor(norm_archor, ksize);
- return Ptr<BaseFilter_GPU>(new LinearFilter_GPU(ksize, anchor, gpu_krnl, GPUFilter2D_callers[CV_MAT_CN(srcType)],
+ return Ptr<BaseFilter_GPU>(new LinearFilter_GPU(ksize, norm_archor, kernel, GPUFilter2D,
borderType));
}
Ptr<FilterEngine_GPU> cv::ocl::createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Point &anchor,
int borderType)
{
- Size ksize = kernel.size();
+ Size ksize = kernel.size(); // TODO remove duplicated parameter
Ptr<BaseFilter_GPU> linearFilter = getLinearFilter_GPU(srcType, dstType, kernel, ksize, anchor, borderType);
return createFilter2D_GPU(linearFilter);
}
-void cv::ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor, int borderType)
+void cv::ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor, double delta, int borderType)
{
+ CV_Assert(delta == 0);
+
if (ddepth < 0)
ddepth = src.depth();
@@ -1222,8 +1274,11 @@ void cv::ocl::Scharr(const oclMat &src, oclMat &dst, int ddepth, int dx, int dy,
sepFilter2D(src, dst, ddepth, kx, ky, Point(-1, -1), delta, bordertype);
}
-void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, double scale)
+void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, double scale,
+ double delta, int borderType)
{
+ CV_Assert(delta == 0);
+
if (!src.clCxt->supportsFeature(FEATURE_CL_DOUBLE) && src.type() == CV_64F)
{
CV_Error(CV_OpenCLDoubleNotSupported, "Selected device doesn't support double");
@@ -1232,17 +1287,17 @@ void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, d
CV_Assert(ksize == 1 || ksize == 3);
- int K[2][9] =
+ double K[2][9] =
{
{0, 1, 0, 1, -4, 1, 0, 1, 0},
{2, 0, 2, 0, -8, 0, 2, 0, 2}
};
- Mat kernel(3, 3, CV_32S, (void *)K[ksize == 3]);
+ Mat kernel(3, 3, CV_64F, (void *)K[ksize == 3 ? 1 : 0]);
if (scale != 1)
kernel *= scale;
- filter2D(src, dst, ddepth, kernel, Point(-1, -1));
+ filter2D(src, dst, ddepth, kernel, Point(-1, -1), 0, borderType);
}
////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/modules/ocl/src/opencl/filtering_filter2D.cl b/modules/ocl/src/opencl/filtering_filter2D.cl
new file mode 100644
index 000000000..f96676689
--- /dev/null
+++ b/modules/ocl/src/opencl/filtering_filter2D.cl
@@ -0,0 +1,370 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+// By downloading, copying, installing or using the software you agree to this license.
+// If you do not agree to this license, do not download, install,
+// copy or use the software.
+//
+//
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2013, Advanced Micro Devices, 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,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifdef BORDER_REPLICATE
+//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr))
+#endif
+
+#ifdef BORDER_REFLECT
+//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
+#endif
+
+#ifdef BORDER_REFLECT_101
+//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
+#endif
+
+//blur function does not support BORDER_WRAP
+#ifdef BORDER_WRAP
+//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
+#endif
+
+#ifdef EXTRA_EXTRAPOLATION // border > src image size
+#ifdef BORDER_CONSTANT
+// None
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ x = max(min(x, maxX - 1), minX); \
+ y = max(min(y, maxY - 1), minY); \
+ }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ if (x < minX) \
+ x -= ((x - maxX + 1) / maxX) * maxX; \
+ if (x >= maxX) \
+ x %= maxX; \
+ if (y < minY) \
+ y -= ((y - maxY + 1) / maxY) * maxY; \
+ if (y >= maxY) \
+ y %= maxY; \
+ }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \
+ { \
+ if (maxX - minX == 1) \
+ x = minX; \
+ else \
+ do \
+ { \
+ if (x < minX) \
+ x = -(x - minX) - 1 + delta; \
+ else \
+ x = maxX - 1 - (x - maxX) - delta; \
+ } \
+ while (x >= maxX || x < minX); \
+ \
+ if (maxY - minY == 1) \
+ y = minY; \
+ else \
+ do \
+ { \
+ if (y < minY) \
+ y = -(y - minY) - 1 + delta; \
+ else \
+ y = maxY - 1 - (y - maxY) - delta; \
+ } \
+ while (y >= maxY || y < minY); \
+ }
+#ifdef BORDER_REFLECT
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0)
+#elif defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1)
+#endif
+#else
+#error No extrapolation method
+#endif
+#else
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ int _row = y - minY, _col = x - minX; \
+ _row = ADDR_H(_row, 0, maxY - minY); \
+ _row = ADDR_B(_row, maxY - minY, _row); \
+ y = _row + minY; \
+ \
+ _col = ADDR_L(_col, 0, maxX - minX); \
+ _col = ADDR_R(_col, maxX - minX, _col); \
+ x = _col + minX; \
+ }
+#endif
+
+#if USE_DOUBLE
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#define FPTYPE double
+#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE)
+#else
+#define FPTYPE float
+#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE)
+#endif
+
+#if DATA_DEPTH == 0
+#define BASE_TYPE uchar
+#elif DATA_DEPTH == 1
+#define BASE_TYPE char
+#elif DATA_DEPTH == 2
+#define BASE_TYPE ushort
+#elif DATA_DEPTH == 3
+#define BASE_TYPE short
+#elif DATA_DEPTH == 4
+#define BASE_TYPE int
+#elif DATA_DEPTH == 5
+#define BASE_TYPE float
+#elif DATA_DEPTH == 6
+#define BASE_TYPE double
+#else
+#error data_depth
+#endif
+
+#define __CAT(x, y) x##y
+#define CAT(x, y) __CAT(x, y)
+
+#define uchar1 uchar
+#define char1 char
+#define ushort1 ushort
+#define short1 short
+#define int1 int
+#define float1 float
+#define double1 double
+
+#define convert_uchar1_sat_rte convert_uchar_sat_rte
+#define convert_char1_sat_rte convert_char_sat_rte
+#define convert_ushort1_sat_rte convert_ushort_sat_rte
+#define convert_short1_sat_rte convert_short_sat_rte
+#define convert_int1_sat_rte convert_int_sat_rte
+#define convert_float1
+#define convert_double1
+
+#if DATA_DEPTH == 5 || DATA_DEPTH == 6
+#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE)
+#else
+#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte)
+#endif
+
+#define VEC_SIZE DATA_CHAN
+
+#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE)
+#define TYPE VEC_TYPE
+
+#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE)
+
+#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE)
+
+struct RectCoords
+{
+ int x1, y1, x2, y2;
+};
+
+//#define DEBUG
+#ifdef DEBUG
+#define DEBUG_ONLY(x) x
+#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0)
+#else
+#define DEBUG_ONLY(x) (void)0
+#define ASSERT(condition) (void)0
+#endif
+
+
+inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords
+#ifdef BORDER_CONSTANT
+ , SCALAR_TYPE borderValue
+#endif
+ )
+{
+#ifdef BORDER_ISOLATED
+ if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+#else
+ if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+#endif
+ {
+ __global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes);
+ return CONVERT_TO_FPTYPE(*ptr);
+ }
+ else
+ {
+#ifdef BORDER_CONSTANT
+ return borderValue;
+#else
+ int selected_col = pos.x;
+ int selected_row = pos.y;
+
+ EXTRAPOLATE(selected_col, selected_row,
+#ifdef BORDER_ISOLATED
+ srcCoords.x1, srcCoords.y1,
+#else
+ 0, 0,
+#endif
+ srcCoords.x2, srcCoords.y2
+ );
+
+ // debug border mapping
+ //printf("pos=%d,%d --> %d, %d\n", pos.x, pos.y, selected_col, selected_row);
+
+ pos = (int2)(selected_col, selected_row);
+ if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+ {
+ __global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes);
+ return CONVERT_TO_FPTYPE(*ptr);
+ }
+ else
+ {
+ // for debug only
+ DEBUG_ONLY(printf("BUG in boxFilter kernel\n"));
+ return (FPTYPE)(0.0f);
+ }
+#endif
+ }
+}
+
+// INPUT PARAMETER: BLOCK_SIZE_Y (via defines)
+
+__kernel
+__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1)))
+void filter2D(__global TYPE *src, const unsigned int srcStepBytes, const int4 srcRC,
+ __global TYPE *dst, const unsigned int dstStepBytes, const int4 dstRC,
+#ifdef BORDER_CONSTANT
+ SCALAR_TYPE borderValue,
+#endif
+ __constant FPTYPE* kernelData // transposed: [KERNEL_SIZE_X][KERNEL_SIZE_Y2_ALIGNED]
+ )
+{
+ const struct RectCoords srcCoords = {srcRC.s0, srcRC.s1, srcRC.s2, srcRC.s3}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY
+ struct RectCoords dstCoords = {dstRC.s0, dstRC.s1, dstRC.s2, dstRC.s3};
+
+ const int local_id = get_local_id(0);
+ const int x = local_id + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X;
+ const int y = get_global_id(1) * BLOCK_SIZE_Y;
+
+ INTERMEDIATE_TYPE data[KERNEL_SIZE_Y];
+ __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE];
+
+ int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y);
+
+ int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y);
+ __global TYPE* dstPtr = (__global TYPE*)((__global char*)dst + pos.x * sizeof(TYPE) + pos.y * dstStepBytes); // Pointer can be out of bounds!
+ bool writeResult = (local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) &&
+ pos.x >= dstCoords.x1 && pos.x < dstCoords.x2);
+
+#if BLOCK_SIZE_Y > 1
+ bool readAllpixels = true;
+ int sy_index = 0; // current index in data[] array
+
+ dstCoords.y2 = min(dstCoords.y2, pos.y + BLOCK_SIZE_Y);
+ for (;
+ pos.y < dstCoords.y2;
+ pos.y++,
+ dstPtr = (__global TYPE*)((__global char*)dstPtr + dstStepBytes))
+#endif
+ {
+ ASSERT(pos.y < dstCoords.y2);
+
+ for (
+#if BLOCK_SIZE_Y > 1
+ int sy = readAllpixels ? 0 : -1; sy < (readAllpixels ? KERNEL_SIZE_Y : 0);
+#else
+ int sy = 0, sy_index = 0; sy < KERNEL_SIZE_Y;
+#endif
+ sy++, srcPos.y++)
+ {
+ data[sy + sy_index] = readSrcPixel(srcPos, src, srcStepBytes, srcCoords
+#ifdef BORDER_CONSTANT
+ , borderValue
+#endif
+ );
+ }
+
+ INTERMEDIATE_TYPE total_sum = 0;
+ for (int sx = 0; sx < KERNEL_SIZE_X; sx++)
+ {
+ {
+ __constant FPTYPE* k = &kernelData[KERNEL_SIZE_Y2_ALIGNED * sx
+#if BLOCK_SIZE_Y > 1
+ + KERNEL_SIZE_Y - sy_index
+#endif
+ ];
+ INTERMEDIATE_TYPE tmp_sum = 0;
+ for (int sy = 0; sy < KERNEL_SIZE_Y; sy++)
+ {
+ tmp_sum += data[sy] * k[sy];
+ }
+
+ sumOfCols[local_id] = tmp_sum;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ int id = local_id + sx - ANCHOR_X;
+ if (id >= 0 && id < LOCAL_SIZE)
+ total_sum += sumOfCols[id];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ if (writeResult)
+ {
+ ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2);
+ *dstPtr = CONVERT_TO_TYPE(total_sum);
+ }
+
+#if BLOCK_SIZE_Y > 1
+ readAllpixels = false;
+#if BLOCK_SIZE_Y > KERNEL_SIZE_Y
+ sy_index = (sy_index + 1 <= KERNEL_SIZE_Y) ? sy_index + 1 : 1;
+#else
+ sy_index++;
+#endif
+#endif // BLOCK_SIZE_Y == 1
+ }
+}
diff --git a/modules/ocl/src/opencl/filtering_laplacian.cl b/modules/ocl/src/opencl/filtering_laplacian.cl
deleted file mode 100644
index ea22967df..000000000
--- a/modules/ocl/src/opencl/filtering_laplacian.cl
+++ /dev/null
@@ -1,381 +0,0 @@
-/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-// By downloading, copying, installing or using the software you agree to this license.
-// If you do not agree to this license, do not download, install,
-// copy or use the software.
-//
-//
-// License Agreement
-// For Open Source Computer Vision Library
-//
-// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
-// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// @Authors
-// Pang Erping, erping@multicorewareinc.com
-// Jia Haipeng, jiahaipeng95@gmail.com
-// Peng Xiao, pengxiao@outlook.com
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-// * Redistribution's of source code must retain the above copyright notice,
-// this list of conditions and the following disclaimer.
-//
-// * Redistribution's in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// * The name of the copyright holders may not be used to endorse or promote products
-// derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors as is and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
-//M*/
-
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////Macro for border type////////////////////////////////////////////
-/////////////////////////////////////////////////////////////////////////////////////////////////
-#ifdef BORDER_REPLICATE
-
-//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
-#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i))
-#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) : (i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr))
-#endif
-
-#ifdef BORDER_REFLECT
-#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? ((l_edge)<<1)-(i)-1 : (i))
-#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? ((t_edge)<<1)-(i)-1 : (i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
-#endif
-
-#ifdef BORDER_REFLECT_101
-//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
-#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? ((l_edge)<<1)-(i) : (i))
-#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
-#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? ((t_edge)<<1)-(i) : (i))
-#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
-#endif
-
-#ifdef IMG_C_1_0
-#define T_IMG uchar
-#define T_IMGx4 uchar4
-#define T_IMG_C1 uchar
-#define CONVERT_TYPE convert_uchar_sat
-#define CONVERT_TYPEx4 convert_uchar4_sat
-#endif
-#ifdef IMG_C_4_0
-#define T_IMG uchar4
-#define T_IMGx4 uchar16
-#define T_IMG_C1 uchar
-#define CONVERT_TYPE convert_uchar4_sat
-#define CONVERT_TYPEx4 convert_uchar16_sat
-#endif
-#ifdef IMG_C_1_5
-#define T_IMG float
-#define T_IMGx4 float4
-#define T_IMG_C1 float
-#define CONVERT_TYPE convert_float
-#define CONVERT_TYPEx4 convert_float4
-#endif
-#ifdef IMG_C_4_5
-#define T_IMG float4
-#define T_IMGx4 float16
-#define T_IMG_C1 float
-#define CONVERT_TYPE convert_float4
-#define CONVERT_TYPEx4 convert_float16
-#endif
-
-#ifndef CN
-#define CN 1
-#endif
-
-#if CN == 1
-#define T_SUM float
-#define T_SUMx4 float4
-#define CONVERT_TYPE_SUM convert_float
-#define CONVERT_TYPE_SUMx4 convert_float4
-#define SUM_ZERO (0.0f)
-#define SUM_ZEROx4 (0.0f, 0.0f, 0.0f, 0.0f)
-#define VLOAD4 vload4
-#define SX x
-#define SY y
-#define SZ z
-#define SW w
-#elif CN == 4
-#define T_SUM float4
-#define T_SUMx4 float16
-#define CONVERT_TYPE_SUM convert_float4
-#define CONVERT_TYPE_SUMx4 convert_float16
-#define SUM_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
-#define SUM_ZEROx4 (0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)
-#define VLOAD4 vload16
-#define SX s0123
-#define SY s4567
-#define SZ s89ab
-#define SW scdef
-#endif
-
-#ifndef FILTER_SIZE
-#define FILTER_SIZE 3
-#endif
-
-#define LOCAL_GROUP_SIZE 16
-
-#define LOCAL_WIDTH ((FILTER_SIZE/2)*2 + LOCAL_GROUP_SIZE)
-#define LOCAL_HEIGHT ((FILTER_SIZE/2)*2 + LOCAL_GROUP_SIZE)
-
-#define FILTER_RADIUS (FILTER_SIZE >> 1)
-
-__kernel void filter2D(
- __global T_IMG *src,
- __global T_IMG *dst,
- int src_step,
- int dst_step,
- __constant float *mat_kernel,
- __local T_IMG *local_data,
- int wholerows,
- int wholecols,
- int src_offset_x,
- int src_offset_y,
- int dst_offset_x,
- int dst_offset_y,
- int cols,
- int rows,
- int operate_cols
-)
-{
- int groupStartCol = get_group_id(0) * get_local_size(0);
- int groupStartRow = get_group_id(1) * get_local_size(1);
-
- int localCol = get_local_id(0);
- int localRow = get_local_id(1);
- int globalCol = groupStartCol + localCol;
- int globalRow = groupStartRow + localRow;
- const int src_offset = mad24(src_offset_y, src_step, src_offset_x);
- const int dst_offset = mad24(dst_offset_y, dst_step, dst_offset_x);
-
-#ifdef BORDER_CONSTANT
- for(int i = localRow; i < LOCAL_HEIGHT; i += get_local_size(1))
- {
- int curRow = groupStartRow + i;
- for(int j = localCol; j < LOCAL_WIDTH; j += get_local_size(0))
- {
- int curCol = groupStartCol + j;
- if(curRow < FILTER_RADIUS - src_offset_y || (curRow - FILTER_RADIUS) >= wholerows - src_offset_y||
- curCol < FILTER_RADIUS - src_offset_x || (curCol - FILTER_RADIUS) >= wholecols - src_offset_x)
- {
- local_data[(i) * LOCAL_WIDTH + j] = 0;
- }
- else
- {
- local_data[(i) * LOCAL_WIDTH + j] = src[(curRow - FILTER_RADIUS) * src_step + curCol - FILTER_RADIUS + src_offset];
- }
- }
- }
-#else
- for(int i = localRow; i < LOCAL_HEIGHT; i += get_local_size(1))
- {
- int curRow = groupStartRow + i;
-
- curRow = ADDR_H(curRow, FILTER_RADIUS - src_offset_y, wholerows - src_offset_y);
-
- curRow = ADDR_B(curRow - FILTER_RADIUS, wholerows - src_offset_y, curRow - FILTER_RADIUS);
-
- for(int j = localCol; j < LOCAL_WIDTH; j += get_local_size(0))
- {
- int curCol = groupStartCol + j;
- curCol = ADDR_L(curCol, FILTER_RADIUS - src_offset_x, wholecols - src_offset_x);
- curCol = ADDR_R(curCol - FILTER_RADIUS, wholecols - src_offset_x, curCol - FILTER_RADIUS);
- if(curRow < wholerows && curCol < wholecols)
- {
- local_data[(i) * LOCAL_WIDTH + j] = src[(curRow) * src_step + curCol + src_offset];
- }
- }
- }
-#endif
-
- barrier(CLK_LOCAL_MEM_FENCE);
- if(globalRow < rows && globalCol < cols)
- {
- T_SUM sum = (T_SUM)(SUM_ZERO);
- int filterIdx = 0;
- for(int i = 0; i < FILTER_SIZE; i++)
- {
- int offset = (i + localRow) * LOCAL_WIDTH;
-
- for(int j = 0; j < FILTER_SIZE; j++)
- {
- sum += CONVERT_TYPE_SUM(local_data[offset + j + localCol]) * mat_kernel[filterIdx++];
- }
- }
- dst[(globalRow)*dst_step + (globalCol) + dst_offset] = CONVERT_TYPE(sum);
- }
-}
-
-/// following is specific for 3x3 kernels
-
-//////////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////Macro for define elements number per thread/////////////////////////////
-////////////////////////////////////////////////////////////////////////////////////////////////////
-
-#define ANX 1
-#define ANY 1
-
-#define ROWS_PER_GROUP 4
-#define ROWS_PER_GROUP_BITS 2
-#define ROWS_FETCH (ROWS_PER_GROUP + ANY + ANY) //(ROWS_PER_GROUP + anY * 2)
-
-#define THREADS_PER_ROW 64
-#define THREADS_PER_ROW_BIT 6
-
-#define ELEMENTS_PER_THREAD 4
-#define ELEMENTS_PER_THREAD_BIT 2
-
-#define LOCAL_MEM_STEP 260 //divup((get_local_size(0) + anX * 2), 4) * 4
-
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////////////8uC1////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////////////////////////
-
-__kernel void filter2D_3x3(
- __global T_IMG *src,
- __global T_IMG *dst,
- int src_step,
- int dst_step,
- __constant float *mat_kernel,
- __local T_IMG *local_data,
- int wholerows,
- int wholecols,
- int src_offset_x,
- int src_offset_y,
- int dst_offset_x,
- int dst_offset_y,
- int cols,
- int rows,
- int operate_cols
-)
-{
- int gX = get_global_id(0);
- int gY = get_global_id(1);
-
- int lX = get_local_id(0);
-
- int groupX_size = get_local_size(0);
- int groupX_id = get_group_id(0);
-
-#define dst_align (dst_offset_x & 3)
- int cols_start_index_group = src_offset_x - dst_align + groupX_size * groupX_id - ANX;
- int rows_start_index = src_offset_y + (gY << ROWS_PER_GROUP_BITS) - ANY;
-
- if((gY << 2) < rows)
- {
- for(int i = 0; i < ROWS_FETCH; ++i)
- {
- if((rows_start_index - src_offset_y) + i < rows + ANY)
- {
-#ifdef BORDER_CONSTANT
- int selected_row = rows_start_index + i;
- int selected_cols = cols_start_index_group + lX;
-
- T_IMG data = src[mad24(selected_row, src_step, selected_cols)];
- int con = selected_row >= 0 && selected_row < wholerows && selected_cols >= 0 && selected_cols < wholecols;
- data = con ? data : (T_IMG)(0);
- local_data[mad24(i, LOCAL_MEM_STEP, lX)] = data;
-
- if(lX < (ANX << 1))
- {
- selected_cols = cols_start_index_group + lX + groupX_size;
-
- data = src[mad24(selected_row, src_step, selected_cols)];
- con = selected_row >= 0 && selected_row < wholerows && selected_cols >= 0 && selected_cols < wholecols;
- data = con ? data : (T_IMG)(0);
- local_data[mad24(i, LOCAL_MEM_STEP, lX) + groupX_size] = data;
- }
-#else
- int selected_row = ADDR_H(rows_start_index + i, 0, wholerows);
- selected_row = ADDR_B(rows_start_index + i, wholerows, selected_row);
-
- int selected_cols = ADDR_L(cols_start_index_group + lX, 0, wholecols);
- selected_cols = ADDR_R(cols_start_index_group + lX, wholecols, selected_cols);
-
- T_IMG data = src[mad24(selected_row, src_step, selected_cols)];
-
- local_data[mad24(i, LOCAL_MEM_STEP, lX)] = data;
-
- if(lX < (ANX << 1))
- {
- selected_cols = cols_start_index_group + lX + groupX_size;
- selected_cols = ADDR_R(selected_cols, wholecols, selected_cols);
-
- data = src[mad24(selected_row, src_step, selected_cols)];
- local_data[mad24(i, LOCAL_MEM_STEP, lX) + groupX_size] = data;
- }
-#endif
- }
- }
- }
- barrier(CLK_LOCAL_MEM_FENCE);
-
- int process_col = groupX_size * groupX_id + ((lX % THREADS_PER_ROW) << 2);
- if(((gY << 2) < rows) && (process_col < operate_cols))
- {
- int dst_cols_start = dst_offset_x;
- int dst_cols_end = dst_offset_x + cols;
- int dst_cols_index = (dst_offset_x + process_col) & 0xfffffffc;
-
- int dst_rows_end = dst_offset_y + rows;
- int dst_rows_index = dst_offset_y + (gY << ROWS_PER_GROUP_BITS) + (lX >> THREADS_PER_ROW_BIT);
- dst = dst + mad24(dst_rows_index, dst_step, dst_cols_index);
-
- T_IMGx4 dst_data = *(__global T_IMGx4 *)dst;
-
- T_SUMx4 sum = (T_SUMx4)SUM_ZEROx4;
- T_IMGx4 data;
-
- for(int i = 0; i < FILTER_SIZE; i++)
- {
-#pragma unroll
- for(int j = 0; j < FILTER_SIZE; j++)
- {
- if(dst_rows_index < dst_rows_end)
- {
- int local_row = (lX >> THREADS_PER_ROW_BIT) + i;
- int local_cols = ((lX % THREADS_PER_ROW) << ELEMENTS_PER_THREAD_BIT) + j;
-
- data = VLOAD4(0, (__local T_IMG_C1 *)(local_data + local_row * LOCAL_MEM_STEP + local_cols));
- sum = sum + (mat_kernel[i * FILTER_SIZE + j] * CONVERT_TYPE_SUMx4(data));
- }
- }
- }
-
- if(dst_rows_index < dst_rows_end)
- {
- T_IMGx4 tmp_dst = CONVERT_TYPEx4(sum);
- tmp_dst.SX = ((dst_cols_index + 0 >= dst_cols_start) && (dst_cols_index + 0 < dst_cols_end)) ?
- tmp_dst.SX : dst_data.SX;
- tmp_dst.SY = ((dst_cols_index + 1 >= dst_cols_start) && (dst_cols_index + 1 < dst_cols_end)) ?
- tmp_dst.SY : dst_data.SY;
- tmp_dst.SZ = ((dst_cols_index + 2 >= dst_cols_start) && (dst_cols_index + 2 < dst_cols_end)) ?
- tmp_dst.SZ : dst_data.SZ;
- tmp_dst.SW = ((dst_cols_index + 3 >= dst_cols_start) && (dst_cols_index + 3 < dst_cols_end)) ?
- tmp_dst.SW : dst_data.SW;
- *(__global T_IMGx4 *)dst = tmp_dst;
- }
- }
-}
diff --git a/modules/ocl/test/test_filters.cpp b/modules/ocl/test/test_filters.cpp
index 3cf7d37b8..a8583b28a 100644
--- a/modules/ocl/test/test_filters.cpp
+++ b/modules/ocl/test/test_filters.cpp
@@ -160,8 +160,8 @@ OCL_TEST_P(LaplacianTest, Accuracy)
{
random_roi();
- Laplacian(src_roi, dst_roi, -1, ksize, scale); // TODO FIXIT , 0, borderType);
- ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale); // TODO FIXIT , 0, borderType);
+ Laplacian(src_roi, dst_roi, -1, ksize, scale, 0, borderType);
+ ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale, 0, borderType);
Near();
}
@@ -298,7 +298,7 @@ OCL_TEST_P(Filter2D, Mat)
kernel *= 1.0 / (double)(ksize * ksize);
cv::filter2D(src_roi, dst_roi, -1, kernel, anchor, 0.0, borderType);
- ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, anchor, /* TODO FIXIT 0.0,*/ borderType);
+ ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, anchor, 0.0, borderType);
Near();
}
From 98f73705dfed3ee4a5f003ce4c284a8c1be25646 Mon Sep 17 00:00:00 2001
From: Alexander Alekhin <alexander.alekhin@itseez.com>
Date: Sat, 26 Oct 2013 11:15:53 +0400
Subject: [PATCH 4/6] ocl: fix morph filters
---
modules/ocl/src/filtering.cpp | 16 +++++-----------
1 file changed, 5 insertions(+), 11 deletions(-)
diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp
index e1255197f..954d1d5aa 100644
--- a/modules/ocl/src/filtering.cpp
+++ b/modules/ocl/src/filtering.cpp
@@ -321,28 +321,22 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
openCLExecuteKernel(clCxt, &filtering_morph, kernelName, globalThreads, localThreads, args, -1, -1, compile_option);
}
-Ptr<BaseFilter_GPU> cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat &kernel, const Size &ksize, Point anchor)
+Ptr<BaseFilter_GPU> cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat &_kernel, const Size &ksize, Point anchor)
{
- static const GPUMorfFilter_t GPUMorfFilter_callers[2][5] =
- {
- {0, GPUErode, 0, GPUErode, GPUErode },
- {0, GPUDilate, 0, GPUDilate, GPUDilate}
- };
-
CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE);
CV_Assert(type == CV_8UC1 || type == CV_8UC3 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC3 || type == CV_32FC4);
normalizeAnchor(anchor, ksize);
Mat kernel8U;
- kernel.convertTo(kernel8U, CV_8U);
- Mat cont_krnl = kernel8U.reshape(1, 1);
+ _kernel.convertTo(kernel8U, CV_8U);
+ Mat kernel = kernel8U.reshape(1, 1);
bool noZero = true;
for(int i = 0; i < kernel.rows * kernel.cols; ++i)
- if(kernel.data[i] != 1)
+ if(kernel.at<uchar>(i) != 1)
noZero = false;
- MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, cont_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]);
+ MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, kernel, op == MORPH_ERODE ? GPUErode : GPUDilate);
if(noZero)
mfgpu->rectKernel = true;
From b10e1e5c7e143b54a2b93fbc170cd3715abe1e7e Mon Sep 17 00:00:00 2001
From: Alexander Alekhin <alexander.alekhin@itseez.com>
Date: Mon, 28 Oct 2013 18:22:31 +0400
Subject: [PATCH 5/6] ocl: filters: update documentation
---
modules/ocl/doc/image_filtering.rst | 40 ++++++++++++++---------------
1 file changed, 19 insertions(+), 21 deletions(-)
diff --git a/modules/ocl/doc/image_filtering.rst b/modules/ocl/doc/image_filtering.rst
index bd929b988..cbec29b11 100644
--- a/modules/ocl/doc/image_filtering.rst
+++ b/modules/ocl/doc/image_filtering.rst
@@ -133,7 +133,7 @@ Creates a normalized 2D box filter.
.. ocv:function:: Ptr<BaseFilter_GPU> ocl::getBoxFilter_GPU(int srcType, int dstType, const Size &ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT)
- :param srcType: Input image type supporting ``CV_8UC1`` and ``CV_8UC4`` .
+ :param srcType: Input image type.
:param dstType: Output image type. It supports only the same values as the source type.
@@ -141,9 +141,7 @@ Creates a normalized 2D box filter.
:param anchor: Anchor point. The default value ``Point(-1, -1)`` means that the anchor is at the kernel center.
- :param borderType: Supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP.
-
-.. note:: This filter does not check out-of-border accesses, so only a proper sub-matrix of a bigger matrix has to be passed to it.
+ :param borderType: Border type.
.. seealso:: :ocv:func:`boxFilter`
@@ -153,21 +151,19 @@ Smooths the image using the normalized box filter.
.. ocv:function:: void ocl::boxFilter(const oclMat &src, oclMat &dst, int ddepth, Size ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT)
- :param src: Input image. ``CV_8UC1`` and ``CV_8UC4`` source types are supported.
+ :param src: Input image.
:param dst: Output image type. The size and type is the same as ``src`` .
- :param ddepth: Output image depth. If -1, the output image has the same depth as the input one. The only values allowed here are ``CV_8U`` and -1.
+ :param ddepth: Desired depth of the destination image. If it is negative, it is the same as ``src.depth()`` . It supports only the same depth as the source image depth.
:param ksize: Kernel size.
:param anchor: Anchor point. The default value ``Point(-1, -1)`` means that the anchor is at the kernel center.
- :param borderType: Supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP.
+ :param borderType: Border type.
-Smoothes image using box filter.Supports data type: CV_8UC1, CV_8UC4, CV_32FC1 and CV_32FC4.
-
-.. note:: This filter does not check out-of-border accesses, so only a proper sub-matrix of a bigger matrix has to be passed to it.
+Smoothes image using box filter.
ocl::blur
-------------
@@ -175,7 +171,7 @@ Acts as a synonym for the normalized box filter.
.. ocv:function:: void ocl::blur(const oclMat &src, oclMat &dst, Size ksize, Point anchor = Point(-1, -1), int borderType = BORDER_CONSTANT)
- :param src: Input image. ``CV_8UC1`` and ``CV_8UC4`` source types are supported.
+ :param src: Input image.
:param dst: Output image type with the same size and type as ``src`` .
@@ -183,9 +179,7 @@ Acts as a synonym for the normalized box filter.
:param anchor: Anchor point. The default value Point(-1, -1) means that the anchor is at the kernel center.
- :param borderType: Supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP.
-
-.. note:: This filter does not check out-of-border accesses, so only a proper sub-matrix of a bigger matrix has to be passed to it.
+ :param borderType: Border type.
.. seealso:: :ocv:func:`blur`, :ocv:func:`ocl::boxFilter`
@@ -217,11 +211,11 @@ Creates a non-separable linear filter.
.. ocv:function:: Ptr<FilterEngine_GPU> ocl::createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT)
- :param srcType: Input image type. Supports ``CV_8U`` , ``CV_16U`` and ``CV_32F`` one and four channel image.
+ :param srcType: Input image type..
:param dstType: Output image type. The same type as ``src`` is supported.
- :param kernel: 2D array of filter coefficients. Floating-point coefficients will be converted to fixed-point representation before the actual processing. Supports size up to 16. For larger kernels use :ocv:func:`ocl::convolve`.
+ :param kernel: 2D array of filter coefficients.
:param anchor: Anchor point. The default value Point(-1, -1) means that the anchor is at the kernel center.
@@ -234,9 +228,9 @@ ocl::filter2D
-----------------
Applies the non-separable 2D linear filter to an image.
-.. ocv:function:: void ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT)
+.. ocv:function:: void ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor = Point(-1, -1), double delta = 0.0, int borderType = BORDER_DEFAULT)
- :param src: Source image. Supports ``CV_8U`` , ``CV_16U`` and ``CV_32F`` one and four channel image.
+ :param src: Source image.
:param dst: Destination image. The size and the number of channels is the same as ``src`` .
@@ -246,9 +240,9 @@ Applies the non-separable 2D linear filter to an image.
:param anchor: Anchor of the kernel that indicates the relative position of a filtered point within the kernel. The anchor resides within the kernel. The special default value (-1,-1) means that the anchor is at the kernel center.
- :param borderType: Pixel extrapolation method. For details, see :ocv:func:`borderInterpolate` .
+ :param delta: optional value added to the filtered pixels before storing them in ``dst``. Value '0' is supported only.
- :param stream: Stream for the asynchronous version.
+ :param borderType: Pixel extrapolation method. For details, see :ocv:func:`borderInterpolate` .
ocl::getLinearRowFilter_GPU
-------------------------------
@@ -447,7 +441,7 @@ ocl::Laplacian
------------------
Returns void
-.. ocv:function:: void ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1)
+.. ocv:function:: void ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1, double delta = 0, int borderType = BORDER_DEFAULT)
:param src: The source image
@@ -459,6 +453,10 @@ Returns void
:param scale: The optional scale factor for the computed Laplacian values (by default, no scaling is applied
+ :param delta: Optional delta value that is added to the results prior to storing them in ``dst`` . Supported value is 0 only.
+
+ :param bordertype: Pixel extrapolation method.
+
The function calculates the Laplacian of the source image by adding up the second x and y derivatives calculated using the Sobel operator.
ocl::convolve
From 58be2546ca6ba3e00ab7f008069f70ff984ba3bd Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Mon, 28 Oct 2013 16:46:41 +0400
Subject: [PATCH 6/6] fixed OpenCL morph operations for case when kernel does
not have zero element
---
modules/ocl/src/filtering.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp
index 954d1d5aa..4a04e2de8 100644
--- a/modules/ocl/src/filtering.cpp
+++ b/modules/ocl/src/filtering.cpp
@@ -416,8 +416,8 @@ void morphOp(int op, const oclMat &src, oclMat &dst, const Mat &_kernel, Point a
else if (iterations > 1 && countNonZero(_kernel) == _kernel.rows * _kernel.cols)
{
anchor = Point(anchor.x * iterations, anchor.y * iterations);
- kernel = getStructuringElement(MORPH_RECT, Size(ksize.width + iterations * (ksize.width - 1),
- ksize.height + iterations * (ksize.height - 1)), anchor);
+ kernel = getStructuringElement(MORPH_RECT, Size(ksize.width + (iterations - 1) * (ksize.width - 1),
+ ksize.height + (iterations - 1) * (ksize.height - 1)), anchor);
iterations = 1;
}
else