From 43d5e2d8b4ea4ff9bf0c83f874a2ac8a79f6ff5e Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 11 Apr 2013 14:53:35 +0400 Subject: [PATCH] removed gpu BORDER_* constants --- modules/core/include/opencv2/core/cuda/common.hpp | 9 -- modules/core/include/opencv2/core/gpu_private.hpp | 4 - modules/core/src/gpumat.cpp | 30 ----- modules/gpuarithm/src/arithm.cpp | 9 +- modules/gpuarithm/src/cuda/copy_make_border.cu | 6 +- modules/gpufilters/src/cuda/column_filter.hpp | 134 ++++++++++----------- modules/gpufilters/src/cuda/filter2d.cu | 6 +- modules/gpufilters/src/cuda/row_filter.hpp | 134 ++++++++++----------- modules/gpufilters/src/filtering.cpp | 15 +-- modules/gpuimgproc/src/bilateral_filter.cpp | 5 +- modules/gpuimgproc/src/cuda/bilateral_filter.cu | 4 +- modules/gpuimgproc/src/cuda/imgproc.cu | 12 +- modules/gpuimgproc/src/imgproc.cpp | 10 +- .../gpuoptflow/src/cuda/optical_flow_farneback.cu | 15 ++- modules/gpuoptflow/src/optical_flow_farneback.cpp | 6 +- modules/gpuwarping/src/cuda/remap.cu | 18 +-- modules/gpuwarping/src/cuda/warp.cu | 18 +-- modules/gpuwarping/src/remap.cpp | 5 +- modules/gpuwarping/src/warp.cpp | 10 +- modules/photo/src/cuda/nlm.cu | 4 +- modules/photo/src/denoising_gpu.cpp | 5 +- 21 files changed, 196 insertions(+), 263 deletions(-) diff --git a/modules/core/include/opencv2/core/cuda/common.hpp b/modules/core/include/opencv2/core/cuda/common.hpp index 774500e..434a3eb 100644 --- a/modules/core/include/opencv2/core/cuda/common.hpp +++ b/modules/core/include/opencv2/core/cuda/common.hpp @@ -87,15 +87,6 @@ namespace cv { namespace gpu namespace cv { namespace gpu { - enum - { - BORDER_REFLECT101_GPU = 0, - BORDER_REPLICATE_GPU, - BORDER_CONSTANT_GPU, - BORDER_REFLECT_GPU, - BORDER_WRAP_GPU - }; - namespace cudev { __host__ __device__ __forceinline__ int divUp(int total, int grain) diff --git a/modules/core/include/opencv2/core/gpu_private.hpp b/modules/core/include/opencv2/core/gpu_private.hpp index be194f5..7692bc2 100644 --- a/modules/core/include/opencv2/core/gpu_private.hpp +++ b/modules/core/include/opencv2/core/gpu_private.hpp @@ -74,10 +74,6 @@ namespace cv { namespace gpu { CV_EXPORTS cv::String getNppErrorMessage(int code); CV_EXPORTS cv::String getCudaDriverApiErrorMessage(int code); - - // Converts CPU border extrapolation mode into GPU internal analogue. - // Returns true if the GPU analogue exists, false otherwise. - CV_EXPORTS bool tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType); }} #ifndef HAVE_CUDA diff --git a/modules/core/src/gpumat.cpp b/modules/core/src/gpumat.cpp index c5d4d7a..11bb419 100644 --- a/modules/core/src/gpumat.cpp +++ b/modules/core/src/gpumat.cpp @@ -1678,33 +1678,3 @@ String cv::gpu::getCudaDriverApiErrorMessage(int code) return getErrorString(code, cu_errors, cu_errors_num); #endif } - -bool cv::gpu::tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType) -{ -#ifndef HAVE_CUDA - (void) cpuBorderType; - (void) gpuBorderType; - return false; -#else - switch (cpuBorderType) - { - case IPL_BORDER_REFLECT_101: - gpuBorderType = cv::gpu::BORDER_REFLECT101_GPU; - return true; - case IPL_BORDER_REPLICATE: - gpuBorderType = cv::gpu::BORDER_REPLICATE_GPU; - return true; - case IPL_BORDER_CONSTANT: - gpuBorderType = cv::gpu::BORDER_CONSTANT_GPU; - return true; - case IPL_BORDER_REFLECT: - gpuBorderType = cv::gpu::BORDER_REFLECT_GPU; - return true; - case IPL_BORDER_WRAP: - gpuBorderType = cv::gpu::BORDER_WRAP_GPU; - return true; - default: - return false; - }; -#endif -} diff --git a/modules/gpuarithm/src/arithm.cpp b/modules/gpuarithm/src/arithm.cpp index cc85cc7..908d963 100644 --- a/modules/gpuarithm/src/arithm.cpp +++ b/modules/gpuarithm/src/arithm.cpp @@ -696,13 +696,13 @@ typedef Npp32s Npp32s_a; void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value, Stream& s) { CV_Assert(src.depth() <= CV_32F && src.channels() <= 4); - CV_Assert(borderType == IPL_BORDER_REFLECT_101 || borderType == IPL_BORDER_REPLICATE || borderType == IPL_BORDER_CONSTANT || borderType == IPL_BORDER_REFLECT || borderType == IPL_BORDER_WRAP); + CV_Assert(borderType == BORDER_REFLECT_101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP); dst.create(src.rows + top + bottom, src.cols + left + right, src.type()); cudaStream_t stream = StreamAccessor::getStream(s); - if (borderType == IPL_BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1)) + if (borderType == BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1)) { NppiSize srcsz; srcsz.width = src.cols; @@ -766,10 +766,7 @@ void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom caller_t func = callers[src.depth()][src.channels() - 1]; CV_Assert(func != 0); - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType)); - - func(src, dst, top, left, gpuBorderType, value, stream); + func(src, dst, top, left, borderType, value, stream); } } diff --git a/modules/gpuarithm/src/cuda/copy_make_border.cu b/modules/gpuarithm/src/cuda/copy_make_border.cu index ed90e9e..d772e09 100644 --- a/modules/gpuarithm/src/cuda/copy_make_border.cu +++ b/modules/gpuarithm/src/cuda/copy_make_border.cu @@ -86,11 +86,11 @@ namespace cv { namespace gpu { namespace cudev static const caller_t callers[5] = { - CopyMakeBorderDispatcher::call, - CopyMakeBorderDispatcher::call, CopyMakeBorderDispatcher::call, + CopyMakeBorderDispatcher::call, CopyMakeBorderDispatcher::call, - CopyMakeBorderDispatcher::call + CopyMakeBorderDispatcher::call, + CopyMakeBorderDispatcher::call }; callers[borderMode](PtrStepSz(src), PtrStepSz(dst), top, left, borderValue, stream); diff --git a/modules/gpufilters/src/cuda/column_filter.hpp b/modules/gpufilters/src/cuda/column_filter.hpp index 39b6d47..6f10c36 100644 --- a/modules/gpufilters/src/cuda/column_filter.hpp +++ b/modules/gpufilters/src/cuda/column_filter.hpp @@ -187,38 +187,38 @@ namespace filter { { 0, - column_filter::caller< 1, T, D, BrdColReflect101>, - column_filter::caller< 2, T, D, BrdColReflect101>, - column_filter::caller< 3, T, D, BrdColReflect101>, - column_filter::caller< 4, T, D, BrdColReflect101>, - column_filter::caller< 5, T, D, BrdColReflect101>, - column_filter::caller< 6, T, D, BrdColReflect101>, - column_filter::caller< 7, T, D, BrdColReflect101>, - column_filter::caller< 8, T, D, BrdColReflect101>, - column_filter::caller< 9, T, D, BrdColReflect101>, - column_filter::caller<10, T, D, BrdColReflect101>, - column_filter::caller<11, T, D, BrdColReflect101>, - column_filter::caller<12, T, D, BrdColReflect101>, - column_filter::caller<13, T, D, BrdColReflect101>, - column_filter::caller<14, T, D, BrdColReflect101>, - column_filter::caller<15, T, D, BrdColReflect101>, - column_filter::caller<16, T, D, BrdColReflect101>, - column_filter::caller<17, T, D, BrdColReflect101>, - column_filter::caller<18, T, D, BrdColReflect101>, - column_filter::caller<19, T, D, BrdColReflect101>, - column_filter::caller<20, T, D, BrdColReflect101>, - column_filter::caller<21, T, D, BrdColReflect101>, - column_filter::caller<22, T, D, BrdColReflect101>, - column_filter::caller<23, T, D, BrdColReflect101>, - column_filter::caller<24, T, D, BrdColReflect101>, - column_filter::caller<25, T, D, BrdColReflect101>, - column_filter::caller<26, T, D, BrdColReflect101>, - column_filter::caller<27, T, D, BrdColReflect101>, - column_filter::caller<28, T, D, BrdColReflect101>, - column_filter::caller<29, T, D, BrdColReflect101>, - column_filter::caller<30, T, D, BrdColReflect101>, - column_filter::caller<31, T, D, BrdColReflect101>, - column_filter::caller<32, T, D, BrdColReflect101> + column_filter::caller< 1, T, D, BrdColConstant>, + column_filter::caller< 2, T, D, BrdColConstant>, + column_filter::caller< 3, T, D, BrdColConstant>, + column_filter::caller< 4, T, D, BrdColConstant>, + column_filter::caller< 5, T, D, BrdColConstant>, + column_filter::caller< 6, T, D, BrdColConstant>, + column_filter::caller< 7, T, D, BrdColConstant>, + column_filter::caller< 8, T, D, BrdColConstant>, + column_filter::caller< 9, T, D, BrdColConstant>, + column_filter::caller<10, T, D, BrdColConstant>, + column_filter::caller<11, T, D, BrdColConstant>, + column_filter::caller<12, T, D, BrdColConstant>, + column_filter::caller<13, T, D, BrdColConstant>, + column_filter::caller<14, T, D, BrdColConstant>, + column_filter::caller<15, T, D, BrdColConstant>, + column_filter::caller<16, T, D, BrdColConstant>, + column_filter::caller<17, T, D, BrdColConstant>, + column_filter::caller<18, T, D, BrdColConstant>, + column_filter::caller<19, T, D, BrdColConstant>, + column_filter::caller<20, T, D, BrdColConstant>, + column_filter::caller<21, T, D, BrdColConstant>, + column_filter::caller<22, T, D, BrdColConstant>, + column_filter::caller<23, T, D, BrdColConstant>, + column_filter::caller<24, T, D, BrdColConstant>, + column_filter::caller<25, T, D, BrdColConstant>, + column_filter::caller<26, T, D, BrdColConstant>, + column_filter::caller<27, T, D, BrdColConstant>, + column_filter::caller<28, T, D, BrdColConstant>, + column_filter::caller<29, T, D, BrdColConstant>, + column_filter::caller<30, T, D, BrdColConstant>, + column_filter::caller<31, T, D, BrdColConstant>, + column_filter::caller<32, T, D, BrdColConstant> }, { 0, @@ -257,41 +257,6 @@ namespace filter }, { 0, - column_filter::caller< 1, T, D, BrdColConstant>, - column_filter::caller< 2, T, D, BrdColConstant>, - column_filter::caller< 3, T, D, BrdColConstant>, - column_filter::caller< 4, T, D, BrdColConstant>, - column_filter::caller< 5, T, D, BrdColConstant>, - column_filter::caller< 6, T, D, BrdColConstant>, - column_filter::caller< 7, T, D, BrdColConstant>, - column_filter::caller< 8, T, D, BrdColConstant>, - column_filter::caller< 9, T, D, BrdColConstant>, - column_filter::caller<10, T, D, BrdColConstant>, - column_filter::caller<11, T, D, BrdColConstant>, - column_filter::caller<12, T, D, BrdColConstant>, - column_filter::caller<13, T, D, BrdColConstant>, - column_filter::caller<14, T, D, BrdColConstant>, - column_filter::caller<15, T, D, BrdColConstant>, - column_filter::caller<16, T, D, BrdColConstant>, - column_filter::caller<17, T, D, BrdColConstant>, - column_filter::caller<18, T, D, BrdColConstant>, - column_filter::caller<19, T, D, BrdColConstant>, - column_filter::caller<20, T, D, BrdColConstant>, - column_filter::caller<21, T, D, BrdColConstant>, - column_filter::caller<22, T, D, BrdColConstant>, - column_filter::caller<23, T, D, BrdColConstant>, - column_filter::caller<24, T, D, BrdColConstant>, - column_filter::caller<25, T, D, BrdColConstant>, - column_filter::caller<26, T, D, BrdColConstant>, - column_filter::caller<27, T, D, BrdColConstant>, - column_filter::caller<28, T, D, BrdColConstant>, - column_filter::caller<29, T, D, BrdColConstant>, - column_filter::caller<30, T, D, BrdColConstant>, - column_filter::caller<31, T, D, BrdColConstant>, - column_filter::caller<32, T, D, BrdColConstant> - }, - { - 0, column_filter::caller< 1, T, D, BrdColReflect>, column_filter::caller< 2, T, D, BrdColReflect>, column_filter::caller< 3, T, D, BrdColReflect>, @@ -359,6 +324,41 @@ namespace filter column_filter::caller<30, T, D, BrdColWrap>, column_filter::caller<31, T, D, BrdColWrap>, column_filter::caller<32, T, D, BrdColWrap> + }, + { + 0, + column_filter::caller< 1, T, D, BrdColReflect101>, + column_filter::caller< 2, T, D, BrdColReflect101>, + column_filter::caller< 3, T, D, BrdColReflect101>, + column_filter::caller< 4, T, D, BrdColReflect101>, + column_filter::caller< 5, T, D, BrdColReflect101>, + column_filter::caller< 6, T, D, BrdColReflect101>, + column_filter::caller< 7, T, D, BrdColReflect101>, + column_filter::caller< 8, T, D, BrdColReflect101>, + column_filter::caller< 9, T, D, BrdColReflect101>, + column_filter::caller<10, T, D, BrdColReflect101>, + column_filter::caller<11, T, D, BrdColReflect101>, + column_filter::caller<12, T, D, BrdColReflect101>, + column_filter::caller<13, T, D, BrdColReflect101>, + column_filter::caller<14, T, D, BrdColReflect101>, + column_filter::caller<15, T, D, BrdColReflect101>, + column_filter::caller<16, T, D, BrdColReflect101>, + column_filter::caller<17, T, D, BrdColReflect101>, + column_filter::caller<18, T, D, BrdColReflect101>, + column_filter::caller<19, T, D, BrdColReflect101>, + column_filter::caller<20, T, D, BrdColReflect101>, + column_filter::caller<21, T, D, BrdColReflect101>, + column_filter::caller<22, T, D, BrdColReflect101>, + column_filter::caller<23, T, D, BrdColReflect101>, + column_filter::caller<24, T, D, BrdColReflect101>, + column_filter::caller<25, T, D, BrdColReflect101>, + column_filter::caller<26, T, D, BrdColReflect101>, + column_filter::caller<27, T, D, BrdColReflect101>, + column_filter::caller<28, T, D, BrdColReflect101>, + column_filter::caller<29, T, D, BrdColReflect101>, + column_filter::caller<30, T, D, BrdColReflect101>, + column_filter::caller<31, T, D, BrdColReflect101>, + column_filter::caller<32, T, D, BrdColReflect101> } }; diff --git a/modules/gpufilters/src/cuda/filter2d.cu b/modules/gpufilters/src/cuda/filter2d.cu index 0bb5fcd..80c93c5 100644 --- a/modules/gpufilters/src/cuda/filter2d.cu +++ b/modules/gpufilters/src/cuda/filter2d.cu @@ -131,11 +131,11 @@ namespace cv { namespace gpu { namespace cudev typedef void (*func_t)(const PtrStepSz srcWhole, int xoff, int yoff, PtrStepSz dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream); static const func_t funcs[] = { - Filter2DCaller::call, - Filter2DCaller::call, Filter2DCaller::call, + Filter2DCaller::call, Filter2DCaller::call, - Filter2DCaller::call + Filter2DCaller::call, + Filter2DCaller::call }; if (stream == 0) diff --git a/modules/gpufilters/src/cuda/row_filter.hpp b/modules/gpufilters/src/cuda/row_filter.hpp index 787da34..3199a02 100644 --- a/modules/gpufilters/src/cuda/row_filter.hpp +++ b/modules/gpufilters/src/cuda/row_filter.hpp @@ -186,38 +186,38 @@ namespace filter { { 0, - row_filter::caller< 1, T, D, BrdRowReflect101>, - row_filter::caller< 2, T, D, BrdRowReflect101>, - row_filter::caller< 3, T, D, BrdRowReflect101>, - row_filter::caller< 4, T, D, BrdRowReflect101>, - row_filter::caller< 5, T, D, BrdRowReflect101>, - row_filter::caller< 6, T, D, BrdRowReflect101>, - row_filter::caller< 7, T, D, BrdRowReflect101>, - row_filter::caller< 8, T, D, BrdRowReflect101>, - row_filter::caller< 9, T, D, BrdRowReflect101>, - row_filter::caller<10, T, D, BrdRowReflect101>, - row_filter::caller<11, T, D, BrdRowReflect101>, - row_filter::caller<12, T, D, BrdRowReflect101>, - row_filter::caller<13, T, D, BrdRowReflect101>, - row_filter::caller<14, T, D, BrdRowReflect101>, - row_filter::caller<15, T, D, BrdRowReflect101>, - row_filter::caller<16, T, D, BrdRowReflect101>, - row_filter::caller<17, T, D, BrdRowReflect101>, - row_filter::caller<18, T, D, BrdRowReflect101>, - row_filter::caller<19, T, D, BrdRowReflect101>, - row_filter::caller<20, T, D, BrdRowReflect101>, - row_filter::caller<21, T, D, BrdRowReflect101>, - row_filter::caller<22, T, D, BrdRowReflect101>, - row_filter::caller<23, T, D, BrdRowReflect101>, - row_filter::caller<24, T, D, BrdRowReflect101>, - row_filter::caller<25, T, D, BrdRowReflect101>, - row_filter::caller<26, T, D, BrdRowReflect101>, - row_filter::caller<27, T, D, BrdRowReflect101>, - row_filter::caller<28, T, D, BrdRowReflect101>, - row_filter::caller<29, T, D, BrdRowReflect101>, - row_filter::caller<30, T, D, BrdRowReflect101>, - row_filter::caller<31, T, D, BrdRowReflect101>, - row_filter::caller<32, T, D, BrdRowReflect101> + row_filter::caller< 1, T, D, BrdRowConstant>, + row_filter::caller< 2, T, D, BrdRowConstant>, + row_filter::caller< 3, T, D, BrdRowConstant>, + row_filter::caller< 4, T, D, BrdRowConstant>, + row_filter::caller< 5, T, D, BrdRowConstant>, + row_filter::caller< 6, T, D, BrdRowConstant>, + row_filter::caller< 7, T, D, BrdRowConstant>, + row_filter::caller< 8, T, D, BrdRowConstant>, + row_filter::caller< 9, T, D, BrdRowConstant>, + row_filter::caller<10, T, D, BrdRowConstant>, + row_filter::caller<11, T, D, BrdRowConstant>, + row_filter::caller<12, T, D, BrdRowConstant>, + row_filter::caller<13, T, D, BrdRowConstant>, + row_filter::caller<14, T, D, BrdRowConstant>, + row_filter::caller<15, T, D, BrdRowConstant>, + row_filter::caller<16, T, D, BrdRowConstant>, + row_filter::caller<17, T, D, BrdRowConstant>, + row_filter::caller<18, T, D, BrdRowConstant>, + row_filter::caller<19, T, D, BrdRowConstant>, + row_filter::caller<20, T, D, BrdRowConstant>, + row_filter::caller<21, T, D, BrdRowConstant>, + row_filter::caller<22, T, D, BrdRowConstant>, + row_filter::caller<23, T, D, BrdRowConstant>, + row_filter::caller<24, T, D, BrdRowConstant>, + row_filter::caller<25, T, D, BrdRowConstant>, + row_filter::caller<26, T, D, BrdRowConstant>, + row_filter::caller<27, T, D, BrdRowConstant>, + row_filter::caller<28, T, D, BrdRowConstant>, + row_filter::caller<29, T, D, BrdRowConstant>, + row_filter::caller<30, T, D, BrdRowConstant>, + row_filter::caller<31, T, D, BrdRowConstant>, + row_filter::caller<32, T, D, BrdRowConstant> }, { 0, @@ -256,41 +256,6 @@ namespace filter }, { 0, - row_filter::caller< 1, T, D, BrdRowConstant>, - row_filter::caller< 2, T, D, BrdRowConstant>, - row_filter::caller< 3, T, D, BrdRowConstant>, - row_filter::caller< 4, T, D, BrdRowConstant>, - row_filter::caller< 5, T, D, BrdRowConstant>, - row_filter::caller< 6, T, D, BrdRowConstant>, - row_filter::caller< 7, T, D, BrdRowConstant>, - row_filter::caller< 8, T, D, BrdRowConstant>, - row_filter::caller< 9, T, D, BrdRowConstant>, - row_filter::caller<10, T, D, BrdRowConstant>, - row_filter::caller<11, T, D, BrdRowConstant>, - row_filter::caller<12, T, D, BrdRowConstant>, - row_filter::caller<13, T, D, BrdRowConstant>, - row_filter::caller<14, T, D, BrdRowConstant>, - row_filter::caller<15, T, D, BrdRowConstant>, - row_filter::caller<16, T, D, BrdRowConstant>, - row_filter::caller<17, T, D, BrdRowConstant>, - row_filter::caller<18, T, D, BrdRowConstant>, - row_filter::caller<19, T, D, BrdRowConstant>, - row_filter::caller<20, T, D, BrdRowConstant>, - row_filter::caller<21, T, D, BrdRowConstant>, - row_filter::caller<22, T, D, BrdRowConstant>, - row_filter::caller<23, T, D, BrdRowConstant>, - row_filter::caller<24, T, D, BrdRowConstant>, - row_filter::caller<25, T, D, BrdRowConstant>, - row_filter::caller<26, T, D, BrdRowConstant>, - row_filter::caller<27, T, D, BrdRowConstant>, - row_filter::caller<28, T, D, BrdRowConstant>, - row_filter::caller<29, T, D, BrdRowConstant>, - row_filter::caller<30, T, D, BrdRowConstant>, - row_filter::caller<31, T, D, BrdRowConstant>, - row_filter::caller<32, T, D, BrdRowConstant> - }, - { - 0, row_filter::caller< 1, T, D, BrdRowReflect>, row_filter::caller< 2, T, D, BrdRowReflect>, row_filter::caller< 3, T, D, BrdRowReflect>, @@ -358,6 +323,41 @@ namespace filter row_filter::caller<30, T, D, BrdRowWrap>, row_filter::caller<31, T, D, BrdRowWrap>, row_filter::caller<32, T, D, BrdRowWrap> + }, + { + 0, + row_filter::caller< 1, T, D, BrdRowReflect101>, + row_filter::caller< 2, T, D, BrdRowReflect101>, + row_filter::caller< 3, T, D, BrdRowReflect101>, + row_filter::caller< 4, T, D, BrdRowReflect101>, + row_filter::caller< 5, T, D, BrdRowReflect101>, + row_filter::caller< 6, T, D, BrdRowReflect101>, + row_filter::caller< 7, T, D, BrdRowReflect101>, + row_filter::caller< 8, T, D, BrdRowReflect101>, + row_filter::caller< 9, T, D, BrdRowReflect101>, + row_filter::caller<10, T, D, BrdRowReflect101>, + row_filter::caller<11, T, D, BrdRowReflect101>, + row_filter::caller<12, T, D, BrdRowReflect101>, + row_filter::caller<13, T, D, BrdRowReflect101>, + row_filter::caller<14, T, D, BrdRowReflect101>, + row_filter::caller<15, T, D, BrdRowReflect101>, + row_filter::caller<16, T, D, BrdRowReflect101>, + row_filter::caller<17, T, D, BrdRowReflect101>, + row_filter::caller<18, T, D, BrdRowReflect101>, + row_filter::caller<19, T, D, BrdRowReflect101>, + row_filter::caller<20, T, D, BrdRowReflect101>, + row_filter::caller<21, T, D, BrdRowReflect101>, + row_filter::caller<22, T, D, BrdRowReflect101>, + row_filter::caller<23, T, D, BrdRowReflect101>, + row_filter::caller<24, T, D, BrdRowReflect101>, + row_filter::caller<25, T, D, BrdRowReflect101>, + row_filter::caller<26, T, D, BrdRowReflect101>, + row_filter::caller<27, T, D, BrdRowReflect101>, + row_filter::caller<28, T, D, BrdRowReflect101>, + row_filter::caller<29, T, D, BrdRowReflect101>, + row_filter::caller<30, T, D, BrdRowReflect101>, + row_filter::caller<31, T, D, BrdRowReflect101>, + row_filter::caller<32, T, D, BrdRowReflect101> } }; diff --git a/modules/gpufilters/src/filtering.cpp b/modules/gpufilters/src/filtering.cpp index 6416325..8232ab8 100644 --- a/modules/gpufilters/src/filtering.cpp +++ b/modules/gpufilters/src/filtering.cpp @@ -783,9 +783,6 @@ Ptr cv::gpu::getLinearFilter_GPU(int srcType, int dstType, const CV_Assert(ksize.width * ksize.height <= 16 * 16); - int gpuBorderType; - CV_Assert( tryConvertToGpuBorderType(brd_type, gpuBorderType) ); - GpuMat gpu_krnl; normalizeKernel(kernel, gpu_krnl, CV_32F); @@ -815,7 +812,7 @@ Ptr cv::gpu::getLinearFilter_GPU(int srcType, int dstType, const break; } - return Ptr(new GpuFilter2D(ksize, anchor, func, gpu_krnl, gpuBorderType)); + return Ptr(new GpuFilter2D(ksize, anchor, func, gpu_krnl, brd_type)); } Ptr cv::gpu::createLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, Point anchor, int borderType) @@ -936,9 +933,6 @@ Ptr cv::gpu::getLinearRowFilter_GPU(int srcType, int bufType, CV_Assert( borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP ); - int gpuBorderType; - CV_Assert( tryConvertToGpuBorderType(borderType, gpuBorderType) ); - const int sdepth = CV_MAT_DEPTH(srcType); const int cn = CV_MAT_CN(srcType); CV_Assert( sdepth <= CV_64F && cn <= 4 ); @@ -955,7 +949,7 @@ Ptr cv::gpu::getLinearRowFilter_GPU(int srcType, int bufType, normalizeAnchor(anchor, ksize); - return Ptr(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, gpuBorderType)); + return Ptr(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, borderType)); } namespace @@ -1041,9 +1035,6 @@ Ptr cv::gpu::getLinearColumnFilter_GPU(int bufType, int ds CV_Assert( borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP ); - int gpuBorderType; - CV_Assert( tryConvertToGpuBorderType(borderType, gpuBorderType) ); - const int ddepth = CV_MAT_DEPTH(dstType); const int cn = CV_MAT_CN(dstType); CV_Assert( ddepth <= CV_64F && cn <= 4 ); @@ -1060,7 +1051,7 @@ Ptr cv::gpu::getLinearColumnFilter_GPU(int bufType, int ds normalizeAnchor(anchor, ksize); - return Ptr(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, gpuBorderType)); + return Ptr(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, borderType)); } Ptr cv::gpu::createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, const Mat& columnKernel, diff --git a/modules/gpuimgproc/src/bilateral_filter.cpp b/modules/gpuimgproc/src/bilateral_filter.cpp index 0c14987..c95dbe4 100644 --- a/modules/gpuimgproc/src/bilateral_filter.cpp +++ b/modules/gpuimgproc/src/bilateral_filter.cpp @@ -89,11 +89,8 @@ void cv::gpu::bilateralFilter(const GpuMat& src, GpuMat& dst, int kernel_size, f 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)); + func(src, dst, kernel_size, sigma_spatial, sigma_color, borderMode, StreamAccessor::getStream(s)); } #endif diff --git a/modules/gpuimgproc/src/cuda/bilateral_filter.cu b/modules/gpuimgproc/src/cuda/bilateral_filter.cu index 4449274..6aa5df2 100644 --- a/modules/gpuimgproc/src/cuda/bilateral_filter.cu +++ b/modules/gpuimgproc/src/cuda/bilateral_filter.cu @@ -150,11 +150,11 @@ namespace cv { namespace gpu { namespace cudev static caller_t funcs[] = { - bilateral_caller, - bilateral_caller, bilateral_caller, + bilateral_caller, bilateral_caller, bilateral_caller, + bilateral_caller }; funcs[borderMode](src, dst, kernel_size, gauss_spatial_coeff, gauss_color_coeff, stream); } diff --git a/modules/gpuimgproc/src/cuda/imgproc.cu b/modules/gpuimgproc/src/cuda/imgproc.cu index c47076f..3f39a43 100644 --- a/modules/gpuimgproc/src/cuda/imgproc.cu +++ b/modules/gpuimgproc/src/cuda/imgproc.cu @@ -269,15 +269,15 @@ namespace cv { namespace gpu { namespace cudev switch (border_type) { - case BORDER_REFLECT101_GPU: + case BORDER_REFLECT101: cornerHarris_kernel<<>>(block_size, k, dst, BrdRowReflect101(Dx.cols), BrdColReflect101(Dx.rows)); break; - case BORDER_REFLECT_GPU: + case BORDER_REFLECT: cornerHarris_kernel<<>>(block_size, k, dst, BrdRowReflect(Dx.cols), BrdColReflect(Dx.rows)); break; - case BORDER_REPLICATE_GPU: + case BORDER_REPLICATE: cornerHarris_kernel<<>>(block_size, k, dst); break; } @@ -381,15 +381,15 @@ namespace cv { namespace gpu { namespace cudev switch (border_type) { - case BORDER_REFLECT101_GPU: + case BORDER_REFLECT101: cornerMinEigenVal_kernel<<>>(block_size, dst, BrdRowReflect101(Dx.cols), BrdColReflect101(Dx.rows)); break; - case BORDER_REFLECT_GPU: + case BORDER_REFLECT: cornerMinEigenVal_kernel<<>>(block_size, dst, BrdRowReflect(Dx.cols), BrdColReflect(Dx.rows)); break; - case BORDER_REPLICATE_GPU: + case BORDER_REPLICATE: cornerMinEigenVal_kernel<<>>(block_size, dst); break; } diff --git a/modules/gpuimgproc/src/imgproc.cpp b/modules/gpuimgproc/src/imgproc.cpp index 939b149..100d091 100644 --- a/modules/gpuimgproc/src/imgproc.cpp +++ b/modules/gpuimgproc/src/imgproc.cpp @@ -552,14 +552,11 @@ void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuMat& D CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT); - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType)); - extractCovData(src, Dx, Dy, buf, blockSize, ksize, borderType, stream); dst.create(src.size(), CV_32F); - cornerHarris_gpu(blockSize, static_cast(k), Dx, Dy, dst, gpuBorderType, StreamAccessor::getStream(stream)); + cornerHarris_gpu(blockSize, static_cast(k), Dx, Dy, dst, borderType, StreamAccessor::getStream(stream)); } void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType) @@ -580,14 +577,11 @@ void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, GpuMat& Dx, GpuM CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT); - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType)); - extractCovData(src, Dx, Dy, buf, blockSize, ksize, borderType, stream); dst.create(src.size(), CV_32F); - cornerMinEigenVal_gpu(blockSize, Dx, Dy, dst, gpuBorderType, StreamAccessor::getStream(stream)); + cornerMinEigenVal_gpu(blockSize, Dx, Dy, dst, borderType, StreamAccessor::getStream(stream)); } diff --git a/modules/gpuoptflow/src/cuda/optical_flow_farneback.cu b/modules/gpuoptflow/src/cuda/optical_flow_farneback.cu index e7ff3a0..68a58c1 100644 --- a/modules/gpuoptflow/src/cuda/optical_flow_farneback.cu +++ b/modules/gpuoptflow/src/cuda/optical_flow_farneback.cu @@ -525,8 +525,11 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback static const caller_t callers[] = { - gaussianBlurCaller >, + 0 /*gaussianBlurCaller >*/, gaussianBlurCaller >, + 0 /*gaussianBlurCaller >*/, + 0 /*gaussianBlurCaller >*/, + gaussianBlurCaller > }; callers[borderMode](src, ksizeHalf, dst, stream); @@ -620,8 +623,11 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback static const caller_t callers[] = { - gaussianBlur5Caller,256>, + 0 /*gaussianBlur5Caller,256>*/, gaussianBlur5Caller,256>, + 0 /*gaussianBlur5Caller,256>*/, + 0 /*gaussianBlur5Caller,256>*/, + gaussianBlur5Caller,256> }; callers[borderMode](src, ksizeHalf, dst, stream); @@ -634,8 +640,11 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback static const caller_t callers[] = { - gaussianBlur5Caller,128>, + 0 /*gaussianBlur5Caller,128>*/, gaussianBlur5Caller,128>, + 0 /*gaussianBlur5Caller,128>*/, + 0 /*gaussianBlur5Caller,128>*/, + gaussianBlur5Caller,128> }; callers[borderMode](src, ksizeHalf, dst, stream); diff --git a/modules/gpuoptflow/src/optical_flow_farneback.cpp b/modules/gpuoptflow/src/optical_flow_farneback.cpp index 8dbf25b..efe2436 100644 --- a/modules/gpuoptflow/src/optical_flow_farneback.cpp +++ b/modules/gpuoptflow/src/optical_flow_farneback.cpp @@ -192,10 +192,10 @@ void cv::gpu::FarnebackOpticalFlow::updateFlow_gaussianBlur( { if (deviceSupports(FEATURE_SET_COMPUTE_12)) cudev::optflow_farneback::gaussianBlur5Gpu( - M, blockSize/2, bufM, BORDER_REPLICATE_GPU, S(streams[0])); + M, blockSize/2, bufM, BORDER_REPLICATE, S(streams[0])); else cudev::optflow_farneback::gaussianBlur5Gpu_CC11( - M, blockSize/2, bufM, BORDER_REPLICATE_GPU, S(streams[0])); + M, blockSize/2, bufM, BORDER_REPLICATE, S(streams[0])); swap(M, bufM); cudev::optflow_farneback::updateFlowGpu(M, flowx, flowy, S(streams[0])); @@ -366,7 +366,7 @@ void cv::gpu::FarnebackOpticalFlow::operator ()( for (int i = 0; i < 2; i++) { cudev::optflow_farneback::gaussianBlurGpu( - frames_[i], smoothSize/2, blurredFrame[i], BORDER_REFLECT101_GPU, S(streams[i])); + frames_[i], smoothSize/2, blurredFrame[i], BORDER_REFLECT101, S(streams[i])); #if ENABLE_GPU_RESIZE resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR, streams[i]); #else diff --git a/modules/gpuwarping/src/cuda/remap.cu b/modules/gpuwarping/src/cuda/remap.cu index dd2c669..c4ea317 100644 --- a/modules/gpuwarping/src/cuda/remap.cu +++ b/modules/gpuwarping/src/cuda/remap.cu @@ -212,25 +212,25 @@ namespace cv { namespace gpu { namespace cudev static const caller_t callers[3][5] = { { - RemapDispatcher::call, - RemapDispatcher::call, RemapDispatcher::call, + RemapDispatcher::call, RemapDispatcher::call, - RemapDispatcher::call + RemapDispatcher::call, + RemapDispatcher::call }, { - RemapDispatcher::call, - RemapDispatcher::call, RemapDispatcher::call, + RemapDispatcher::call, RemapDispatcher::call, - RemapDispatcher::call + RemapDispatcher::call, + RemapDispatcher::call }, { - RemapDispatcher::call, - RemapDispatcher::call, RemapDispatcher::call, + RemapDispatcher::call, RemapDispatcher::call, - RemapDispatcher::call + RemapDispatcher::call, + RemapDispatcher::call } }; diff --git a/modules/gpuwarping/src/cuda/warp.cu b/modules/gpuwarping/src/cuda/warp.cu index 8c5a067..83db79e 100644 --- a/modules/gpuwarping/src/cuda/warp.cu +++ b/modules/gpuwarping/src/cuda/warp.cu @@ -281,25 +281,25 @@ namespace cv { namespace gpu { namespace cudev static const func_t funcs[3][5] = { { - WarpDispatcher::call, - WarpDispatcher::call, WarpDispatcher::call, + WarpDispatcher::call, WarpDispatcher::call, - WarpDispatcher::call + WarpDispatcher::call, + WarpDispatcher::call }, { - WarpDispatcher::call, - WarpDispatcher::call, WarpDispatcher::call, + WarpDispatcher::call, WarpDispatcher::call, - WarpDispatcher::call + WarpDispatcher::call, + WarpDispatcher::call }, { - WarpDispatcher::call, - WarpDispatcher::call, WarpDispatcher::call, + WarpDispatcher::call, WarpDispatcher::call, - WarpDispatcher::call + WarpDispatcher::call, + WarpDispatcher::call } }; diff --git a/modules/gpuwarping/src/remap.cpp b/modules/gpuwarping/src/remap.cpp index 3157665..131f937 100644 --- a/modules/gpuwarping/src/remap.cpp +++ b/modules/gpuwarping/src/remap.cpp @@ -83,9 +83,6 @@ void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const Gp const func_t func = funcs[src.depth()][src.channels() - 1]; CV_Assert(func != 0); - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType)); - dst.create(xmap.size(), src.type()); Scalar_ borderValueFloat; @@ -96,7 +93,7 @@ void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const Gp src.locateROI(wholeSize, ofs); func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, xmap, ymap, - dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(stream), deviceSupports(FEATURE_SET_COMPUTE_20)); + dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(stream), deviceSupports(FEATURE_SET_COMPUTE_20)); } #endif // HAVE_CUDA diff --git a/modules/gpuwarping/src/warp.cpp b/modules/gpuwarping/src/warp.cpp index e15c11b..b3c44e5 100644 --- a/modules/gpuwarping/src/warp.cpp +++ b/modules/gpuwarping/src/warp.cpp @@ -289,9 +289,6 @@ void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsiz const func_t func = funcs[src.depth()][src.channels() - 1]; CV_Assert(func != 0); - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType)); - float coeffs[2 * 3]; Mat coeffsMat(2, 3, CV_32F, (void*)coeffs); @@ -308,7 +305,7 @@ void cv::gpu::warpAffine(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsiz borderValueFloat = borderValue; func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, coeffs, - dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20)); + dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20)); } } @@ -427,9 +424,6 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size const func_t func = funcs[src.depth()][src.channels() - 1]; CV_Assert(func != 0); - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType)); - float coeffs[3 * 3]; Mat coeffsMat(3, 3, CV_32F, (void*)coeffs); @@ -446,7 +440,7 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size borderValueFloat = borderValue; func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, coeffs, - dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20)); + dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20)); } } diff --git a/modules/photo/src/cuda/nlm.cu b/modules/photo/src/cuda/nlm.cu index 92bfccf..0304469 100644 --- a/modules/photo/src/cuda/nlm.cu +++ b/modules/photo/src/cuda/nlm.cu @@ -161,11 +161,11 @@ namespace cv { namespace gpu { namespace cudev static func_t funcs[] = { - nlm_caller, - nlm_caller, nlm_caller, + nlm_caller, nlm_caller, nlm_caller, + nlm_caller }; funcs[borderMode](src, dst, search_radius, block_radius, h, stream); } diff --git a/modules/photo/src/denoising_gpu.cpp b/modules/photo/src/denoising_gpu.cpp index 2164731..65d6f81 100644 --- a/modules/photo/src/denoising_gpu.cpp +++ b/modules/photo/src/denoising_gpu.cpp @@ -85,11 +85,8 @@ void cv::gpu::nonLocalMeans(const GpuMat& src, GpuMat& dst, float h, int search_ 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)); - dst.create(src.size(), src.type()); - func(src, dst, search_window/2, block_window/2, h, gpuBorderType, StreamAccessor::getStream(s)); + func(src, dst, search_window/2, block_window/2, h, borderMode, StreamAccessor::getStream(s)); } namespace cv { namespace gpu { namespace cudev -- 2.7.4