From d3e24f3cbff8a5f1503ff45cdf8d1d8118e4c049 Mon Sep 17 00:00:00 2001 From: Nghia Ho Date: Fri, 3 Jan 2014 14:18:07 +1100 Subject: [PATCH 01/13] Improved documentation for neural network --- modules/ml/doc/neural_networks.rst | 3 +++ 1 file changed, 3 insertions(+) diff --git a/modules/ml/doc/neural_networks.rst b/modules/ml/doc/neural_networks.rst index 0496e2201a..776bf243bd 100644 --- a/modules/ml/doc/neural_networks.rst +++ b/modules/ml/doc/neural_networks.rst @@ -240,6 +240,7 @@ This method applies the specified training algorithm to computing/adjusting the The RPROP training algorithm is parallelized with the TBB library. +If you are using the default ``cvANN_MLP::SIGMOID_SYM`` activation function then the output should be in the range [-1,1], instead of [0,1], for optimal results. CvANN_MLP::predict ------------------ @@ -257,6 +258,8 @@ Predicts responses for input samples. The method returns a dummy value which should be ignored. +If you are using the default ``cvANN_MLP::SIGMOID_SYM`` activation function with the default parameter values fparam1=0 and fparam2=0 then the function used is y = 1.7159*tanh(2/3 * x), so the output will range from [-1.7159, 1.7159], instead of [0,1]. + CvANN_MLP::get_layer_count -------------------------- Returns the number of layers in the MLP. From 6bf599b1bca8a58c7a656ddc169f7be0fc3094c6 Mon Sep 17 00:00:00 2001 From: Drew Jetter Date: Sat, 18 Jan 2014 16:39:50 -0700 Subject: [PATCH 02/13] Fixed bug #3489: The code assumed that two global variables would be constructed in a particular order, but global variable initialization order is compiler-dependent. --- modules/core/src/system.cpp | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) diff --git a/modules/core/src/system.cpp b/modules/core/src/system.cpp index ee0799ce05..ef45b9b5b6 100644 --- a/modules/core/src/system.cpp +++ b/modules/core/src/system.cpp @@ -1131,17 +1131,24 @@ public: } } }; -static TLSContainerStorage tlsContainerStorage; + +// This is a wrapper function that will ensure 'tlsContainerStorage' is constructed on first use. +// For more information: http://www.parashift.com/c++-faq/static-init-order-on-first-use.html +static TLSContainerStorage& getTLSContainerStorage() +{ + static TLSContainerStorage *tlsContainerStorage = new TLSContainerStorage(); + return *tlsContainerStorage; +} TLSDataContainer::TLSDataContainer() : key_(-1) { - key_ = tlsContainerStorage.allocateKey(this); + key_ = getTLSContainerStorage().allocateKey(this); } TLSDataContainer::~TLSDataContainer() { - tlsContainerStorage.releaseKey(key_, this); + getTLSContainerStorage().releaseKey(key_, this); key_ = -1; } @@ -1166,7 +1173,7 @@ TLSStorage::~TLSStorage() void*& data = tlsData_[i]; if (data) { - tlsContainerStorage.destroyData(i, data); + getTLSContainerStorage().destroyData(i, data); data = NULL; } } From 6cb90c0e976ff24af85167e71d88c9d9b2ed06c7 Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Mon, 20 Jan 2014 10:23:46 +0100 Subject: [PATCH 03/13] fix cross-compilation issue with Numpy --- cmake/OpenCVDetectPython.cmake | 37 ++++++++++++++++++++++++++-------- 1 file changed, 29 insertions(+), 8 deletions(-) diff --git a/cmake/OpenCVDetectPython.cmake b/cmake/OpenCVDetectPython.cmake index 3326bcd98e..8c488e6561 100644 --- a/cmake/OpenCVDetectPython.cmake +++ b/cmake/OpenCVDetectPython.cmake @@ -80,14 +80,29 @@ if(PYTHON_EXECUTABLE) endif() SET(PYTHON_PACKAGES_PATH "${_PYTHON_PACKAGES_PATH}" CACHE PATH "Where to install the python packages.") - if(NOT PYTHON_NUMPY_INCLUDE_DIR) - # Attempt to discover the NumPy include directory. If this succeeds, then build python API with NumPy - execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import os; os.environ['DISTUTILS_USE_SDK']='1'; import numpy.distutils; print numpy.distutils.misc_util.get_numpy_include_dirs()[0]" - RESULT_VARIABLE PYTHON_NUMPY_PROCESS - OUTPUT_VARIABLE PYTHON_NUMPY_INCLUDE_DIR - OUTPUT_STRIP_TRAILING_WHITESPACE) + if(NOT CMAKE_CROSSCOMPILING) + if(NOT PYTHON_NUMPY_INCLUDE_DIR) + # Attempt to discover the NumPy include directory. If this succeeds, then build python API with NumPy + execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import os; os.environ['DISTUTILS_USE_SDK']='1'; import numpy.distutils; print numpy.distutils.misc_util.get_numpy_include_dirs()[0]" + RESULT_VARIABLE PYTHON_NUMPY_PROCESS + OUTPUT_VARIABLE PYTHON_NUMPY_INCLUDE_DIR + OUTPUT_STRIP_TRAILING_WHITESPACE) - if(PYTHON_NUMPY_PROCESS EQUAL 0) + if(PYTHON_NUMPY_PROCESS EQUAL 0) + file(TO_CMAKE_PATH "${PYTHON_NUMPY_INCLUDE_DIR}" _PYTHON_NUMPY_INCLUDE_DIR) + set(PYTHON_NUMPY_INCLUDE_DIR ${_PYTHON_NUMPY_INCLUDE_DIR} CACHE PATH "Path to numpy headers") + endif() + endif() + else() + if(NOT PYTHON_NUMPY_INCLUDE_DIR) + message(STATUS "Cannot probe for Python/Numpy support (because we are cross-compiling OpenCV)") + message(STATUS "If you want to enable Python/Numpy support, set the following variables:") + message(STATUS " PYTHON_EXECUTABLE") + message(STATUS " PYTHON_INCLUDE_DIR") + message(STATUS " PYTHON_LIBRARY") + message(STATUS " PYTHON_NUMPY_INCLUDE_DIR") + message(STATUS " PYTHON_NUMPY_VERSION") + else() file(TO_CMAKE_PATH "${PYTHON_NUMPY_INCLUDE_DIR}" _PYTHON_NUMPY_INCLUDE_DIR) set(PYTHON_NUMPY_INCLUDE_DIR ${_PYTHON_NUMPY_INCLUDE_DIR} CACHE PATH "Path to numpy headers") endif() @@ -95,10 +110,16 @@ if(PYTHON_EXECUTABLE) if(PYTHON_NUMPY_INCLUDE_DIR) set(PYTHON_USE_NUMPY TRUE) - execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import numpy; print numpy.version.version" + if(NOT CMAKE_CROSSCOMPILING) + execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import numpy; print numpy.version.version" RESULT_VARIABLE PYTHON_NUMPY_PROCESS OUTPUT_VARIABLE PYTHON_NUMPY_VERSION OUTPUT_STRIP_TRAILING_WHITESPACE) + else() + if(NOT PYTHON_NUMPY_VERSION) + set(PYTHON_NUMPY_VERSION "undefined - cannot be probed because of the cross-compilation") + endif() + endif() endif() endif(NOT ANDROID AND NOT IOS) From eabcfa56527de1117f09665f9c6c18b8d4212d9d Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Tue, 21 Jan 2014 10:01:52 +0100 Subject: [PATCH 04/13] reorder the if's for clarity --- cmake/OpenCVDetectPython.cmake | 40 +++++++++++++++------------------- 1 file changed, 18 insertions(+), 22 deletions(-) diff --git a/cmake/OpenCVDetectPython.cmake b/cmake/OpenCVDetectPython.cmake index 8c488e6561..618b0fd222 100644 --- a/cmake/OpenCVDetectPython.cmake +++ b/cmake/OpenCVDetectPython.cmake @@ -80,21 +80,8 @@ if(PYTHON_EXECUTABLE) endif() SET(PYTHON_PACKAGES_PATH "${_PYTHON_PACKAGES_PATH}" CACHE PATH "Where to install the python packages.") - if(NOT CMAKE_CROSSCOMPILING) - if(NOT PYTHON_NUMPY_INCLUDE_DIR) - # Attempt to discover the NumPy include directory. If this succeeds, then build python API with NumPy - execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import os; os.environ['DISTUTILS_USE_SDK']='1'; import numpy.distutils; print numpy.distutils.misc_util.get_numpy_include_dirs()[0]" - RESULT_VARIABLE PYTHON_NUMPY_PROCESS - OUTPUT_VARIABLE PYTHON_NUMPY_INCLUDE_DIR - OUTPUT_STRIP_TRAILING_WHITESPACE) - - if(PYTHON_NUMPY_PROCESS EQUAL 0) - file(TO_CMAKE_PATH "${PYTHON_NUMPY_INCLUDE_DIR}" _PYTHON_NUMPY_INCLUDE_DIR) - set(PYTHON_NUMPY_INCLUDE_DIR ${_PYTHON_NUMPY_INCLUDE_DIR} CACHE PATH "Path to numpy headers") - endif() - endif() - else() - if(NOT PYTHON_NUMPY_INCLUDE_DIR) + if(NOT PYTHON_NUMPY_INCLUDE_DIR) + if(CMAKE_CROSSCOMPILING) message(STATUS "Cannot probe for Python/Numpy support (because we are cross-compiling OpenCV)") message(STATUS "If you want to enable Python/Numpy support, set the following variables:") message(STATUS " PYTHON_EXECUTABLE") @@ -103,22 +90,31 @@ if(PYTHON_EXECUTABLE) message(STATUS " PYTHON_NUMPY_INCLUDE_DIR") message(STATUS " PYTHON_NUMPY_VERSION") else() - file(TO_CMAKE_PATH "${PYTHON_NUMPY_INCLUDE_DIR}" _PYTHON_NUMPY_INCLUDE_DIR) - set(PYTHON_NUMPY_INCLUDE_DIR ${_PYTHON_NUMPY_INCLUDE_DIR} CACHE PATH "Path to numpy headers") + # Attempt to discover the NumPy include directory. If this succeeds, then build python API with NumPy + execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import os; os.environ['DISTUTILS_USE_SDK']='1'; import numpy.distutils; print numpy.distutils.misc_util.get_numpy_include_dirs()[0]" + RESULT_VARIABLE PYTHON_NUMPY_PROCESS + OUTPUT_VARIABLE PYTHON_NUMPY_INCLUDE_DIR + OUTPUT_STRIP_TRAILING_WHITESPACE) + + if(NOT PYTHON_NUMPY_PROCESS EQUAL 0) + unset(PYTHON_NUMPY_INCLUDE_DIR) + endif() endif() endif() if(PYTHON_NUMPY_INCLUDE_DIR) + file(TO_CMAKE_PATH "${PYTHON_NUMPY_INCLUDE_DIR}" _PYTHON_NUMPY_INCLUDE_DIR) + set(PYTHON_NUMPY_INCLUDE_DIR ${_PYTHON_NUMPY_INCLUDE_DIR} CACHE PATH "Path to numpy headers") set(PYTHON_USE_NUMPY TRUE) - if(NOT CMAKE_CROSSCOMPILING) + if(CMAKE_CROSSCOMPILING) + if(NOT PYTHON_NUMPY_VERSION) + set(PYTHON_NUMPY_VERSION "undefined - cannot be probed because of the cross-compilation") + endif() + else() execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import numpy; print numpy.version.version" RESULT_VARIABLE PYTHON_NUMPY_PROCESS OUTPUT_VARIABLE PYTHON_NUMPY_VERSION OUTPUT_STRIP_TRAILING_WHITESPACE) - else() - if(NOT PYTHON_NUMPY_VERSION) - set(PYTHON_NUMPY_VERSION "undefined - cannot be probed because of the cross-compilation") - endif() endif() endif() endif(NOT ANDROID AND NOT IOS) From d84738769422aad33038d90681c47486e47a0380 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 21 Jan 2014 14:35:34 +0400 Subject: [PATCH 05/13] split CUDA Hough sources --- modules/gpu/src/cuda/build_point_list.cu | 138 ++++ .../cuda/{hough.cu => generalized_hough.cu} | 638 +----------------- modules/gpu/src/cuda/hough_circles.cu | 254 +++++++ modules/gpu/src/cuda/hough_lines.cu | 212 ++++++ modules/gpu/src/cuda/hough_segments.cu | 249 +++++++ .../src/{hough.cpp => generalized_hough.cpp} | 303 --------- modules/gpu/src/hough_circles.cpp | 223 ++++++ modules/gpu/src/hough_lines.cpp | 142 ++++ modules/gpu/src/hough_segments.cpp | 110 +++ 9 files changed, 1330 insertions(+), 939 deletions(-) create mode 100644 modules/gpu/src/cuda/build_point_list.cu rename modules/gpu/src/cuda/{hough.cu => generalized_hough.cu} (65%) create mode 100644 modules/gpu/src/cuda/hough_circles.cu create mode 100644 modules/gpu/src/cuda/hough_lines.cu create mode 100644 modules/gpu/src/cuda/hough_segments.cu rename modules/gpu/src/{hough.cpp => generalized_hough.cpp} (79%) create mode 100644 modules/gpu/src/hough_circles.cpp create mode 100644 modules/gpu/src/hough_lines.cpp create mode 100644 modules/gpu/src/hough_segments.cpp diff --git a/modules/gpu/src/cuda/build_point_list.cu b/modules/gpu/src/cuda/build_point_list.cu new file mode 100644 index 0000000000..8a9c73b3f1 --- /dev/null +++ b/modules/gpu/src/cuda/build_point_list.cu @@ -0,0 +1,138 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#if !defined CUDA_DISABLER + +#include "opencv2/gpu/device/common.hpp" +#include "opencv2/gpu/device/emulation.hpp" + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + __device__ static int g_counter; + + template + __global__ void buildPointList(const PtrStepSzb src, unsigned int* list) + { + __shared__ unsigned int s_queues[4][32 * PIXELS_PER_THREAD]; + __shared__ int s_qsize[4]; + __shared__ int s_globStart[4]; + + const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (threadIdx.x == 0) + s_qsize[threadIdx.y] = 0; + __syncthreads(); + + if (y < src.rows) + { + // fill the queue + const uchar* srcRow = src.ptr(y); + for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x) + { + if (srcRow[xx]) + { + const unsigned int val = (y << 16) | xx; + const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1); + s_queues[threadIdx.y][qidx] = val; + } + } + } + + __syncthreads(); + + // let one thread reserve the space required in the global list + if (threadIdx.x == 0 && threadIdx.y == 0) + { + // find how many items are stored in each list + int totalSize = 0; + for (int i = 0; i < blockDim.y; ++i) + { + s_globStart[i] = totalSize; + totalSize += s_qsize[i]; + } + + // calculate the offset in the global list + const int globalOffset = atomicAdd(&g_counter, totalSize); + for (int i = 0; i < blockDim.y; ++i) + s_globStart[i] += globalOffset; + } + + __syncthreads(); + + // copy local queues to global queue + const int qsize = s_qsize[threadIdx.y]; + int gidx = s_globStart[threadIdx.y] + threadIdx.x; + for(int i = threadIdx.x; i < qsize; i += blockDim.x, gidx += blockDim.x) + list[gidx] = s_queues[threadIdx.y][i]; + } + + int buildPointList_gpu(PtrStepSzb src, unsigned int* list) + { + const int PIXELS_PER_THREAD = 16; + + void* counterPtr; + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + + const dim3 block(32, 4); + const dim3 grid(divUp(src.cols, block.x * PIXELS_PER_THREAD), divUp(src.rows, block.y)); + + cudaSafeCall( cudaFuncSetCacheConfig(buildPointList, cudaFuncCachePreferShared) ); + + buildPointList<<>>(src, list); + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + + int totalCount; + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + + return totalCount; + } + } +}}} + +#endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/hough.cu b/modules/gpu/src/cuda/generalized_hough.cu similarity index 65% rename from modules/gpu/src/cuda/hough.cu rename to modules/gpu/src/cuda/generalized_hough.cu index 59eba26081..5e2041eae4 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/generalized_hough.cu @@ -43,651 +43,18 @@ #if !defined CUDA_DISABLER #include -#include +#include #include "opencv2/gpu/device/common.hpp" #include "opencv2/gpu/device/emulation.hpp" #include "opencv2/gpu/device/vec_math.hpp" #include "opencv2/gpu/device/functional.hpp" -#include "opencv2/gpu/device/limits.hpp" -#include "opencv2/gpu/device/dynamic_smem.hpp" namespace cv { namespace gpu { namespace device { namespace hough { - __device__ int g_counter; - - //////////////////////////////////////////////////////////////////////// - // buildPointList - - template - __global__ void buildPointList(const PtrStepSzb src, unsigned int* list) - { - __shared__ unsigned int s_queues[4][32 * PIXELS_PER_THREAD]; - __shared__ int s_qsize[4]; - __shared__ int s_globStart[4]; - - const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x; - const int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (threadIdx.x == 0) - s_qsize[threadIdx.y] = 0; - __syncthreads(); - - if (y < src.rows) - { - // fill the queue - const uchar* srcRow = src.ptr(y); - for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x) - { - if (srcRow[xx]) - { - const unsigned int val = (y << 16) | xx; - const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1); - s_queues[threadIdx.y][qidx] = val; - } - } - } - - __syncthreads(); - - // let one thread reserve the space required in the global list - if (threadIdx.x == 0 && threadIdx.y == 0) - { - // find how many items are stored in each list - int totalSize = 0; - for (int i = 0; i < blockDim.y; ++i) - { - s_globStart[i] = totalSize; - totalSize += s_qsize[i]; - } - - // calculate the offset in the global list - const int globalOffset = atomicAdd(&g_counter, totalSize); - for (int i = 0; i < blockDim.y; ++i) - s_globStart[i] += globalOffset; - } - - __syncthreads(); - - // copy local queues to global queue - const int qsize = s_qsize[threadIdx.y]; - int gidx = s_globStart[threadIdx.y] + threadIdx.x; - for(int i = threadIdx.x; i < qsize; i += blockDim.x, gidx += blockDim.x) - list[gidx] = s_queues[threadIdx.y][i]; - } - - int buildPointList_gpu(PtrStepSzb src, unsigned int* list) - { - const int PIXELS_PER_THREAD = 16; - - void* counterPtr; - cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - - cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); - - const dim3 block(32, 4); - const dim3 grid(divUp(src.cols, block.x * PIXELS_PER_THREAD), divUp(src.rows, block.y)); - - cudaSafeCall( cudaFuncSetCacheConfig(buildPointList, cudaFuncCachePreferShared) ); - - buildPointList<<>>(src, list); - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - - int totalCount; - cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - - return totalCount; - } - - //////////////////////////////////////////////////////////////////////// - // linesAccum - - __global__ void linesAccumGlobal(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho) - { - const int n = blockIdx.x; - const float ang = n * theta; - - float sinVal; - float cosVal; - sincosf(ang, &sinVal, &cosVal); - sinVal *= irho; - cosVal *= irho; - - const int shift = (numrho - 1) / 2; - - int* accumRow = accum.ptr(n + 1); - for (int i = threadIdx.x; i < count; i += blockDim.x) - { - const unsigned int val = list[i]; - - const int x = (val & 0xFFFF); - const int y = (val >> 16) & 0xFFFF; - - int r = __float2int_rn(x * cosVal + y * sinVal); - r += shift; - - ::atomicAdd(accumRow + r + 1, 1); - } - } - - __global__ void linesAccumShared(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho) - { - int* smem = DynamicSharedMem(); - - for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x) - smem[i] = 0; - - __syncthreads(); - - const int n = blockIdx.x; - const float ang = n * theta; - - float sinVal; - float cosVal; - sincosf(ang, &sinVal, &cosVal); - sinVal *= irho; - cosVal *= irho; - - const int shift = (numrho - 1) / 2; - - for (int i = threadIdx.x; i < count; i += blockDim.x) - { - const unsigned int val = list[i]; - - const int x = (val & 0xFFFF); - const int y = (val >> 16) & 0xFFFF; - - int r = __float2int_rn(x * cosVal + y * sinVal); - r += shift; - - Emulation::smem::atomicAdd(&smem[r + 1], 1); - } - - __syncthreads(); - - int* accumRow = accum.ptr(n + 1); - for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x) - accumRow[i] = smem[i]; - } - - void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20) - { - const dim3 block(has20 ? 1024 : 512); - const dim3 grid(accum.rows - 2); - - size_t smemSize = (accum.cols - 1) * sizeof(int); - - if (smemSize < sharedMemPerBlock - 1000) - linesAccumShared<<>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2); - else - linesAccumGlobal<<>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2); - - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - } - - //////////////////////////////////////////////////////////////////////// - // linesGetResult - - __global__ void linesGetResult(const PtrStepSzi accum, float2* out, int* votes, const int maxSize, const float rho, const float theta, const int threshold, const int numrho) - { - const int r = blockIdx.x * blockDim.x + threadIdx.x; - const int n = blockIdx.y * blockDim.y + threadIdx.y; - - if (r >= accum.cols - 2 || n >= accum.rows - 2) - return; - - const int curVotes = accum(n + 1, r + 1); - - if (curVotes > threshold && - curVotes > accum(n + 1, r) && - curVotes >= accum(n + 1, r + 2) && - curVotes > accum(n, r + 1) && - curVotes >= accum(n + 2, r + 1)) - { - const float radius = (r - (numrho - 1) * 0.5f) * rho; - const float angle = n * theta; - - const int ind = ::atomicAdd(&g_counter, 1); - if (ind < maxSize) - { - out[ind] = make_float2(radius, angle); - votes[ind] = curVotes; - } - } - } - - int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort) - { - void* counterPtr; - cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - - cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); - - const dim3 block(32, 8); - const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); - - cudaSafeCall( cudaFuncSetCacheConfig(linesGetResult, cudaFuncCachePreferL1) ); - - linesGetResult<<>>(accum, out, votes, maxSize, rho, theta, threshold, accum.cols - 2); - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - - int totalCount; - cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - - totalCount = ::min(totalCount, maxSize); - - if (doSort && totalCount > 0) - { - thrust::device_ptr outPtr(out); - thrust::device_ptr votesPtr(votes); - thrust::sort_by_key(votesPtr, votesPtr + totalCount, outPtr, thrust::greater()); - } - - return totalCount; - } - - //////////////////////////////////////////////////////////////////////// - // houghLinesProbabilistic - - texture tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp); - - __global__ void houghLinesProbabilistic(const PtrStepSzi accum, - int4* out, const int maxSize, - const float rho, const float theta, - const int lineGap, const int lineLength, - const int rows, const int cols) - { - const int r = blockIdx.x * blockDim.x + threadIdx.x; - const int n = blockIdx.y * blockDim.y + threadIdx.y; - - if (r >= accum.cols - 2 || n >= accum.rows - 2) - return; - - const int curVotes = accum(n + 1, r + 1); - - if (curVotes >= lineLength && - curVotes > accum(n, r) && - curVotes > accum(n, r + 1) && - curVotes > accum(n, r + 2) && - curVotes > accum(n + 1, r) && - curVotes > accum(n + 1, r + 2) && - curVotes > accum(n + 2, r) && - curVotes > accum(n + 2, r + 1) && - curVotes > accum(n + 2, r + 2)) - { - const float radius = (r - (accum.cols - 2 - 1) * 0.5f) * rho; - const float angle = n * theta; - - float cosa; - float sina; - sincosf(angle, &sina, &cosa); - - float2 p0 = make_float2(cosa * radius, sina * radius); - float2 dir = make_float2(-sina, cosa); - - float2 pb[4] = {make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1)}; - float a; - - if (dir.x != 0) - { - a = -p0.x / dir.x; - pb[0].x = 0; - pb[0].y = p0.y + a * dir.y; - - a = (cols - 1 - p0.x) / dir.x; - pb[1].x = cols - 1; - pb[1].y = p0.y + a * dir.y; - } - if (dir.y != 0) - { - a = -p0.y / dir.y; - pb[2].x = p0.x + a * dir.x; - pb[2].y = 0; - - a = (rows - 1 - p0.y) / dir.y; - pb[3].x = p0.x + a * dir.x; - pb[3].y = rows - 1; - } - - if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < rows)) - { - p0 = pb[0]; - if (dir.x < 0) - dir = -dir; - } - else if (pb[1].x == cols - 1 && (pb[0].y >= 0 && pb[0].y < rows)) - { - p0 = pb[1]; - if (dir.x > 0) - dir = -dir; - } - else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < cols)) - { - p0 = pb[2]; - if (dir.y < 0) - dir = -dir; - } - else if (pb[3].y == rows - 1 && (pb[3].x >= 0 && pb[3].x < cols)) - { - p0 = pb[3]; - if (dir.y > 0) - dir = -dir; - } - - float2 d; - if (::fabsf(dir.x) > ::fabsf(dir.y)) - { - d.x = dir.x > 0 ? 1 : -1; - d.y = dir.y / ::fabsf(dir.x); - } - else - { - d.x = dir.x / ::fabsf(dir.y); - d.y = dir.y > 0 ? 1 : -1; - } - - float2 line_end[2]; - int gap; - bool inLine = false; - - float2 p1 = p0; - if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows) - return; - - for (;;) - { - if (tex2D(tex_mask, p1.x, p1.y)) - { - gap = 0; - - if (!inLine) - { - line_end[0] = p1; - line_end[1] = p1; - inLine = true; - } - else - { - line_end[1] = p1; - } - } - else if (inLine) - { - if (++gap > lineGap) - { - bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength || - ::abs(line_end[1].y - line_end[0].y) >= lineLength; - - if (good_line) - { - const int ind = ::atomicAdd(&g_counter, 1); - if (ind < maxSize) - out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); - } - - gap = 0; - inLine = false; - } - } - - p1 = p1 + d; - if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows) - { - if (inLine) - { - bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength || - ::abs(line_end[1].y - line_end[0].y) >= lineLength; - - if (good_line) - { - const int ind = ::atomicAdd(&g_counter, 1); - if (ind < maxSize) - out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); - } - - } - break; - } - } - } - } - - int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength) - { - void* counterPtr; - cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - - cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); - - const dim3 block(32, 8); - const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); - - bindTexture(&tex_mask, mask); - - houghLinesProbabilistic<<>>(accum, - out, maxSize, - rho, theta, - lineGap, lineLength, - mask.rows, mask.cols); - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - - int totalCount; - cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - - totalCount = ::min(totalCount, maxSize); - - return totalCount; - } - - //////////////////////////////////////////////////////////////////////// - // circlesAccumCenters - - __global__ void circlesAccumCenters(const unsigned int* list, const int count, const PtrStepi dx, const PtrStepi dy, - PtrStepi accum, const int width, const int height, const int minRadius, const int maxRadius, const float idp) - { - const int SHIFT = 10; - const int ONE = 1 << SHIFT; - - const int tid = blockIdx.x * blockDim.x + threadIdx.x; - - if (tid >= count) - return; - - const unsigned int val = list[tid]; - - const int x = (val & 0xFFFF); - const int y = (val >> 16) & 0xFFFF; - - const int vx = dx(y, x); - const int vy = dy(y, x); - - if (vx == 0 && vy == 0) - return; - - const float mag = ::sqrtf(vx * vx + vy * vy); - - const int x0 = __float2int_rn((x * idp) * ONE); - const int y0 = __float2int_rn((y * idp) * ONE); - - int sx = __float2int_rn((vx * idp) * ONE / mag); - int sy = __float2int_rn((vy * idp) * ONE / mag); - - // Step from minRadius to maxRadius in both directions of the gradient - for (int k1 = 0; k1 < 2; ++k1) - { - int x1 = x0 + minRadius * sx; - int y1 = y0 + minRadius * sy; - - for (int r = minRadius; r <= maxRadius; x1 += sx, y1 += sy, ++r) - { - const int x2 = x1 >> SHIFT; - const int y2 = y1 >> SHIFT; - - if (x2 < 0 || x2 >= width || y2 < 0 || y2 >= height) - break; - - ::atomicAdd(accum.ptr(y2 + 1) + x2 + 1, 1); - } - - sx = -sx; - sy = -sy; - } - } - - void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp) - { - const dim3 block(256); - const dim3 grid(divUp(count, block.x)); - - cudaSafeCall( cudaFuncSetCacheConfig(circlesAccumCenters, cudaFuncCachePreferL1) ); - - circlesAccumCenters<<>>(list, count, dx, dy, accum, accum.cols - 2, accum.rows - 2, minRadius, maxRadius, idp); - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - } - - //////////////////////////////////////////////////////////////////////// - // buildCentersList - - __global__ void buildCentersList(const PtrStepSzi accum, unsigned int* centers, const int threshold) - { - const int x = blockIdx.x * blockDim.x + threadIdx.x; - const int y = blockIdx.y * blockDim.y + threadIdx.y; - - if (x < accum.cols - 2 && y < accum.rows - 2) - { - const int top = accum(y, x + 1); - - const int left = accum(y + 1, x); - const int cur = accum(y + 1, x + 1); - const int right = accum(y + 1, x + 2); - - const int bottom = accum(y + 2, x + 1); - - if (cur > threshold && cur > top && cur >= bottom && cur > left && cur >= right) - { - const unsigned int val = (y << 16) | x; - const int idx = ::atomicAdd(&g_counter, 1); - centers[idx] = val; - } - } - } - - int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold) - { - void* counterPtr; - cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - - cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); - - const dim3 block(32, 8); - const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); - - cudaSafeCall( cudaFuncSetCacheConfig(buildCentersList, cudaFuncCachePreferL1) ); - - buildCentersList<<>>(accum, centers, threshold); - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - - int totalCount; - cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - - return totalCount; - } - - //////////////////////////////////////////////////////////////////////// - // circlesAccumRadius - - __global__ void circlesAccumRadius(const unsigned int* centers, const unsigned int* list, const int count, - float3* circles, const int maxCircles, const float dp, - const int minRadius, const int maxRadius, const int histSize, const int threshold) - { - int* smem = DynamicSharedMem(); - - for (int i = threadIdx.x; i < histSize + 2; i += blockDim.x) - smem[i] = 0; - __syncthreads(); - - unsigned int val = centers[blockIdx.x]; - - float cx = (val & 0xFFFF); - float cy = (val >> 16) & 0xFFFF; - - cx = (cx + 0.5f) * dp; - cy = (cy + 0.5f) * dp; - - for (int i = threadIdx.x; i < count; i += blockDim.x) - { - val = list[i]; - - const int x = (val & 0xFFFF); - const int y = (val >> 16) & 0xFFFF; - - const float rad = ::sqrtf((cx - x) * (cx - x) + (cy - y) * (cy - y)); - if (rad >= minRadius && rad <= maxRadius) - { - const int r = __float2int_rn(rad - minRadius); - - Emulation::smem::atomicAdd(&smem[r + 1], 1); - } - } - - __syncthreads(); - - for (int i = threadIdx.x; i < histSize; i += blockDim.x) - { - const int curVotes = smem[i + 1]; - - if (curVotes >= threshold && curVotes > smem[i] && curVotes >= smem[i + 2]) - { - const int ind = ::atomicAdd(&g_counter, 1); - if (ind < maxCircles) - circles[ind] = make_float3(cx, cy, i + minRadius); - } - } - } - - int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count, - float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20) - { - void* counterPtr; - cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); - - cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); - - const dim3 block(has20 ? 1024 : 512); - const dim3 grid(centersCount); - - const int histSize = maxRadius - minRadius + 1; - size_t smemSize = (histSize + 2) * sizeof(int); - - circlesAccumRadius<<>>(centers, list, count, circles, maxCircles, dp, minRadius, maxRadius, histSize, threshold); - cudaSafeCall( cudaGetLastError() ); - - cudaSafeCall( cudaDeviceSynchronize() ); - - int totalCount; - cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); - - totalCount = ::min(totalCount, maxCircles); - - return totalCount; - } - - //////////////////////////////////////////////////////////////////////// - // Generalized Hough + __device__ static int g_counter; template __global__ void buildEdgePointList(const PtrStepSzb edges, const PtrStep dx, const PtrStep dy, unsigned int* coordList, float* thetaList) @@ -1706,5 +1073,4 @@ namespace cv { namespace gpu { namespace device } }}} - #endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/hough_circles.cu b/modules/gpu/src/cuda/hough_circles.cu new file mode 100644 index 0000000000..5df7887b0b --- /dev/null +++ b/modules/gpu/src/cuda/hough_circles.cu @@ -0,0 +1,254 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#if !defined CUDA_DISABLER + +#include "opencv2/gpu/device/common.hpp" +#include "opencv2/gpu/device/emulation.hpp" +#include "opencv2/gpu/device/dynamic_smem.hpp" + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + __device__ static int g_counter; + + //////////////////////////////////////////////////////////////////////// + // circlesAccumCenters + + __global__ void circlesAccumCenters(const unsigned int* list, const int count, const PtrStepi dx, const PtrStepi dy, + PtrStepi accum, const int width, const int height, const int minRadius, const int maxRadius, const float idp) + { + const int SHIFT = 10; + const int ONE = 1 << SHIFT; + + const int tid = blockIdx.x * blockDim.x + threadIdx.x; + + if (tid >= count) + return; + + const unsigned int val = list[tid]; + + const int x = (val & 0xFFFF); + const int y = (val >> 16) & 0xFFFF; + + const int vx = dx(y, x); + const int vy = dy(y, x); + + if (vx == 0 && vy == 0) + return; + + const float mag = ::sqrtf(vx * vx + vy * vy); + + const int x0 = __float2int_rn((x * idp) * ONE); + const int y0 = __float2int_rn((y * idp) * ONE); + + int sx = __float2int_rn((vx * idp) * ONE / mag); + int sy = __float2int_rn((vy * idp) * ONE / mag); + + // Step from minRadius to maxRadius in both directions of the gradient + for (int k1 = 0; k1 < 2; ++k1) + { + int x1 = x0 + minRadius * sx; + int y1 = y0 + minRadius * sy; + + for (int r = minRadius; r <= maxRadius; x1 += sx, y1 += sy, ++r) + { + const int x2 = x1 >> SHIFT; + const int y2 = y1 >> SHIFT; + + if (x2 < 0 || x2 >= width || y2 < 0 || y2 >= height) + break; + + ::atomicAdd(accum.ptr(y2 + 1) + x2 + 1, 1); + } + + sx = -sx; + sy = -sy; + } + } + + void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp) + { + const dim3 block(256); + const dim3 grid(divUp(count, block.x)); + + cudaSafeCall( cudaFuncSetCacheConfig(circlesAccumCenters, cudaFuncCachePreferL1) ); + + circlesAccumCenters<<>>(list, count, dx, dy, accum, accum.cols - 2, accum.rows - 2, minRadius, maxRadius, idp); + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + } + + //////////////////////////////////////////////////////////////////////// + // buildCentersList + + __global__ void buildCentersList(const PtrStepSzi accum, unsigned int* centers, const int threshold) + { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x < accum.cols - 2 && y < accum.rows - 2) + { + const int top = accum(y, x + 1); + + const int left = accum(y + 1, x); + const int cur = accum(y + 1, x + 1); + const int right = accum(y + 1, x + 2); + + const int bottom = accum(y + 2, x + 1); + + if (cur > threshold && cur > top && cur >= bottom && cur > left && cur >= right) + { + const unsigned int val = (y << 16) | x; + const int idx = ::atomicAdd(&g_counter, 1); + centers[idx] = val; + } + } + } + + int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold) + { + void* counterPtr; + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + + const dim3 block(32, 8); + const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); + + cudaSafeCall( cudaFuncSetCacheConfig(buildCentersList, cudaFuncCachePreferL1) ); + + buildCentersList<<>>(accum, centers, threshold); + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + + int totalCount; + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + + return totalCount; + } + + //////////////////////////////////////////////////////////////////////// + // circlesAccumRadius + + __global__ void circlesAccumRadius(const unsigned int* centers, const unsigned int* list, const int count, + float3* circles, const int maxCircles, const float dp, + const int minRadius, const int maxRadius, const int histSize, const int threshold) + { + int* smem = DynamicSharedMem(); + + for (int i = threadIdx.x; i < histSize + 2; i += blockDim.x) + smem[i] = 0; + __syncthreads(); + + unsigned int val = centers[blockIdx.x]; + + float cx = (val & 0xFFFF); + float cy = (val >> 16) & 0xFFFF; + + cx = (cx + 0.5f) * dp; + cy = (cy + 0.5f) * dp; + + for (int i = threadIdx.x; i < count; i += blockDim.x) + { + val = list[i]; + + const int x = (val & 0xFFFF); + const int y = (val >> 16) & 0xFFFF; + + const float rad = ::sqrtf((cx - x) * (cx - x) + (cy - y) * (cy - y)); + if (rad >= minRadius && rad <= maxRadius) + { + const int r = __float2int_rn(rad - minRadius); + + Emulation::smem::atomicAdd(&smem[r + 1], 1); + } + } + + __syncthreads(); + + for (int i = threadIdx.x; i < histSize; i += blockDim.x) + { + const int curVotes = smem[i + 1]; + + if (curVotes >= threshold && curVotes > smem[i] && curVotes >= smem[i + 2]) + { + const int ind = ::atomicAdd(&g_counter, 1); + if (ind < maxCircles) + circles[ind] = make_float3(cx, cy, i + minRadius); + } + } + } + + int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count, + float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20) + { + void* counterPtr; + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + + const dim3 block(has20 ? 1024 : 512); + const dim3 grid(centersCount); + + const int histSize = maxRadius - minRadius + 1; + size_t smemSize = (histSize + 2) * sizeof(int); + + circlesAccumRadius<<>>(centers, list, count, circles, maxCircles, dp, minRadius, maxRadius, histSize, threshold); + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + + int totalCount; + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + + totalCount = ::min(totalCount, maxCircles); + + return totalCount; + } + } +}}} + +#endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/hough_lines.cu b/modules/gpu/src/cuda/hough_lines.cu new file mode 100644 index 0000000000..2f2fe9a7e2 --- /dev/null +++ b/modules/gpu/src/cuda/hough_lines.cu @@ -0,0 +1,212 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#if !defined CUDA_DISABLER + +#include +#include + +#include "opencv2/gpu/device/common.hpp" +#include "opencv2/gpu/device/emulation.hpp" +#include "opencv2/gpu/device/dynamic_smem.hpp" + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + __device__ static int g_counter; + + //////////////////////////////////////////////////////////////////////// + // linesAccum + + __global__ void linesAccumGlobal(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho) + { + const int n = blockIdx.x; + const float ang = n * theta; + + float sinVal; + float cosVal; + sincosf(ang, &sinVal, &cosVal); + sinVal *= irho; + cosVal *= irho; + + const int shift = (numrho - 1) / 2; + + int* accumRow = accum.ptr(n + 1); + for (int i = threadIdx.x; i < count; i += blockDim.x) + { + const unsigned int val = list[i]; + + const int x = (val & 0xFFFF); + const int y = (val >> 16) & 0xFFFF; + + int r = __float2int_rn(x * cosVal + y * sinVal); + r += shift; + + ::atomicAdd(accumRow + r + 1, 1); + } + } + + __global__ void linesAccumShared(const unsigned int* list, const int count, PtrStepi accum, const float irho, const float theta, const int numrho) + { + int* smem = DynamicSharedMem(); + + for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x) + smem[i] = 0; + + __syncthreads(); + + const int n = blockIdx.x; + const float ang = n * theta; + + float sinVal; + float cosVal; + sincosf(ang, &sinVal, &cosVal); + sinVal *= irho; + cosVal *= irho; + + const int shift = (numrho - 1) / 2; + + for (int i = threadIdx.x; i < count; i += blockDim.x) + { + const unsigned int val = list[i]; + + const int x = (val & 0xFFFF); + const int y = (val >> 16) & 0xFFFF; + + int r = __float2int_rn(x * cosVal + y * sinVal); + r += shift; + + Emulation::smem::atomicAdd(&smem[r + 1], 1); + } + + __syncthreads(); + + int* accumRow = accum.ptr(n + 1); + for (int i = threadIdx.x; i < numrho + 1; i += blockDim.x) + accumRow[i] = smem[i]; + } + + void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20) + { + const dim3 block(has20 ? 1024 : 512); + const dim3 grid(accum.rows - 2); + + size_t smemSize = (accum.cols - 1) * sizeof(int); + + if (smemSize < sharedMemPerBlock - 1000) + linesAccumShared<<>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2); + else + linesAccumGlobal<<>>(list, count, accum, 1.0f / rho, theta, accum.cols - 2); + + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + } + + //////////////////////////////////////////////////////////////////////// + // linesGetResult + + __global__ void linesGetResult(const PtrStepSzi accum, float2* out, int* votes, const int maxSize, const float rho, const float theta, const int threshold, const int numrho) + { + const int r = blockIdx.x * blockDim.x + threadIdx.x; + const int n = blockIdx.y * blockDim.y + threadIdx.y; + + if (r >= accum.cols - 2 || n >= accum.rows - 2) + return; + + const int curVotes = accum(n + 1, r + 1); + + if (curVotes > threshold && + curVotes > accum(n + 1, r) && + curVotes >= accum(n + 1, r + 2) && + curVotes > accum(n, r + 1) && + curVotes >= accum(n + 2, r + 1)) + { + const float radius = (r - (numrho - 1) * 0.5f) * rho; + const float angle = n * theta; + + const int ind = ::atomicAdd(&g_counter, 1); + if (ind < maxSize) + { + out[ind] = make_float2(radius, angle); + votes[ind] = curVotes; + } + } + } + + int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort) + { + void* counterPtr; + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + + const dim3 block(32, 8); + const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); + + cudaSafeCall( cudaFuncSetCacheConfig(linesGetResult, cudaFuncCachePreferL1) ); + + linesGetResult<<>>(accum, out, votes, maxSize, rho, theta, threshold, accum.cols - 2); + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + + int totalCount; + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + + totalCount = ::min(totalCount, maxSize); + + if (doSort && totalCount > 0) + { + thrust::device_ptr outPtr(out); + thrust::device_ptr votesPtr(votes); + thrust::sort_by_key(votesPtr, votesPtr + totalCount, outPtr, thrust::greater()); + } + + return totalCount; + } + } +}}} + + +#endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/hough_segments.cu b/modules/gpu/src/cuda/hough_segments.cu new file mode 100644 index 0000000000..f55bb4de9f --- /dev/null +++ b/modules/gpu/src/cuda/hough_segments.cu @@ -0,0 +1,249 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#if !defined CUDA_DISABLER + +#include "opencv2/gpu/device/common.hpp" +#include "opencv2/gpu/device/vec_math.hpp" + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + __device__ int g_counter; + + texture tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp); + + __global__ void houghLinesProbabilistic(const PtrStepSzi accum, + int4* out, const int maxSize, + const float rho, const float theta, + const int lineGap, const int lineLength, + const int rows, const int cols) + { + const int r = blockIdx.x * blockDim.x + threadIdx.x; + const int n = blockIdx.y * blockDim.y + threadIdx.y; + + if (r >= accum.cols - 2 || n >= accum.rows - 2) + return; + + const int curVotes = accum(n + 1, r + 1); + + if (curVotes >= lineLength && + curVotes > accum(n, r) && + curVotes > accum(n, r + 1) && + curVotes > accum(n, r + 2) && + curVotes > accum(n + 1, r) && + curVotes > accum(n + 1, r + 2) && + curVotes > accum(n + 2, r) && + curVotes > accum(n + 2, r + 1) && + curVotes > accum(n + 2, r + 2)) + { + const float radius = (r - (accum.cols - 2 - 1) * 0.5f) * rho; + const float angle = n * theta; + + float cosa; + float sina; + sincosf(angle, &sina, &cosa); + + float2 p0 = make_float2(cosa * radius, sina * radius); + float2 dir = make_float2(-sina, cosa); + + float2 pb[4] = {make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1)}; + float a; + + if (dir.x != 0) + { + a = -p0.x / dir.x; + pb[0].x = 0; + pb[0].y = p0.y + a * dir.y; + + a = (cols - 1 - p0.x) / dir.x; + pb[1].x = cols - 1; + pb[1].y = p0.y + a * dir.y; + } + if (dir.y != 0) + { + a = -p0.y / dir.y; + pb[2].x = p0.x + a * dir.x; + pb[2].y = 0; + + a = (rows - 1 - p0.y) / dir.y; + pb[3].x = p0.x + a * dir.x; + pb[3].y = rows - 1; + } + + if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < rows)) + { + p0 = pb[0]; + if (dir.x < 0) + dir = -dir; + } + else if (pb[1].x == cols - 1 && (pb[0].y >= 0 && pb[0].y < rows)) + { + p0 = pb[1]; + if (dir.x > 0) + dir = -dir; + } + else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < cols)) + { + p0 = pb[2]; + if (dir.y < 0) + dir = -dir; + } + else if (pb[3].y == rows - 1 && (pb[3].x >= 0 && pb[3].x < cols)) + { + p0 = pb[3]; + if (dir.y > 0) + dir = -dir; + } + + float2 d; + if (::fabsf(dir.x) > ::fabsf(dir.y)) + { + d.x = dir.x > 0 ? 1 : -1; + d.y = dir.y / ::fabsf(dir.x); + } + else + { + d.x = dir.x / ::fabsf(dir.y); + d.y = dir.y > 0 ? 1 : -1; + } + + float2 line_end[2]; + int gap; + bool inLine = false; + + float2 p1 = p0; + if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows) + return; + + for (;;) + { + if (tex2D(tex_mask, p1.x, p1.y)) + { + gap = 0; + + if (!inLine) + { + line_end[0] = p1; + line_end[1] = p1; + inLine = true; + } + else + { + line_end[1] = p1; + } + } + else if (inLine) + { + if (++gap > lineGap) + { + bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength || + ::abs(line_end[1].y - line_end[0].y) >= lineLength; + + if (good_line) + { + const int ind = ::atomicAdd(&g_counter, 1); + if (ind < maxSize) + out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); + } + + gap = 0; + inLine = false; + } + } + + p1 = p1 + d; + if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows) + { + if (inLine) + { + bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength || + ::abs(line_end[1].y - line_end[0].y) >= lineLength; + + if (good_line) + { + const int ind = ::atomicAdd(&g_counter, 1); + if (ind < maxSize) + out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y); + } + + } + break; + } + } + } + } + + int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength) + { + void* counterPtr; + cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) ); + + cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) ); + + const dim3 block(32, 8); + const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y)); + + bindTexture(&tex_mask, mask); + + houghLinesProbabilistic<<>>(accum, + out, maxSize, + rho, theta, + lineGap, lineLength, + mask.rows, mask.cols); + cudaSafeCall( cudaGetLastError() ); + + cudaSafeCall( cudaDeviceSynchronize() ); + + int totalCount; + cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) ); + + totalCount = ::min(totalCount, maxSize); + + return totalCount; + } + } +}}} + + +#endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/hough.cpp b/modules/gpu/src/generalized_hough.cpp similarity index 79% rename from modules/gpu/src/hough.cpp rename to modules/gpu/src/generalized_hough.cpp index 09cf01850e..a92c37d1a5 100644 --- a/modules/gpu/src/hough.cpp +++ b/modules/gpu/src/generalized_hough.cpp @@ -48,16 +48,6 @@ using namespace cv::gpu; #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) -void cv::gpu::HoughLines(const GpuMat&, GpuMat&, float, float, int, bool, int) { throw_nogpu(); } -void cv::gpu::HoughLines(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, bool, int) { throw_nogpu(); } -void cv::gpu::HoughLinesDownload(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); } - -void cv::gpu::HoughLinesP(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, int, int) { throw_nogpu(); } - -void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, int, float, float, int, int, int, int, int) { throw_nogpu(); } -void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, HoughCirclesBuf&, int, float, float, int, int, int, int, int) { throw_nogpu(); } -void cv::gpu::HoughCirclesDownload(const GpuMat&, OutputArray) { throw_nogpu(); } - Ptr cv::gpu::GeneralizedHough_GPU::create(int) { throw_nogpu(); return Ptr(); } cv::gpu::GeneralizedHough_GPU::~GeneralizedHough_GPU() {} void cv::gpu::GeneralizedHough_GPU::setTemplate(const GpuMat&, int, Point) { throw_nogpu(); } @@ -77,299 +67,6 @@ namespace cv { namespace gpu { namespace device } }}} -////////////////////////////////////////////////////////// -// HoughLines - -namespace cv { namespace gpu { namespace device -{ - namespace hough - { - void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20); - int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort); - } -}}} - -void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, float rho, float theta, int threshold, bool doSort, int maxLines) -{ - HoughLinesBuf buf; - HoughLines(src, lines, buf, rho, theta, threshold, doSort, maxLines); -} - -void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort, int maxLines) -{ - using namespace cv::gpu::device::hough; - - CV_Assert(src.type() == CV_8UC1); - CV_Assert(src.cols < std::numeric_limits::max()); - CV_Assert(src.rows < std::numeric_limits::max()); - - ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list); - unsigned int* srcPoints = buf.list.ptr(); - - const int pointsCount = buildPointList_gpu(src, srcPoints); - if (pointsCount == 0) - { - lines.release(); - return; - } - - const int numangle = cvRound(CV_PI / theta); - const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); - CV_Assert(numangle > 0 && numrho > 0); - - ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum); - buf.accum.setTo(Scalar::all(0)); - - DeviceInfo devInfo; - linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20)); - - ensureSizeIsEnough(2, maxLines, CV_32FC2, lines); - - int linesCount = linesGetResult_gpu(buf.accum, lines.ptr(0), lines.ptr(1), maxLines, rho, theta, threshold, doSort); - if (linesCount > 0) - lines.cols = linesCount; - else - lines.release(); -} - -void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_votes_) -{ - if (d_lines.empty()) - { - h_lines_.release(); - if (h_votes_.needed()) - h_votes_.release(); - return; - } - - CV_Assert(d_lines.rows == 2 && d_lines.type() == CV_32FC2); - - h_lines_.create(1, d_lines.cols, CV_32FC2); - Mat h_lines = h_lines_.getMat(); - d_lines.row(0).download(h_lines); - - if (h_votes_.needed()) - { - h_votes_.create(1, d_lines.cols, CV_32SC1); - Mat h_votes = h_votes_.getMat(); - GpuMat d_votes(1, d_lines.cols, CV_32SC1, const_cast(d_lines.ptr(1))); - d_votes.download(h_votes); - } -} - -////////////////////////////////////////////////////////// -// HoughLinesP - -namespace cv { namespace gpu { namespace device -{ - namespace hough - { - int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength); - } -}}} - -void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int minLineLength, int maxLineGap, int maxLines) -{ - using namespace cv::gpu::device::hough; - - CV_Assert( src.type() == CV_8UC1 ); - CV_Assert( src.cols < std::numeric_limits::max() ); - CV_Assert( src.rows < std::numeric_limits::max() ); - - ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list); - unsigned int* srcPoints = buf.list.ptr(); - - const int pointsCount = buildPointList_gpu(src, srcPoints); - if (pointsCount == 0) - { - lines.release(); - return; - } - - const int numangle = cvRound(CV_PI / theta); - const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); - CV_Assert( numangle > 0 && numrho > 0 ); - - ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum); - buf.accum.setTo(Scalar::all(0)); - - DeviceInfo devInfo; - linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20)); - - ensureSizeIsEnough(1, maxLines, CV_32SC4, lines); - - int linesCount = houghLinesProbabilistic_gpu(src, buf.accum, lines.ptr(), maxLines, rho, theta, maxLineGap, minLineLength); - - if (linesCount > 0) - lines.cols = linesCount; - else - lines.release(); -} - -////////////////////////////////////////////////////////// -// HoughCircles - -namespace cv { namespace gpu { namespace device -{ - namespace hough - { - void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp); - int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold); - int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count, - float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20); - } -}}} - -void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, int method, float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles) -{ - HoughCirclesBuf buf; - HoughCircles(src, circles, buf, method, dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles); -} - -void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, HoughCirclesBuf& buf, int method, - float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles) -{ - using namespace cv::gpu::device::hough; - - CV_Assert(src.type() == CV_8UC1); - CV_Assert(src.cols < std::numeric_limits::max()); - CV_Assert(src.rows < std::numeric_limits::max()); - CV_Assert(method == CV_HOUGH_GRADIENT); - CV_Assert(dp > 0); - CV_Assert(minRadius > 0 && maxRadius > minRadius); - CV_Assert(cannyThreshold > 0); - CV_Assert(votesThreshold > 0); - CV_Assert(maxCircles > 0); - - const float idp = 1.0f / dp; - - cv::gpu::Canny(src, buf.cannyBuf, buf.edges, std::max(cannyThreshold / 2, 1), cannyThreshold); - - ensureSizeIsEnough(2, src.size().area(), CV_32SC1, buf.list); - unsigned int* srcPoints = buf.list.ptr(0); - unsigned int* centers = buf.list.ptr(1); - - const int pointsCount = buildPointList_gpu(buf.edges, srcPoints); - if (pointsCount == 0) - { - circles.release(); - return; - } - - ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, buf.accum); - buf.accum.setTo(Scalar::all(0)); - - circlesAccumCenters_gpu(srcPoints, pointsCount, buf.cannyBuf.dx, buf.cannyBuf.dy, buf.accum, minRadius, maxRadius, idp); - - int centersCount = buildCentersList_gpu(buf.accum, centers, votesThreshold); - if (centersCount == 0) - { - circles.release(); - return; - } - - if (minDist > 1) - { - cv::AutoBuffer oldBuf_(centersCount); - cv::AutoBuffer newBuf_(centersCount); - int newCount = 0; - - ushort2* oldBuf = oldBuf_; - ushort2* newBuf = newBuf_; - - cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) ); - - const int cellSize = cvRound(minDist); - const int gridWidth = (src.cols + cellSize - 1) / cellSize; - const int gridHeight = (src.rows + cellSize - 1) / cellSize; - - std::vector< std::vector > grid(gridWidth * gridHeight); - - const float minDist2 = minDist * minDist; - - for (int i = 0; i < centersCount; ++i) - { - ushort2 p = oldBuf[i]; - - bool good = true; - - int xCell = static_cast(p.x / cellSize); - int yCell = static_cast(p.y / cellSize); - - int x1 = xCell - 1; - int y1 = yCell - 1; - int x2 = xCell + 1; - int y2 = yCell + 1; - - // boundary check - x1 = std::max(0, x1); - y1 = std::max(0, y1); - x2 = std::min(gridWidth - 1, x2); - y2 = std::min(gridHeight - 1, y2); - - for (int yy = y1; yy <= y2; ++yy) - { - for (int xx = x1; xx <= x2; ++xx) - { - vector& m = grid[yy * gridWidth + xx]; - - for(size_t j = 0; j < m.size(); ++j) - { - float dx = (float)(p.x - m[j].x); - float dy = (float)(p.y - m[j].y); - - if (dx * dx + dy * dy < minDist2) - { - good = false; - goto break_out; - } - } - } - } - - break_out: - - if(good) - { - grid[yCell * gridWidth + xCell].push_back(p); - - newBuf[newCount++] = p; - } - } - - cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) ); - centersCount = newCount; - } - - ensureSizeIsEnough(1, maxCircles, CV_32FC3, circles); - - const int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, circles.ptr(), maxCircles, - dp, minRadius, maxRadius, votesThreshold, deviceSupports(FEATURE_SET_COMPUTE_20)); - - if (circlesCount > 0) - circles.cols = circlesCount; - else - circles.release(); -} - -void cv::gpu::HoughCirclesDownload(const GpuMat& d_circles, cv::OutputArray h_circles_) -{ - if (d_circles.empty()) - { - h_circles_.release(); - return; - } - - CV_Assert(d_circles.rows == 1 && d_circles.type() == CV_32FC3); - - h_circles_.create(1, d_circles.cols, CV_32FC3); - Mat h_circles = h_circles_.getMat(); - d_circles.download(h_circles); -} - -////////////////////////////////////////////////////////// -// GeneralizedHough - namespace cv { namespace gpu { namespace device { namespace hough diff --git a/modules/gpu/src/hough_circles.cpp b/modules/gpu/src/hough_circles.cpp new file mode 100644 index 0000000000..74aa7394f2 --- /dev/null +++ b/modules/gpu/src/hough_circles.cpp @@ -0,0 +1,223 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" + +using namespace std; +using namespace cv; +using namespace cv::gpu; + +#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) + +void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, int, float, float, int, int, int, int, int) { throw_nogpu(); } +void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, HoughCirclesBuf&, int, float, float, int, int, int, int, int) { throw_nogpu(); } +void cv::gpu::HoughCirclesDownload(const GpuMat&, OutputArray) { throw_nogpu(); } + +#else /* !defined (HAVE_CUDA) */ + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + int buildPointList_gpu(PtrStepSzb src, unsigned int* list); + } +}}} + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp); + int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold); + int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count, + float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20); + } +}}} + +void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, int method, float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles) +{ + HoughCirclesBuf buf; + HoughCircles(src, circles, buf, method, dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles); +} + +void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, HoughCirclesBuf& buf, int method, + float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles) +{ + using namespace cv::gpu::device::hough; + + CV_Assert(src.type() == CV_8UC1); + CV_Assert(src.cols < std::numeric_limits::max()); + CV_Assert(src.rows < std::numeric_limits::max()); + CV_Assert(method == CV_HOUGH_GRADIENT); + CV_Assert(dp > 0); + CV_Assert(minRadius > 0 && maxRadius > minRadius); + CV_Assert(cannyThreshold > 0); + CV_Assert(votesThreshold > 0); + CV_Assert(maxCircles > 0); + + const float idp = 1.0f / dp; + + cv::gpu::Canny(src, buf.cannyBuf, buf.edges, std::max(cannyThreshold / 2, 1), cannyThreshold); + + ensureSizeIsEnough(2, src.size().area(), CV_32SC1, buf.list); + unsigned int* srcPoints = buf.list.ptr(0); + unsigned int* centers = buf.list.ptr(1); + + const int pointsCount = buildPointList_gpu(buf.edges, srcPoints); + if (pointsCount == 0) + { + circles.release(); + return; + } + + ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, buf.accum); + buf.accum.setTo(Scalar::all(0)); + + circlesAccumCenters_gpu(srcPoints, pointsCount, buf.cannyBuf.dx, buf.cannyBuf.dy, buf.accum, minRadius, maxRadius, idp); + + int centersCount = buildCentersList_gpu(buf.accum, centers, votesThreshold); + if (centersCount == 0) + { + circles.release(); + return; + } + + if (minDist > 1) + { + cv::AutoBuffer oldBuf_(centersCount); + cv::AutoBuffer newBuf_(centersCount); + int newCount = 0; + + ushort2* oldBuf = oldBuf_; + ushort2* newBuf = newBuf_; + + cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) ); + + const int cellSize = cvRound(minDist); + const int gridWidth = (src.cols + cellSize - 1) / cellSize; + const int gridHeight = (src.rows + cellSize - 1) / cellSize; + + std::vector< std::vector > grid(gridWidth * gridHeight); + + const float minDist2 = minDist * minDist; + + for (int i = 0; i < centersCount; ++i) + { + ushort2 p = oldBuf[i]; + + bool good = true; + + int xCell = static_cast(p.x / cellSize); + int yCell = static_cast(p.y / cellSize); + + int x1 = xCell - 1; + int y1 = yCell - 1; + int x2 = xCell + 1; + int y2 = yCell + 1; + + // boundary check + x1 = std::max(0, x1); + y1 = std::max(0, y1); + x2 = std::min(gridWidth - 1, x2); + y2 = std::min(gridHeight - 1, y2); + + for (int yy = y1; yy <= y2; ++yy) + { + for (int xx = x1; xx <= x2; ++xx) + { + vector& m = grid[yy * gridWidth + xx]; + + for(size_t j = 0; j < m.size(); ++j) + { + float dx = (float)(p.x - m[j].x); + float dy = (float)(p.y - m[j].y); + + if (dx * dx + dy * dy < minDist2) + { + good = false; + goto break_out; + } + } + } + } + + break_out: + + if(good) + { + grid[yCell * gridWidth + xCell].push_back(p); + + newBuf[newCount++] = p; + } + } + + cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) ); + centersCount = newCount; + } + + ensureSizeIsEnough(1, maxCircles, CV_32FC3, circles); + + const int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, circles.ptr(), maxCircles, + dp, minRadius, maxRadius, votesThreshold, deviceSupports(FEATURE_SET_COMPUTE_20)); + + if (circlesCount > 0) + circles.cols = circlesCount; + else + circles.release(); +} + +void cv::gpu::HoughCirclesDownload(const GpuMat& d_circles, cv::OutputArray h_circles_) +{ + if (d_circles.empty()) + { + h_circles_.release(); + return; + } + + CV_Assert(d_circles.rows == 1 && d_circles.type() == CV_32FC3); + + h_circles_.create(1, d_circles.cols, CV_32FC3); + Mat h_circles = h_circles_.getMat(); + d_circles.download(h_circles); +} + +#endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/hough_lines.cpp b/modules/gpu/src/hough_lines.cpp new file mode 100644 index 0000000000..4cc4067b20 --- /dev/null +++ b/modules/gpu/src/hough_lines.cpp @@ -0,0 +1,142 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" + +using namespace std; +using namespace cv; +using namespace cv::gpu; + +#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) + +void cv::gpu::HoughLines(const GpuMat&, GpuMat&, float, float, int, bool, int) { throw_nogpu(); } +void cv::gpu::HoughLines(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, bool, int) { throw_nogpu(); } +void cv::gpu::HoughLinesDownload(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); } + +#else /* !defined (HAVE_CUDA) */ + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + int buildPointList_gpu(PtrStepSzb src, unsigned int* list); + } +}}} + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20); + int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort); + } +}}} + +void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, float rho, float theta, int threshold, bool doSort, int maxLines) +{ + HoughLinesBuf buf; + HoughLines(src, lines, buf, rho, theta, threshold, doSort, maxLines); +} + +void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort, int maxLines) +{ + using namespace cv::gpu::device::hough; + + CV_Assert(src.type() == CV_8UC1); + CV_Assert(src.cols < std::numeric_limits::max()); + CV_Assert(src.rows < std::numeric_limits::max()); + + ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list); + unsigned int* srcPoints = buf.list.ptr(); + + const int pointsCount = buildPointList_gpu(src, srcPoints); + if (pointsCount == 0) + { + lines.release(); + return; + } + + const int numangle = cvRound(CV_PI / theta); + const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); + CV_Assert(numangle > 0 && numrho > 0); + + ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum); + buf.accum.setTo(Scalar::all(0)); + + DeviceInfo devInfo; + linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20)); + + ensureSizeIsEnough(2, maxLines, CV_32FC2, lines); + + int linesCount = linesGetResult_gpu(buf.accum, lines.ptr(0), lines.ptr(1), maxLines, rho, theta, threshold, doSort); + if (linesCount > 0) + lines.cols = linesCount; + else + lines.release(); +} + +void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_votes_) +{ + if (d_lines.empty()) + { + h_lines_.release(); + if (h_votes_.needed()) + h_votes_.release(); + return; + } + + CV_Assert(d_lines.rows == 2 && d_lines.type() == CV_32FC2); + + h_lines_.create(1, d_lines.cols, CV_32FC2); + Mat h_lines = h_lines_.getMat(); + d_lines.row(0).download(h_lines); + + if (h_votes_.needed()) + { + h_votes_.create(1, d_lines.cols, CV_32SC1); + Mat h_votes = h_votes_.getMat(); + GpuMat d_votes(1, d_lines.cols, CV_32SC1, const_cast(d_lines.ptr(1))); + d_votes.download(h_votes); + } +} + +#endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/hough_segments.cpp b/modules/gpu/src/hough_segments.cpp new file mode 100644 index 0000000000..c34f33a626 --- /dev/null +++ b/modules/gpu/src/hough_segments.cpp @@ -0,0 +1,110 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" + +using namespace std; +using namespace cv; +using namespace cv::gpu; + +#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) + +void cv::gpu::HoughLinesP(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, int, int) { throw_nogpu(); } + +#else /* !defined (HAVE_CUDA) */ + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + int buildPointList_gpu(PtrStepSzb src, unsigned int* list); + } +}}} + +namespace cv { namespace gpu { namespace device +{ + namespace hough + { + void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20); + int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength); + } +}}} + +void cv::gpu::HoughLinesP(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int minLineLength, int maxLineGap, int maxLines) +{ + using namespace cv::gpu::device::hough; + + CV_Assert( src.type() == CV_8UC1 ); + CV_Assert( src.cols < std::numeric_limits::max() ); + CV_Assert( src.rows < std::numeric_limits::max() ); + + ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list); + unsigned int* srcPoints = buf.list.ptr(); + + const int pointsCount = buildPointList_gpu(src, srcPoints); + if (pointsCount == 0) + { + lines.release(); + return; + } + + const int numangle = cvRound(CV_PI / theta); + const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); + CV_Assert( numangle > 0 && numrho > 0 ); + + ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum); + buf.accum.setTo(Scalar::all(0)); + + DeviceInfo devInfo; + linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20)); + + ensureSizeIsEnough(1, maxLines, CV_32SC4, lines); + + int linesCount = houghLinesProbabilistic_gpu(src, buf.accum, lines.ptr(), maxLines, rho, theta, maxLineGap, minLineLength); + + if (linesCount > 0) + lines.cols = linesCount; + else + lines.release(); +} + +#endif /* !defined (HAVE_CUDA) */ From 33d42b740c6fe938b63a0b25c9ad51741aba48c3 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 21 Jan 2014 14:35:51 +0400 Subject: [PATCH 06/13] disable CUDA generalized Hough Transform --- modules/gpu/src/cuda/generalized_hough.cu | 2 ++ modules/gpu/src/generalized_hough.cpp | 2 ++ modules/gpu/test/test_hough.cpp | 2 +- 3 files changed, 5 insertions(+), 1 deletion(-) diff --git a/modules/gpu/src/cuda/generalized_hough.cu b/modules/gpu/src/cuda/generalized_hough.cu index 5e2041eae4..35451e7e23 100644 --- a/modules/gpu/src/cuda/generalized_hough.cu +++ b/modules/gpu/src/cuda/generalized_hough.cu @@ -40,6 +40,8 @@ // //M*/ +#define CUDA_DISABLER + #if !defined CUDA_DISABLER #include diff --git a/modules/gpu/src/generalized_hough.cpp b/modules/gpu/src/generalized_hough.cpp index a92c37d1a5..867ba7ee71 100644 --- a/modules/gpu/src/generalized_hough.cpp +++ b/modules/gpu/src/generalized_hough.cpp @@ -40,6 +40,8 @@ // //M*/ +#define CUDA_DISABLER + #include "precomp.hpp" using namespace std; diff --git a/modules/gpu/test/test_hough.cpp b/modules/gpu/test/test_hough.cpp index f876a7a2b0..6441fc69ab 100644 --- a/modules/gpu/test/test_hough.cpp +++ b/modules/gpu/test/test_hough.cpp @@ -189,7 +189,7 @@ PARAM_TEST_CASE(GeneralizedHough, cv::gpu::DeviceInfo, UseRoi) { }; -GPU_TEST_P(GeneralizedHough, POSITION) +GPU_TEST_P(GeneralizedHough, DISABLED_POSITION) { const cv::gpu::DeviceInfo devInfo = GET_PARAM(0); cv::gpu::setDevice(devInfo.deviceID()); From b75cbfde45c00fc956f033d0af7fe6d63312fd33 Mon Sep 17 00:00:00 2001 From: Alexander Smorkalov Date: Fri, 17 Jan 2014 16:30:31 +0400 Subject: [PATCH 07/13] All installed files marked with component names for install customization. --- 3rdparty/libjasper/CMakeLists.txt | 2 +- 3rdparty/libjpeg/CMakeLists.txt | 2 +- 3rdparty/libpng/CMakeLists.txt | 2 +- 3rdparty/libtiff/CMakeLists.txt | 2 +- 3rdparty/openexr/CMakeLists.txt | 2 +- 3rdparty/tbb/CMakeLists.txt | 6 ++--- 3rdparty/zlib/CMakeLists.txt | 2 +- apps/haartraining/CMakeLists.txt | 12 ++++----- apps/traincascade/CMakeLists.txt | 4 +-- cmake/OpenCVDetectAndroidSDK.cmake | 10 +++---- cmake/OpenCVGenAndroidMK.cmake | 2 +- cmake/OpenCVGenConfig.cmake | 26 +++++++++---------- cmake/OpenCVGenHeaders.cmake | 2 +- cmake/OpenCVGenPkgconfig.cmake | 2 +- cmake/OpenCVModule.cmake | 12 ++++----- data/CMakeLists.txt | 8 +++--- doc/CMakeLists.txt | 4 +-- include/CMakeLists.txt | 4 +-- modules/androidcamera/CMakeLists.txt | 2 +- .../camera_wrapper/CMakeLists.txt | 2 +- modules/gpu/CMakeLists.txt | 2 +- modules/highgui/CMakeLists.txt | 2 +- modules/java/CMakeLists.txt | 26 +++++++++---------- modules/python/CMakeLists.txt | 12 ++++----- platforms/android/libinfo/CMakeLists.txt | 2 +- platforms/android/package/CMakeLists.txt | 2 +- platforms/android/service/CMakeLists.txt | 2 +- samples/c/CMakeLists.txt | 4 +-- samples/cpp/CMakeLists.txt | 4 +-- samples/gpu/CMakeLists.txt | 4 +-- samples/gpu/performance/CMakeLists.txt | 2 +- samples/ocl/CMakeLists.txt | 4 +-- 32 files changed, 87 insertions(+), 87 deletions(-) diff --git a/3rdparty/libjasper/CMakeLists.txt b/3rdparty/libjasper/CMakeLists.txt index dda9cd2556..7b3dcb08a5 100644 --- a/3rdparty/libjasper/CMakeLists.txt +++ b/3rdparty/libjasper/CMakeLists.txt @@ -46,5 +46,5 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(NOT BUILD_SHARED_LIBS) - ocv_install_target(${JASPER_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main) + ocv_install_target(${JASPER_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev) endif() diff --git a/3rdparty/libjpeg/CMakeLists.txt b/3rdparty/libjpeg/CMakeLists.txt index 8d622f24ff..02d71ade2b 100644 --- a/3rdparty/libjpeg/CMakeLists.txt +++ b/3rdparty/libjpeg/CMakeLists.txt @@ -39,5 +39,5 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(NOT BUILD_SHARED_LIBS) - ocv_install_target(${JPEG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main) + ocv_install_target(${JPEG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev) endif() diff --git a/3rdparty/libpng/CMakeLists.txt b/3rdparty/libpng/CMakeLists.txt index 42b6263e52..8d3d5f4976 100644 --- a/3rdparty/libpng/CMakeLists.txt +++ b/3rdparty/libpng/CMakeLists.txt @@ -55,5 +55,5 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(NOT BUILD_SHARED_LIBS) - ocv_install_target(${PNG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main) + ocv_install_target(${PNG_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev) endif() diff --git a/3rdparty/libtiff/CMakeLists.txt b/3rdparty/libtiff/CMakeLists.txt index 6c34fb5e34..addbb5551c 100644 --- a/3rdparty/libtiff/CMakeLists.txt +++ b/3rdparty/libtiff/CMakeLists.txt @@ -115,5 +115,5 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(NOT BUILD_SHARED_LIBS) - ocv_install_target(${TIFF_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main) + ocv_install_target(${TIFF_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev) endif() diff --git a/3rdparty/openexr/CMakeLists.txt b/3rdparty/openexr/CMakeLists.txt index 2b11436e1d..c4facad2fc 100644 --- a/3rdparty/openexr/CMakeLists.txt +++ b/3rdparty/openexr/CMakeLists.txt @@ -62,7 +62,7 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(NOT BUILD_SHARED_LIBS) - ocv_install_target(IlmImf EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main) + ocv_install_target(IlmImf EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev) endif() set(OPENEXR_INCLUDE_PATHS ${OPENEXR_INCLUDE_PATHS} PARENT_SCOPE) diff --git a/3rdparty/tbb/CMakeLists.txt b/3rdparty/tbb/CMakeLists.txt index 272c195b36..e16f6cd38d 100644 --- a/3rdparty/tbb/CMakeLists.txt +++ b/3rdparty/tbb/CMakeLists.txt @@ -248,9 +248,9 @@ if(ENABLE_SOLUTION_FOLDERS) endif() ocv_install_target(tbb EXPORT OpenCVModules - RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main - LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main - ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main + RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs + LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs + ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev ) # get TBB version diff --git a/3rdparty/zlib/CMakeLists.txt b/3rdparty/zlib/CMakeLists.txt index f1b28fd396..410f2420bf 100644 --- a/3rdparty/zlib/CMakeLists.txt +++ b/3rdparty/zlib/CMakeLists.txt @@ -95,5 +95,5 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(NOT BUILD_SHARED_LIBS) - ocv_install_target(${ZLIB_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT main) + ocv_install_target(${ZLIB_LIBRARY} EXPORT OpenCVModules ARCHIVE DESTINATION ${OPENCV_3P_LIB_INSTALL_PATH} COMPONENT dev) endif() diff --git a/apps/haartraining/CMakeLists.txt b/apps/haartraining/CMakeLists.txt index cdc280556e..d8a3c55c8d 100644 --- a/apps/haartraining/CMakeLists.txt +++ b/apps/haartraining/CMakeLists.txt @@ -71,14 +71,14 @@ set_target_properties(opencv_performance PROPERTIES if(INSTALL_CREATE_DISTRIB) if(BUILD_SHARED_LIBS) - install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main) - install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main) - install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main) + install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev) + install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev) + install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev) endif() else() - install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main) - install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main) - install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main) + install(TARGETS opencv_haartraining RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev) + install(TARGETS opencv_createsamples RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev) + install(TARGETS opencv_performance RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev) endif() if(ENABLE_SOLUTION_FOLDERS) diff --git a/apps/traincascade/CMakeLists.txt b/apps/traincascade/CMakeLists.txt index 8f6fbe0348..941c0ec711 100644 --- a/apps/traincascade/CMakeLists.txt +++ b/apps/traincascade/CMakeLists.txt @@ -35,8 +35,8 @@ endif() if(INSTALL_CREATE_DISTRIB) if(BUILD_SHARED_LIBS) - install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT main) + install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} CONFIGURATIONS Release COMPONENT dev) endif() else() - install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main) + install(TARGETS ${the_target} RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev) endif() diff --git a/cmake/OpenCVDetectAndroidSDK.cmake b/cmake/OpenCVDetectAndroidSDK.cmake index 393dbb62d2..7a37051e49 100644 --- a/cmake/OpenCVDetectAndroidSDK.cmake +++ b/cmake/OpenCVDetectAndroidSDK.cmake @@ -344,20 +344,20 @@ macro(add_android_project target path) add_custom_command(TARGET ${target} POST_BUILD COMMAND ${CMAKE_COMMAND} -E copy "${android_proj_bin_dir}/bin/${target}-debug.apk" "${OpenCV_BINARY_DIR}/bin/${target}.apk") if(INSTALL_ANDROID_EXAMPLES AND "${target}" MATCHES "^example-") #apk - install(FILES "${OpenCV_BINARY_DIR}/bin/${target}.apk" DESTINATION "samples" COMPONENT main) + install(FILES "${OpenCV_BINARY_DIR}/bin/${target}.apk" DESTINATION "samples" COMPONENT samples) get_filename_component(sample_dir "${path}" NAME) #java part list(REMOVE_ITEM android_proj_files ${ANDROID_MANIFEST_FILE}) foreach(f ${android_proj_files} ${ANDROID_MANIFEST_FILE}) get_filename_component(install_subdir "${f}" PATH) - install(FILES "${android_proj_bin_dir}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT main) + install(FILES "${android_proj_bin_dir}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT samples) endforeach() #jni part + eclipse files file(GLOB_RECURSE jni_files RELATIVE "${path}" "${path}/jni/*" "${path}/.cproject") ocv_list_filterout(jni_files "\\\\.svn") foreach(f ${jni_files} ".classpath" ".project" ".settings/org.eclipse.jdt.core.prefs") get_filename_component(install_subdir "${f}" PATH) - install(FILES "${path}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT main) + install(FILES "${path}/${f}" DESTINATION "samples/${sample_dir}/${install_subdir}" COMPONENT samples) endforeach() #update proj if(android_proj_lib_deps_commands) @@ -365,9 +365,9 @@ 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 main) + )" COMPONENT dev) #empty 'gen' - install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}/gen\")" COMPONENT main) + install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/samples/${sample_dir}/gen\")" COMPONENT samples) endif() endif() endmacro() diff --git a/cmake/OpenCVGenAndroidMK.cmake b/cmake/OpenCVGenAndroidMK.cmake index eed47652b4..45193a2732 100644 --- a/cmake/OpenCVGenAndroidMK.cmake +++ b/cmake/OpenCVGenAndroidMK.cmake @@ -116,5 +116,5 @@ if(ANDROID) set(OPENCV_3RDPARTY_LIBS_DIR_CONFIGCMAKE "\$(OPENCV_THIS_DIR)/../3rdparty/libs/\$(OPENCV_TARGET_ARCH_ABI)") configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCV.mk.in" "${CMAKE_BINARY_DIR}/unix-install/OpenCV.mk" IMMEDIATE @ONLY) - install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCV.mk DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}) + install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCV.mk DESTINATION ${OPENCV_CONFIG_INSTALL_PATH} COMPONENT dev) endif(ANDROID) diff --git a/cmake/OpenCVGenConfig.cmake b/cmake/OpenCVGenConfig.cmake index 411d22582d..18411e8786 100644 --- a/cmake/OpenCVGenConfig.cmake +++ b/cmake/OpenCVGenConfig.cmake @@ -109,18 +109,18 @@ if(UNIX) # ANDROID configuration is created here also # /(share|lib)/*/ (U) # /(share|lib)/*/(cmake|CMake)/ (U) if(INSTALL_TO_MANGLED_PATHS) - install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/) - install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/) - install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ FILE OpenCVModules${modules_file_suffix}.cmake) + install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ COMPONENT dev) + install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ COMPONENT dev) + install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}-${OPENCV_VERSION}/ FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev) else() - install(FILES "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/) - install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/) - install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ FILE OpenCVModules${modules_file_suffix}.cmake) + install(FILES "${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ COMPONENT dev) + install(FILES ${CMAKE_BINARY_DIR}/unix-install/OpenCVConfig-version.cmake DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ COMPONENT dev) + install(EXPORT OpenCVModules DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev) endif() endif() if(ANDROID) - install(FILES "${OpenCV_SOURCE_DIR}/platforms/android/android.toolchain.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/) + install(FILES "${OpenCV_SOURCE_DIR}/platforms/android/android.toolchain.cmake" DESTINATION ${OPENCV_CONFIG_INSTALL_PATH}/ COMPONENT dev) endif() # -------------------------------------------------------------------------------------------- @@ -134,12 +134,12 @@ if(WIN32) configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig.cmake.in" "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" IMMEDIATE @ONLY) configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/OpenCVConfig-version.cmake.in" "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" IMMEDIATE @ONLY) if(BUILD_SHARED_LIBS) - install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib") - install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib" FILE OpenCVModules${modules_file_suffix}.cmake) + install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib" COMPONENT dev) + install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}lib" FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev) else() - install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib") - install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib" FILE OpenCVModules${modules_file_suffix}.cmake) + install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig.cmake" DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib" COMPONENT dev) + install(EXPORT OpenCVModules DESTINATION "${OpenCV_INSTALL_BINARIES_PREFIX}staticlib" FILE OpenCVModules${modules_file_suffix}.cmake COMPONENT dev) endif() - install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}") - install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}/") + install(FILES "${CMAKE_BINARY_DIR}/win-install/OpenCVConfig-version.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}" COMPONENT dev) + install(FILES "${OpenCV_SOURCE_DIR}/cmake/OpenCVConfig.cmake" DESTINATION "${CMAKE_INSTALL_PREFIX}/" COMPONENT dev) endif() diff --git a/cmake/OpenCVGenHeaders.cmake b/cmake/OpenCVGenHeaders.cmake index 35da0fb4be..c892a929c2 100644 --- a/cmake/OpenCVGenHeaders.cmake +++ b/cmake/OpenCVGenHeaders.cmake @@ -23,4 +23,4 @@ set(OPENCV_MODULE_DEFINITIONS_CONFIGMAKE "${OPENCV_MODULE_DEFINITIONS_CONFIGMAKE #endforeach() configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/opencv_modules.hpp.in" "${OPENCV_CONFIG_FILE_INCLUDE_DIR}/opencv2/opencv_modules.hpp") -install(FILES "${OPENCV_CONFIG_FILE_INCLUDE_DIR}/opencv2/opencv_modules.hpp" DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2 COMPONENT main) +install(FILES "${OPENCV_CONFIG_FILE_INCLUDE_DIR}/opencv2/opencv_modules.hpp" DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2 COMPONENT dev) diff --git a/cmake/OpenCVGenPkgconfig.cmake b/cmake/OpenCVGenPkgconfig.cmake index cd54f11bf3..13b1e44970 100644 --- a/cmake/OpenCVGenPkgconfig.cmake +++ b/cmake/OpenCVGenPkgconfig.cmake @@ -81,5 +81,5 @@ configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/opencv-XXX.pc.in" @ONLY IMMEDIATE) if(UNIX AND NOT ANDROID) - install(FILES ${CMAKE_BINARY_DIR}/unix-install/${OPENCV_PC_FILE_NAME} DESTINATION ${OPENCV_LIB_INSTALL_PATH}/pkgconfig) + install(FILES ${CMAKE_BINARY_DIR}/unix-install/${OPENCV_PC_FILE_NAME} DESTINATION ${OPENCV_LIB_INSTALL_PATH}/pkgconfig COMPONENT dev) endif() diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake index 3dd749b053..2f90a97f36 100644 --- a/cmake/OpenCVModule.cmake +++ b/cmake/OpenCVModule.cmake @@ -577,9 +577,9 @@ macro(ocv_create_module) endif() ocv_install_target(${the_module} EXPORT OpenCVModules - RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main - LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main - ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main + RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev + LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs + ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT dev ) # only "public" headers need to be installed @@ -587,7 +587,7 @@ macro(ocv_create_module) foreach(hdr ${OPENCV_MODULE_${the_module}_HEADERS}) string(REGEX REPLACE "^.*opencv2/" "opencv2/" hdr2 "${hdr}") if(hdr2 MATCHES "^(opencv2/.*)/[^/]+.h(..)?$") - install(FILES ${hdr} DESTINATION "${OPENCV_INCLUDE_INSTALL_PATH}/${CMAKE_MATCH_1}" COMPONENT main) + install(FILES ${hdr} DESTINATION "${OPENCV_INCLUDE_INSTALL_PATH}/${CMAKE_MATCH_1}" COMPONENT dev) endif() endforeach() endif() @@ -795,7 +795,7 @@ function(ocv_add_samples) endif() if(WIN32) - install(TARGETS ${the_target} RUNTIME DESTINATION "samples/${module_id}" COMPONENT main) + install(TARGETS ${the_target} RUNTIME DESTINATION "samples/${module_id}" COMPONENT samples) endif() endforeach() endif() @@ -805,7 +805,7 @@ function(ocv_add_samples) file(GLOB sample_files "${samples_path}/*") install(FILES ${sample_files} DESTINATION share/OpenCV/samples/${module_id} - PERMISSIONS OWNER_READ GROUP_READ WORLD_READ) + PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples) endif() endfunction() diff --git a/data/CMakeLists.txt b/data/CMakeLists.txt index 70efd6fd03..48094df402 100644 --- a/data/CMakeLists.txt +++ b/data/CMakeLists.txt @@ -2,9 +2,9 @@ file(GLOB HAAR_CASCADES haarcascades/*.xml) file(GLOB LBP_CASCADES lbpcascades/*.xml) if(ANDROID) - install(FILES ${HAAR_CASCADES} DESTINATION sdk/etc/haarcascades COMPONENT main) - install(FILES ${LBP_CASCADES} DESTINATION sdk/etc/lbpcascades COMPONENT main) + install(FILES ${HAAR_CASCADES} DESTINATION sdk/etc/haarcascades COMPONENT libs) + install(FILES ${LBP_CASCADES} DESTINATION sdk/etc/lbpcascades COMPONENT libs) elseif(NOT WIN32) - install(FILES ${HAAR_CASCADES} DESTINATION share/OpenCV/haarcascades COMPONENT main) - install(FILES ${LBP_CASCADES} DESTINATION share/OpenCV/lbpcascades COMPONENT main) + install(FILES ${HAAR_CASCADES} DESTINATION share/OpenCV/haarcascades COMPONENT libs) + install(FILES ${LBP_CASCADES} DESTINATION share/OpenCV/lbpcascades COMPONENT libs) endif() diff --git a/doc/CMakeLists.txt b/doc/CMakeLists.txt index 5793f7298b..3557a2bd2c 100644 --- a/doc/CMakeLists.txt +++ b/doc/CMakeLists.txt @@ -143,11 +143,11 @@ if(BUILD_DOCS AND HAVE_SPHINX) endif() foreach(f ${DOC_LIST}) - install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" COMPONENT main) + install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" COMPONENT docs) endforeach() foreach(f ${OPTIONAL_DOC_LIST}) - install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" OPTIONAL) + install(FILES "${f}" DESTINATION "${OPENCV_DOC_INSTALL_PATH}" OPTIONAL COMPONENT docs) endforeach() endif() diff --git a/include/CMakeLists.txt b/include/CMakeLists.txt index ed3b85a8fc..b4e48e6fa7 100644 --- a/include/CMakeLists.txt +++ b/include/CMakeLists.txt @@ -1,7 +1,7 @@ file(GLOB old_hdrs "opencv/*.h*") install(FILES ${old_hdrs} DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv - COMPONENT main) + COMPONENT dev) install(FILES "opencv2/opencv.hpp" DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2 - COMPONENT main) + COMPONENT dev) diff --git a/modules/androidcamera/CMakeLists.txt b/modules/androidcamera/CMakeLists.txt index 8ac8ced88e..3858ba9f6d 100644 --- a/modules/androidcamera/CMakeLists.txt +++ b/modules/androidcamera/CMakeLists.txt @@ -40,6 +40,6 @@ else() get_filename_component(wrapper_name "${wrapper}" NAME) install(FILES "${LIBRARY_OUTPUT_PATH}/${wrapper_name}" DESTINATION ${OPENCV_LIB_INSTALL_PATH} - COMPONENT main) + COMPONENT libs) endforeach() endif() diff --git a/modules/androidcamera/camera_wrapper/CMakeLists.txt b/modules/androidcamera/camera_wrapper/CMakeLists.txt index 21b9ee1ad2..bc5585a7a8 100644 --- a/modules/androidcamera/camera_wrapper/CMakeLists.txt +++ b/modules/androidcamera/camera_wrapper/CMakeLists.txt @@ -63,4 +63,4 @@ if (NOT (CMAKE_BUILD_TYPE MATCHES "debug")) endif() -install(TARGETS ${the_target} LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main) +install(TARGETS ${the_target} LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs) diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 9171febc74..6de3e0efab 100644 --- a/modules/gpu/CMakeLists.txt +++ b/modules/gpu/CMakeLists.txt @@ -82,7 +82,7 @@ ocv_create_module(${cuda_link_libs}) if(HAVE_CUDA) install(FILES src/nvidia/NPP_staging/NPP_staging.hpp src/nvidia/core/NCV.hpp DESTINATION ${OPENCV_INCLUDE_INSTALL_PATH}/opencv2/${name} - COMPONENT main) + COMPONENT dev) endif() ocv_add_precompiled_headers(${the_module}) diff --git a/modules/highgui/CMakeLists.txt b/modules/highgui/CMakeLists.txt index fd2eec6a1e..7d2547bcff 100644 --- a/modules/highgui/CMakeLists.txt +++ b/modules/highgui/CMakeLists.txt @@ -315,7 +315,7 @@ if(WIN32 AND WITH_FFMPEG) COMMENT "Copying ${ffmpeg_path} to the output directory") endif() - install(FILES "${ffmpeg_path}" DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT main RENAME "${ffmpeg_bare_name_ver}") + install(FILES "${ffmpeg_path}" DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs RENAME "${ffmpeg_bare_name_ver}") endif() ocv_add_accuracy_tests() diff --git a/modules/java/CMakeLists.txt b/modules/java/CMakeLists.txt index 3a6ebe8362..198048ebe2 100644 --- a/modules/java/CMakeLists.txt +++ b/modules/java/CMakeLists.txt @@ -175,7 +175,7 @@ foreach(java_file ${step3_input_files}) if(ANDROID) get_filename_component(install_subdir "${java_file_name}" PATH) - install(FILES "${output_name}" DESTINATION "${JAVA_INSTALL_ROOT}/src/org/opencv/${install_subdir}" COMPONENT main) + install(FILES "${output_name}" DESTINATION "${JAVA_INSTALL_ROOT}/src/org/opencv/${install_subdir}" COMPONENT java) endif() endforeach() @@ -189,7 +189,7 @@ if(ANDROID) if(NOT file MATCHES "jni/.+") get_filename_component(install_subdir "${file}" PATH) - install(FILES "${OpenCV_BINARY_DIR}/${file}" DESTINATION "${JAVA_INSTALL_ROOT}/${install_subdir}" COMPONENT main) + install(FILES "${OpenCV_BINARY_DIR}/${file}" DESTINATION "${JAVA_INSTALL_ROOT}/${install_subdir}" COMPONENT java) endif() endforeach() @@ -225,11 +225,11 @@ if(ANDROID AND ANDROID_EXECUTABLE) list(APPEND copied_files ${lib_target_files} "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}") list(APPEND step3_input_files "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}") - install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_PROJECT_PROPERTIES_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main) - install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main) + install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_PROJECT_PROPERTIES_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT java) + install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT java) # creating empty 'gen' and 'res' folders - install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/gen\")" COMPONENT main) - install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/res\")" COMPONENT main) + install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/gen\")" COMPONENT java) + install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/res\")" COMPONENT java) endif(ANDROID AND ANDROID_EXECUTABLE) set(step3_depends ${step2_depends} ${step3_input_files} ${copied_files}) @@ -282,7 +282,7 @@ else(ANDROID) else(WIN32) set(JAR_INSTALL_DIR share/OpenCV/java) endif(WIN32) - install(FILES ${JAR_FILE} DESTINATION ${JAR_INSTALL_DIR} COMPONENT main) + install(FILES ${JAR_FILE} DESTINATION ${JAR_INSTALL_DIR} COMPONENT java) endif(ANDROID) # step 5: build native part @@ -353,17 +353,17 @@ endif() if(ANDROID) ocv_install_target(${the_module} EXPORT OpenCVModules - LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main - ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main) + LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT java + ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT java) else() if(NOT INSTALL_CREATE_DISTRIB) ocv_install_target(${the_module} EXPORT OpenCVModules - RUNTIME DESTINATION ${JAR_INSTALL_DIR} COMPONENT main - LIBRARY DESTINATION ${JAR_INSTALL_DIR} COMPONENT main) + RUNTIME DESTINATION ${JAR_INSTALL_DIR} COMPONENT java + LIBRARY DESTINATION ${JAR_INSTALL_DIR} COMPONENT java) else() ocv_install_target(${the_module} EXPORT OpenCVModules - RUNTIME DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT main - LIBRARY DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT main) + RUNTIME DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT java + LIBRARY DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT java) endif() endif() diff --git a/modules/python/CMakeLists.txt b/modules/python/CMakeLists.txt index 2c44a3906d..bab8b061b3 100644 --- a/modules/python/CMakeLists.txt +++ b/modules/python/CMakeLists.txt @@ -108,17 +108,17 @@ endif() if(WIN32) set(PYTHON_INSTALL_ARCHIVE "") else() - set(PYTHON_INSTALL_ARCHIVE ARCHIVE DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main) + set(PYTHON_INSTALL_ARCHIVE ARCHIVE DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python) endif() if(NOT INSTALL_CREATE_DISTRIB) install(TARGETS ${the_module} ${PYTHON_INSTALL_CONFIGURATIONS} - RUNTIME DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main - LIBRARY DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main + RUNTIME DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python + LIBRARY DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python ${PYTHON_INSTALL_ARCHIVE} ) - install(FILES src2/cv.py ${PYTHON_INSTALL_CONFIGURATIONS} DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT main) + install(FILES src2/cv.py ${PYTHON_INSTALL_CONFIGURATIONS} DESTINATION ${PYTHON_PACKAGES_PATH} COMPONENT python) else() if(DEFINED PYTHON_VERSION_MAJOR) set(__ver "${PYTHON_VERSION_MAJOR}.${PYTHON_VERSION_MINOR}") @@ -127,7 +127,7 @@ else() endif() install(TARGETS ${the_module} CONFIGURATIONS Release - RUNTIME DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT main - LIBRARY DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT main + RUNTIME DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT python + LIBRARY DESTINATION python/${__ver}/${OpenCV_ARCH} COMPONENT python ) endif() diff --git a/platforms/android/libinfo/CMakeLists.txt b/platforms/android/libinfo/CMakeLists.txt index 028413ec6e..55dd278594 100644 --- a/platforms/android/libinfo/CMakeLists.txt +++ b/platforms/android/libinfo/CMakeLists.txt @@ -36,4 +36,4 @@ set_target_properties(${the_module} PROPERTIES ) get_filename_component(lib_name "libopencv_info.so" NAME) -install(FILES "${LIBRARY_OUTPUT_PATH}/${lib_name}" DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main) +install(FILES "${LIBRARY_OUTPUT_PATH}/${lib_name}" DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs) diff --git a/platforms/android/package/CMakeLists.txt b/platforms/android/package/CMakeLists.txt index 1382a078cf..b48a55a6a1 100644 --- a/platforms/android/package/CMakeLists.txt +++ b/platforms/android/package/CMakeLists.txt @@ -89,6 +89,6 @@ add_custom_command( DEPENDS "${OpenCV_BINARY_DIR}/bin/classes.jar.dephelper" "${PACKAGE_DIR}/res/values/strings.xml" "${PACKAGE_DIR}/res/drawable/icon.png" ${camera_wrappers} opencv_java ) -install(FILES "${APK_NAME}" DESTINATION "apk/" COMPONENT main) +install(FILES "${APK_NAME}" DESTINATION "apk/" COMPONENT libs) add_custom_target(android_package ALL SOURCES "${APK_NAME}" ) add_dependencies(android_package opencv_java) diff --git a/platforms/android/service/CMakeLists.txt b/platforms/android/service/CMakeLists.txt index dde1455138..c99b71392f 100644 --- a/platforms/android/service/CMakeLists.txt +++ b/platforms/android/service/CMakeLists.txt @@ -3,4 +3,4 @@ if(BUILD_ANDROID_SERVICE) #add_subdirectory(engine_test) endif() -install(FILES "readme.txt" DESTINATION "apk/" COMPONENT main) +install(FILES "readme.txt" DESTINATION "apk/" COMPONENT libs) diff --git a/samples/c/CMakeLists.txt b/samples/c/CMakeLists.txt index 77a42949d0..ab6e15dbf2 100644 --- a/samples/c/CMakeLists.txt +++ b/samples/c/CMakeLists.txt @@ -39,7 +39,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND) set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG") endif() install(TARGETS ${the_target} - RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/c" COMPONENT main) + RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/c" COMPONENT samples) endif() ENDMACRO() @@ -55,5 +55,5 @@ if (INSTALL_C_EXAMPLES AND NOT WIN32) file(GLOB C_SAMPLES *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd ) install(FILES ${C_SAMPLES} DESTINATION share/OpenCV/samples/c - PERMISSIONS OWNER_READ GROUP_READ WORLD_READ) + PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples) endif () diff --git a/samples/cpp/CMakeLists.txt b/samples/cpp/CMakeLists.txt index ebee5bd0a8..e0842d9e4a 100644 --- a/samples/cpp/CMakeLists.txt +++ b/samples/cpp/CMakeLists.txt @@ -68,7 +68,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND) set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG") endif() install(TARGETS ${the_target} - RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${sample_subfolder}" COMPONENT main) + RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${sample_subfolder}" COMPONENT samples) endif() ENDMACRO() @@ -92,5 +92,5 @@ if (INSTALL_C_EXAMPLES AND NOT WIN32) file(GLOB C_SAMPLES *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd ) install(FILES ${C_SAMPLES} DESTINATION share/OpenCV/samples/cpp - PERMISSIONS OWNER_READ GROUP_READ WORLD_READ) + PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples) endif() diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt index 732a9172a5..7093cf5d19 100644 --- a/samples/gpu/CMakeLists.txt +++ b/samples/gpu/CMakeLists.txt @@ -65,7 +65,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND) if(MSVC AND NOT BUILD_SHARED_LIBS) set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG") endif() - install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT main) + install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT samples) endif() ENDMACRO() @@ -84,5 +84,5 @@ if (INSTALL_C_EXAMPLES AND NOT WIN32) file(GLOB install_list *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd ) install(FILES ${install_list} DESTINATION share/OpenCV/samples/${project} - PERMISSIONS OWNER_READ GROUP_READ WORLD_READ) + PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples) endif() diff --git a/samples/gpu/performance/CMakeLists.txt b/samples/gpu/performance/CMakeLists.txt index 8f3caac5b3..0b2346f91f 100644 --- a/samples/gpu/performance/CMakeLists.txt +++ b/samples/gpu/performance/CMakeLists.txt @@ -23,7 +23,7 @@ if(ENABLE_SOLUTION_FOLDERS) endif() if(WIN32) - install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/gpu" COMPONENT main) + install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/gpu" COMPONENT samples) endif() if(INSTALL_C_EXAMPLES AND NOT WIN32) diff --git a/samples/ocl/CMakeLists.txt b/samples/ocl/CMakeLists.txt index 8db77d52c8..4139211000 100644 --- a/samples/ocl/CMakeLists.txt +++ b/samples/ocl/CMakeLists.txt @@ -38,7 +38,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND) if(MSVC AND NOT BUILD_SHARED_LIBS) set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG") endif() - install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT main) + install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT samples) endif() ENDMACRO() @@ -55,5 +55,5 @@ if (INSTALL_C_EXAMPLES AND NOT WIN32) file(GLOB install_list *.c *.cpp *.jpg *.png *.data makefile.* build_all.sh *.dsp *.cmd ) install(FILES ${install_list} DESTINATION share/OpenCV/samples/${project} - PERMISSIONS OWNER_READ GROUP_READ WORLD_READ) + PERMISSIONS OWNER_READ GROUP_READ WORLD_READ COMPONENT samples) endif() From 7821fe2bde331c6b1abd612315ca9fc59da58619 Mon Sep 17 00:00:00 2001 From: Alexander Smorkalov Date: Sat, 18 Jan 2014 23:50:07 +0400 Subject: [PATCH 08/13] Initial Linux packages build rools for CPack. --- CMakeLists.txt | 6 +++ LICENSE | 33 ++++++++++++++ cmake/OpenCVModule.cmake | 2 +- cmake/OpenCVPackaging.cmake | 89 +++++++++++++++++++++++++++++++++++++ 4 files changed, 129 insertions(+), 1 deletion(-) create mode 100644 LICENSE create mode 100644 cmake/OpenCVPackaging.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index d7db8fd130..eb25cd3474 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -975,3 +975,9 @@ ocv_finalize_status() if("${CMAKE_CURRENT_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_BINARY_DIR}") message(WARNING "The source directory is the same as binary directory. \"make clean\" may damage the source tree") endif() + +# ---------------------------------------------------------------------------- +# CPack stuff +# ---------------------------------------------------------------------------- + +include(cmake/OpenCVPackaging.cmake) diff --git a/LICENSE b/LICENSE new file mode 100644 index 0000000000..5e32d88b47 --- /dev/null +++ b/LICENSE @@ -0,0 +1,33 @@ +By downloading, copying, installing or using the software you agree to this license. +If you do not agree to this license, do not download, install, +copy or use the software. + + + License Agreement + For Open Source Computer Vision Library + (3-clause BSD License) + +Redistribution and use in source and binary forms, with or without modification, +are permitted provided that the following conditions are met: + + * Redistributions of source code must retain the above copyright notice, + this list of conditions and the following disclaimer. + + * Redistributions in binary form must reproduce the above copyright notice, + this list of conditions and the following disclaimer in the documentation + and/or other materials provided with the distribution. + + * Neither the names of the copyright holders nor the names of the contributors + may be used to endorse or promote products derived from this software + without specific prior written permission. + +This software is provided by the copyright holders and contributors "as is" and +any express or implied warranties, including, but not limited to, the implied +warranties of merchantability and fitness for a particular purpose are disclaimed. +In no event shall copyright holders or contributors be liable for any direct, +indirect, incidental, special, exemplary, or consequential damages +(including, but not limited to, procurement of substitute goods or services; +loss of use, data, or profits; or business interruption) however caused +and on any theory of liability, whether in contract, strict liability, +or tort (including negligence or otherwise) arising in any way out of +the use of this software, even if advised of the possibility of such damage. diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake index 2f90a97f36..6734462fc2 100644 --- a/cmake/OpenCVModule.cmake +++ b/cmake/OpenCVModule.cmake @@ -577,7 +577,7 @@ macro(ocv_create_module) endif() ocv_install_target(${the_module} EXPORT OpenCVModules - RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT dev + RUNTIME DESTINATION ${OPENCV_BIN_INSTALL_PATH} COMPONENT libs LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT libs ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT dev ) diff --git a/cmake/OpenCVPackaging.cmake b/cmake/OpenCVPackaging.cmake new file mode 100644 index 0000000000..4a81c255bf --- /dev/null +++ b/cmake/OpenCVPackaging.cmake @@ -0,0 +1,89 @@ +if(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake") +set(CPACK_set_DESTDIR "on") + +if(NOT OPENCV_CUSTOM_PACKAGE_INFO) + set(CPACK_PACKAGE_DESCRIPTION "Open Computer Vision Library") + set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "OpenCV") + set(CPACK_PACKAGE_VENDOR "OpenCV Foundation") + set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE") + set(CPACK_PACKAGE_CONTACT "admin@opencv.org") +endif(NOT OPENCV_CUSTOM_PACKAGE_INFO) + +set(CPACK_PACKAGE_VERSION_MAJOR "${OPENCV_VERSION_MAJOR}") +set(CPACK_PACKAGE_VERSION_MINOR "${OPENCV_VERSION_MINOR}") +set(CPACK_PACKAGE_VERSION_PATCH "${OPENCV_VERSION_PATCH}") + +#arch +if(X86) + set(CPACK_DEBIAN_ARCHITECTURE "i386") + set(CPACK_RPM_PACKAGE_ARCHITECTURE "i686") +elseif(X86_64) + set(CPACK_DEBIAN_ARCHITECTURE "amd64") + set(CPACK_RPM_PACKAGE_ARCHITECTURE "amd64") +elseif(ARM) + set(CPACK_DEBIAN_ARCHITECTURE "armhf") + set(CPACK_RPM_PACKAGE_ARCHITECTURE "armhf") +else() + set(CPACK_DEBIAN_ARCHITECTURE ${CMAKE_SYSTEM_PROCESSOR}) + set(CPACK_RPM_PACKAGE_ARCHITECTURE ${CMAKE_SYSTEM_PROCESSOR}) +endif() + +if(CPACK_GENERATOR STREQUAL "DEB") + set(OPENCV_PACKAGE_ARCH_SUFFIX ${CPACK_DEBIAN_ARCHITECTURE}) +elseif(CPACK_GENERATOR STREQUAL "RPM") + set(OPENCV_PACKAGE_ARCH_SUFFIX ${CPACK_RPM_PACKAGE_ARCHITECTURE}) +else() + set(OPENCV_PACKAGE_ARCH_SUFFIX ${CMAKE_SYSTEM_PROCESSOR}) +endif() + +set(CPACK_PACKAGE_FILE_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VCSVERSION}-${OPENCV_PACKAGE_ARCH_SUFFIX}") +set(CPACK_SOURCE_PACKAGE_FILE_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VCSVERSION}-${OPENCV_PACKAGE_ARCH_SUFFIX}") + +#rpm options +set(CPACK_RPM_COMPONENT_INSTALL TRUE) +set(CPACK_RPM_PACKAGE_LICENSE ${CPACK_RESOURCE_FILE_LICENSE}) + +#deb options +set(CPACK_DEB_COMPONENT_INSTALL TRUE) +set(CPACK_DEBIAN_PACKAGE_PRIORITY "extra") + +#depencencies +set(CPACK_DEBIAN_PACKAGE_SHLIBDEPS TRUE) +set(CPACK_COMPONENT_samples_DEPENDS libs) +set(CPACK_COMPONENT_dev_DEPENDS libs) +set(CPACK_COMPONENT_docs_DEPENDS libs) +set(CPACK_COMPONENT_java_DEPENDS libs) +set(CPACK_COMPONENT_python_DEPENDS libs) + +if(NOT OPENCV_CUSTOM_PACKAGE_INFO) + set(CPACK_COMPONENT_libs_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}") + set(CPACK_COMPONENT_libs_DESCRIPTION "Open Computer Vision Library") + + set(CPACK_COMPONENT_python_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-python") + set(CPACK_COMPONENT_python_DESCRIPTION "Python bindings for Open Computer Vision Library") + + set(CPACK_COMPONENT_java_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-java") + set(CPACK_COMPONENT_java_DESCRIPTION "Java bindings for Open Computer Vision Library") + + set(CPACK_COMPONENT_dev_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-dev") + set(CPACK_COMPONENT_dev_DESCRIPTION "Development files for Open Computer Vision Library") + + set(CPACK_COMPONENT_docs_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-docs") + set(CPACK_COMPONENT_docs_DESCRIPTION "Documentation for Open Computer Vision Library") + + set(CPACK_COMPONENT_samples_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}-samples") + set(CPACK_COMPONENT_samples_DESCRIPTION "Samples for Open Computer Vision Library") +endif(NOT OPENCV_CUSTOM_PACKAGE_INFO) + +if(NOT OPENCV_CUSTOM_PACKAGE_LAYOUT) + set(CPACK_libs_COMPONENT_INSTALL TRUE) + set(CPACK_dev_COMPONENT_INSTALL TRUE) + set(CPACK_docs_COMPONENT_INSTALL TRUE) + set(CPACK_python_COMPONENT_INSTALL TRUE) + set(CPACK_java_COMPONENT_INSTALL TRUE) + set(CPACK_samples_COMPONENT_INSTALL TRUE) +endif(NOT OPENCV_CUSTOM_PACKAGE_LAYOUT) + +include(CPack) + +ENDif(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake") \ No newline at end of file From dda999545c9dd1cca56081a4b2d56755210b840a Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Wed, 22 Jan 2014 10:40:14 +0400 Subject: [PATCH 09/13] fix GpuMat::copyTo method with mask: fill destination matrix with zeros if it was reallocated --- modules/core/src/gpumat.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/modules/core/src/gpumat.cpp b/modules/core/src/gpumat.cpp index ec26801ddc..96691919fd 100644 --- a/modules/core/src/gpumat.cpp +++ b/modules/core/src/gpumat.cpp @@ -620,11 +620,19 @@ void cv::gpu::GpuMat::copyTo(GpuMat& m) const void cv::gpu::GpuMat::copyTo(GpuMat& mat, const GpuMat& mask) const { if (mask.empty()) + { copyTo(mat); + } else { + uchar* data0 = mat.data; + mat.create(size(), type()); + // do not leave dst uninitialized + if (mat.data != data0) + mat.setTo(Scalar::all(0)); + gpuFuncTable()->copyWithMask(*this, mat, mask); } } From 167a26642e58486a06517bdf6c1b3cdeae7152ac Mon Sep 17 00:00:00 2001 From: Vincent Rabaud Date: Wed, 22 Jan 2014 15:26:14 +0100 Subject: [PATCH 10/13] fix message sent to user during cross_compilation --- cmake/OpenCVDetectPython.cmake | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/cmake/OpenCVDetectPython.cmake b/cmake/OpenCVDetectPython.cmake index 618b0fd222..d02b7596a6 100644 --- a/cmake/OpenCVDetectPython.cmake +++ b/cmake/OpenCVDetectPython.cmake @@ -84,11 +84,9 @@ if(PYTHON_EXECUTABLE) if(CMAKE_CROSSCOMPILING) message(STATUS "Cannot probe for Python/Numpy support (because we are cross-compiling OpenCV)") message(STATUS "If you want to enable Python/Numpy support, set the following variables:") - message(STATUS " PYTHON_EXECUTABLE") - message(STATUS " PYTHON_INCLUDE_DIR") - message(STATUS " PYTHON_LIBRARY") + message(STATUS " PYTHON_INCLUDE_PATH") + message(STATUS " PYTHON_LIBRARIES") message(STATUS " PYTHON_NUMPY_INCLUDE_DIR") - message(STATUS " PYTHON_NUMPY_VERSION") else() # Attempt to discover the NumPy include directory. If this succeeds, then build python API with NumPy execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import os; os.environ['DISTUTILS_USE_SDK']='1'; import numpy.distutils; print numpy.distutils.misc_util.get_numpy_include_dirs()[0]" From dca56841459dadfa01ac74eb1ac01cd7e05d8522 Mon Sep 17 00:00:00 2001 From: Andrey Pavlenko Date: Wed, 22 Jan 2014 21:47:50 +0400 Subject: [PATCH 11/13] removing duplicated legacy license, the actual instance is in 'opencv/LICENSE' --- doc/license.txt | 37 ------------------- .../android_binary_package/O4A_SDK.rst | 2 +- 2 files changed, 1 insertion(+), 38 deletions(-) delete mode 100644 doc/license.txt diff --git a/doc/license.txt b/doc/license.txt deleted file mode 100644 index 8824228d03..0000000000 --- a/doc/license.txt +++ /dev/null @@ -1,37 +0,0 @@ -IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. - - By downloading, copying, installing or using the software you agree to this license. - If you do not agree to this license, do not download, install, - copy or use the software. - - - License Agreement - For Open Source Computer Vision Library - -Copyright (C) 2000-2008, Intel Corporation, all rights reserved. -Copyright (C) 2008-2011, Willow Garage Inc., all rights reserved. -Third party copyrights are property of their respective owners. - -Redistribution and use in source and binary forms, with or without modification, -are permitted provided that the following conditions are met: - - * Redistributions of source code must retain the above copyright notice, - this list of conditions and the following disclaimer. - - * Redistributions in binary form must reproduce the above copyright notice, - this list of conditions and the following disclaimer in the documentation - and/or other materials provided with the distribution. - - * The name of the copyright holders may not be used to endorse or promote products - derived from this software without specific prior written permission. - -This software is provided by the copyright holders and contributors "as is" and -any express or implied warranties, including, but not limited to, the implied -warranties of merchantability and fitness for a particular purpose are disclaimed. -In no event shall the Intel Corporation or contributors be liable for any direct, -indirect, incidental, special, exemplary, or consequential damages -(including, but not limited to, procurement of substitute goods or services; -loss of use, data, or profits; or business interruption) however caused -and on any theory of liability, whether in contract, strict liability, -or tort (including negligence or otherwise) arising in any way out of -the use of this software, even if advised of the possibility of such damage. diff --git a/doc/tutorials/introduction/android_binary_package/O4A_SDK.rst b/doc/tutorials/introduction/android_binary_package/O4A_SDK.rst index 9a683ea496..ef9337aae2 100644 --- a/doc/tutorials/introduction/android_binary_package/O4A_SDK.rst +++ b/doc/tutorials/introduction/android_binary_package/O4A_SDK.rst @@ -66,7 +66,7 @@ The structure of package contents looks as follows: | |_ armeabi-v7a | |_ x86 | - |_ license.txt + |_ LICENSE |_ README.android * :file:`sdk` folder contains OpenCV API and libraries for Android: From 0a4a1d7526ba893e818365875cd5aa916fc9f92b Mon Sep 17 00:00:00 2001 From: Andrey Pavlenko Date: Fri, 24 Jan 2014 10:07:21 +0400 Subject: [PATCH 12/13] temporary disabling hanging test --- modules/ocl/perf/perf_haar.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/ocl/perf/perf_haar.cpp b/modules/ocl/perf/perf_haar.cpp index 86890a8911..3aedd88868 100644 --- a/modules/ocl/perf/perf_haar.cpp +++ b/modules/ocl/perf/perf_haar.cpp @@ -97,8 +97,8 @@ PERF_TEST_P( OCL_Cascade_Image_MinSize, CascadeClassifier, testing::Combine( testing::Values( string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt.xml") ), testing::Values( string("cv/shared/lena.png"), - string("cv/cascadeandhog/images/bttf301.png"), - string("cv/cascadeandhog/images/class57.png") ), + string("cv/cascadeandhog/images/bttf301.png")/*, + string("cv/cascadeandhog/images/class57.png")*/ ), testing::Values(30, 64, 90) ) ) { const string cascasePath = get<0>(GetParam()); From 086792ec06ff78a45fdd5e1b2fa7f72fe3b7c9a2 Mon Sep 17 00:00:00 2001 From: Alexander Smorkalov Date: Thu, 23 Jan 2014 20:57:30 +0400 Subject: [PATCH 13/13] Improvements in package build. --- cmake/OpenCVPackaging.cmake | 35 ++++++++++++++++++++++++++--------- 1 file changed, 26 insertions(+), 9 deletions(-) diff --git a/cmake/OpenCVPackaging.cmake b/cmake/OpenCVPackaging.cmake index 4a81c255bf..3b8d4db547 100644 --- a/cmake/OpenCVPackaging.cmake +++ b/cmake/OpenCVPackaging.cmake @@ -2,24 +2,29 @@ if(EXISTS "${CMAKE_ROOT}/Modules/CPack.cmake") set(CPACK_set_DESTDIR "on") if(NOT OPENCV_CUSTOM_PACKAGE_INFO) - set(CPACK_PACKAGE_DESCRIPTION "Open Computer Vision Library") - set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "OpenCV") + set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "Open Computer Vision Library") + set(CPACK_PACKAGE_DESCRIPTION +"OpenCV (Open Source Computer Vision Library) is an open source computer vision +and machine learning software library. OpenCV was built to provide a common +infrastructure for computer vision applications and to accelerate the use of +machine perception in the commercial products. Being a BSD-licensed product, +OpenCV makes it easy for businesses to utilize and modify the code.") set(CPACK_PACKAGE_VENDOR "OpenCV Foundation") set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE") set(CPACK_PACKAGE_CONTACT "admin@opencv.org") + set(CPACK_PACKAGE_VERSION_MAJOR "${OPENCV_VERSION_MAJOR}") + set(CPACK_PACKAGE_VERSION_MINOR "${OPENCV_VERSION_MINOR}") + set(CPACK_PACKAGE_VERSION_PATCH "${OPENCV_VERSION_PATCH}") + set(CPACK_PACKAGE_VERSION "${OPENCV_VCSVERSION}") endif(NOT OPENCV_CUSTOM_PACKAGE_INFO) -set(CPACK_PACKAGE_VERSION_MAJOR "${OPENCV_VERSION_MAJOR}") -set(CPACK_PACKAGE_VERSION_MINOR "${OPENCV_VERSION_MINOR}") -set(CPACK_PACKAGE_VERSION_PATCH "${OPENCV_VERSION_PATCH}") - #arch if(X86) set(CPACK_DEBIAN_ARCHITECTURE "i386") set(CPACK_RPM_PACKAGE_ARCHITECTURE "i686") elseif(X86_64) set(CPACK_DEBIAN_ARCHITECTURE "amd64") - set(CPACK_RPM_PACKAGE_ARCHITECTURE "amd64") + set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64") elseif(ARM) set(CPACK_DEBIAN_ARCHITECTURE "armhf") set(CPACK_RPM_PACKAGE_ARCHITECTURE "armhf") @@ -41,11 +46,16 @@ set(CPACK_SOURCE_PACKAGE_FILE_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VCSVERSION}-$ #rpm options set(CPACK_RPM_COMPONENT_INSTALL TRUE) -set(CPACK_RPM_PACKAGE_LICENSE ${CPACK_RESOURCE_FILE_LICENSE}) +set(CPACK_RPM_PACKAGE_SUMMARY ${CPACK_PACKAGE_DESCRIPTION_SUMMARY}) +set(CPACK_RPM_PACKAGE_DESCRIPTION ${CPACK_PACKAGE_DESCRIPTION}) +set(CPACK_RPM_PACKAGE_URL "http://opencv.org") +set(CPACK_RPM_PACKAGE_LICENSE "BSD") #deb options set(CPACK_DEB_COMPONENT_INSTALL TRUE) -set(CPACK_DEBIAN_PACKAGE_PRIORITY "extra") +set(CPACK_DEBIAN_PACKAGE_PRIORITY "optional") +set(CPACK_DEBIAN_PACKAGE_SECTION "libs") +set(CPACK_DEBIAN_PACKAGE_HOMEPAGE "http://opencv.org") #depencencies set(CPACK_DEBIAN_PACKAGE_SHLIBDEPS TRUE) @@ -55,6 +65,13 @@ set(CPACK_COMPONENT_docs_DEPENDS libs) set(CPACK_COMPONENT_java_DEPENDS libs) set(CPACK_COMPONENT_python_DEPENDS libs) +if(HAVE_CUDA) + string(REPLACE "." "-" cuda_version_suffix ${CUDA_VERSION}) + set(CPACK_DEB_libs_PACKAGE_DEPENDS "cuda-core-libs-${cuda_version_suffix}, cuda-extra-libs-${cuda_version_suffix}") + set(CPACK_COMPONENT_dev_DEPENDS libs) + set(CPACK_DEB_dev_PACKAGE_DEPENDS "cuda-headers-${cuda_version_suffix}") +endif() + if(NOT OPENCV_CUSTOM_PACKAGE_INFO) set(CPACK_COMPONENT_libs_DISPLAY_NAME "lib${CMAKE_PROJECT_NAME}") set(CPACK_COMPONENT_libs_DESCRIPTION "Open Computer Vision Library")