Merge pull request #836 from jet47:gpu-modules

cmake/OpenCVDetectCUDA.cmake

@@ -28,6 +28,9 @@ if(CUDA_FOUND)
 
   if(WITH_NVCUVID)
     find_cuda_helper_libs(nvcuvid)
+    if(WIN32)
+      find_cuda_helper_libs(nvcuvenc)
+    endif()
     set(HAVE_NVCUVID 1)
   endif()
 

modules/core/include/opencv2/core/cuda/border_interpolate.hpp

@@ -73,8 +73,8 @@ namespace cv { namespace gpu { namespace cudev
             return (x >= 0 && x < width) ? saturate_cast<D>(data[x]) : val;
         }
 
-        const int width;
-        const D val;
+        int width;
+        D val;
     };
 
     template <typename D> struct BrdColConstant
@@ -98,8 +98,8 @@ namespace cv { namespace gpu { namespace cudev
             return (y >= 0 && y < height) ? saturate_cast<D>(*(const T*)((const char*)data + y * step)) : val;
         }
 
-        const int height;
-        const D val;
+        int height;
+        D val;
     };
 
     template <typename D> struct BrdConstant
@@ -120,9 +120,9 @@ namespace cv { namespace gpu { namespace cudev
             return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val;
         }
 
-        const int height;
-        const int width;
-        const D val;
+        int height;
+        int width;
+        D val;
     };
 
     //////////////////////////////////////////////////////////////
@@ -165,7 +165,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(data[idx_col(x)]);
         }
 
-        const int last_col;
+        int last_col;
     };
 
     template <typename D> struct BrdColReplicate
@@ -205,7 +205,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(*(const T*)((const char*)data + idx_row(y) * step));
         }
 
-        const int last_row;
+        int last_row;
     };
 
     template <typename D> struct BrdReplicate
@@ -255,8 +255,8 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(src(idx_row(y), idx_col(x)));
         }
 
-        const int last_row;
-        const int last_col;
+        int last_row;
+        int last_col;
     };
 
     //////////////////////////////////////////////////////////////
@@ -299,7 +299,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(data[idx_col(x)]);
         }
 
-        const int last_col;
+        int last_col;
     };
 
     template <typename D> struct BrdColReflect101
@@ -339,7 +339,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step));
         }
 
-        const int last_row;
+        int last_row;
     };
 
     template <typename D> struct BrdReflect101
@@ -389,8 +389,8 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(src(idx_row(y), idx_col(x)));
         }
 
-        const int last_row;
-        const int last_col;
+        int last_row;
+        int last_col;
     };
 
     //////////////////////////////////////////////////////////////
@@ -433,7 +433,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(data[idx_col(x)]);
         }
 
-        const int last_col;
+        int last_col;
     };
 
     template <typename D> struct BrdColReflect
@@ -473,7 +473,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step));
         }
 
-        const int last_row;
+        int last_row;
     };
 
     template <typename D> struct BrdReflect
@@ -523,8 +523,8 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(src(idx_row(y), idx_col(x)));
         }
 
-        const int last_row;
-        const int last_col;
+        int last_row;
+        int last_col;
     };
 
     //////////////////////////////////////////////////////////////
@@ -567,7 +567,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(data[idx_col(x)]);
         }
 
-        const int width;
+        int width;
     };
 
     template <typename D> struct BrdColWrap
@@ -607,7 +607,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(*(const D*)((const char*)data + idx_row(y) * step));
         }
 
-        const int height;
+        int height;
     };
 
     template <typename D> struct BrdWrap
@@ -664,8 +664,8 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<D>(src(idx_row(y), idx_col(x)));
         }
 
-        const int height;
-        const int width;
+        int height;
+        int width;
     };
 
     //////////////////////////////////////////////////////////////
@@ -683,8 +683,8 @@ namespace cv { namespace gpu { namespace cudev
             return b.at(y, x, ptr);
         }
 
-        const Ptr2D ptr;
-        const B b;
+        Ptr2D ptr;
+        B b;
     };
 
     // under win32 there is some bug with templated types that passed as kernel parameters
@@ -704,10 +704,10 @@ namespace cv { namespace gpu { namespace cudev
             return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val;
         }
 
-        const Ptr2D src;
-        const int height;
-        const int width;
-        const D val;
+        Ptr2D src;
+        int height;
+        int width;
+        D val;
     };
 }}} // namespace cv { namespace gpu { namespace cudev
 

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)

modules/core/include/opencv2/core/cuda/emulation.hpp

@@ -43,6 +43,7 @@
 #ifndef OPENCV_GPU_EMULATION_HPP_
 #define OPENCV_GPU_EMULATION_HPP_
 
+#include "common.hpp"
 #include "warp_reduce.hpp"
 
 namespace cv { namespace gpu { namespace cudev
@@ -131,8 +132,130 @@ namespace cv { namespace gpu { namespace cudev
                 return ::atomicMin(address, val);
 #endif
             }
+        }; // struct cmem
+
+        struct glob
+        {
+            static __device__ __forceinline__ int atomicAdd(int* address, int val)
+            {
+                return ::atomicAdd(address, val);
+            }
+            static __device__ __forceinline__ unsigned int atomicAdd(unsigned int* address, unsigned int val)
+            {
+                return ::atomicAdd(address, val);
+            }
+            static __device__ __forceinline__ float atomicAdd(float* address, float val)
+            {
+            #if __CUDA_ARCH__ >= 200
+                return ::atomicAdd(address, val);
+            #else
+                int* address_as_i = (int*) address;
+                int old = *address_as_i, assumed;
+                do {
+                    assumed = old;
+                    old = ::atomicCAS(address_as_i, assumed,
+                        __float_as_int(val + __int_as_float(assumed)));
+                } while (assumed != old);
+                return __int_as_float(old);
+            #endif
+            }
+            static __device__ __forceinline__ double atomicAdd(double* address, double val)
+            {
+            #if __CUDA_ARCH__ >= 130
+                unsigned long long int* address_as_ull = (unsigned long long int*) address;
+                unsigned long long int old = *address_as_ull, assumed;
+                do {
+                    assumed = old;
+                    old = ::atomicCAS(address_as_ull, assumed,
+                        __double_as_longlong(val + __longlong_as_double(assumed)));
+                } while (assumed != old);
+                return __longlong_as_double(old);
+            #else
+                (void) address;
+                (void) val;
+                return 0.0;
+            #endif
+            }
+
+            static __device__ __forceinline__ int atomicMin(int* address, int val)
+            {
+                return ::atomicMin(address, val);
+            }
+            static __device__ __forceinline__ float atomicMin(float* address, float val)
+            {
+            #if __CUDA_ARCH__ >= 120
+                int* address_as_i = (int*) address;
+                int old = *address_as_i, assumed;
+                do {
+                    assumed = old;
+                    old = ::atomicCAS(address_as_i, assumed,
+                        __float_as_int(::fminf(val, __int_as_float(assumed))));
+                } while (assumed != old);
+                return __int_as_float(old);
+            #else
+                (void) address;
+                (void) val;
+                return 0.0f;
+            #endif
+            }
+            static __device__ __forceinline__ double atomicMin(double* address, double val)
+            {
+            #if __CUDA_ARCH__ >= 130
+                unsigned long long int* address_as_ull = (unsigned long long int*) address;
+                unsigned long long int old = *address_as_ull, assumed;
+                do {
+                    assumed = old;
+                    old = ::atomicCAS(address_as_ull, assumed,
+                        __double_as_longlong(::fmin(val, __longlong_as_double(assumed))));
+                } while (assumed != old);
+                return __longlong_as_double(old);
+            #else
+                (void) address;
+                (void) val;
+                return 0.0;
+            #endif
+            }
+
+            static __device__ __forceinline__ int atomicMax(int* address, int val)
+            {
+                return ::atomicMax(address, val);
+            }
+            static __device__ __forceinline__ float atomicMax(float* address, float val)
+            {
+            #if __CUDA_ARCH__ >= 120
+                int* address_as_i = (int*) address;
+                int old = *address_as_i, assumed;
+                do {
+                    assumed = old;
+                    old = ::atomicCAS(address_as_i, assumed,
+                        __float_as_int(::fmaxf(val, __int_as_float(assumed))));
+                } while (assumed != old);
+                return __int_as_float(old);
+            #else
+                (void) address;
+                (void) val;
+                return 0.0f;
+            #endif
+            }
+            static __device__ __forceinline__ double atomicMax(double* address, double val)
+            {
+            #if __CUDA_ARCH__ >= 130
+                unsigned long long int* address_as_ull = (unsigned long long int*) address;
+                unsigned long long int old = *address_as_ull, assumed;
+                do {
+                    assumed = old;
+                    old = ::atomicCAS(address_as_ull, assumed,
+                        __double_as_longlong(::fmax(val, __longlong_as_double(assumed))));
+                } while (assumed != old);
+                return __longlong_as_double(old);
+            #else
+                (void) address;
+                (void) val;
+                return 0.0;
+            #endif
+            }
         };
-    };
+    }; //struct Emulation
 }}} // namespace cv { namespace gpu { namespace cudev
 
 #endif /* OPENCV_GPU_EMULATION_HPP_ */

modules/core/include/opencv2/core/cuda/filters.hpp

@@ -67,7 +67,7 @@ namespace cv { namespace gpu { namespace cudev
             return src(__float2int_rz(y), __float2int_rz(x));
         }
 
-        const Ptr2D src;
+        Ptr2D src;
     };
 
     template <typename Ptr2D> struct LinearFilter
@@ -107,7 +107,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<elem_type>(out);
         }
 
-        const Ptr2D src;
+        Ptr2D src;
     };
 
     template <typename Ptr2D> struct CubicFilter
@@ -166,7 +166,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<elem_type>(res);
         }
 
-        const Ptr2D src;
+        Ptr2D src;
     };
     // for integer scaling
     template <typename Ptr2D> struct IntegerAreaFilter
@@ -203,7 +203,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<elem_type>(out);
         }
 
-        const Ptr2D src;
+        Ptr2D src;
         float scale_x, scale_y ,scale;
     };
 
@@ -269,7 +269,7 @@ namespace cv { namespace gpu { namespace cudev
             return saturate_cast<elem_type>(out);
         }
 
-        const Ptr2D src;
+        Ptr2D src;
         float scale_x, scale_y;
         int width, haight;
     };

modules/core/include/opencv2/core/cuda/functional.hpp

@@ -552,8 +552,8 @@ namespace cv { namespace gpu { namespace cudev
 
         __device__ __forceinline__ thresh_binary_func():unary_function<T, T>(){}
 
-        const T thresh;
-        const T maxVal;
+        T thresh;
+        T maxVal;
     };
 
     template <typename T> struct thresh_binary_inv_func : unary_function<T, T>
@@ -570,8 +570,8 @@ namespace cv { namespace gpu { namespace cudev
 
         __device__ __forceinline__ thresh_binary_inv_func():unary_function<T, T>(){}
 
-        const T thresh;
-        const T maxVal;
+        T thresh;
+        T maxVal;
     };
 
     template <typename T> struct thresh_trunc_func : unary_function<T, T>
@@ -588,7 +588,7 @@ namespace cv { namespace gpu { namespace cudev
 
         __device__ __forceinline__ thresh_trunc_func():unary_function<T, T>(){}
 
-        const T thresh;
+        T thresh;
     };
 
     template <typename T> struct thresh_to_zero_func : unary_function<T, T>
@@ -604,7 +604,7 @@ namespace cv { namespace gpu { namespace cudev
 
         __device__ __forceinline__ thresh_to_zero_func():unary_function<T, T>(){}
 
-        const T thresh;
+        T thresh;
     };
 
     template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T>
@@ -620,7 +620,7 @@ namespace cv { namespace gpu { namespace cudev
 
         __device__ __forceinline__ thresh_to_zero_inv_func():unary_function<T, T>(){}
 
-        const T thresh;
+        T thresh;
     };
 //bound!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! ============>
     // Function Object Adaptors
@@ -636,7 +636,7 @@ namespace cv { namespace gpu { namespace cudev
         __device__ __forceinline__ unary_negate(const unary_negate& other) : unary_function<typename Predicate::argument_type, bool>(){}
         __device__ __forceinline__ unary_negate() : unary_function<typename Predicate::argument_type, bool>(){}
 
-      const Predicate pred;
+      Predicate pred;
     };
 
     template <typename Predicate> __host__ __device__ __forceinline__ unary_negate<Predicate> not1(const Predicate& pred)
@@ -659,7 +659,7 @@ namespace cv { namespace gpu { namespace cudev
         __device__ __forceinline__ binary_negate() :
         binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool>(){}
 
-        const Predicate pred;
+        Predicate pred;
     };
 
     template <typename BinaryPredicate> __host__ __device__ __forceinline__ binary_negate<BinaryPredicate> not2(const BinaryPredicate& pred)
@@ -679,8 +679,8 @@ namespace cv { namespace gpu { namespace cudev
         __device__ __forceinline__ binder1st(const binder1st& other) :
         unary_function<typename Op::second_argument_type, typename Op::result_type>(){}
 
-        const Op op;
-        const typename Op::first_argument_type arg1;
+        Op op;
+        typename Op::first_argument_type arg1;
     };
 
     template <typename Op, typename T> __host__ __device__ __forceinline__ binder1st<Op> bind1st(const Op& op, const T& x)
@@ -700,8 +700,8 @@ namespace cv { namespace gpu { namespace cudev
          __device__ __forceinline__ binder2nd(const binder2nd& other) :
         unary_function<typename Op::first_argument_type, typename Op::result_type>(), op(other.op), arg2(other.arg2){}
 
-        const Op op;
-        const typename Op::second_argument_type arg2;
+        Op op;
+        typename Op::second_argument_type arg2;
     };
 
     template <typename Op, typename T> __host__ __device__ __forceinline__ binder2nd<Op> bind2nd(const Op& op, const T& x)

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

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
-}

modules/gpu/CMakeLists.txt

@@ -3,101 +3,10 @@ if(ANDROID OR IOS)
 endif()
 
 set(the_description "GPU-accelerated Computer Vision")
-ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_photo opencv_legacy)
 
-ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda")
+ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4100 /wd4324 /wd4512 -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter)
 
-file(GLOB lib_hdrs               "include/opencv2/*.hpp"                       "include/opencv2/${name}/*.hpp"               "include/opencv2/${name}/*.h")
-file(GLOB lib_int_hdrs           "src/*.hpp"      "src/*.h")
-file(GLOB lib_cuda_hdrs          "src/cuda/*.hpp" "src/cuda/*.h")
-file(GLOB lib_srcs               "src/*.cpp")
-file(GLOB lib_cuda               "src/cuda/*.cu*")
-
-source_group("Include"        FILES ${lib_hdrs})
-source_group("Src\\Host"      FILES ${lib_srcs} ${lib_int_hdrs})
-source_group("Src\\Cuda"      FILES ${lib_cuda} ${lib_cuda_hdrs})
-
-if(HAVE_CUDA)
-  file(GLOB_RECURSE ncv_srcs "src/nvidia/*.cpp" "src/nvidia/*.h*")
-  file(GLOB_RECURSE ncv_cuda "src/nvidia/*.cu")
-  set(ncv_files ${ncv_srcs} ${ncv_cuda})
-
-  source_group("Src\\NVidia" FILES ${ncv_files})
-  ocv_include_directories("src/nvidia" "src/nvidia/core" "src/nvidia/NPP_staging" ${CUDA_INCLUDE_DIRS})
-  ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter /wd4211 /wd4201 /wd4100 /wd4505 /wd4408)
-
-  if(MSVC)
-    if(NOT ENABLE_NOISY_WARNINGS)
-      foreach(var CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_RELEASE CMAKE_CXX_FLAGS_DEBUG)
-        string(REPLACE "/W4" "/W3" ${var} "${${var}}")
-      endforeach()
-
-      set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} -Xcompiler /wd4251)
-    endif()
-  endif()
-
-  ocv_cuda_compile(cuda_objs ${lib_cuda} ${ncv_cuda})
-
-  set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
-
-  if(WITH_NVCUVID)
-    set(cuda_link_libs ${cuda_link_libs} ${CUDA_nvcuvid_LIBRARY})
-  endif()
-
-  if(WIN32)
-    find_cuda_helper_libs(nvcuvenc)
-    set(cuda_link_libs ${cuda_link_libs} ${CUDA_nvcuvenc_LIBRARY})
-  endif()
-
-  if(WITH_FFMPEG)
-    set(cuda_link_libs ${cuda_link_libs} ${HIGHGUI_LIBRARIES})
-  endif()
-else()
-  set(lib_cuda "")
-  set(cuda_objs "")
-  set(cuda_link_libs "")
-  set(ncv_files "")
-endif()
-
-ocv_set_module_sources(
-  HEADERS ${lib_hdrs}
-  SOURCES ${lib_int_hdrs} ${lib_cuda_hdrs} ${lib_srcs} ${lib_cuda} ${ncv_files} ${cuda_objs}
-  )
-
-ocv_create_module(${cuda_link_libs})
-
-if(HAVE_CUDA)
-  if(HAVE_CUFFT)
-    CUDA_ADD_CUFFT_TO_TARGET(${the_module})
-  endif()
-
-  if(HAVE_CUBLAS)
-    CUDA_ADD_CUBLAS_TO_TARGET(${the_module})
-  endif()
-
-  install(FILES src/nvidia/NPP_staging/NPP_staging.hpp  src/nvidia/core/NCV.hpp
-    DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2/${name}
-    COMPONENT main)
-endif()
-
-ocv_add_precompiled_headers(${the_module})
-
-################################################################################################################
-################################      GPU Module Tests     #####################################################
-################################################################################################################
-file(GLOB test_srcs "test/*.cpp")
-file(GLOB test_hdrs "test/*.hpp" "test/*.h")
-
-set(nvidia "")
-if(HAVE_CUDA)
-  file(GLOB nvidia "test/nvidia/*.cpp" "test/nvidia/*.hpp" "test/nvidia/*.h")
-  set(nvidia FILES "Src\\\\\\\\NVidia" ${nvidia}) # 8 ugly backslashes :'(
-endif()
-
-ocv_add_accuracy_tests(FILES "Include" ${test_hdrs}
-                       FILES "Src" ${test_srcs}
-                       ${nvidia})
-ocv_add_perf_tests()
+ocv_define_module(gpu opencv_calib3d opencv_objdetect opencv_gpuarithm opencv_gpuwarping OPTIONAL opencv_gpulegacy)
 
 if(HAVE_CUDA)
   add_subdirectory(perf4au)

modules/gpu/app/nv_perf_test/CMakeLists.txt

-cmake_minimum_required(VERSION 2.8.3)
-
-project(nv_perf_test)
-
-find_package(OpenCV REQUIRED)
-include_directories(${OpenCV_INCLUDE_DIR})
-
-add_executable(${PROJECT_NAME} main.cpp)
-
-target_link_libraries(${PROJECT_NAME} ${OpenCV_LIBS})

modules/gpu/app/nv_perf_test/main.cpp

-#include <cstdio>
-#define HAVE_CUDA 1
-#include <opencv2/core.hpp>
-#include <opencv2/gpu.hpp>
-#include <opencv2/highgui.hpp>
-#include <opencv2/video.hpp>
-#include <opencv2/ts.hpp>
-
-static void printOsInfo()
-{
-#if defined _WIN32
-#   if defined _WIN64
-        printf("[----------]\n[ GPU INFO ] \tRun on OS Windows x64.\n[----------]\n"); fflush(stdout);
-#   else
-        printf("[----------]\n[ GPU INFO ] \tRun on OS Windows x32.\n[----------]\n"); fflush(stdout);
-#   endif
-#elif defined linux
-#   if defined _LP64
-        printf("[----------]\n[ GPU INFO ] \tRun on OS Linux x64.\n[----------]\n"); fflush(stdout);
-#   else
-        printf("[----------]\n[ GPU INFO ] \tRun on OS Linux x32.\n[----------]\n"); fflush(stdout);
-#   endif
-#elif defined __APPLE__
-#   if defined _LP64
-        printf("[----------]\n[ GPU INFO ] \tRun on OS Apple x64.\n[----------]\n"); fflush(stdout);
-#   else
-        printf("[----------]\n[ GPU INFO ] \tRun on OS Apple x32.\n[----------]\n"); fflush(stdout);
-#   endif
-#endif
-}
-
-static void printCudaInfo()
-{
-    const int deviceCount = cv::gpu::getCudaEnabledDeviceCount();
-
-    printf("[----------]\n"); fflush(stdout);
-    printf("[ GPU INFO ] \tCUDA device count:: %d.\n", deviceCount); fflush(stdout);
-    printf("[----------]\n"); fflush(stdout);
-
-    for (int i = 0; i < deviceCount; ++i)
-    {
-        cv::gpu::DeviceInfo info(i);
-
-        printf("[----------]\n"); fflush(stdout);
-        printf("[ DEVICE   ] \t# %d %s.\n", i, info.name().c_str()); fflush(stdout);
-        printf("[          ] \tCompute capability: %d.%d\n", info.majorVersion(), info.minorVersion()); fflush(stdout);
-        printf("[          ] \tMulti Processor Count:  %d\n", info.multiProcessorCount()); fflush(stdout);
-        printf("[          ] \tTotal memory: %d Mb\n", static_cast<int>(static_cast<int>(info.totalMemory() / 1024.0) / 1024.0)); fflush(stdout);
-        printf("[          ] \tFree  memory: %d Mb\n", static_cast<int>(static_cast<int>(info.freeMemory()  / 1024.0) / 1024.0)); fflush(stdout);
-        if (!info.isCompatible())
-            printf("[ GPU INFO ] \tThis device is NOT compatible with current GPU module build\n");
-        printf("[----------]\n"); fflush(stdout);
-    }
-}
-
-int main(int argc, char* argv[])
-{
-    printOsInfo();
-    printCudaInfo();
-
-    perf::Regression::Init("nv_perf_test");
-    perf::TestBase::Init(argc, argv);
-    testing::InitGoogleTest(&argc, argv);
-
-    return RUN_ALL_TESTS();
-}
-
-#define DEF_PARAM_TEST(name, ...) typedef ::perf::TestBaseWithParam< std::tr1::tuple< __VA_ARGS__ > > name
-#define DEF_PARAM_TEST_1(name, param_type) typedef ::perf::TestBaseWithParam< param_type > name
-
-//////////////////////////////////////////////////////////
-// HoughLinesP
-
-DEF_PARAM_TEST_1(Image, std::string);
-
-GPU_PERF_TEST_P(Image, HoughLinesP, testing::Values(std::string("im1_1280x800.jpg")))
-{
-    declare.time(30.0);
-
-    std::string fileName = GetParam();
-
-    const float rho = 1.f;
-    const float theta = 1.f;
-    const int threshold = 40;
-    const int minLineLenght = 20;
-    const int maxLineGap = 5;
-
-    cv::Mat image = cv::imread(fileName, cv::IMREAD_GRAYSCALE);
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::GpuMat d_image(image);
-        cv::gpu::GpuMat d_lines;
-        cv::gpu::HoughLinesBuf d_buf;
-
-        cv::gpu::HoughLinesP(d_image, d_lines, d_buf, rho, theta, minLineLenght, maxLineGap);
-
-        TEST_CYCLE()
-        {
-            cv::gpu::HoughLinesP(d_image, d_lines, d_buf, rho, theta, minLineLenght, maxLineGap);
-        }
-    }
-    else
-    {
-        cv::Mat mask;
-        cv::Canny(image, mask, 50, 100);
-
-        std::vector<cv::Vec4i> lines;
-        cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLenght, maxLineGap);
-
-        TEST_CYCLE()
-        {
-            cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLenght, maxLineGap);
-        }
-    }
-
-    SANITY_CHECK(0);
-}
-
-//////////////////////////////////////////////////////////
-// GoodFeaturesToTrack
-
-DEF_PARAM_TEST(Image_Depth, std::string, perf::MatDepth);
-
-GPU_PERF_TEST_P(Image_Depth, GoodFeaturesToTrack,
-                testing::Combine(
-                testing::Values(std::string("im1_1280x800.jpg")),
-                testing::Values(CV_8U, CV_16U)
-                ))
-{
-    declare.time(60);
-
-    const std::string fileName = std::tr1::get<0>(GetParam());
-    const int depth = std::tr1::get<1>(GetParam());
-
-    const int maxCorners = 5000;
-    const double qualityLevel = 0.05;
-    const int minDistance = 5;
-    const int blockSize = 3;
-    const bool useHarrisDetector = true;
-    const double k = 0.05;
-
-    cv::Mat src = cv::imread(fileName, cv::IMREAD_GRAYSCALE);
-    if (src.empty())
-        FAIL() << "Unable to load source image [" << fileName << "]";
-
-    if (depth != CV_8U)
-        src.convertTo(src, depth);
-
-    cv::Mat mask(src.size(), CV_8UC1, cv::Scalar::all(1));
-    mask(cv::Rect(0, 0, 100, 100)).setTo(cv::Scalar::all(0));
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::GoodFeaturesToTrackDetector_GPU d_detector(maxCorners, qualityLevel, minDistance, blockSize, useHarrisDetector, k);
-
-        cv::gpu::GpuMat d_src(src);
-        cv::gpu::GpuMat d_mask(mask);
-        cv::gpu::GpuMat d_pts;
-
-        d_detector(d_src, d_pts, d_mask);
-
-        TEST_CYCLE()
-        {
-            d_detector(d_src, d_pts, d_mask);
-        }
-    }
-    else
-    {
-        if (depth != CV_8U)
-            FAIL() << "Unsupported depth";
-
-        cv::Mat pts;
-
-        cv::goodFeaturesToTrack(src, pts, maxCorners, qualityLevel, minDistance, mask, blockSize, useHarrisDetector, k);
-
-        TEST_CYCLE()
-        {
-            cv::goodFeaturesToTrack(src, pts, maxCorners, qualityLevel, minDistance, mask, blockSize, useHarrisDetector, k);
-        }
-    }
-
-    SANITY_CHECK(0);
-}
-
-//////////////////////////////////////////////////////////
-// OpticalFlowPyrLKSparse
-
-typedef std::pair<std::string, std::string> string_pair;
-
-DEF_PARAM_TEST(ImagePair_Depth_GraySource, string_pair, perf::MatDepth, bool);
-
-GPU_PERF_TEST_P(ImagePair_Depth_GraySource, OpticalFlowPyrLKSparse,
-                testing::Combine(
-                    testing::Values(string_pair("im1_1280x800.jpg", "im2_1280x800.jpg")),
-                    testing::Values(CV_8U, CV_16U),
-                    testing::Bool()
-                    ))
-{
-    declare.time(60);
-
-    const string_pair fileNames = std::tr1::get<0>(GetParam());
-    const int depth = std::tr1::get<1>(GetParam());
-    const bool graySource = std::tr1::get<2>(GetParam());
-
-    // PyrLK params
-    const cv::Size winSize(15, 15);
-    const int maxLevel = 5;
-    const cv::TermCriteria criteria(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, 30, 0.01);
-
-    // GoodFeaturesToTrack params
-    const int maxCorners = 5000;
-    const double qualityLevel = 0.05;
-    const int minDistance = 5;
-    const int blockSize = 3;
-    const bool useHarrisDetector = true;
-    const double k = 0.05;
-
-    cv::Mat src1 = cv::imread(fileNames.first, graySource ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
-    if (src1.empty())
-        FAIL() << "Unable to load source image [" << fileNames.first << "]";
-
-    cv::Mat src2 = cv::imread(fileNames.second, graySource ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
-    if (src2.empty())
-        FAIL() << "Unable to load source image [" << fileNames.second << "]";
-
-    cv::Mat gray_src;
-    if (graySource)
-        gray_src = src1;
-    else
-        cv::cvtColor(src1, gray_src, cv::COLOR_BGR2GRAY);
-
-    cv::Mat pts;
-    cv::goodFeaturesToTrack(gray_src, pts, maxCorners, qualityLevel, minDistance, cv::noArray(), blockSize, useHarrisDetector, k);
-
-    if (depth != CV_8U)
-    {
-        src1.convertTo(src1, depth);
-        src2.convertTo(src2, depth);
-    }
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::GpuMat d_src1(src1);
-        cv::gpu::GpuMat d_src2(src2);
-        cv::gpu::GpuMat d_pts(pts.reshape(2, 1));
-        cv::gpu::GpuMat d_nextPts;
-        cv::gpu::GpuMat d_status;
-
-        cv::gpu::PyrLKOpticalFlow d_pyrLK;
-        d_pyrLK.winSize = winSize;
-        d_pyrLK.maxLevel = maxLevel;
-        d_pyrLK.iters = criteria.maxCount;
-        d_pyrLK.useInitialFlow = false;
-
-        d_pyrLK.sparse(d_src1, d_src2, d_pts, d_nextPts, d_status);
-
-        TEST_CYCLE()
-        {
-            d_pyrLK.sparse(d_src1, d_src2, d_pts, d_nextPts, d_status);
-        }
-    }
-    else
-    {
-        if (depth != CV_8U)
-            FAIL() << "Unsupported depth";
-
-        cv::Mat nextPts;
-        cv::Mat status;
-
-        cv::calcOpticalFlowPyrLK(src1, src2, pts, nextPts, status, cv::noArray(), winSize, maxLevel, criteria);
-
-        TEST_CYCLE()
-        {
-            cv::calcOpticalFlowPyrLK(src1, src2, pts, nextPts, status, cv::noArray(), winSize, maxLevel, criteria);
-        }
-    }
-
-    SANITY_CHECK(0);
-}
-
-//////////////////////////////////////////////////////////
-// OpticalFlowFarneback
-
-DEF_PARAM_TEST(ImagePair_Depth, string_pair, perf::MatDepth);
-
-GPU_PERF_TEST_P(ImagePair_Depth, OpticalFlowFarneback,
-                testing::Combine(
-                    testing::Values(string_pair("im1_1280x800.jpg", "im2_1280x800.jpg")),
-                    testing::Values(CV_8U, CV_16U)
-                    ))
-{
-    declare.time(500);
-
-    const string_pair fileNames = std::tr1::get<0>(GetParam());
-    const int depth = std::tr1::get<1>(GetParam());
-
-    const double pyrScale = 0.5;
-    const int numLevels = 6;
-    const int winSize = 7;
-    const int numIters = 15;
-    const int polyN = 7;
-    const double polySigma = 1.5;
-    const int flags = cv::OPTFLOW_USE_INITIAL_FLOW;
-
-    cv::Mat src1 = cv::imread(fileNames.first, cv::IMREAD_GRAYSCALE);
-    if (src1.empty())
-        FAIL() << "Unable to load source image [" << fileNames.first << "]";
-
-    cv::Mat src2 = cv::imread(fileNames.second, cv::IMREAD_GRAYSCALE);
-    if (src2.empty())
-        FAIL() << "Unable to load source image [" << fileNames.second << "]";
-
-    if (depth != CV_8U)
-    {
-        src1.convertTo(src1, depth);
-        src2.convertTo(src2, depth);
-    }
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::GpuMat d_src1(src1);
-        cv::gpu::GpuMat d_src2(src2);
-        cv::gpu::GpuMat d_u(src1.size(), CV_32FC1, cv::Scalar::all(0));
-        cv::gpu::GpuMat d_v(src1.size(), CV_32FC1, cv::Scalar::all(0));
-
-        cv::gpu::FarnebackOpticalFlow d_farneback;
-        d_farneback.pyrScale = pyrScale;
-        d_farneback.numLevels = numLevels;
-        d_farneback.winSize = winSize;
-        d_farneback.numIters = numIters;
-        d_farneback.polyN = polyN;
-        d_farneback.polySigma = polySigma;
-        d_farneback.flags = flags;
-
-        d_farneback(d_src1, d_src2, d_u, d_v);
-
-        TEST_CYCLE_N(10)
-        {
-            d_farneback(d_src1, d_src2, d_u, d_v);
-        }
-    }
-    else
-    {
-        if (depth != CV_8U)
-            FAIL() << "Unsupported depth";
-
-        cv::Mat flow(src1.size(), CV_32FC2, cv::Scalar::all(0));
-
-        cv::calcOpticalFlowFarneback(src1, src2, flow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags);
-
-        TEST_CYCLE_N(10)
-        {
-            cv::calcOpticalFlowFarneback(src1, src2, flow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags);
-        }
-    }
-
-    SANITY_CHECK(0);
-}
-
-//////////////////////////////////////////////////////////
-// OpticalFlowBM
-
-void calcOpticalFlowBM(const cv::Mat& prev, const cv::Mat& curr,
-                       cv::Size bSize, cv::Size shiftSize, cv::Size maxRange, int usePrevious,
-                       cv::Mat& velx, cv::Mat& vely)
-{
-    cv::Size sz((curr.cols - bSize.width + shiftSize.width)/shiftSize.width, (curr.rows - bSize.height + shiftSize.height)/shiftSize.height);
-
-    velx.create(sz, CV_32FC1);
-    vely.create(sz, CV_32FC1);
-
-    CvMat cvprev = prev;
-    CvMat cvcurr = curr;
-
-    CvMat cvvelx = velx;
-    CvMat cvvely = vely;
-
-    cvCalcOpticalFlowBM(&cvprev, &cvcurr, bSize, shiftSize, maxRange, usePrevious, &cvvelx, &cvvely);
-}
-
-DEF_PARAM_TEST(ImagePair_BlockSize_ShiftSize_MaxRange, string_pair, cv::Size, cv::Size, cv::Size);
-
-GPU_PERF_TEST_P(ImagePair_BlockSize_ShiftSize_MaxRange, OpticalFlowBM,
-                testing::Combine(
-                    testing::Values(string_pair("im1_1280x800.jpg", "im2_1280x800.jpg")),
-                    testing::Values(cv::Size(16, 16)),
-                    testing::Values(cv::Size(2, 2)),
-                    testing::Values(cv::Size(16, 16))
-                    ))
-{
-    declare.time(3000);
-
-    const string_pair fileNames = std::tr1::get<0>(GetParam());
-    const cv::Size block_size = std::tr1::get<1>(GetParam());
-    const cv::Size shift_size = std::tr1::get<2>(GetParam());
-    const cv::Size max_range = std::tr1::get<3>(GetParam());
-
-    cv::Mat src1 = cv::imread(fileNames.first, cv::IMREAD_GRAYSCALE);
-    if (src1.empty())
-        FAIL() << "Unable to load source image [" << fileNames.first << "]";
-
-    cv::Mat src2 = cv::imread(fileNames.second, cv::IMREAD_GRAYSCALE);
-    if (src2.empty())
-        FAIL() << "Unable to load source image [" << fileNames.second << "]";
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::GpuMat d_src1(src1);
-        cv::gpu::GpuMat d_src2(src2);
-        cv::gpu::GpuMat d_velx, d_vely, buf;
-
-        cv::gpu::calcOpticalFlowBM(d_src1, d_src2, block_size, shift_size, max_range, false, d_velx, d_vely, buf);
-
-        TEST_CYCLE_N(10)
-        {
-            cv::gpu::calcOpticalFlowBM(d_src1, d_src2, block_size, shift_size, max_range, false, d_velx, d_vely, buf);
-        }
-    }
-    else
-    {
-        cv::Mat velx, vely;
-
-        calcOpticalFlowBM(src1, src2, block_size, shift_size, max_range, false, velx, vely);
-
-        TEST_CYCLE_N(10)
-        {
-            calcOpticalFlowBM(src1, src2, block_size, shift_size, max_range, false, velx, vely);
-        }
-    }
-
-    SANITY_CHECK(0);
-}
-
-GPU_PERF_TEST_P(ImagePair_BlockSize_ShiftSize_MaxRange, FastOpticalFlowBM,
-                testing::Combine(
-                    testing::Values(string_pair("im1_1280x800.jpg", "im2_1280x800.jpg")),
-                    testing::Values(cv::Size(16, 16)),
-                    testing::Values(cv::Size(1, 1)),
-                    testing::Values(cv::Size(16, 16))
-                    ))
-{
-    declare.time(3000);
-
-    const string_pair fileNames = std::tr1::get<0>(GetParam());
-    const cv::Size block_size = std::tr1::get<1>(GetParam());
-    const cv::Size shift_size = std::tr1::get<2>(GetParam());
-    const cv::Size max_range = std::tr1::get<3>(GetParam());
-
-    cv::Mat src1 = cv::imread(fileNames.first, cv::IMREAD_GRAYSCALE);
-    if (src1.empty())
-        FAIL() << "Unable to load source image [" << fileNames.first << "]";
-
-    cv::Mat src2 = cv::imread(fileNames.second, cv::IMREAD_GRAYSCALE);
-    if (src2.empty())
-        FAIL() << "Unable to load source image [" << fileNames.second << "]";
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::GpuMat d_src1(src1);
-        cv::gpu::GpuMat d_src2(src2);
-        cv::gpu::GpuMat d_velx, d_vely;
-
-        cv::gpu::FastOpticalFlowBM fastBM;
-
-        fastBM(d_src1, d_src2, d_velx, d_vely, max_range.width, block_size.width);
-
-        TEST_CYCLE_N(10)
-        {
-            fastBM(d_src1, d_src2, d_velx, d_vely, max_range.width, block_size.width);
-        }
-    }
-    else
-    {
-        cv::Mat velx, vely;
-
-        calcOpticalFlowBM(src1, src2, block_size, shift_size, max_range, false, velx, vely);
-
-        TEST_CYCLE_N(10)
-        {
-            calcOpticalFlowBM(src1, src2, block_size, shift_size, max_range, false, velx, vely);
-        }
-    }
-
-    SANITY_CHECK(0);
-}

modules/gpu/doc/calib3d.rst

+Camera Calibration and 3D Reconstruction
+========================================
+
+.. highlight:: cpp
+
+
+
+gpu::solvePnPRansac
+-------------------
+Finds the object pose from 3D-2D point correspondences.
+
+.. ocv:function:: void gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat, const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess=false, int num_iters=100, float max_dist=8.0, int min_inlier_count=100, vector<int>* inliers=NULL)
+
+    :param object: Single-row matrix of object points.
+
+    :param image: Single-row matrix of image points.
+
+    :param camera_mat: 3x3 matrix of intrinsic camera parameters.
+
+    :param dist_coef: Distortion coefficients. See :ocv:func:`undistortPoints` for details.
+
+    :param rvec: Output 3D rotation vector.
+
+    :param tvec: Output 3D translation vector.
+
+    :param use_extrinsic_guess: Flag to indicate that the function must use ``rvec`` and ``tvec`` as an initial transformation guess. It is not supported for now.
+
+    :param num_iters: Maximum number of RANSAC iterations.
+
+    :param max_dist: Euclidean distance threshold to detect whether point is inlier or not.
+
+    :param min_inlier_count: Flag to indicate that the function must stop if greater or equal number of inliers is achieved. It is not supported for now.
+
+    :param inliers: Output vector of inlier indices.
+
+.. seealso:: :ocv:func:`solvePnPRansac`

modules/gpu/doc/gpu.rst

@@ -8,12 +8,5 @@ gpu. GPU-accelerated Computer Vision
     introduction
     initalization_and_information
     data_structures
-    operations_on_matrices
-    per_element_operations
-    image_processing
-    matrix_reductions
     object_detection
-    feature_detection_and_description
-    image_filtering
-    camera_calibration_and_3d_reconstruction
-    video
+    calib3d

modules/gpu/perf/perf_calib3d.cpp

@@ -46,181 +46,8 @@ using namespace std;
 using namespace testing;
 using namespace perf;
 
-//////////////////////////////////////////////////////////////////////
-// StereoBM
-
-typedef std::tr1::tuple<string, string> pair_string;
-DEF_PARAM_TEST_1(ImagePair, pair_string);
-
-PERF_TEST_P(ImagePair, Calib3D_StereoBM,
-            Values(pair_string("gpu/perf/aloe.png", "gpu/perf/aloeR.png")))
-{
-    declare.time(300.0);
-
-    const cv::Mat imgLeft = readImage(GET_PARAM(0), cv::IMREAD_GRAYSCALE);
-    ASSERT_FALSE(imgLeft.empty());
-
-    const cv::Mat imgRight = readImage(GET_PARAM(1), cv::IMREAD_GRAYSCALE);
-    ASSERT_FALSE(imgRight.empty());
-
-    const int preset = 0;
-    const int ndisp = 256;
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::StereoBM_GPU d_bm(preset, ndisp);
-
-        const cv::gpu::GpuMat d_imgLeft(imgLeft);
-        const cv::gpu::GpuMat d_imgRight(imgRight);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() d_bm(d_imgLeft, d_imgRight, dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Ptr<cv::StereoBM> bm = cv::createStereoBM(ndisp);
-
-        cv::Mat dst;
-
-        TEST_CYCLE() bm->compute(imgLeft, imgRight, dst);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// StereoBeliefPropagation
-
-PERF_TEST_P(ImagePair, Calib3D_StereoBeliefPropagation,
-            Values(pair_string("gpu/stereobp/aloe-L.png", "gpu/stereobp/aloe-R.png")))
-{
-    declare.time(300.0);
-
-    const cv::Mat imgLeft = readImage(GET_PARAM(0));
-    ASSERT_FALSE(imgLeft.empty());
-
-    const cv::Mat imgRight = readImage(GET_PARAM(1));
-    ASSERT_FALSE(imgRight.empty());
-
-    const int ndisp = 64;
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::StereoBeliefPropagation d_bp(ndisp);
-
-        const cv::gpu::GpuMat d_imgLeft(imgLeft);
-        const cv::gpu::GpuMat d_imgRight(imgRight);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() d_bp(d_imgLeft, d_imgRight, dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// StereoConstantSpaceBP
-
-PERF_TEST_P(ImagePair, Calib3D_StereoConstantSpaceBP,
-            Values(pair_string("gpu/stereobm/aloe-L.png", "gpu/stereobm/aloe-R.png")))
-{
-    declare.time(300.0);
-
-    const cv::Mat imgLeft = readImage(GET_PARAM(0), cv::IMREAD_GRAYSCALE);
-    ASSERT_FALSE(imgLeft.empty());
-
-    const cv::Mat imgRight = readImage(GET_PARAM(1), cv::IMREAD_GRAYSCALE);
-    ASSERT_FALSE(imgRight.empty());
-
-    const int ndisp = 128;
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::StereoConstantSpaceBP d_csbp(ndisp);
-
-        const cv::gpu::GpuMat d_imgLeft(imgLeft);
-        const cv::gpu::GpuMat d_imgRight(imgRight);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() d_csbp(d_imgLeft, d_imgRight, dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// DisparityBilateralFilter
-
-PERF_TEST_P(ImagePair, Calib3D_DisparityBilateralFilter,
-            Values(pair_string("gpu/stereobm/aloe-L.png", "gpu/stereobm/aloe-disp.png")))
-{
-    const cv::Mat img = readImage(GET_PARAM(0), cv::IMREAD_GRAYSCALE);
-    ASSERT_FALSE(img.empty());
-
-    const cv::Mat disp = readImage(GET_PARAM(1), cv::IMREAD_GRAYSCALE);
-    ASSERT_FALSE(disp.empty());
-
-    const int ndisp = 128;
-
-    if (PERF_RUN_GPU())
-    {
-        cv::gpu::DisparityBilateralFilter d_filter(ndisp);
-
-        const cv::gpu::GpuMat d_img(img);
-        const cv::gpu::GpuMat d_disp(disp);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() d_filter(d_disp, d_img, dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// TransformPoints
-
 DEF_PARAM_TEST_1(Count, int);
 
-PERF_TEST_P(Count, Calib3D_TransformPoints,
-            Values(5000, 10000, 20000))
-{
-    const int count = GetParam();
-
-    cv::Mat src(1, count, CV_32FC3);
-    declare.in(src, WARMUP_RNG);
-
-    const cv::Mat rvec = cv::Mat::ones(1, 3, CV_32FC1);
-    const cv::Mat tvec = cv::Mat::ones(1, 3, CV_32FC1);
-
-    if (PERF_RUN_GPU())
-    {
-        const cv::gpu::GpuMat d_src(src);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() cv::gpu::transformPoints(d_src, rvec, tvec, dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}
-
 //////////////////////////////////////////////////////////////////////
 // ProjectPoints
 
@@ -306,66 +133,3 @@ PERF_TEST_P(Count, Calib3D_SolvePnPRansac,
         CPU_SANITY_CHECK(tvec, 1e-6);
     }
 }
-
-//////////////////////////////////////////////////////////////////////
-// ReprojectImageTo3D
-
-PERF_TEST_P(Sz_Depth, Calib3D_ReprojectImageTo3D,
-            Combine(GPU_TYPICAL_MAT_SIZES,
-                    Values(CV_8U, CV_16S)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int depth = GET_PARAM(1);
-
-    cv::Mat src(size, depth);
-    declare.in(src, WARMUP_RNG);
-
-    cv::Mat Q(4, 4, CV_32FC1);
-    cv::randu(Q, 0.1, 1.0);
-
-    if (PERF_RUN_GPU())
-    {
-        const cv::gpu::GpuMat d_src(src);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() cv::gpu::reprojectImageTo3D(d_src, dst, Q);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::reprojectImageTo3D(src, dst, Q);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// DrawColorDisp
-
-PERF_TEST_P(Sz_Depth, Calib3D_DrawColorDisp,
-            Combine(GPU_TYPICAL_MAT_SIZES,
-                    Values(CV_8U, CV_16S)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int type = GET_PARAM(1);
-
-    cv::Mat src(size, type);
-    declare.in(src, WARMUP_RNG);
-
-    if (PERF_RUN_GPU())
-    {
-        const cv::gpu::GpuMat d_src(src);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() cv::gpu::drawColorDisp(d_src, dst, 255);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}

modules/gpu/perf/perf_precomp.hpp

@@ -51,21 +51,12 @@
 #ifndef __OPENCV_PERF_PRECOMP_HPP__
 #define __OPENCV_PERF_PRECOMP_HPP__
 
-#include <cstdio>
-#include <iostream>
-
 #include "opencv2/ts.hpp"
 #include "opencv2/ts/gpu_perf.hpp"
 
-#include "opencv2/core.hpp"
-#include "opencv2/highgui.hpp"
 #include "opencv2/gpu.hpp"
 #include "opencv2/calib3d.hpp"
-#include "opencv2/imgproc.hpp"
-#include "opencv2/video.hpp"
-#include "opencv2/photo.hpp"
-
-#include "opencv2/core/gpu_private.hpp"
+#include "opencv2/objdetect.hpp"
 
 #ifdef GTEST_CREATE_SHARED_LIBRARY
 #error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined

modules/gpu/perf4au/CMakeLists.txt

-set(PERF4AU_REQUIRED_DEPS opencv_core opencv_imgproc opencv_highgui opencv_video opencv_legacy opencv_gpu opencv_ts)
+set(PERF4AU_REQUIRED_DEPS opencv_core opencv_imgproc opencv_highgui opencv_video opencv_legacy opencv_ml opencv_ts opencv_gpufilters opencv_gpuimgproc opencv_gpuoptflow)
 
 ocv_check_dependencies(${PERF4AU_REQUIRED_DEPS})
 
@@ -25,4 +25,3 @@ if(WIN32)
         set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG")
     endif()
 endif()
-

modules/gpu/perf4au/main.cpp

@@ -40,18 +40,15 @@
 //
 //M*/
 
-#include <cstdio>
+#include "opencv2/ts.hpp"
+#include "opencv2/ts/gpu_perf.hpp"
+
+#include "opencv2/gpuimgproc.hpp"
+#include "opencv2/gpuoptflow.hpp"
 
-#ifdef HAVE_CVCONFIG_H
-#include "cvconfig.h"
-#endif
-#include "opencv2/core.hpp"
-#include "opencv2/gpu.hpp"
 #include "opencv2/highgui.hpp"
 #include "opencv2/video.hpp"
 #include "opencv2/legacy.hpp"
-#include "opencv2/ts.hpp"
-#include "opencv2/ts/gpu_perf.hpp"
 
 int main(int argc, char* argv[])
 {