removed gpu BORDER_* constants
authorVladislav Vinogradov <vlad.vinogradov@itseez.com>
Thu, 11 Apr 2013 10:53:35 +0000 (14:53 +0400)
committerVladislav Vinogradov <vlad.vinogradov@itseez.com>
Thu, 18 Apr 2013 07:33:34 +0000 (11:33 +0400)
21 files changed:
modules/core/include/opencv2/core/cuda/common.hpp
modules/core/include/opencv2/core/gpu_private.hpp
modules/core/src/gpumat.cpp
modules/gpuarithm/src/arithm.cpp
modules/gpuarithm/src/cuda/copy_make_border.cu
modules/gpufilters/src/cuda/column_filter.hpp
modules/gpufilters/src/cuda/filter2d.cu
modules/gpufilters/src/cuda/row_filter.hpp
modules/gpufilters/src/filtering.cpp
modules/gpuimgproc/src/bilateral_filter.cpp
modules/gpuimgproc/src/cuda/bilateral_filter.cu
modules/gpuimgproc/src/cuda/imgproc.cu
modules/gpuimgproc/src/imgproc.cpp
modules/gpuoptflow/src/cuda/optical_flow_farneback.cu
modules/gpuoptflow/src/optical_flow_farneback.cpp
modules/gpuwarping/src/cuda/remap.cu
modules/gpuwarping/src/cuda/warp.cu
modules/gpuwarping/src/remap.cpp
modules/gpuwarping/src/warp.cpp
modules/photo/src/cuda/nlm.cu
modules/photo/src/denoising_gpu.cpp

index 774500e..434a3eb 100644 (file)
@@ -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)
index be194f5..7692bc2 100644 (file)
 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
index c5d4d7a..11bb419 100644 (file)
@@ -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
-}
index cc85cc7..908d963 100644 (file)
@@ -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);
     }
 }
 
index ed90e9e..d772e09 100644 (file)
@@ -86,11 +86,11 @@ namespace cv { namespace gpu { namespace cudev
 
             static const caller_t callers[5] =
             {
-                CopyMakeBorderDispatcher<BrdReflect101, vec_type>::call,
-                CopyMakeBorderDispatcher<BrdReplicate, vec_type>::call,
                 CopyMakeBorderDispatcher<BrdConstant, vec_type>::call,
+                CopyMakeBorderDispatcher<BrdReplicate, vec_type>::call,
                 CopyMakeBorderDispatcher<BrdReflect, vec_type>::call,
-                CopyMakeBorderDispatcher<BrdWrap, vec_type>::call
+                CopyMakeBorderDispatcher<BrdWrap, vec_type>::call,
+                CopyMakeBorderDispatcher<BrdReflect101, vec_type>::call
             };
 
             callers[borderMode](PtrStepSz<vec_type>(src), PtrStepSz<vec_type>(dst), top, left, borderValue, stream);
index 39b6d47..6f10c36 100644 (file)
@@ -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>
             }
         };
 
index 0bb5fcd..80c93c5 100644 (file)
@@ -131,11 +131,11 @@ namespace cv { namespace gpu { namespace cudev
             typedef void (*func_t)(const PtrStepSz<T> srcWhole, int xoff, int yoff, PtrStepSz<D> dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream);
             static const func_t funcs[] =
             {
-                Filter2DCaller<T, D, BrdReflect101>::call,
-                Filter2DCaller<T, D, BrdReplicate>::call,
                 Filter2DCaller<T, D, BrdConstant>::call,
+                Filter2DCaller<T, D, BrdReplicate>::call,
                 Filter2DCaller<T, D, BrdReflect>::call,
-                Filter2DCaller<T, D, BrdWrap>::call
+                Filter2DCaller<T, D, BrdWrap>::call,
+                Filter2DCaller<T, D, BrdReflect101>::call
             };
 
             if (stream == 0)
index 787da34..3199a02 100644 (file)
@@ -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>
             }
         };
 
index 6416325..8232ab8 100644 (file)
@@ -783,9 +783,6 @@ Ptr<BaseFilter_GPU> 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<BaseFilter_GPU> cv::gpu::getLinearFilter_GPU(int srcType, int dstType, const
         break;
     }
 
-    return Ptr<BaseFilter_GPU>(new GpuFilter2D(ksize, anchor, func, gpu_krnl, gpuBorderType));
+    return Ptr<BaseFilter_GPU>(new GpuFilter2D(ksize, anchor, func, gpu_krnl, brd_type));
 }
 
 Ptr<FilterEngine_GPU> cv::gpu::createLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, Point anchor, int borderType)
@@ -936,9 +933,6 @@ Ptr<BaseRowFilter_GPU> 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<BaseRowFilter_GPU> cv::gpu::getLinearRowFilter_GPU(int srcType, int bufType,
 
     normalizeAnchor(anchor, ksize);
 
-    return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, gpuBorderType));
+    return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, borderType));
 }
 
 namespace
@@ -1041,9 +1035,6 @@ Ptr<BaseColumnFilter_GPU> 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<BaseColumnFilter_GPU> cv::gpu::getLinearColumnFilter_GPU(int bufType, int ds
 
     normalizeAnchor(anchor, ksize);
 
-    return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, gpuBorderType));
+    return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, borderType));
 }
 
 Ptr<FilterEngine_GPU> cv::gpu::createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, const Mat& columnKernel,
index 0c14987..c95dbe4 100644 (file)
@@ -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
index 4449274..6aa5df2 100644 (file)
@@ -150,11 +150,11 @@ namespace cv { namespace gpu { namespace cudev
 
             static caller_t funcs[] =
             {
-                bilateral_caller<T, BrdReflect101>,
-                bilateral_caller<T, BrdReplicate>,
                 bilateral_caller<T, BrdConstant>,
+                bilateral_caller<T, BrdReplicate>,
                 bilateral_caller<T, BrdReflect>,
                 bilateral_caller<T, BrdWrap>,
+                bilateral_caller<T, BrdReflect101>
             };
             funcs[borderMode](src, dst, kernel_size, gauss_spatial_coeff, gauss_color_coeff, stream);
         }
index c47076f..3f39a43 100644 (file)
@@ -269,15 +269,15 @@ namespace cv { namespace gpu { namespace cudev
 
             switch (border_type)
             {
-            case BORDER_REFLECT101_GPU:
+            case BORDER_REFLECT101:
                 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows));
                 break;
 
-            case BORDER_REFLECT_GPU:
+            case BORDER_REFLECT:
                 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows));
                 break;
 
-            case BORDER_REPLICATE_GPU:
+            case BORDER_REPLICATE:
                 cornerHarris_kernel<<<grid, block, 0, stream>>>(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<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows));
                 break;
 
-            case BORDER_REFLECT_GPU:
+            case BORDER_REFLECT:
                 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows));
                 break;
 
-            case BORDER_REPLICATE_GPU:
+            case BORDER_REPLICATE:
                 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst);
                 break;
             }
index 939b149..100d091 100644 (file)
@@ -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<float>(k), Dx, Dy, dst, gpuBorderType, StreamAccessor::getStream(stream));
+    cornerHarris_gpu(blockSize, static_cast<float>(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));
 }
 
 
index e7ff3a0..68a58c1 100644 (file)
@@ -525,8 +525,11 @@ namespace cv { namespace gpu { namespace cudev { namespace optflow_farneback
 
         static const caller_t callers[] =
         {
-            gaussianBlurCaller<BrdReflect101<float> >,
+            0 /*gaussianBlurCaller<BrdConstant<float> >*/,
             gaussianBlurCaller<BrdReplicate<float> >,
+            0 /*gaussianBlurCaller<BrdReflect<float> >*/,
+            0 /*gaussianBlurCaller<BrdWrap<float> >*/,
+            gaussianBlurCaller<BrdReflect101<float> >
         };
 
         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<BrdReflect101<float>,256>,
+            0 /*gaussianBlur5Caller<BrdConstant<float>,256>*/,
             gaussianBlur5Caller<BrdReplicate<float>,256>,
+            0 /*gaussianBlur5Caller<BrdReflect<float>,256>*/,
+            0 /*gaussianBlur5Caller<BrdWrap<float>,256>*/,
+            gaussianBlur5Caller<BrdReflect101<float>,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<BrdReflect101<float>,128>,
+            0 /*gaussianBlur5Caller<BrdConstant<float>,128>*/,
             gaussianBlur5Caller<BrdReplicate<float>,128>,
+            0 /*gaussianBlur5Caller<BrdReflect<float>,128>*/,
+            0 /*gaussianBlur5Caller<BrdWrap<float>,128>*/,
+            gaussianBlur5Caller<BrdReflect101<float>,128>
         };
 
         callers[borderMode](src, ksizeHalf, dst, stream);
index 8dbf25b..efe2436 100644 (file)
@@ -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
index dd2c669..c4ea317 100644 (file)
@@ -212,25 +212,25 @@ namespace cv { namespace gpu { namespace cudev
             static const caller_t callers[3][5] =
             {
                 {
-                    RemapDispatcher<PointFilter, BrdReflect101, T>::call,
-                    RemapDispatcher<PointFilter, BrdReplicate, T>::call,
                     RemapDispatcher<PointFilter, BrdConstant, T>::call,
+                    RemapDispatcher<PointFilter, BrdReplicate, T>::call,
                     RemapDispatcher<PointFilter, BrdReflect, T>::call,
-                    RemapDispatcher<PointFilter, BrdWrap, T>::call
+                    RemapDispatcher<PointFilter, BrdWrap, T>::call,
+                    RemapDispatcher<PointFilter, BrdReflect101, T>::call
                 },
                 {
-                    RemapDispatcher<LinearFilter, BrdReflect101, T>::call,
-                    RemapDispatcher<LinearFilter, BrdReplicate, T>::call,
                     RemapDispatcher<LinearFilter, BrdConstant, T>::call,
+                    RemapDispatcher<LinearFilter, BrdReplicate, T>::call,
                     RemapDispatcher<LinearFilter, BrdReflect, T>::call,
-                    RemapDispatcher<LinearFilter, BrdWrap, T>::call
+                    RemapDispatcher<LinearFilter, BrdWrap, T>::call,
+                    RemapDispatcher<LinearFilter, BrdReflect101, T>::call
                 },
                 {
-                    RemapDispatcher<CubicFilter, BrdReflect101, T>::call,
-                    RemapDispatcher<CubicFilter, BrdReplicate, T>::call,
                     RemapDispatcher<CubicFilter, BrdConstant, T>::call,
+                    RemapDispatcher<CubicFilter, BrdReplicate, T>::call,
                     RemapDispatcher<CubicFilter, BrdReflect, T>::call,
-                    RemapDispatcher<CubicFilter, BrdWrap, T>::call
+                    RemapDispatcher<CubicFilter, BrdWrap, T>::call,
+                    RemapDispatcher<CubicFilter, BrdReflect101, T>::call
                 }
             };
 
index 8c5a067..83db79e 100644 (file)
@@ -281,25 +281,25 @@ namespace cv { namespace gpu { namespace cudev
             static const func_t funcs[3][5] =
             {
                 {
-                    WarpDispatcher<Transform, PointFilter, BrdReflect101, T>::call,
-                    WarpDispatcher<Transform, PointFilter, BrdReplicate, T>::call,
                     WarpDispatcher<Transform, PointFilter, BrdConstant, T>::call,
+                    WarpDispatcher<Transform, PointFilter, BrdReplicate, T>::call,
                     WarpDispatcher<Transform, PointFilter, BrdReflect, T>::call,
-                    WarpDispatcher<Transform, PointFilter, BrdWrap, T>::call
+                    WarpDispatcher<Transform, PointFilter, BrdWrap, T>::call,
+                    WarpDispatcher<Transform, PointFilter, BrdReflect101, T>::call
                 },
                 {
-                    WarpDispatcher<Transform, LinearFilter, BrdReflect101, T>::call,
-                    WarpDispatcher<Transform, LinearFilter, BrdReplicate, T>::call,
                     WarpDispatcher<Transform, LinearFilter, BrdConstant, T>::call,
+                    WarpDispatcher<Transform, LinearFilter, BrdReplicate, T>::call,
                     WarpDispatcher<Transform, LinearFilter, BrdReflect, T>::call,
-                    WarpDispatcher<Transform, LinearFilter, BrdWrap, T>::call
+                    WarpDispatcher<Transform, LinearFilter, BrdWrap, T>::call,
+                    WarpDispatcher<Transform, LinearFilter, BrdReflect101, T>::call
                 },
                 {
-                    WarpDispatcher<Transform, CubicFilter, BrdReflect101, T>::call,
-                    WarpDispatcher<Transform, CubicFilter, BrdReplicate, T>::call,
                     WarpDispatcher<Transform, CubicFilter, BrdConstant, T>::call,
+                    WarpDispatcher<Transform, CubicFilter, BrdReplicate, T>::call,
                     WarpDispatcher<Transform, CubicFilter, BrdReflect, T>::call,
-                    WarpDispatcher<Transform, CubicFilter, BrdWrap, T>::call
+                    WarpDispatcher<Transform, CubicFilter, BrdWrap, T>::call,
+                    WarpDispatcher<Transform, CubicFilter, BrdReflect101, T>::call
                 }
             };
 
index 3157665..131f937 100644 (file)
@@ -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_<float> 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
index e15c11b..b3c44e5 100644 (file)
@@ -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));
     }
 }
 
index 92bfccf..0304469 100644 (file)
@@ -161,11 +161,11 @@ namespace cv { namespace gpu { namespace cudev
 
             static func_t funcs[] =
             {
-                nlm_caller<T, BrdReflect101>,
-                nlm_caller<T, BrdReplicate>,
                 nlm_caller<T, BrdConstant>,
+                nlm_caller<T, BrdReplicate>,
                 nlm_caller<T, BrdReflect>,
                 nlm_caller<T, BrdWrap>,
+                nlm_caller<T, BrdReflect101>
             };
             funcs[borderMode](src, dst, search_radius, block_radius, h, stream);
         }
index 2164731..65d6f81 100644 (file)
@@ -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