Merge pull request #3158 from arielelkin:master
authorVadim Pisarevsky <vadim.pisarevsky@gmail.com>
Tue, 26 Aug 2014 16:07:10 +0000 (16:07 +0000)
committerVadim Pisarevsky <vadim.pisarevsky@gmail.com>
Tue, 26 Aug 2014 16:07:10 +0000 (16:07 +0000)
22 files changed:
modules/calib3d/src/dls.h
modules/calib3d/src/epnp.h
modules/calib3d/src/p3p.h
modules/core/src/copy.cpp
modules/core/src/ocl.cpp
modules/cudaoptflow/src/tvl1flow.cpp
modules/imgproc/src/deriv.cpp
modules/imgproc/src/opencl/laplacian5.cl
modules/stitching/include/opencv2/stitching/detail/matchers.hpp
modules/stitching/src/matchers.cpp
modules/stitching/src/precomp.hpp
modules/stitching/src/stitcher.cpp
modules/stitching/test/test_matchers.cpp
modules/videoio/src/cap.cpp
modules/videoio/src/cap_dshow.cpp
modules/world/src/precomp.hpp
modules/world/src/world_init.cpp
samples/cpp/stitching_detailed.cpp
samples/gpu/CMakeLists.txt
samples/gpu/performance/CMakeLists.txt
samples/gpu/performance/tests.cpp
samples/gpu/surf_keypoint_matcher.cpp

index 30321fc..d379844 100644 (file)
@@ -24,9 +24,9 @@ private:
     {
         for(int i = 0; i < N; i++)
         {
-            p.at<double>(0,i) = opoints.at<OpointType>(0,i).x;
-            p.at<double>(1,i) = opoints.at<OpointType>(0,i).y;
-            p.at<double>(2,i) = opoints.at<OpointType>(0,i).z;
+            p.at<double>(0,i) = opoints.at<OpointType>(i).x;
+            p.at<double>(1,i) = opoints.at<OpointType>(i).y;
+            p.at<double>(2,i) = opoints.at<OpointType>(i).z;
 
             // compute mean of object points
             mn.at<double>(0) += p.at<double>(0,i);
@@ -34,12 +34,12 @@ private:
             mn.at<double>(2) += p.at<double>(2,i);
 
             // make z into unit vectors from normalized pixel coords
-            double sr = std::pow(ipoints.at<IpointType>(0,i).x, 2) +
-                        std::pow(ipoints.at<IpointType>(0,i).y, 2) + (double)1;
+            double sr = std::pow(ipoints.at<IpointType>(i).x, 2) +
+                        std::pow(ipoints.at<IpointType>(i).y, 2) + (double)1;
                    sr = std::sqrt(sr);
 
-            z.at<double>(0,i) = ipoints.at<IpointType>(0,i).x / sr;
-            z.at<double>(1,i) = ipoints.at<IpointType>(0,i).y / sr;
+            z.at<double>(0,i) = ipoints.at<IpointType>(i).x / sr;
+            z.at<double>(1,i) = ipoints.at<IpointType>(i).y / sr;
             z.at<double>(2,i) = (double)1 / sr;
         }
 
index dd42b01..2619f75 100644 (file)
@@ -27,12 +27,12 @@ class epnp {
   {
       for(int i = 0; i < number_of_correspondences; i++)
       {
-          pws[3 * i    ] = opoints.at<OpointType>(0,i).x;
-          pws[3 * i + 1] = opoints.at<OpointType>(0,i).y;
-          pws[3 * i + 2] = opoints.at<OpointType>(0,i).z;
+          pws[3 * i    ] = opoints.at<OpointType>(i).x;
+          pws[3 * i + 1] = opoints.at<OpointType>(i).y;
+          pws[3 * i + 2] = opoints.at<OpointType>(i).z;
 
-          us[2 * i    ] = ipoints.at<IpointType>(0,i).x*fu + uc;
-          us[2 * i + 1] = ipoints.at<IpointType>(0,i).y*fv + vc;
+          us[2 * i    ] = ipoints.at<IpointType>(i).x*fu + uc;
+          us[2 * i + 1] = ipoints.at<IpointType>(i).y*fv + vc;
       }
   }
   double reprojection_error(const double R[3][3], const double t[3]);
index 57f8d7d..00d99ae 100644 (file)
@@ -37,11 +37,11 @@ class p3p
       points.resize(20);
       for(int i = 0; i < 4; i++)
       {
-          points[i*5] = ipoints.at<IpointType>(0,i).x*fx + cx;
-          points[i*5+1] = ipoints.at<IpointType>(0,i).y*fy + cy;
-          points[i*5+2] = opoints.at<OpointType>(0,i).x;
-          points[i*5+3] = opoints.at<OpointType>(0,i).y;
-          points[i*5+4] = opoints.at<OpointType>(0,i).z;
+          points[i*5] = ipoints.at<IpointType>(i).x*fx + cx;
+          points[i*5+1] = ipoints.at<IpointType>(i).y*fy + cy;
+          points[i*5+2] = opoints.at<OpointType>(i).x;
+          points[i*5+3] = opoints.at<OpointType>(i).y;
+          points[i*5+4] = opoints.at<OpointType>(i).z;
       }
   }
   void init_inverse_parameters();
index 82b2000..c98e3dc 100644 (file)
@@ -615,9 +615,15 @@ enum { FLIP_COLS = 1 << 0, FLIP_ROWS = 1 << 1, FLIP_BOTH = FLIP_ROWS | FLIP_COLS
 static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
 {
     CV_Assert(flipCode >= -1 && flipCode <= 1);
+
+    const ocl::Device & dev = ocl::Device::getDefault();
     int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
             flipType, kercn = std::min(ocl::predictOptimalVectorWidth(_src, _dst), 4);
 
+    bool doubleSupport = dev.doubleFPConfig() > 0;
+    if (!doubleSupport && depth == CV_64F)
+        kercn = cn;
+
     if (cn > 4)
         return false;
 
@@ -629,14 +635,13 @@ static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
     else
         kernelName = "arithm_flip_rows_cols", flipType = FLIP_BOTH;
 
-    ocl::Device dev = ocl::Device::getDefault();
     int pxPerWIy = (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU)) ? 4 : 1;
     kercn = (cn!=3 || flipType == FLIP_ROWS) ? std::max(kercn, cn) : cn;
 
     ocl::Kernel k(kernelName, ocl::core::flip_oclsrc,
         format( "-D T=%s -D T1=%s -D cn=%d -D PIX_PER_WI_Y=%d -D kercn=%d",
-                ocl::memopTypeToStr(CV_MAKE_TYPE(depth, kercn)),
-                ocl::memopTypeToStr(depth), cn, pxPerWIy, kercn));
+                kercn != cn ? ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)) : ocl::memopTypeToStr(CV_MAKE_TYPE(depth, kercn)),
+                kercn != cn ? ocl::typeToStr(depth) : ocl::memopTypeToStr(depth), cn, pxPerWIy, kercn));
     if (k.empty())
         return false;
 
index d279c02..fc36ee0 100644 (file)
@@ -4360,10 +4360,10 @@ const char* memopTypeToStr(int type)
 {
     static const char* tab[] =
     {
-        "uchar", "uchar2", "uchar3", "uchar4", 0, 0, 0, "uchar8", 0, 0, 0, 0, 0, 0, 0, "uchar16",
-        "char", "char2", "char3", "char4", 0, 0, 0, "char8", 0, 0, 0, 0, 0, 0, 0, "char16",
-        "ushort", "ushort2", "ushort3", "ushort4",0, 0, 0, "ushort8", 0, 0, 0, 0, 0, 0, 0, "ushort16",
-        "short", "short2", "short3", "short4", 0, 0, 0, "short8", 0, 0, 0, 0, 0, 0, 0, "short16",
+        "uchar", "short", "uchar3", "int", 0, 0, 0, "int2", 0, 0, 0, 0, 0, 0, 0, "int4",
+        "char", "short", "char3", "int", 0, 0, 0, "int2", 0, 0, 0, 0, 0, 0, 0, "int4",
+        "ushort", "int", "ushort3", "int2",0, 0, 0, "int4", 0, 0, 0, 0, 0, 0, 0, "int8",
+        "short", "int", "short3", "int2", 0, 0, 0, "int4", 0, 0, 0, 0, 0, 0, 0, "int8",
         "int", "int2", "int3", "int4", 0, 0, 0, "int8", 0, 0, 0, 0, 0, 0, 0, "int16",
         "int", "int2", "int3", "int4", 0, 0, 0, "int8", 0, 0, 0, 0, 0, 0, 0, "int16",
         "ulong", "ulong2", "ulong3", "ulong4", 0, 0, 0, "ulong8", 0, 0, 0, 0, 0, 0, 0, "ulong16",
index a54e212..d0efffa 100644 (file)
@@ -47,7 +47,7 @@
 cv::cuda::OpticalFlowDual_TVL1_CUDA::OpticalFlowDual_TVL1_CUDA() { throw_no_cuda(); }
 void cv::cuda::OpticalFlowDual_TVL1_CUDA::operator ()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
 void cv::cuda::OpticalFlowDual_TVL1_CUDA::collectGarbage() {}
-void cv::cuda::OpticalFlowDual_TVL1_CUDA::procOneScale(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
+void cv::cuda::OpticalFlowDual_TVL1_CUDA::procOneScale(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
 
 #else
 
index c3579fe..dd8f124 100644 (file)
@@ -643,21 +643,109 @@ void cv::Scharr( InputArray _src, OutputArray _dst, int ddepth, int dx, int dy,
 
 namespace cv {
 
+#define LAPLACIAN_LOCAL_MEM(tileX, tileY, ksize, elsize) (((tileX) + 2 * (int)((ksize) / 2)) * (3 * (tileY) + 2 * (int)((ksize) / 2)) * elsize)
+
 static bool ocl_Laplacian5(InputArray _src, OutputArray _dst,
                            const Mat & kd, const Mat & ks, double scale, double delta,
                            int borderType, int depth, int ddepth)
 {
+    const size_t tileSizeX = 16;
+    const size_t tileSizeYmin = 8;
+
+    const ocl::Device dev = ocl::Device::getDefault();
+
+    int stype = _src.type();
+    int sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype), esz = CV_ELEM_SIZE(stype);
+
+    bool doubleSupport = dev.doubleFPConfig() > 0;
+    if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F))
+        return false;
+
+    Mat kernelX = kd.reshape(1, 1);
+    if (kernelX.cols % 2 != 1)
+        return false;
+    Mat kernelY = ks.reshape(1, 1);
+    if (kernelY.cols % 2 != 1)
+        return false;
+    CV_Assert(kernelX.cols == kernelY.cols);
+
+    size_t wgs = dev.maxWorkGroupSize();
+    size_t lmsz = dev.localMemSize();
+
+    size_t src_step = _src.step(), src_offset = _src.offset();
+    if (((src_offset % src_step) % esz == 0) &&
+        (
+         (borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE) ||
+         ((borderType == BORDER_REFLECT || borderType == BORDER_WRAP || borderType == BORDER_REFLECT_101) &&
+          (_src.cols() >= kernelX.cols && _src.rows() >= kernelY.cols))
+        ) &&
+        (tileSizeX * tileSizeYmin  <= wgs) &&
+        (LAPLACIAN_LOCAL_MEM(tileSizeX, tileSizeYmin, kernelX.cols, cn * 4) <= lmsz)
+       )
+    {
+        Size size = _src.size(), wholeSize;
+        Point origin;
+        int dtype = CV_MAKE_TYPE(ddepth, cn);
+        int wdepth = CV_32F;
+
+        size_t tileSizeY = wgs / tileSizeX;
+        while ((tileSizeX * tileSizeY > wgs) || (LAPLACIAN_LOCAL_MEM(tileSizeX, tileSizeY, kernelX.cols, cn * 4) > lmsz))
+        {
+            tileSizeY /= 2;
+        }
+        size_t lt2[2] = { tileSizeX, tileSizeY};
+        size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1] };
+
+        char cvt[2][40];
+        const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP",
+                                           "BORDER_REFLECT_101" };
+
+        String opts = cv::format("-D BLK_X=%d -D BLK_Y=%d -D RADIUS=%d%s%s"
+                                 " -D convertToWT=%s -D convertToDT=%s"
+                                 " -D %s -D srcT1=%s -D dstT1=%s -D WT1=%s"
+                                 " -D srcT=%s -D dstT=%s -D WT=%s"
+                                 " -D CN=%d ",
+                                 (int)lt2[0], (int)lt2[1], kernelX.cols / 2,
+                                 ocl::kernelToStr(kernelX, wdepth, "KERNEL_MATRIX_X").c_str(),
+                                 ocl::kernelToStr(kernelY, wdepth, "KERNEL_MATRIX_Y").c_str(),
+                                 ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
+                                 ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]),
+                                 borderMap[borderType],
+                                 ocl::typeToStr(sdepth), ocl::typeToStr(ddepth), ocl::typeToStr(wdepth),
+                                 ocl::typeToStr(CV_MAKETYPE(sdepth, cn)),
+                                 ocl::typeToStr(CV_MAKETYPE(ddepth, cn)),
+                                 ocl::typeToStr(CV_MAKETYPE(wdepth, cn)),
+                                 cn);
+
+        ocl::Kernel k("laplacian", ocl::imgproc::laplacian5_oclsrc, opts);
+        if (k.empty())
+            return false;
+        UMat src = _src.getUMat();
+        _dst.create(size, dtype);
+        UMat dst = _dst.getUMat();
+
+        int src_offset_x = static_cast<int>((src_offset % src_step) / esz);
+        int src_offset_y = static_cast<int>(src_offset / src_step);
+
+        src.locateROI(wholeSize, origin);
+
+        k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y,
+               wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst),
+               static_cast<float>(scale), static_cast<float>(delta));
+
+        return k.run(2, gt2, lt2, false);
+    }
     int iscale = cvRound(scale), idelta = cvRound(delta);
-    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0,
-            floatCoeff = std::fabs(delta - idelta) > DBL_EPSILON || std::fabs(scale - iscale) > DBL_EPSILON;
-    int cn = _src.channels(), wdepth = std::max(depth, floatCoeff ? CV_32F : CV_32S), kercn = 1;
+    bool floatCoeff = std::fabs(delta - idelta) > DBL_EPSILON || std::fabs(scale - iscale) > DBL_EPSILON;
+    int wdepth = std::max(depth, floatCoeff ? CV_32F : CV_32S), kercn = 1;
 
     if (!doubleSupport && wdepth == CV_64F)
         return false;
 
     char cvt[2][40];
     ocl::Kernel k("sumConvert", ocl::imgproc::laplacian5_oclsrc,
-                  format("-D srcT=%s -D WT=%s -D dstT=%s -D coeffT=%s -D wdepth=%d "
+                  format("-D ONLY_SUM_CONVERT "
+                         "-D srcT=%s -D WT=%s -D dstT=%s -D coeffT=%s -D wdepth=%d "
                          "-D convertToWT=%s -D convertToDT=%s%s",
                          ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)),
                          ocl::typeToStr(CV_MAKE_TYPE(wdepth, kercn)),
index 3e15e09..ed6a0d4 100644 (file)
@@ -5,8 +5,11 @@
 // Copyright (C) 2014, Itseez, Inc., all rights reserved.
 // Third party copyrights are property of their respective owners.
 
+
 #define noconvert
 
+#ifdef ONLY_SUM_CONVERT
+
 __kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1_offset,
                          __global const uchar * src2ptr, int src2_step, int src2_offset,
                          __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
@@ -32,3 +35,172 @@ __kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1
 #endif
     }
 }
+
+#else
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+/////////////////////////////////Macro for border type////////////////////////////////////////////
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+#ifdef BORDER_CONSTANT
+// CCCCCC|abcdefgh|CCCCCCC
+#define EXTRAPOLATE(x, maxV)
+#elif defined BORDER_REPLICATE
+// aaaaaa|abcdefgh|hhhhhhh
+#define EXTRAPOLATE(x, maxV) \
+    { \
+        (x) = max(min((x), (maxV) - 1), 0); \
+    }
+#elif defined BORDER_WRAP
+// cdefgh|abcdefgh|abcdefg
+#define EXTRAPOLATE(x, maxV) \
+    { \
+        (x) = ( (x) + (maxV) ) % (maxV); \
+    }
+#elif defined BORDER_REFLECT
+// fedcba|abcdefgh|hgfedcb
+#define EXTRAPOLATE(x, maxV) \
+    { \
+        (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \
+    }
+#elif defined BORDER_REFLECT_101
+// gfedcb|abcdefgh|gfedcba
+#define EXTRAPOLATE(x, maxV) \
+    { \
+        (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \
+    }
+#else
+#error No extrapolation method
+#endif
+
+#if CN != 3
+#define loadpix(addr) *(__global const srcT *)(addr)
+#define storepix(val, addr)  *(__global dstT *)(addr) = val
+#define SRCSIZE (int)sizeof(srcT)
+#define DSTSIZE (int)sizeof(dstT)
+#else
+#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
+#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
+#define SRCSIZE (int)sizeof(srcT1)*3
+#define DSTSIZE (int)sizeof(dstT1)*3
+#endif
+
+#define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x)))
+
+#ifdef BORDER_CONSTANT
+// CCCCCC|abcdefgh|CCCCCCC
+#define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y))
+#else
+#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y))
+#endif
+
+// horizontal and vertical filter kernels
+// should be defined on host during compile time to avoid overhead
+#define DIG(a) a,
+__constant WT1 mat_kernelX[] = { KERNEL_MATRIX_X };
+__constant WT1 mat_kernelY[] = { KERNEL_MATRIX_Y };
+
+__kernel void laplacian(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width,
+                         __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+                         WT1 scale, WT1 delta)
+{
+    __local WT lsmem[BLK_Y + 2 * RADIUS][BLK_X + 2 * RADIUS];
+    __local WT lsmemDy1[BLK_Y][BLK_X + 2 * RADIUS];
+    __local WT lsmemDy2[BLK_Y][BLK_X + 2 * RADIUS];
+
+    int lix = get_local_id(0);
+    int liy = get_local_id(1);
+
+    int x = get_global_id(0);
+
+    int srcX = x + srcOffsetX - RADIUS;
+
+    int clocY = liy;
+    do
+    {
+        int yb = clocY + srcOffsetY - RADIUS;
+        EXTRAPOLATE(yb, (height));
+
+        int clocX = lix;
+        int cSrcX = srcX;
+        do
+        {
+            int xb = cSrcX;
+            EXTRAPOLATE(xb,(width));
+            lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 );
+
+            clocX += BLK_X;
+            cSrcX += BLK_X;
+        }
+        while(clocX < BLK_X+(RADIUS*2));
+
+        clocY += BLK_Y;
+    }
+    while (clocY < BLK_Y+(RADIUS*2));
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    WT scale_v = (WT)scale;
+    WT delta_v = (WT)delta;
+    for (int y = 0; y < dst_rows; y+=BLK_Y)
+    {
+        int i, clocX = lix;
+        WT sum1 = (WT) 0;
+        WT sum2 = (WT) 0;
+        do
+        {
+            sum1 = (WT) 0;
+            sum2 = (WT) 0;
+            for (i=0; i<=2*RADIUS; i++)
+            {
+                sum1 = mad(lsmem[liy + i][clocX], mat_kernelY[i], sum1);
+                sum2 = mad(lsmem[liy + i][clocX], mat_kernelX[i], sum2);
+            }
+            lsmemDy1[liy][clocX] = sum1;
+            lsmemDy2[liy][clocX] = sum2;
+            clocX += BLK_X;
+        }
+        while(clocX < BLK_X+(RADIUS*2));
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+        if ((x < dst_cols) && (y + liy < dst_rows))
+        {
+            sum1 = (WT) 0;
+            sum2 = (WT) 0;
+            for (i=0; i<=2*RADIUS; i++)
+            {
+                sum1 = mad(lsmemDy1[liy][lix+i], mat_kernelX[i], sum1);
+                sum2 = mad(lsmemDy2[liy][lix+i], mat_kernelY[i], sum2);
+            }
+
+            WT sum = mad(scale_v, (sum1 + sum2), delta_v);
+            storepix(convertToDT(sum), Dst + mad24(y + liy, dst_step, mad24(x, DSTSIZE, dst_offset)));
+        }
+
+        for (int i = liy * BLK_X + lix; i < (RADIUS*2) * (BLK_X+(RADIUS*2)); i += BLK_X * BLK_Y)
+        {
+            int clocX = i % (BLK_X+(RADIUS*2));
+            int clocY = i / (BLK_X+(RADIUS*2));
+            lsmem[clocY][clocX] = lsmem[clocY + BLK_Y][clocX];
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+        int yb = y + liy + BLK_Y + srcOffsetY + RADIUS;
+        EXTRAPOLATE(yb, (height));
+
+        clocX = lix;
+        int cSrcX = x + srcOffsetX - RADIUS;
+        do
+        {
+            int xb = cSrcX;
+            EXTRAPOLATE(xb,(width));
+            lsmem[liy + 2*RADIUS][clocX] = ELEM(xb, yb, (width), (height), 0 );
+
+            clocX += BLK_X;
+            cSrcX += BLK_X;
+        }
+        while(clocX < BLK_X+(RADIUS*2));
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+}
+
+#endif
\ No newline at end of file
index d87a1ff..c0fb5d9 100644 (file)
@@ -48,8 +48,8 @@
 
 #include "opencv2/opencv_modules.hpp"
 
-#ifdef HAVE_OPENCV_NONFREE
-#  include "opencv2/nonfree/cuda.hpp"
+#ifdef HAVE_OPENCV_XFEATURES2D
+#  include "opencv2/xfeatures2d/cuda.hpp"
 #endif
 
 namespace cv {
@@ -104,7 +104,7 @@ private:
 };
 
 
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
 class CV_EXPORTS SurfFeaturesFinderGpu : public FeaturesFinder
 {
 public:
index a512e35..45e9725 100644 (file)
@@ -46,10 +46,10 @@ using namespace cv;
 using namespace cv::detail;
 using namespace cv::cuda;
 
-#ifdef HAVE_OPENCV_NONFREE
-#include "opencv2/nonfree.hpp"
+#ifdef HAVE_OPENCV_XFEATURES2D
+#include "opencv2/xfeatures2d.hpp"
 
-static bool makeUseOfNonfree = initModule_nonfree();
+static bool makeUseOfNonfree = initModule_xfeatures2d();
 #endif
 
 namespace {
@@ -443,7 +443,7 @@ void OrbFeaturesFinder::find(InputArray image, ImageFeatures &features)
     }
 }
 
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
 SurfFeaturesFinderGpu::SurfFeaturesFinderGpu(double hess_thresh, int num_octaves, int num_layers,
                                              int num_octaves_descr, int num_layers_descr)
 {
index 759d036..c70d9f9 100644 (file)
@@ -87,8 +87,8 @@
 #  include "opencv2/cuda.hpp"
 #endif
 
-#ifdef HAVE_OPENCV_NONFREE
-#  include "opencv2/nonfree/cuda.hpp"
+#ifdef HAVE_OPENCV_XFEATURES2D
+#  include "opencv2/xfeatures2d/cuda.hpp"
 #endif
 
 #include "../../imgproc/src/gcgraph.hpp"
index 0d4623a..f3bbe38 100644 (file)
@@ -59,7 +59,7 @@ Stitcher Stitcher::createDefault(bool try_use_gpu)
 #ifdef HAVE_OPENCV_CUDA
     if (try_use_gpu && cuda::getCudaEnabledDeviceCount() > 0)
     {
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
         stitcher.setFeaturesFinder(makePtr<detail::SurfFeaturesFinderGpu>());
 #else
         stitcher.setFeaturesFinder(makePtr<detail::OrbFeaturesFinder>());
@@ -70,7 +70,7 @@ Stitcher Stitcher::createDefault(bool try_use_gpu)
     else
 #endif
     {
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
         stitcher.setFeaturesFinder(makePtr<detail::SurfFeaturesFinder>());
 #else
         stitcher.setFeaturesFinder(makePtr<detail::OrbFeaturesFinder>());
index c7f068b..6deed7b 100644 (file)
@@ -42,7 +42,7 @@
 #include "test_precomp.hpp"
 #include "opencv2/opencv_modules.hpp"
 
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
 
 using namespace cv;
 using namespace std;
index e36dc05..30ec7ec 100644 (file)
@@ -621,15 +621,15 @@ Ptr<IVideoCapture> VideoCapture::createCameraCapture(int index)
         {
 #ifdef HAVE_DSHOW
         case CV_CAP_DSHOW:
-            capture = Ptr<IVideoCapture>(new cv::VideoCapture_DShow(index));
-            if (capture)
+            capture = makePtr<VideoCapture_DShow>(index);
+            if (capture && capture.dynamicCast<VideoCapture_DShow>()->isOpened())
                 return capture;
             break; // CV_CAP_DSHOW
 #endif
 #ifdef HAVE_INTELPERC
         case CV_CAP_INTELPERC:
-            capture = Ptr<IVideoCapture>(new cv::VideoCapture_IntelPerC());
-            if (capture)
+            capture = makePtr<VideoCapture_IntelPerC>();
+            if (capture && capture.dynamicCast<VideoCapture_IntelPerC>()->isOpened())
                 return capture;
             break; // CV_CAP_INTEL_PERC
 #endif
index bc9f05e..efbbd08 100644 (file)
@@ -1590,7 +1590,7 @@ bool videoInput::isFrameNew(int id){
 
 bool videoInput::isDeviceSetup(int id){
 
-    if(id<devicesFound && VDList[id]->readyToCapture)return true;
+    if(id>=0 && id<devicesFound && VDList[id]->readyToCapture)return true;
     else return false;
 
 }
index a46e470..4e8ab39 100644 (file)
@@ -53,8 +53,8 @@
 #ifdef HAVE_OPENCV_FEATURES2D
 #include "opencv2/features2d.hpp"
 #endif
-#ifdef HAVE_OPENCV_NONFREE
-#include "opencv2/nonfree.hpp"
+#ifdef HAVE_OPENCV_XFEATURES2D
+#include "opencv2/xfeatures2d/nonfree.hpp"
 #endif
 
 #include "opencv2/world.hpp"
index 685e44b..c8b31eb 100644 (file)
@@ -51,8 +51,8 @@ bool cv::initAll()
 #ifdef HAVE_OPENCV_FEATURES2D
     && initModule_features2d()
 #endif
-#ifdef HAVE_OPENCV_NONFREE
-    && initModule_nonfree()
+#ifdef HAVE_OPENCV_XFEATURES2D
+    && initModule_xfeatures2d()
 #endif
     ;
 }
index b6fd1e9..0d9b4a6 100644 (file)
@@ -384,7 +384,7 @@ int main(int argc, char* argv[])
     Ptr<FeaturesFinder> finder;
     if (features_type == "surf")
     {
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
         if (try_cuda && cuda::getCudaEnabledDeviceCount() > 0)
             finder = makePtr<SurfFeaturesFinderGpu>();
         else
index 0995295..65fe4ef 100644 (file)
@@ -19,8 +19,8 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
     "${OpenCV_SOURCE_DIR}/modules/gpu/src/nvidia/core"
     )
 
-  if(HAVE_opencv_nonfree)
-    ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/nonfree/include")
+  if(HAVE_opencv_xfeatures2d)
+    ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/xfeatures2d/include")
   endif()
 
   if(HAVE_opencv_cudacodec)
@@ -52,8 +52,8 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
       ocv_target_link_libraries(${the_target} ${CUDA_CUDA_LIBRARY})
     endif()
 
-    if(HAVE_opencv_nonfree)
-      ocv_target_link_libraries(${the_target} opencv_nonfree)
+    if(HAVE_opencv_xfeatures2d)
+      ocv_target_link_libraries(${the_target} opencv_xfeatures2d)
     endif()
     if(HAVE_opencv_cudacodec)
       ocv_target_link_libraries(${the_target} opencv_cudacodec)
index 07125c2..0b25663 100644 (file)
@@ -3,15 +3,15 @@ set(the_target "example_gpu_performance")
 file(GLOB sources "performance/*.cpp")
 file(GLOB headers "performance/*.h")
 
-if(HAVE_opencv_nonfree)
-  ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/nonfree/include")
+if(HAVE_opencv_xfeatures2d)
+  ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/xfeatures2d/include")
 endif()
 
 add_executable(${the_target} ${sources} ${headers})
 ocv_target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${OPENCV_CUDA_SAMPLES_REQUIRED_DEPS})
 
-if(HAVE_opencv_nonfree)
-  ocv_target_link_libraries(${the_target} opencv_nonfree)
+if(HAVE_opencv_xfeatures2d)
+  ocv_target_link_libraries(${the_target} opencv_xfeatures2d)
 endif()
 
 set_target_properties(${the_target} PROPERTIES
index 3ca9e6a..4d6f778 100644 (file)
@@ -17,9 +17,9 @@
 
 #include "opencv2/opencv_modules.hpp"
 
-#ifdef HAVE_OPENCV_NONFREE
-#include "opencv2/nonfree/cuda.hpp"
-#include "opencv2/nonfree/nonfree.hpp"
+#ifdef HAVE_OPENCV_XFEATURES2D
+#include "opencv2/xfeatures2d/cuda.hpp"
+#include "opencv2/xfeatures2d/nonfree.hpp"
 #endif
 
 using namespace std;
@@ -274,7 +274,7 @@ TEST(meanShift)
     }
 }
 
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
 
 TEST(SURF)
 {
index c267dd4..21cbc9d 100644 (file)
@@ -2,13 +2,13 @@
 
 #include "opencv2/opencv_modules.hpp"
 
-#ifdef HAVE_OPENCV_NONFREE
+#ifdef HAVE_OPENCV_XFEATURES2D
 
 #include "opencv2/core/core.hpp"
 #include "opencv2/features2d/features2d.hpp"
 #include "opencv2/highgui/highgui.hpp"
 #include "opencv2/cudafeatures2d.hpp"
-#include "opencv2/nonfree/cuda.hpp"
+#include "opencv2/xfeatures2d/cuda.hpp"
 
 using namespace std;
 using namespace cv;