diff --git a/CMakeLists.txt b/CMakeLists.txt index 7a61aeaca0b5..04158f753485 100644 --- a/CMakeLists.txt +++ b/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() diff --git a/cmake/OpenCVDetectAndroidSDK.cmake b/cmake/OpenCVDetectAndroidSDK.cmake index e14bcf65555f..7fc45108cb19 100644 --- a/cmake/OpenCVDetectAndroidSDK.cmake +++ b/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() diff --git a/cmake/OpenCVUtils.cmake b/cmake/OpenCVUtils.cmake index 677d7f5d5c40..f2a0197f829f 100644 --- a/cmake/OpenCVUtils.cmake +++ b/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}) diff --git a/cmake/templates/OpenCV.mk.in b/cmake/templates/OpenCV.mk.in index 078e02039fae..33d36601a160 100644 --- a/cmake/templates/OpenCV.mk.in +++ b/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) diff --git a/cmake/templates/opencv_run_all_tests.sh.in b/cmake/templates/opencv_run_all_tests.sh.in index b614900026ca..77dc1191a9b0 100644 --- a/cmake/templates/opencv_run_all_tests.sh.in +++ b/cmake/templates/opencv_run_all_tests.sh.in @@ -1,6 +1,6 @@ #!/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 diff --git a/data/CMakeLists.txt b/data/CMakeLists.txt index 2f10c82f64ac..998e78520cc1 100644 --- a/data/CMakeLists.txt +++ b/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 diff --git a/modules/core/include/opencv2/core/affine.hpp b/modules/core/include/opencv2/core/affine.hpp index 8a54762924ba..7284525c81ea 100644 --- a/modules/core/include/opencv2/core/affine.hpp +++ b/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 diff --git a/modules/cuda/test/test_gpumat.cpp b/modules/cuda/test/test_gpumat.cpp index 207635273d18..dcd368c085bf 100644 --- a/modules/cuda/test/test_gpumat.cpp +++ b/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); } } diff --git a/modules/cudaarithm/test/test_reductions.cpp b/modules/cudaarithm/test/test_reductions.cpp index 1d1594eb530b..5fd7e2dec98c 100644 --- a/modules/cudaarithm/test/test_reductions.cpp +++ b/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) diff --git a/modules/cudaimgproc/src/canny.cpp b/modules/cudaimgproc/src/canny.cpp index b22094d603de..eed4a284e5a6 100644 --- a/modules/cudaimgproc/src/canny.cpp +++ b/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::max() && image_size.height < std::numeric_limits::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(low_thresh_), static_cast(high_thresh_)); - canny::edgesHysteresisLocal(map_, st1_.ptr()); + canny::edgesHysteresisLocal(map_, st1_.ptr()); - canny::edgesHysteresisGlobal(map_, st1_.ptr(), st2_.ptr()); + canny::edgesHysteresisGlobal(map_, st1_.ptr(), st2_.ptr()); canny::getEdges(map_, edges); } diff --git a/modules/cudaimgproc/src/cuda/canny.cu b/modules/cudaimgproc/src/cuda/canny.cu index 9b691e404394..043d6e5ff764 100644 --- a/modules/cudaimgproc/src/cuda/canny.cu +++ b/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); } } diff --git a/modules/cudaimgproc/test/test_color.cpp b/modules/cudaimgproc/test/test_color.cpp index e439450079be..449444277354 100644 --- a/modules/cudaimgproc/test/test_color.cpp +++ b/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) diff --git a/modules/java/generator/src/java/android+OpenCVLoader.java b/modules/java/generator/src/java/android+OpenCVLoader.java index 46e62eb34736..0892e3af3c82 100644 --- a/modules/java/generator/src/java/android+OpenCVLoader.java +++ b/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); } /** diff --git a/modules/java/generator/src/java/android+StaticHelper.java b/modules/java/generator/src/java/android+StaticHelper.java index 8d0629c8d365..10442c904d7b 100644 --- a/modules/java/generator/src/java/android+StaticHelper.java +++ b/modules/java/generator/src/java/android+StaticHelper.java @@ -7,11 +7,21 @@ 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 @@ private static boolean loadLibrary(String Name) try { System.loadLibrary(Name); - Log.d(TAG, "OpenCV libs init was ok!"); + Log.d(TAG, "Library " + Name + " loaded"); } catch(UnsatisfiedLinkError e) { diff --git a/modules/stitching/src/blenders.cpp b/modules/stitching/src/blenders.cpp index 3e9cfb7c7dc6..446bfc131ff7 100644 --- a/modules/stitching/src/blenders.cpp +++ b/modules/stitching/src/blenders.cpp @@ -512,6 +512,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, std::vector &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 &pyr) gpu_pyr[0].download(pyr[0]); #else (void)pyr; + CV_Error(Error::StsNotImplemented, "CUDA optimization is unavailable"); #endif }