Merge pull request #2310 from SpecLad:merge-2.4

CMakeLists.txt

@@ -253,13 +253,24 @@ if(WIN32)
     message(STATUS "Can't detect runtime and/or arch")
     set(OpenCV_INSTALL_BINARIES_PREFIX "")
   endif()
+elseif(ANDROID)
+  set(OpenCV_INSTALL_BINARIES_PREFIX "sdk/native/")
 else()
   set(OpenCV_INSTALL_BINARIES_PREFIX "")
 endif()
 
-set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples")
+if(ANDROID)
+  set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples/${ANDROID_NDK_ABI_NAME}")
+else()
+  set(OPENCV_SAMPLES_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}samples")
+endif()
+
+if(ANDROID)
+  set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin/${ANDROID_NDK_ABI_NAME}")
+else()
+  set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin")
+endif()
 
-set(OPENCV_BIN_INSTALL_PATH "${OpenCV_INSTALL_BINARIES_PREFIX}bin")
 if(NOT OPENCV_TEST_INSTALL_PATH)
   set(OPENCV_TEST_INSTALL_PATH "${OPENCV_BIN_INSTALL_PATH}")
 endif()

cmake/OpenCVDetectAndroidSDK.cmake

@@ -365,7 +365,7 @@ macro(add_android_project target path)
       endif()
       install(CODE "EXECUTE_PROCESS(COMMAND ${ANDROID_EXECUTABLE} --silent update project --path . --target \"${android_proj_sdk_target}\" --name \"${target}\" ${inst_lib_opt}
                                     WORKING_DIRECTORY \"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}\"
-                                   )"  COMPONENT dev)
+                                   )"  COMPONENT samples)
       #empty 'gen'
       install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}/gen\")" COMPONENT samples)
     endif()

cmake/OpenCVUtils.cmake

@@ -467,6 +467,20 @@ macro(ocv_convert_to_full_paths VAR)
 endmacro()
 
 
+# convert list of paths to libraries names without lib prefix
+macro(ocv_convert_to_lib_name var)
+  set(__tmp "")
+  foreach(path ${ARGN})
+    get_filename_component(__tmp_name "${path}" NAME_WE)
+    string(REGEX REPLACE "^lib" "" __tmp_name ${__tmp_name})
+    list(APPEND __tmp "${__tmp_name}")
+  endforeach()
+  set(${var} ${__tmp})
+  unset(__tmp)
+  unset(__tmp_name)
+endmacro()
+
+
 # add install command
 function(ocv_install_target)
   install(TARGETS ${ARGN})

cmake/templates/OpenCV.mk.in

@@ -2,6 +2,13 @@
 # you might need to define NDK_USE_CYGPATH=1 before calling the ndk-build
 
 USER_LOCAL_PATH:=$(LOCAL_PATH)
+
+USER_LOCAL_C_INCLUDES:=$(LOCAL_C_INCLUDES)
+USER_LOCAL_CFLAGS:=$(LOCAL_CFLAGS)
+USER_LOCAL_STATIC_LIBRARIES:=$(LOCAL_STATIC_LIBRARIES)
+USER_LOCAL_SHARED_LIBRARIES:=$(LOCAL_SHARED_LIBRARIES)
+USER_LOCAL_LDLIBS:=$(LOCAL_LDLIBS)
+
 LOCAL_PATH:=$(subst ?,,$(firstword ?$(subst \, ,$(subst /, ,$(call my-dir)))))
 
 OPENCV_TARGET_ARCH_ABI:=$(TARGET_ARCH_ABI)
@@ -47,7 +54,7 @@ else
     endif
 endif
 
-ifeq (${OPENCV_CAMERA_MODULES},on)
+ifeq ($(OPENCV_CAMERA_MODULES),on)
     ifeq ($(TARGET_ARCH_ABI),armeabi)
         OPENCV_CAMERA_MODULES:=@OPENCV_CAMERA_LIBS_ARMEABI_CONFIGCMAKE@
     endif
@@ -113,6 +120,13 @@ ifeq ($(OPENCV_LOCAL_CFLAGS),)
 endif
 
 include $(CLEAR_VARS)
+
+LOCAL_C_INCLUDES:=$(USER_LOCAL_C_INCLUDES)
+LOCAL_CFLAGS:=$(USER_LOCAL_CFLAGS)
+LOCAL_STATIC_LIBRARIES:=$(USER_LOCAL_STATIC_LIBRARIES)
+LOCAL_SHARED_LIBRARIES:=$(USER_LOCAL_SHARED_LIBRARIES)
+LOCAL_LDLIBS:=$(USER_LOCAL_LDLIBS)
+
 LOCAL_C_INCLUDES += $(OPENCV_LOCAL_C_INCLUDES)
 LOCAL_CFLAGS     += $(OPENCV_LOCAL_CFLAGS)
 

cmake/templates/opencv_run_all_tests.sh.in

 #!/bin/sh
 
-OPENCV_TEST_PATH=@OPENCV_TEST_INSTALL_PATH@
+OPENCV_TEST_PATH=@CMAKE_INSTALL_PREFIX@/@OPENCV_TEST_INSTALL_PATH@
 export OPENCV_TEST_DATA_PATH=@CMAKE_INSTALL_PREFIX@/share/OpenCV/testdata
 
 SUMMARY_STATUS=0

data/CMakeLists.txt

@@ -13,6 +13,10 @@ if(INSTALL_TESTS AND OPENCV_TEST_DATA_PATH)
   if(ANDROID)
     install(DIRECTORY ${OPENCV_TEST_DATA_PATH} DESTINATION sdk/etc/testdata COMPONENT tests)
   elseif(NOT WIN32)
-    install(DIRECTORY ${OPENCV_TEST_DATA_PATH} DESTINATION share/OpenCV/testdata COMPONENT tests)
+    # CPack does not set correct permissions by default, so we do it explicitly.
+    install(DIRECTORY ${OPENCV_TEST_DATA_PATH}
+            DIRECTORY_PERMISSIONS OWNER_WRITE OWNER_READ OWNER_EXECUTE
+              GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE
+            DESTINATION share/OpenCV/testdata COMPONENT tests)
   endif()
 endif()
\ No newline at end of file

modules/core/include/opencv2/core/affine.hpp

@@ -70,7 +70,7 @@ namespace cv
         //Rodrigues vector
         Affine3(const Vec3& rvec, const Vec3& t = Vec3::all(0));
 
-        //Combines all contructors above. Supports 4x4, 3x4, 3x3, 1x3, 3x1 sizes of data matrix
+        //Combines all contructors above. Supports 4x4, 4x3, 3x3, 1x3, 3x1 sizes of data matrix
         explicit Affine3(const Mat& data, const Vec3& t = Vec3::all(0));
 
         //From 16th element array

modules/cuda/test/test_gpumat.cpp

@@ -281,7 +281,7 @@ CUDA_TEST_P(ConvertTo, WithOutScaling)
         cv::Mat dst_gold;
         src.convertTo(dst_gold, depth2);
 
-        EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+        EXPECT_MAT_NEAR(dst_gold, dst, depth2 < CV_32F ? 1.0 : 1e-4);
     }
 }
 

modules/cudaarithm/test/test_reductions.cpp

@@ -734,7 +734,7 @@ CUDA_TEST_P(Normalize, WithOutMask)
     cv::Mat dst_gold;
     cv::normalize(src, dst_gold, alpha, beta, norm_type, type);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, type < CV_32F ? 1.0 : 1e-4);
 }
 
 CUDA_TEST_P(Normalize, WithMask)

modules/cudaimgproc/src/canny.cpp

@@ -58,9 +58,9 @@ namespace canny
 
     void calcMap(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, PtrStepSzi map, float low_thresh, float high_thresh);
 
-    void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1);
+    void edgesHysteresisLocal(PtrStepSzi map, short2* st1);
 
-    void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2);
+    void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2);
 
     void getEdges(PtrStepSzi map, PtrStepSzb dst);
 }
@@ -194,6 +194,8 @@ namespace
 
     void CannyImpl::createBuf(Size image_size)
     {
+        CV_Assert(image_size.width < std::numeric_limits<short>::max() && image_size.height < std::numeric_limits<short>::max());
+
         ensureSizeIsEnough(image_size, CV_32SC1, dx_);
         ensureSizeIsEnough(image_size, CV_32SC1, dy_);
 
@@ -209,8 +211,8 @@ namespace
         ensureSizeIsEnough(image_size, CV_32FC1, mag_);
         ensureSizeIsEnough(image_size, CV_32SC1, map_);
 
-        ensureSizeIsEnough(1, image_size.area(), CV_16UC2, st1_);
-        ensureSizeIsEnough(1, image_size.area(), CV_16UC2, st2_);
+        ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st1_);
+        ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st2_);
     }
 
     void CannyImpl::CannyCaller(GpuMat& edges)
@@ -218,9 +220,9 @@ namespace
         map_.setTo(Scalar::all(0));
         canny::calcMap(dx_, dy_, mag_, map_, static_cast<float>(low_thresh_), static_cast<float>(high_thresh_));
 
-        canny::edgesHysteresisLocal(map_, st1_.ptr<ushort2>());
+        canny::edgesHysteresisLocal(map_, st1_.ptr<short2>());
 
-        canny::edgesHysteresisGlobal(map_, st1_.ptr<ushort2>(), st2_.ptr<ushort2>());
+        canny::edgesHysteresisGlobal(map_, st1_.ptr<short2>(), st2_.ptr<short2>());
 
         canny::getEdges(map_, edges);
     }

modules/cudaimgproc/src/cuda/canny.cu

@@ -239,30 +239,35 @@ namespace canny
 {
     __device__ int counter = 0;
 
-    __global__ void edgesHysteresisLocalKernel(PtrStepSzi map, ushort2* st)
+    __device__ __forceinline__ bool checkIdx(int y, int x, int rows, int cols)
+    {
+        return (y >= 0) && (y < rows) && (x >= 0) && (x < cols);
+    }
+
+    __global__ void edgesHysteresisLocalKernel(PtrStepSzi map, short2* st)
     {
         __shared__ volatile int smem[18][18];
 
         const int x = blockIdx.x * blockDim.x + threadIdx.x;
         const int y = blockIdx.y * blockDim.y + threadIdx.y;
 
-        smem[threadIdx.y + 1][threadIdx.x + 1] = x < map.cols && y < map.rows ? map(y, x) : 0;
+        smem[threadIdx.y + 1][threadIdx.x + 1] = checkIdx(y, x, map.rows, map.cols) ? map(y, x) : 0;
         if (threadIdx.y == 0)
-            smem[0][threadIdx.x + 1] = y > 0 ? map(y - 1, x) : 0;
+            smem[0][threadIdx.x + 1] = checkIdx(y - 1, x, map.rows, map.cols) ? map(y - 1, x) : 0;
         if (threadIdx.y == blockDim.y - 1)
-            smem[blockDim.y + 1][threadIdx.x + 1] = y + 1 < map.rows ? map(y + 1, x) : 0;
+            smem[blockDim.y + 1][threadIdx.x + 1] = checkIdx(y + 1, x, map.rows, map.cols) ? map(y + 1, x) : 0;
         if (threadIdx.x == 0)
-            smem[threadIdx.y + 1][0] = x > 0 ? map(y, x - 1) : 0;
+            smem[threadIdx.y + 1][0] = checkIdx(y, x - 1, map.rows, map.cols) ? map(y, x - 1) : 0;
         if (threadIdx.x == blockDim.x - 1)
-            smem[threadIdx.y + 1][blockDim.x + 1] = x + 1 < map.cols ? map(y, x + 1) : 0;
+            smem[threadIdx.y + 1][blockDim.x + 1] = checkIdx(y, x + 1, map.rows, map.cols) ? map(y, x + 1) : 0;
         if (threadIdx.x == 0 && threadIdx.y == 0)
-            smem[0][0] = y > 0 && x > 0 ? map(y - 1, x - 1) : 0;
+            smem[0][0] = checkIdx(y - 1, x - 1, map.rows, map.cols) ? map(y - 1, x - 1) : 0;
         if (threadIdx.x == blockDim.x - 1 && threadIdx.y == 0)
-            smem[0][blockDim.x + 1] = y > 0 && x + 1 < map.cols ? map(y - 1, x + 1) : 0;
+            smem[0][blockDim.x + 1] = checkIdx(y - 1, x + 1, map.rows, map.cols) ? map(y - 1, x + 1) : 0;
         if (threadIdx.x == 0 && threadIdx.y == blockDim.y - 1)
-            smem[blockDim.y + 1][0] = y + 1 < map.rows && x > 0 ? map(y + 1, x - 1) : 0;
+            smem[blockDim.y + 1][0] = checkIdx(y + 1, x - 1, map.rows, map.cols) ? map(y + 1, x - 1) : 0;
         if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1)
-            smem[blockDim.y + 1][blockDim.x + 1] = y + 1 < map.rows && x + 1 < map.cols ? map(y + 1, x + 1) : 0;
+            smem[blockDim.y + 1][blockDim.x + 1] = checkIdx(y + 1, x + 1, map.rows, map.cols) ? map(y + 1, x + 1) : 0;
 
         __syncthreads();
 
@@ -317,11 +322,11 @@ namespace canny
         if (n > 0)
         {
             const int ind =  ::atomicAdd(&counter, 1);
-            st[ind] = make_ushort2(x, y);
+            st[ind] = make_short2(x, y);
         }
     }
 
-    void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1)
+    void edgesHysteresisLocal(PtrStepSzi map, short2* st1)
     {
         void* counter_ptr;
         cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) );
@@ -345,13 +350,13 @@ namespace canny
     __constant__ int c_dx[8] = {-1,  0,  1, -1, 1, -1, 0, 1};
     __constant__ int c_dy[8] = {-1, -1, -1,  0, 0,  1, 1, 1};
 
-    __global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, ushort2* st1, ushort2* st2, const int count)
+    __global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, short2* st1, short2* st2, const int count)
     {
         const int stack_size = 512;
 
         __shared__ int s_counter;
         __shared__ int s_ind;
-        __shared__ ushort2 s_st[stack_size];
+        __shared__ short2 s_st[stack_size];
 
         if (threadIdx.x == 0)
             s_counter = 0;
@@ -363,14 +368,14 @@ namespace canny
         if (ind >= count)
             return;
 
-        ushort2 pos = st1[ind];
+        short2 pos = st1[ind];
 
         if (threadIdx.x < 8)
         {
             pos.x += c_dx[threadIdx.x];
             pos.y += c_dy[threadIdx.x];
 
-            if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
+            if (pos.x > 0 && pos.x < map.cols - 1 && pos.y > 0 && pos.y < map.rows - 1 && map(pos.y, pos.x) == 1)
             {
                 map(pos.y, pos.x) = 2;
 
@@ -402,7 +407,7 @@ namespace canny
                 pos.x += c_dx[threadIdx.x & 7];
                 pos.y += c_dy[threadIdx.x & 7];
 
-                if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
+                if (pos.x > 0 && pos.x < map.cols - 1 && pos.y > 0 && pos.y < map.rows - 1 && map(pos.y, pos.x) == 1)
                 {
                     map(pos.y, pos.x) = 2;
 
@@ -419,8 +424,10 @@ namespace canny
         {
             if (threadIdx.x == 0)
             {
-                ind = ::atomicAdd(&counter, s_counter);
-                s_ind = ind - s_counter;
+                s_ind = ::atomicAdd(&counter, s_counter);
+
+                if (s_ind + s_counter > map.cols * map.rows)
+                    s_counter = 0;
             }
 
             __syncthreads();
@@ -432,7 +439,7 @@ namespace canny
         }
     }
 
-    void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2)
+    void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2)
     {
         void* counter_ptr;
         cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) );
@@ -454,6 +461,8 @@ namespace canny
 
             cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) );
 
+            count = min(count, map.cols * map.rows);
+
             std::swap(st1, st2);
         }
     }

modules/cudaimgproc/test/test_color.cpp

@@ -715,7 +715,7 @@ CUDA_TEST_P(CvtColor, BGR2YCrCb)
     cv::Mat dst_gold;
     cv::cvtColor(src, dst_gold, cv::COLOR_BGR2YCrCb);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 CUDA_TEST_P(CvtColor, RGB2YCrCb)
@@ -728,7 +728,7 @@ CUDA_TEST_P(CvtColor, RGB2YCrCb)
     cv::Mat dst_gold;
     cv::cvtColor(src, dst_gold, cv::COLOR_RGB2YCrCb);
 
-    EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 CUDA_TEST_P(CvtColor, BGR2YCrCb4)
@@ -749,7 +749,7 @@ CUDA_TEST_P(CvtColor, BGR2YCrCb4)
     cv::split(h_dst, channels);
     cv::merge(channels, 3, h_dst);
 
-    EXPECT_MAT_NEAR(dst_gold, h_dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 CUDA_TEST_P(CvtColor, RGBA2YCrCb4)
@@ -771,7 +771,7 @@ CUDA_TEST_P(CvtColor, RGBA2YCrCb4)
     cv::split(h_dst, channels);
     cv::merge(channels, 3, h_dst);
 
-    EXPECT_MAT_NEAR(dst_gold, h_dst, 1.0);
+    EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
 }
 
 CUDA_TEST_P(CvtColor, YCrCb2BGR)

modules/java/generator/src/java/android+OpenCVLoader.java

@@ -48,7 +48,17 @@ public class OpenCVLoader
      */
     public static boolean initDebug()
     {
-        return StaticHelper.initOpenCV();
+        return StaticHelper.initOpenCV(false);
+    }
+
+    /**
+     * Loads and initializes OpenCV library from current application package. Roughly, it's an analog of system.loadLibrary("opencv_java").
+     * @param InitCuda load and initialize CUDA runtime libraries.
+     * @return Returns true is initialization of OpenCV was successful.
+     */
+    public static boolean initDebug(boolean InitCuda)
+    {
+        return StaticHelper.initOpenCV(InitCuda);
     }
 
     /**

modules/java/generator/src/java/android+StaticHelper.java

@@ -7,11 +7,21 @@ import android.util.Log;
 
 class StaticHelper {
 
-    public static boolean initOpenCV()
+    public static boolean initOpenCV(boolean InitCuda)
     {
         boolean result;
         String libs = "";
 
+        if(InitCuda)
+        {
+            loadLibrary("cudart");
+            loadLibrary("nppc");
+            loadLibrary("nppi");
+            loadLibrary("npps");
+            loadLibrary("cufft");
+            loadLibrary("cublas");
+        }
+
         Log.d(TAG, "Trying to get library list");
 
         try
@@ -52,7 +62,7 @@ class StaticHelper {
         try
         {
             System.loadLibrary(Name);
-            Log.d(TAG, "OpenCV libs init was ok!");
+            Log.d(TAG, "Library " + Name + " loaded");
         }
         catch(UnsatisfiedLinkError e)
         {

modules/stitching/src/blenders.cpp

@@ -512,6 +512,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, std::vector<Mat> &pyr)
     (void)img;
     (void)num_levels;
     (void)pyr;
+    CV_Error(Error::StsNotImplemented, "CUDA optimization is unavailable");
 #endif
 }
 
@@ -549,6 +550,7 @@ void restoreImageFromLaplacePyrGpu(std::vector<Mat> &pyr)
     gpu_pyr[0].download(pyr[0]);
 #else
     (void)pyr;
+    CV_Error(Error::StsNotImplemented, "CUDA optimization is unavailable");
 #endif
 }