added cv::bilateralFilter to T-API
This commit is contained in:
Ilya Lavrenov
parent
24465bbd05
commit
93c6fe66ce
59
modules/imgproc/src/opencl/bilateral.cl
Normal file
59
modules/imgproc/src/opencl/bilateral.cl
Normal file
@ -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);
|
||||
}
|
||||
}
|
||||
@ -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<float> _color_weight(cn * 256);
|
||||
std::vector<float> _space_weight(d * d);
|
||||
std::vector<int> _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;
|
||||
@ -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 );
|
||||
|
||||
176
modules/imgproc/test/ocl/test_filters.cpp
Normal file
176
modules/imgproc/test/ocl/test_filters.cpp
Normal file
@ -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
|
||||
Loading…
Reference in New Issue
Block a user