mirror of https://github.com/opencv/opencv.git
noob
12 years ago
commit
fba081992d
85 changed files with 1560 additions and 500 deletions
Binary file not shown.
Binary file not shown.
@ -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,219 @@ |
||||
/*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 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 "internal_shared.hpp" |
||||
#include "opencv2/gpu/device/limits.hpp" |
||||
|
||||
namespace cv { namespace gpu { namespace device |
||||
{ |
||||
namespace disp_bilateral_filter |
||||
{ |
||||
__constant__ float* ctable_color; |
||||
__constant__ float* ctable_space; |
||||
__constant__ size_t ctable_space_step; |
||||
|
||||
__constant__ int cndisp; |
||||
__constant__ int cradius; |
||||
|
||||
__constant__ short cedge_disc; |
||||
__constant__ short cmax_disc; |
||||
|
||||
void disp_load_constants(float* table_color, PtrStepSzf table_space, int ndisp, int radius, short edge_disc, short max_disc) |
||||
{ |
||||
cudaSafeCall( cudaMemcpyToSymbol(ctable_color, &table_color, sizeof(table_color)) ); |
||||
cudaSafeCall( cudaMemcpyToSymbol(ctable_space, &table_space.data, sizeof(table_space.data)) ); |
||||
size_t table_space_step = table_space.step / sizeof(float); |
||||
cudaSafeCall( cudaMemcpyToSymbol(ctable_space_step, &table_space_step, sizeof(size_t)) ); |
||||
|
||||
cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); |
||||
cudaSafeCall( cudaMemcpyToSymbol(cradius, &radius, sizeof(int)) ); |
||||
|
||||
cudaSafeCall( cudaMemcpyToSymbol(cedge_disc, &edge_disc, sizeof(short)) ); |
||||
cudaSafeCall( cudaMemcpyToSymbol(cmax_disc, &max_disc, sizeof(short)) ); |
||||
} |
||||
|
||||
template <int channels> |
||||
struct DistRgbMax |
||||
{ |
||||
static __device__ __forceinline__ uchar calc(const uchar* a, const uchar* b) |
||||
{ |
||||
uchar x = ::abs(a[0] - b[0]); |
||||
uchar y = ::abs(a[1] - b[1]); |
||||
uchar z = ::abs(a[2] - b[2]); |
||||
return (::max(::max(x, y), z)); |
||||
} |
||||
}; |
||||
|
||||
template <> |
||||
struct DistRgbMax<1> |
||||
{ |
||||
static __device__ __forceinline__ uchar calc(const uchar* a, const uchar* b) |
||||
{ |
||||
return ::abs(a[0] - b[0]); |
||||
} |
||||
}; |
||||
|
||||
template <int channels, typename T> |
||||
__global__ void disp_bilateral_filter(int t, T* disp, size_t disp_step, const uchar* img, size_t img_step, int h, int w) |
||||
{ |
||||
const int y = blockIdx.y * blockDim.y + threadIdx.y; |
||||
const int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + t) & 1); |
||||
|
||||
T dp[5]; |
||||
|
||||
if (y > 0 && y < h - 1 && x > 0 && x < w - 1) |
||||
{ |
||||
dp[0] = *(disp + (y ) * disp_step + x + 0); |
||||
dp[1] = *(disp + (y-1) * disp_step + x + 0); |
||||
dp[2] = *(disp + (y ) * disp_step + x - 1); |
||||
dp[3] = *(disp + (y+1) * disp_step + x + 0); |
||||
dp[4] = *(disp + (y ) * disp_step + x + 1); |
||||
|
||||
if(::abs(dp[1] - dp[0]) >= cedge_disc || ::abs(dp[2] - dp[0]) >= cedge_disc || ::abs(dp[3] - dp[0]) >= cedge_disc || ::abs(dp[4] - dp[0]) >= cedge_disc) |
||||
{ |
||||
const int ymin = ::max(0, y - cradius); |
||||
const int xmin = ::max(0, x - cradius); |
||||
const int ymax = ::min(h - 1, y + cradius); |
||||
const int xmax = ::min(w - 1, x + cradius); |
||||
|
||||
float cost[] = {0.0f, 0.0f, 0.0f, 0.0f, 0.0f}; |
||||
|
||||
const uchar* ic = img + y * img_step + channels * x; |
||||
|
||||
for(int yi = ymin; yi <= ymax; yi++) |
||||
{ |
||||
const T* disp_y = disp + yi * disp_step; |
||||
|
||||
for(int xi = xmin; xi <= xmax; xi++) |
||||
{ |
||||
const uchar* in = img + yi * img_step + channels * xi; |
||||
|
||||
uchar dist_rgb = DistRgbMax<channels>::calc(in, ic); |
||||
|
||||
const float weight = ctable_color[dist_rgb] * (ctable_space + ::abs(y-yi)* ctable_space_step)[::abs(x-xi)]; |
||||
|
||||
const T disp_reg = disp_y[xi]; |
||||
|
||||
cost[0] += ::min(cmax_disc, ::abs(disp_reg - dp[0])) * weight; |
||||
cost[1] += ::min(cmax_disc, ::abs(disp_reg - dp[1])) * weight; |
||||
cost[2] += ::min(cmax_disc, ::abs(disp_reg - dp[2])) * weight; |
||||
cost[3] += ::min(cmax_disc, ::abs(disp_reg - dp[3])) * weight; |
||||
cost[4] += ::min(cmax_disc, ::abs(disp_reg - dp[4])) * weight; |
||||
} |
||||
} |
||||
|
||||
float minimum = numeric_limits<float>::max(); |
||||
int id = 0; |
||||
|
||||
if (cost[0] < minimum) |
||||
{ |
||||
minimum = cost[0]; |
||||
id = 0; |
||||
} |
||||
if (cost[1] < minimum) |
||||
{ |
||||
minimum = cost[1]; |
||||
id = 1; |
||||
} |
||||
if (cost[2] < minimum) |
||||
{ |
||||
minimum = cost[2]; |
||||
id = 2; |
||||
} |
||||
if (cost[3] < minimum) |
||||
{ |
||||
minimum = cost[3]; |
||||
id = 3; |
||||
} |
||||
if (cost[4] < minimum) |
||||
{ |
||||
minimum = cost[4]; |
||||
id = 4; |
||||
} |
||||
|
||||
*(disp + y * disp_step + x) = dp[id]; |
||||
} |
||||
} |
||||
} |
||||
|
||||
template <typename T> |
||||
void disp_bilateral_filter(PtrStepSz<T> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream) |
||||
{ |
||||
dim3 threads(32, 8, 1); |
||||
dim3 grid(1, 1, 1); |
||||
grid.x = divUp(disp.cols, threads.x << 1); |
||||
grid.y = divUp(disp.rows, threads.y); |
||||
|
||||
switch (channels) |
||||
{ |
||||
case 1: |
||||
for (int i = 0; i < iters; ++i) |
||||
{ |
||||
disp_bilateral_filter<1><<<grid, threads, 0, stream>>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
|
||||
disp_bilateral_filter<1><<<grid, threads, 0, stream>>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
} |
||||
break; |
||||
case 3: |
||||
for (int i = 0; i < iters; ++i) |
||||
{ |
||||
disp_bilateral_filter<3><<<grid, threads, 0, stream>>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
|
||||
disp_bilateral_filter<3><<<grid, threads, 0, stream>>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
} |
||||
break; |
||||
default: |
||||
cv::gpu::error("Unsupported channels count", __FILE__, __LINE__, "disp_bilateral_filter"); |
||||
} |
||||
|
||||
if (stream == 0) |
||||
cudaSafeCall( cudaDeviceSynchronize() ); |
||||
} |
||||
|
||||
template void disp_bilateral_filter<uchar>(PtrStepSz<uchar> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); |
||||
template void disp_bilateral_filter<short>(PtrStepSz<short> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); |
||||
} // namespace bilateral_filter |
||||
}}} // namespace cv { namespace gpu { namespace device |
||||
@ -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
|
||||
@ -1,21 +1,21 @@ |
||||
package org.opencv.android; |
||||
|
||||
/** |
||||
* Installation callback interface
|
||||
* Installation callback interface. |
||||
*/ |
||||
public interface InstallCallbackInterface |
||||
{ |
||||
/** |
||||
* Target package name |
||||
* @return Return target package name |
||||
* Target package name. |
||||
* @return Return target package name. |
||||
*/ |
||||
public String getPackageName(); |
||||
/** |
||||
* Installation of package is approved |
||||
* Installation is approved. |
||||
*/ |
||||
public void install(); |
||||
/** |
||||
* Installation canceled |
||||
* Installation is canceled. |
||||
*/ |
||||
public void cancel(); |
||||
}; |
||||
|
||||
@ -1,44 +1,44 @@ |
||||
package org.opencv.android; |
||||
|
||||
/** |
||||
* Interface for callback object in case of asynchronous initialization of OpenCV |
||||
* Interface for callback object in case of asynchronous initialization of OpenCV. |
||||
*/ |
||||
public interface LoaderCallbackInterface |
||||
{ |
||||
/** |
||||
* OpenCV initialization finished successfully |
||||
* OpenCV initialization finished successfully. |
||||
*/ |
||||
static final int SUCCESS = 0; |
||||
/** |
||||
* OpenCV library installation via Google Play service was initialized. Application restart is required |
||||
* OpenCV library installation via Google Play service has been initialized. Restart the application. |
||||
*/ |
||||
static final int RESTART_REQUIRED = 1; |
||||
/** |
||||
* Google Play (Android Market) cannot be invoked |
||||
* Google Play Market cannot be invoked. |
||||
*/ |
||||
static final int MARKET_ERROR = 2; |
||||
/** |
||||
* OpenCV library installation was canceled by user |
||||
* OpenCV library installation has been canceled by the user. |
||||
*/ |
||||
static final int INSTALL_CANCELED = 3; |
||||
/** |
||||
* Version of OpenCV Manager Service is incompatible with this app. Service update is needed |
||||
* This version of OpenCV Manager Service is incompatible with the app. Possibly, a service update is required. |
||||
*/ |
||||
static final int INCOMPATIBLE_MANAGER_VERSION = 4; |
||||
/** |
||||
* OpenCV library initialization failed |
||||
* OpenCV library initialization has failed. |
||||
*/ |
||||
static final int INIT_FAILED = 0xff; |
||||
|
||||
/** |
||||
* Callback method that is called after OpenCV library initialization |
||||
* @param status Status of initialization. See Initialization status constants |
||||
* Callback method, called after OpenCV library initialization. |
||||
* @param status status of initialization (see initialization status constants). |
||||
*/ |
||||
public void onManagerConnected(int status); |
||||
|
||||
/** |
||||
* Callback method that is called in case when package installation is needed |
||||
* @param callback Answer object with approve and cancel methods and package description |
||||
* Callback method, called in case the package installation is needed. |
||||
* @param callback answer object with approve and cancel methods and the package description. |
||||
*/ |
||||
public void onPackageInstall(InstallCallbackInterface callback); |
||||
}; |
||||
|
||||
@ -1,33 +1,33 @@ |
||||
package org.opencv.engine; |
||||
|
||||
/** |
||||
* Class provides Java interface to OpenCV Engine Service. Is synchronous with native OpenCVEngine class. |
||||
* Class provides a Java interface for OpenCV Engine Service. It's synchronous with native OpenCVEngine class. |
||||
*/ |
||||
interface OpenCVEngineInterface |
||||
{ |
||||
/** |
||||
* @return Return service version |
||||
* @return Returns service version. |
||||
*/ |
||||
int getEngineVersion(); |
||||
|
||||
/** |
||||
* Find installed OpenCV library |
||||
* @param OpenCV version |
||||
* @return Return path to OpenCV native libs or empty string if OpenCV was not found |
||||
* Finds an installed OpenCV library. |
||||
* @param OpenCV version. |
||||
* @return Returns path to OpenCV native libs or an empty string if OpenCV can not be found. |
||||
*/ |
||||
String getLibPathByVersion(String version); |
||||
|
||||
/** |
||||
* Try to install defined version of OpenCV from Google Play (Android Market). |
||||
* @param OpenCV version |
||||
* @return Return true if installation was successful or OpenCV package has been already installed |
||||
* Tries to install defined version of OpenCV from Google Play Market. |
||||
* @param OpenCV version. |
||||
* @return Returns true if installation was successful or OpenCV package has been already installed. |
||||
*/ |
||||
boolean installVersion(String version); |
||||
|
||||
/** |
||||
* Return list of libraries in loading order separated by ";" symbol |
||||
* @param OpenCV version |
||||
* @return Return OpenCV libraries names separated by symbol ";" in loading order |
||||
* Returns list of libraries in loading order, separated by semicolon. |
||||
* @param OpenCV version. |
||||
* @return Returns names of OpenCV libraries, separated by semicolon. |
||||
*/ |
||||
String getLibraryList(String version); |
||||
} |
||||
Loading…
Reference in new issue