diff --git a/modules/imgproc/src/opencl/bilateral.cl b/modules/imgproc/src/opencl/bilateral.cl new file mode 100644 index 0000000000..f459cfc850 --- /dev/null +++ b/modules/imgproc/src/opencl/bilateral.cl @@ -0,0 +1,59 @@ +// 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 +// Rock Li, Rock.li@amd.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. + +__kernel void bilateral(__global const uchar * src, int src_step, int src_offset, + __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, + __constant float * color_weight, __constant float * space_weight, __constant int * space_ofs) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (y < dst_rows && x < dst_cols) + { + int src_index = mad24(y + radius, src_step, x + radius + src_offset); + int dst_index = mad24(y, dst_step, x + dst_offset); + float sum = 0.f, wsum = 0.f; + int val0 = convert_int(src[src_index]); + + #pragma unroll + for (int k = 0; k < maxk; k++ ) + { + int val = convert_int(src[src_index + space_ofs[k]]); + float w = space_weight[k] * color_weight[abs(val - val0)]; + sum += (float)(val) * w; + wsum += w; + } + dst[dst_index] = convert_uchar_rtz(sum / wsum + 0.5f); + } +} diff --git a/modules/imgproc/src/smooth.cpp b/modules/imgproc/src/smooth.cpp index 2e2eabf4e1..7ef850863f 100644 --- a/modules/imgproc/src/smooth.cpp +++ b/modules/imgproc/src/smooth.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" /* * This file includes the code, contributed by Simon Perreault @@ -1914,19 +1915,91 @@ private: }; #endif +static bool ocl_bilateralFilter_8u(InputArray _src, OutputArray _dst, int d, + double sigma_color, double sigma_space, + int borderType) +{ + int type = _src.type(), cn = CV_MAT_CN(type); + int i, j, maxk, radius; + + if ( type != CV_8UC1 ) + return false; + + if (sigma_color <= 0) + sigma_color = 1; + if (sigma_space <= 0) + sigma_space = 1; + + double gauss_color_coeff = -0.5 / (sigma_color * sigma_color); + double gauss_space_coeff = -0.5 / (sigma_space * sigma_space); + + if ( d <= 0 ) + radius = cvRound(sigma_space * 1.5); + else + radius = d / 2; + radius = MAX(radius, 1); + d = radius * 2 + 1; + + UMat src = _src.getUMat(), dst = _dst.getUMat(), temp; + if (src.u == dst.u) + return false; + + copyMakeBorder(src, temp, radius, radius, radius, radius, borderType); + + std::vector _color_weight(cn * 256); + std::vector _space_weight(d * d); + std::vector _space_ofs(d * d); + float *color_weight = &_color_weight[0]; + float *space_weight = &_space_weight[0]; + int *space_ofs = &_space_ofs[0]; + + // initialize color-related bilateral filter coefficients + for( i = 0; i < 256 * cn; i++ ) + color_weight[i] = (float)std::exp(i * i * gauss_color_coeff); + + // initialize space-related bilateral filter coefficients + for( i = -radius, maxk = 0; i <= radius; i++ ) + for( j = -radius; j <= radius; j++ ) + { + double r = std::sqrt((double)i * i + (double)j * j); + if ( r > radius ) + continue; + space_weight[maxk] = (float)std::exp(r * r * gauss_space_coeff); + space_ofs[maxk++] = (int)(i * temp.step + j); + } + + ocl::Kernel k("bilateral", ocl::imgproc::bilateral_oclsrc, + format("-D radius=%d -D maxk=%d", radius, maxk)); + if (k.empty()) + return false; + + Mat mcolor_weight(1, cn * 256, CV_32FC1, color_weight); + Mat mspace_weight(1, d * d, CV_32FC1, space_weight); + Mat mspace_ofs(1, d * d, CV_32SC1, space_ofs); + UMat ucolor_weight, uspace_weight, uspace_ofs; + mcolor_weight.copyTo(ucolor_weight); + mspace_weight.copyTo(uspace_weight); + mspace_ofs.copyTo(uspace_ofs); + + k.args(ocl::KernelArg::ReadOnlyNoSize(temp), ocl::KernelArg::WriteOnly(dst), + ocl::KernelArg::PtrReadOnly(ucolor_weight), + ocl::KernelArg::PtrReadOnly(uspace_weight), + ocl::KernelArg::PtrReadOnly(uspace_ofs)); + + size_t globalsize[2] = { dst.cols, dst.rows }; + return k.run(2, globalsize, NULL, false); +} + static void bilateralFilter_8u( const Mat& src, Mat& dst, int d, double sigma_color, double sigma_space, int borderType ) { - int cn = src.channels(); int i, j, maxk, radius; Size size = src.size(); - CV_Assert( (src.type() == CV_8UC1 || src.type() == CV_8UC3) && - src.type() == dst.type() && src.size() == dst.size() && - src.data != dst.data ); + CV_Assert( (src.type() == CV_8UC1 || src.type() == CV_8UC3) && src.data != dst.data ); if( sigma_color <= 0 ) sigma_color = 1; @@ -1973,7 +2046,7 @@ bilateralFilter_8u( const Mat& src, Mat& dst, int d, { j = -radius; - for( ;j <= radius; j++ ) + for( ; j <= radius; j++ ) { double r = std::sqrt((double)i*i + (double)j*j); if( r > radius ) @@ -2184,9 +2257,7 @@ bilateralFilter_32f( const Mat& src, Mat& dst, int d, float len, scale_index; Size size = src.size(); - CV_Assert( (src.type() == CV_32FC1 || src.type() == CV_32FC3) && - src.type() == dst.type() && src.size() == dst.size() && - src.data != dst.data ); + CV_Assert( (src.type() == CV_32FC1 || src.type() == CV_32FC3) && src.data != dst.data ); if( sigma_color <= 0 ) sigma_color = 1; @@ -2267,9 +2338,13 @@ void cv::bilateralFilter( InputArray _src, OutputArray _dst, int d, double sigmaColor, double sigmaSpace, int borderType ) { - Mat src = _src.getMat(); - _dst.create( src.size(), src.type() ); - Mat dst = _dst.getMat(); + _dst.create( _src.size(), _src.type() ); + + if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() && + ocl_bilateralFilter_8u(_src, _dst, d, sigmaColor, sigmaSpace, borderType)) + return; + + Mat src = _src.getMat(), dst = _dst.getMat(); if( src.depth() == CV_8U ) bilateralFilter_8u( src, dst, d, sigmaColor, sigmaSpace, borderType ); diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp new file mode 100644 index 0000000000..58d23e8358 --- /dev/null +++ b/modules/imgproc/test/ocl/test_filters.cpp @@ -0,0 +1,176 @@ +/*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. +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Niko Li, newlife20080214@gmail.com +// Jia Haipeng, jiahaipeng95@gmail.com +// Zero Lin, Zero.Lin@amd.com +// Zhang Ying, zhangying913@gmail.com +// Yao Wang, bitwangyaoyao@gmail.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*/ + +#include "test_precomp.hpp" +#include "cvconfig.h" +#include "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +PARAM_TEST_CASE(FilterTestBase, MatType, + int, // kernel size + Size, // dx, dy + int, // border type + double, // optional parameter + bool) // roi or not +{ + int type, borderType, ksize; + Size size; + double param; + bool useRoi; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + type = GET_PARAM(0); + ksize = GET_PARAM(1); + size = GET_PARAM(2); + borderType = GET_PARAM(3); + param = GET_PARAM(4); + useRoi = GET_PARAM(5); + } + + void random_roi(int minSize = 1) + { + 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); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -60, 70); + + UMAT_UPLOAD_INPUT_PARAMETER(src) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } + + void Near() + { + int depth = CV_MAT_DEPTH(type); + bool isFP = depth >= CV_32F; + + if (isFP) + Near(1e-6, true); + else + Near(1, false); + } + + void Near(double threshold, bool relative) + { + if (relative) + { + EXPECT_MAT_NEAR_RELATIVE(dst, udst, threshold); + EXPECT_MAT_NEAR_RELATIVE(dst_roi, udst_roi, threshold); + } + else + { + EXPECT_MAT_NEAR(dst, udst, threshold); + EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold); + } + } +}; + +//////////////////////////////////////////////////////////////////////////////////////////////////// +// Bilateral + +typedef FilterTestBase Bilateral; + +OCL_TEST_P(Bilateral, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + double sigmacolor = rng.uniform(20, 100); + double sigmaspace = rng.uniform(10, 40); + + OCL_OFF(cv::bilateralFilter(src_roi, dst_roi, ksize, sigmacolor, sigmaspace, borderType)); + OCL_ON(cv::bilateralFilter(usrc_roi, udst_roi, ksize, sigmacolor, sigmaspace, borderType)); + + Near(); + } +} + +////////////////////////////////////////////////////////////////////////////////////////////////////////////// + +#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 + +#define FILTER_DATATYPES Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4, \ + CV_32FC1, CV_32FC3, CV_32FC4, \ + CV_64FC1, CV_64FC3, CV_64FC4) + +OCL_INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine( + Values((MatType)CV_8UC1), + Values(5, 9), + Values(Size(0, 0)), // not used + FILTER_BORDER_SET_NO_ISOLATED, + Values(0.0), // not used + Bool())); + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL