mirror of https://github.com/opencv/opencv.git
added GPU non local means brute force filter + testspull/26/head
Anatoly Baksheev
12 years ago
parent
a09679044e
commit
0ba01afd83
19 changed files with 744 additions and 190 deletions
@ -0,0 +1,98 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace testing; |
||||
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BilateralFilter
|
||||
|
||||
DEF_PARAM_TEST(Sz_Depth_Cn_KernelSz, cv::Size, MatDepth , int, int); |
||||
|
||||
PERF_TEST_P(Sz_Depth_Cn_KernelSz, Denoising_BilateralFilter,
|
||||
Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U, CV_16U, CV_32F), GPU_CHANNELS_1_3_4, Values(3, 5, 9))) |
||||
{ |
||||
declare.time(30.0); |
||||
|
||||
cv::Size size = GET_PARAM(0); |
||||
int depth = GET_PARAM(1); |
||||
int channels = GET_PARAM(2); |
||||
int kernel_size = GET_PARAM(3); |
||||
|
||||
float sigma_color = 7; |
||||
float sigma_spatial = 5; |
||||
int borderMode = cv::BORDER_REFLECT101; |
||||
|
||||
int type = CV_MAKE_TYPE(depth, channels); |
||||
|
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
cv::gpu::bilateralFilter(d_src, d_dst, kernel_size, sigma_color, sigma_spatial, borderMode); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::bilateralFilter(d_src, d_dst, kernel_size, sigma_color, sigma_spatial, borderMode); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::Mat dst; |
||||
|
||||
cv::bilateralFilter(src, dst, kernel_size, sigma_color, sigma_spatial, borderMode); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::bilateralFilter(src, dst, kernel_size, sigma_color, sigma_spatial, borderMode); |
||||
} |
||||
} |
||||
} |
||||
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// nonLocalMeans
|
||||
|
||||
DEF_PARAM_TEST(Sz_Depth_Cn_WinSz_BlockSz, cv::Size, MatDepth , int, int, int); |
||||
|
||||
PERF_TEST_P(Sz_Depth_Cn_WinSz_BlockSz, Denoising_NonLocalMeans,
|
||||
Combine(GPU_TYPICAL_MAT_SIZES, Values<MatDepth>(CV_8U), Values(1), Values(21), Values(5, 7))) |
||||
{ |
||||
declare.time(30.0); |
||||
|
||||
cv::Size size = GET_PARAM(0); |
||||
int depth = GET_PARAM(1); |
||||
int channels = GET_PARAM(2); |
||||
|
||||
int search_widow_size = GET_PARAM(3); |
||||
int block_size = GET_PARAM(4); |
||||
|
||||
float h = 10; |
||||
int borderMode = cv::BORDER_REFLECT101; |
||||
|
||||
int type = CV_MAKE_TYPE(depth, channels); |
||||
|
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
cv::gpu::nonLocalMeans(d_src, d_dst, h, search_widow_size, block_size, borderMode); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::nonLocalMeans(d_src, d_dst, h, search_widow_size, block_size, borderMode); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
FAIL(); |
||||
} |
||||
} |
||||
@ -0,0 +1,143 @@ |
||||
/*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) 2000-2008, Intel Corporation, all rights reserved. |
||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved. |
||||
// Copyright (C) 1993-2011, NVIDIA Corporation, 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 bpied warranties, including, but not limited to, the bpied |
||||
// 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 "internal_shared.hpp" |
||||
|
||||
#include "opencv2/gpu/device/vec_traits.hpp" |
||||
#include "opencv2/gpu/device/vec_math.hpp" |
||||
#include "opencv2/gpu/device/border_interpolate.hpp" |
||||
|
||||
using namespace cv::gpu; |
||||
|
||||
typedef unsigned char uchar; |
||||
typedef unsigned short ushort; |
||||
|
||||
////////////////////////////////////////////////////////////////////////////////// |
||||
/// Non local means denosings |
||||
|
||||
namespace cv { namespace gpu { namespace device |
||||
{ |
||||
namespace imgproc |
||||
{ |
||||
__device__ __forceinline__ float norm2(const float& v) { return v*v; } |
||||
__device__ __forceinline__ float norm2(const float2& v) { return v.x*v.x + v.y*v.y; } |
||||
__device__ __forceinline__ float norm2(const float3& v) { return v.x*v.x + v.y*v.y + v.z*v.z; } |
||||
__device__ __forceinline__ float norm2(const float4& v) { return v.x*v.x + v.y*v.y + v.z*v.z + v.w*v.w; } |
||||
|
||||
template<typename T, typename B> |
||||
__global__ void nlm_kernel(const PtrStepSz<T> src, PtrStep<T> dst, const B b, int search_radius, int block_radius, float h2_inv_half) |
||||
{ |
||||
typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type value_type; |
||||
|
||||
const int x = blockDim.x * blockIdx.x + threadIdx.x; |
||||
const int y = blockDim.y * blockIdx.y + threadIdx.y; |
||||
|
||||
if (x >= src.cols || y >= src.rows) |
||||
return; |
||||
|
||||
float block_radius2_inv = -1.f/(block_radius * block_radius); |
||||
|
||||
value_type sum1 = VecTraits<value_type>::all(0); |
||||
float sum2 = 0.f; |
||||
|
||||
for(float cy = -search_radius; cy <= search_radius; ++cy) |
||||
for(float cx = -search_radius; cx <= search_radius; ++cx) |
||||
{ |
||||
float color2 = 0; |
||||
for(float by = -block_radius; by <= block_radius; ++by) |
||||
for(float bx = -block_radius; bx <= block_radius; ++bx) |
||||
{ |
||||
value_type v1 = saturate_cast<value_type>(src(y + by, x + bx)); |
||||
value_type v2 = saturate_cast<value_type>(src(y + cy + by, x + cx + bx)); |
||||
color2 += norm2(v1 - v2); |
||||
} |
||||
|
||||
float dist2 = cx * cx + cy * cy; |
||||
float w = __expf(color2 * h2_inv_half + dist2 * block_radius2_inv); |
||||
|
||||
sum1 = sum1 + saturate_cast<value_type>(src(y + cy, x + cy)) * w; |
||||
sum2 += w; |
||||
} |
||||
|
||||
dst(y, x) = saturate_cast<T>(sum1 / sum2); |
||||
|
||||
} |
||||
|
||||
template<typename T, template <typename> class B> |
||||
void nlm_caller(const PtrStepSzb src, PtrStepSzb dst, int search_radius, int block_radius, float h, cudaStream_t stream) |
||||
{ |
||||
dim3 block (32, 8); |
||||
dim3 grid (divUp (src.cols, block.x), divUp (src.rows, block.y)); |
||||
|
||||
B<T> b(src.rows, src.cols); |
||||
|
||||
float h2_inv_half = -0.5f/(h * h * VecTraits<T>::cn); |
||||
|
||||
cudaSafeCall( cudaFuncSetCacheConfig (nlm_kernel<T, B<T> >, cudaFuncCachePreferL1) ); |
||||
nlm_kernel<<<grid, block>>>((PtrStepSz<T>)src, (PtrStepSz<T>)dst, b, search_radius, block_radius, h2_inv_half); |
||||
cudaSafeCall ( cudaGetLastError () ); |
||||
|
||||
if (stream == 0) |
||||
cudaSafeCall( cudaDeviceSynchronize() ); |
||||
} |
||||
|
||||
template<typename T> |
||||
void nlm_bruteforce_gpu(const PtrStepSzb& src, PtrStepSzb dst, int search_radius, int block_radius, float h, int borderMode, cudaStream_t stream) |
||||
{ |
||||
typedef void (*func_t)(const PtrStepSzb src, PtrStepSzb dst, int search_radius, int block_radius, float h, cudaStream_t stream); |
||||
|
||||
static func_t funcs[] = |
||||
{ |
||||
nlm_caller<T, BrdReflect101>, |
||||
nlm_caller<T, BrdReplicate>, |
||||
nlm_caller<T, BrdConstant>, |
||||
nlm_caller<T, BrdReflect>, |
||||
nlm_caller<T, BrdWrap>, |
||||
}; |
||||
funcs[borderMode](src, dst, search_radius, block_radius, h, stream); |
||||
} |
||||
|
||||
template void nlm_bruteforce_gpu<uchar>(const PtrStepSzb&, PtrStepSzb, int, int, float, int, cudaStream_t); |
||||
template void nlm_bruteforce_gpu<uchar3>(const PtrStepSzb&, PtrStepSzb, int, int, float, int, cudaStream_t); |
||||
} |
||||
}}} |
||||
@ -0,0 +1,135 @@ |
||||
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||
// Copyright (C) 2009, Willow Garage 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 GpuMaterials 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 bpied warranties, including, but not limited to, the bpied
|
||||
// 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 "precomp.hpp" |
||||
|
||||
using namespace cv; |
||||
using namespace cv::gpu; |
||||
|
||||
#if !defined (HAVE_CUDA) |
||||
|
||||
cv::gpu::bilateralFilter(const GpuMat&, GpuMat&, int, float, float, int, Stream&) { throw_nogpu(); } |
||||
|
||||
#else |
||||
|
||||
|
||||
namespace cv { namespace gpu { namespace device |
||||
{ |
||||
namespace imgproc |
||||
{ |
||||
template<typename T> |
||||
void bilateral_filter_gpu(const PtrStepSzb& src, PtrStepSzb dst, int kernel_size, float sigma_spatial, float sigma_color, int borderMode, cudaStream_t stream); |
||||
|
||||
template<typename T> |
||||
void nlm_bruteforce_gpu(const PtrStepSzb& src, PtrStepSzb dst, int search_radius, int block_radius, float h, int borderMode, cudaStream_t stream); |
||||
} |
||||
}}} |
||||
|
||||
void cv::gpu::bilateralFilter(const GpuMat& src, GpuMat& dst, int kernel_size, float sigma_color, float sigma_spatial, int borderMode, Stream& s) |
||||
{ |
||||
using cv::gpu::device::imgproc::bilateral_filter_gpu; |
||||
|
||||
typedef void (*func_t)(const PtrStepSzb& src, PtrStepSzb dst, int kernel_size, float sigma_spatial, float sigma_color, int borderMode, cudaStream_t s); |
||||
|
||||
static const func_t funcs[6][4] = |
||||
{ |
||||
{bilateral_filter_gpu<uchar> , 0 /*bilateral_filter_gpu<uchar2>*/ , bilateral_filter_gpu<uchar3> , bilateral_filter_gpu<uchar4> }, |
||||
{0 /*bilateral_filter_gpu<schar>*/, 0 /*bilateral_filter_gpu<schar2>*/ , 0 /*bilateral_filter_gpu<schar3>*/, 0 /*bilateral_filter_gpu<schar4>*/}, |
||||
{bilateral_filter_gpu<ushort> , 0 /*bilateral_filter_gpu<ushort2>*/, bilateral_filter_gpu<ushort3> , bilateral_filter_gpu<ushort4> }, |
||||
{bilateral_filter_gpu<short> , 0 /*bilateral_filter_gpu<short2>*/ , bilateral_filter_gpu<short3> , bilateral_filter_gpu<short4> }, |
||||
{0 /*bilateral_filter_gpu<int>*/ , 0 /*bilateral_filter_gpu<int2>*/ , 0 /*bilateral_filter_gpu<int3>*/ , 0 /*bilateral_filter_gpu<int4>*/ }, |
||||
{bilateral_filter_gpu<float> , 0 /*bilateral_filter_gpu<float2>*/ , bilateral_filter_gpu<float3> , bilateral_filter_gpu<float4> } |
||||
}; |
||||
|
||||
sigma_color = (sigma_color <= 0 ) ? 1 : sigma_color; |
||||
sigma_spatial = (sigma_spatial <= 0 ) ? 1 : sigma_spatial; |
||||
|
||||
|
||||
int radius = (kernel_size <= 0) ? cvRound(sigma_spatial*1.5) : kernel_size/2; |
||||
kernel_size = std::max(radius, 1)*2 + 1; |
||||
|
||||
CV_Assert(src.depth() <= CV_32F && src.channels() <= 4); |
||||
const func_t func = funcs[src.depth()][src.channels() - 1]; |
||||
CV_Assert(func != 0); |
||||
|
||||
CV_Assert(borderMode == BORDER_REFLECT101 || borderMode == BORDER_REPLICATE || borderMode == BORDER_CONSTANT || borderMode == BORDER_REFLECT || borderMode == BORDER_WRAP); |
||||
|
||||
int gpuBorderType; |
||||
CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType)); |
||||
|
||||
dst.create(src.size(), src.type()); |
||||
func(src, dst, kernel_size, sigma_spatial, sigma_color, gpuBorderType, StreamAccessor::getStream(s)); |
||||
} |
||||
|
||||
void cv::gpu::nonLocalMeans(const GpuMat& src, GpuMat& dst, float h, int search_window_size, int block_size, int borderMode, Stream& s) |
||||
{ |
||||
using cv::gpu::device::imgproc::nlm_bruteforce_gpu; |
||||
typedef void (*func_t)(const PtrStepSzb& src, PtrStepSzb dst, int search_radius, int block_radius, float h, int borderMode, cudaStream_t stream); |
||||
|
||||
static const func_t funcs[4] = { nlm_bruteforce_gpu<uchar>, 0 /*nlm_bruteforce_gpu<uchar2>*/ , nlm_bruteforce_gpu<uchar3>, 0/*nlm_bruteforce_gpu<uchar4>,*/ }; |
||||
|
||||
CV_Assert(src.type() == CV_8U || src.type() == CV_8UC3); |
||||
|
||||
const func_t func = funcs[src.channels() - 1]; |
||||
CV_Assert(func != 0); |
||||
|
||||
int b = borderMode; |
||||
CV_Assert(b == BORDER_REFLECT101 || b == BORDER_REPLICATE || b == BORDER_CONSTANT || b == BORDER_REFLECT || b == BORDER_WRAP); |
||||
|
||||
int gpuBorderType; |
||||
CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType)); |
||||
|
||||
int search_radius = search_window_size/2; |
||||
int block_radius = block_size/2; |
||||
|
||||
dst.create(src.size(), src.type()); |
||||
func(src, dst, search_radius, block_radius, h, gpuBorderType, StreamAccessor::getStream(s)); |
||||
} |
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
#endif |
||||
@ -0,0 +1,140 @@ |
||||
/*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.
|
||||
//
|
||||
//
|
||||
// Intel License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2000, Intel Corporation, 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 Intel Corporation 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" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
////////////////////////////////////////////////////////
|
||||
// BilateralFilter
|
||||
|
||||
PARAM_TEST_CASE(BilateralFilter, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo; |
||||
cv::Size size; |
||||
int type; |
||||
int kernel_size; |
||||
float sigma_color; |
||||
float sigma_spatial; |
||||
|
||||
virtual void SetUp() |
||||
{ |
||||
devInfo = GET_PARAM(0); |
||||
size = GET_PARAM(1); |
||||
type = GET_PARAM(2); |
||||
|
||||
kernel_size = 5; |
||||
sigma_color = 10.f; |
||||
sigma_spatial = 3.5f; |
||||
|
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
} |
||||
}; |
||||
|
||||
TEST_P(BilateralFilter, Accuracy) |
||||
{ |
||||
cv::Mat src = randomMat(size, type); |
||||
//cv::Mat src = readImage("hog/road.png", cv::IMREAD_GRAYSCALE);
|
||||
//cv::Mat src = readImage("csstereobp/aloe-R.png", cv::IMREAD_GRAYSCALE);
|
||||
|
||||
src.convertTo(src, type); |
||||
cv::gpu::GpuMat dst; |
||||
|
||||
cv::gpu::bilateralFilter(loadMat(src), dst, kernel_size, sigma_color, sigma_spatial); |
||||
|
||||
cv::Mat dst_gold; |
||||
cv::bilateralFilter(src, dst_gold, kernel_size, sigma_color, sigma_spatial); |
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, src.depth() == CV_32F ? 1e-3 : 1.0); |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, BilateralFilter, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(cv::Size(128, 128), cv::Size(113, 113), cv::Size(639, 481)), |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_32FC1), MatType(CV_32FC3)) |
||||
)); |
||||
|
||||
|
||||
////////////////////////////////////////////////////////
|
||||
// Brute Force Non local means
|
||||
|
||||
struct NonLocalMeans: testing::TestWithParam<cv::gpu::DeviceInfo> |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo; |
||||
|
||||
virtual void SetUp() |
||||
{ |
||||
devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
} |
||||
}; |
||||
|
||||
TEST_P(NonLocalMeans, Regression) |
||||
{ |
||||
using cv::gpu::GpuMat; |
||||
|
||||
cv::Mat bgr = readImage("denoising/lena_noised_gaussian_sigma=20_multi_0.png", cv::IMREAD_COLOR); |
||||
ASSERT_FALSE(bgr.empty()); |
||||
|
||||
cv::Mat gray; |
||||
cv::cvtColor(bgr, gray, CV_BGR2GRAY); |
||||
|
||||
GpuMat dbgr, dgray; |
||||
cv::gpu::nonLocalMeans(GpuMat(bgr), dbgr, 10); |
||||
cv::gpu::nonLocalMeans(GpuMat(gray), dgray, 10); |
||||
|
||||
#if 0 |
||||
dumpImage("denoising/denoised_lena_bgr.png", cv::Mat(dbgr)); |
||||
dumpImage("denoising/denoised_lena_gray.png", cv::Mat(dgray)); |
||||
#endif |
||||
|
||||
cv::Mat bgr_gold = readImage("denoising/denoised_lena_bgr.png", cv::IMREAD_COLOR); |
||||
cv::Mat gray_gold = readImage("denoising/denoised_lena_gray.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(bgr_gold.empty() || gray_gold.empty()); |
||||
|
||||
EXPECT_MAT_NEAR(bgr_gold, dbgr, 1e-4); |
||||
EXPECT_MAT_NEAR(gray_gold, dgray, 1e-4); |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, NonLocalMeans, ALL_DEVICES); |
||||
|
||||
|
||||
#endif // HAVE_CUDA
|
||||
Loading…
Reference in new issue