From d3ac1bc314136d0654fad24a9ee5b01c29e9f5af Mon Sep 17 00:00:00 2001
From: Pierre-Emmanuel Viel
Date: Fri, 20 Dec 2013 01:00:55 +0100
Subject: [PATCH 01/47] When a cluster is empty for KMeans, it's better to give
it the point from another cluster j that is the furthest one from center j.
---
modules/flann/include/opencv2/flann/kmeans_index.h | 11 +++++++----
1 file changed, 7 insertions(+), 4 deletions(-)
diff --git a/modules/flann/include/opencv2/flann/kmeans_index.h b/modules/flann/include/opencv2/flann/kmeans_index.h
index 3fea956a74..489ed80565 100644
--- a/modules/flann/include/opencv2/flann/kmeans_index.h
+++ b/modules/flann/include/opencv2/flann/kmeans_index.h
@@ -759,10 +759,13 @@ private:
for (int k=0; k
Date: Tue, 7 Jan 2014 19:38:57 -0800
Subject: [PATCH 02/47] added epsilon value to weights in the MergeMertens in
order to avoid zero weights for pixels from uniformly filled areas of image
---
modules/photo/src/merge.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/modules/photo/src/merge.cpp b/modules/photo/src/merge.cpp
index 7adfb5ec68..295e03c95f 100644
--- a/modules/photo/src/merge.cpp
+++ b/modules/photo/src/merge.cpp
@@ -208,7 +208,7 @@ public:
if(channels == 3) {
weights[i] = weights[i].mul(saturation);
}
- weights[i] = weights[i].mul(wellexp);
+ weights[i] = weights[i].mul(wellexp) + 1e-12f;
weight_sum += weights[i];
}
int maxlevel = static_cast(logf(static_cast(min(size.width, size.height))) / logf(2.0f));
From 89dd828e3cd4f4b46922185133e6296801b218fb Mon Sep 17 00:00:00 2001
From: Volodymyr Kysenko
Date: Thu, 9 Jan 2014 14:10:00 -0800
Subject: [PATCH 03/47] added test for correct handling of uniforma areas in
the MergeMertens
---
modules/photo/test/test_hdr.cpp | 10 ++++++++++
1 file changed, 10 insertions(+)
diff --git a/modules/photo/test/test_hdr.cpp b/modules/photo/test/test_hdr.cpp
index 82ae25f525..27773fb384 100644
--- a/modules/photo/test/test_hdr.cpp
+++ b/modules/photo/test/test_hdr.cpp
@@ -166,6 +166,16 @@ TEST(Photo_MergeMertens, regression)
merge->process(images, result);
result.convertTo(result, CV_8UC3, 255);
checkEqual(expected, result, 3, "Mertens");
+
+ Mat uniform(100, 100, CV_8UC3);
+ uniform = Scalar(0, 255, 0);
+
+ images.clear();
+ images.push_back(uniform);
+
+ merge->process(images, result);
+ result.convertTo(result, CV_8UC3, 255);
+ checkEqual(uniform, result, 1e-2f, "Mertens");
}
TEST(Photo_MergeDebevec, regression)
From 8f6ebc2427da0a2db870b32fc351642d65deac5d Mon Sep 17 00:00:00 2001
From: Martin Dlouhy
Date: Mon, 24 Feb 2014 07:54:08 +0100
Subject: [PATCH 04/47] fixed rotated rectangle (center instead of corner)
---
.../py_contours/py_contour_features/py_contour_features.rst | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst b/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst
index 6b7c661cc5..8220fb501e 100644
--- a/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst
+++ b/doc/py_tutorials/py_imgproc/py_contours/py_contour_features/py_contour_features.rst
@@ -123,7 +123,7 @@ Let (x,y) be the top-left coordinate of the rectangle and (w,h) be its width and
7.b. Rotated Rectangle
-----------------------
-Here, bounding rectangle is drawn with minimum area, so it considers the rotation also. The function used is **cv2.minAreaRect()**. It returns a Box2D structure which contains following detals - ( top-left corner(x,y), (width, height), angle of rotation ). But to draw this rectangle, we need 4 corners of the rectangle. It is obtained by the function **cv2.boxPoints()**
+Here, bounding rectangle is drawn with minimum area, so it considers the rotation also. The function used is **cv2.minAreaRect()**. It returns a Box2D structure which contains following detals - ( center (x,y), (width, height), angle of rotation ). But to draw this rectangle, we need 4 corners of the rectangle. It is obtained by the function **cv2.boxPoints()**
::
rect = cv2.minAreaRect(cnt)
From b3e18d23a31313b4885f0c2fb8ca8afc0ff8199c Mon Sep 17 00:00:00 2001
From: Alexander Smorkalov
Date: Wed, 5 Mar 2014 11:25:27 +0400
Subject: [PATCH 05/47] Implicit CUDA and OpenCL control for module definition
added.
Feature allows to exclude CUDA or OpenCL optimizations at all even CUDA is used
on build. Exclusion of CUDA or OpenCL cut unwanted dependencies.
---
cmake/OpenCVModule.cmake | 66 +++++++++++++++++--------
modules/nonfree/CMakeLists.txt | 2 +-
modules/superres/src/cuda/btv_l1_gpu.cu | 2 +-
modules/ts/CMakeLists.txt | 2 +-
4 files changed, 49 insertions(+), 23 deletions(-)
diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake
index 03818018d9..c9c351113c 100644
--- a/cmake/OpenCVModule.cmake
+++ b/cmake/OpenCVModule.cmake
@@ -479,39 +479,49 @@ endmacro()
# finds and sets headers and sources for the standard OpenCV module
# Usage:
# ocv_glob_module_sources()
-macro(ocv_glob_module_sources)
+macro(ocv_glob_module_sources EXCLUDE_CUDA EXCLUDE_OPENCL)
file(GLOB_RECURSE lib_srcs "src/*.cpp")
file(GLOB_RECURSE lib_int_hdrs "src/*.hpp" "src/*.h")
file(GLOB lib_hdrs "include/opencv2/${name}/*.hpp" "include/opencv2/${name}/*.h")
file(GLOB lib_hdrs_detail "include/opencv2/${name}/detail/*.hpp" "include/opencv2/${name}/detail/*.h")
- file(GLOB lib_cuda_srcs "src/cuda/*.cu")
- set(cuda_objs "")
- set(lib_cuda_hdrs "")
- if(HAVE_CUDA)
- ocv_include_directories(${CUDA_INCLUDE_DIRS})
- file(GLOB lib_cuda_hdrs "src/cuda/*.hpp")
+ if (NOT ${EXCLUDE_CUDA})
+ file(GLOB lib_cuda_srcs "src/cuda/*.cu")
+ set(cuda_objs "")
+ set(lib_cuda_hdrs "")
+ if(HAVE_CUDA)
+ ocv_include_directories(${CUDA_INCLUDE_DIRS})
+ file(GLOB lib_cuda_hdrs "src/cuda/*.hpp")
- ocv_cuda_compile(cuda_objs ${lib_cuda_srcs} ${lib_cuda_hdrs})
- source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs})
+ ocv_cuda_compile(cuda_objs ${lib_cuda_srcs} ${lib_cuda_hdrs})
+ source_group("Src\\Cuda" FILES ${lib_cuda_srcs} ${lib_cuda_hdrs})
+ endif()
+ else()
+ set(cuda_objs "")
+ set(lib_cuda_srcs "")
+ set(lib_cuda_hdrs "")
endif()
source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs})
- file(GLOB cl_kernels "src/opencl/*.cl")
- if(HAVE_opencv_ocl AND cl_kernels)
- ocv_include_directories(${OPENCL_INCLUDE_DIRS})
- add_custom_command(
- OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp"
- COMMAND ${CMAKE_COMMAND} -DCL_DIR="${CMAKE_CURRENT_SOURCE_DIR}/src/opencl" -DOUTPUT="${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" -P "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake"
- DEPENDS ${cl_kernels} "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake")
- source_group("OpenCL" FILES ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
- list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
+ if (NOT ${EXCLUDE_OPENCL})
+ file(GLOB cl_kernels "src/opencl/*.cl")
+ if(HAVE_opencv_ocl AND cl_kernels)
+ ocv_include_directories(${OPENCL_INCLUDE_DIRS})
+ add_custom_command(
+ OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp"
+ COMMAND ${CMAKE_COMMAND} -DCL_DIR="${CMAKE_CURRENT_SOURCE_DIR}/src/opencl" -DOUTPUT="${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" -P "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake"
+ DEPENDS ${cl_kernels} "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake")
+ source_group("OpenCL" FILES ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
+ list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
+ endif()
endif()
source_group("Include" FILES ${lib_hdrs})
source_group("Include\\detail" FILES ${lib_hdrs_detail})
+ message(":${EXCLUDE_CUDA}: ${lib_cuda_srcs}")
+
ocv_set_module_sources(${ARGN} HEADERS ${lib_hdrs} ${lib_hdrs_detail}
SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs})
endmacro()
@@ -614,9 +624,25 @@ endmacro()
# Usage:
# ocv_define_module(module_name [INTERNAL] [REQUIRED] [] [OPTIONAL ])
macro(ocv_define_module module_name)
- ocv_add_module(${module_name} ${ARGN})
+ set(_tmp_argn ${ARGN})
+ set(exclude_cuda 0)
+ set(exclude_opencl 0)
+ set(argv0 ${ARGV1})
+ set(argv1 ${ARGV2})
+ set(argv2 ${ARGV3})
+ foreach(i RANGE 0 2)
+ if("${argv${i}}" STREQUAL "EXCLUDE_CUDA")
+ set(exclude_cuda 1)
+ list(REMOVE_AT _tmp_argn ${i})
+ elseif ("${argv${i}}" STREQUAL "EXCLUDE_OPENCL")
+ set(exclude_opencl 1)
+ list(REMOVE_AT _tmp_argn ${i})
+ endif()
+ endforeach()
+
+ ocv_add_module(${module_name} ${_tmp_argn})
ocv_module_include_directories()
- ocv_glob_module_sources()
+ ocv_glob_module_sources(${exclude_cuda} ${exclude_opencl})
ocv_create_module()
ocv_add_precompiled_headers(${the_module})
diff --git a/modules/nonfree/CMakeLists.txt b/modules/nonfree/CMakeLists.txt
index b43273bc80..571614be0d 100644
--- a/modules/nonfree/CMakeLists.txt
+++ b/modules/nonfree/CMakeLists.txt
@@ -6,7 +6,7 @@ set(the_description "Functionality with possible limitations on the use")
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wshadow)
if(ENABLE_DYNAMIC_CUDA)
add_definitions(-DDYNAMIC_CUDA_SUPPORT)
- ocv_define_module(nonfree opencv_imgproc opencv_features2d opencv_calib3d OPTIONAL opencv_ocl)
+ ocv_define_module(nonfree EXCLUDE_CUDA opencv_imgproc opencv_features2d opencv_calib3d OPTIONAL opencv_ocl)
else()
ocv_define_module(nonfree opencv_imgproc opencv_features2d opencv_calib3d OPTIONAL opencv_gpu opencv_ocl)
endif()
diff --git a/modules/superres/src/cuda/btv_l1_gpu.cu b/modules/superres/src/cuda/btv_l1_gpu.cu
index b4d96190ae..4b0ebdc592 100644
--- a/modules/superres/src/cuda/btv_l1_gpu.cu
+++ b/modules/superres/src/cuda/btv_l1_gpu.cu
@@ -42,7 +42,7 @@
#include "opencv2/opencv_modules.hpp"
-#ifdef HAVE_OPENCV_GPU
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
#include "opencv2/gpu/device/common.hpp"
#include "opencv2/gpu/device/transform.hpp"
diff --git a/modules/ts/CMakeLists.txt b/modules/ts/CMakeLists.txt
index bb56da2d98..dcd3e1563b 100644
--- a/modules/ts/CMakeLists.txt
+++ b/modules/ts/CMakeLists.txt
@@ -11,7 +11,7 @@ ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
ocv_add_module(ts opencv_core opencv_features2d)
-ocv_glob_module_sources()
+ocv_glob_module_sources(0 0)
ocv_module_include_directories()
ocv_create_module()
From 0eaeff06418223c89f86fc2fdcf48fbc90f4c4cb Mon Sep 17 00:00:00 2001
From: kurodash
Date: Fri, 7 Mar 2014 19:02:37 +0900
Subject: [PATCH 06/47] fix: use "cvAlloc" wrapper function for malloc.
---
modules/core/src/persistence.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/modules/core/src/persistence.cpp b/modules/core/src/persistence.cpp
index 7759f708b6..4a6e0c9ec5 100644
--- a/modules/core/src/persistence.cpp
+++ b/modules/core/src/persistence.cpp
@@ -4855,7 +4855,7 @@ cvRegisterType( const CvTypeInfo* _info )
"Type name should contain only letters, digits, - and _" );
}
- info = (CvTypeInfo*)malloc( sizeof(*info) + len + 1 );
+ info = (CvTypeInfo*)cvAlloc( sizeof(*info) + len + 1 );
*info = *_info;
info->type_name = (char*)(info + 1);
@@ -4893,7 +4893,7 @@ cvUnregisterType( const char* type_name )
if( !CvType::first || !CvType::last )
CvType::first = CvType::last = 0;
- free( info );
+ cvFree( info );
}
}
From a87607e3ef6d30a2ccd8bd3b516fde4647dce16d Mon Sep 17 00:00:00 2001
From: Firat Kalaycilar
Date: Wed, 12 Mar 2014 16:14:59 +0200
Subject: [PATCH 07/47] Fixed an issue with weight assignment causing the
resulting GMM weights to be unsorted in BackgroundSubtractorMOG2
---
modules/video/src/bgfg_gaussmix2.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/modules/video/src/bgfg_gaussmix2.cpp b/modules/video/src/bgfg_gaussmix2.cpp
index 6bbb960482..b14bc8e1e2 100644
--- a/modules/video/src/bgfg_gaussmix2.cpp
+++ b/modules/video/src/bgfg_gaussmix2.cpp
@@ -319,7 +319,7 @@ struct MOG2Invoker : ParallelLoopBody
for( int mode = 0; mode < nmodes; mode++, mean_m += nchannels )
{
float weight = alpha1*gmm[mode].weight + prune;//need only weight if fit is found
-
+ int swap_count = 0;
////
//fit not found yet
if( !fitsPDF )
@@ -384,6 +384,7 @@ struct MOG2Invoker : ParallelLoopBody
if( weight < gmm[i-1].weight )
break;
+ swap_count++;
//swap one up
std::swap(gmm[i], gmm[i-1]);
for( int c = 0; c < nchannels; c++ )
@@ -401,7 +402,7 @@ struct MOG2Invoker : ParallelLoopBody
nmodes--;
}
- gmm[mode].weight = weight;//update weight by the calculated value
+ gmm[mode-swap_count].weight = weight;//update weight by the calculated value
totalWeight += weight;
}
//go through all modes
From f9484bae8ab748132e753b8c217c5ded36a5c9dd Mon Sep 17 00:00:00 2001
From: kuroda sho
Date: Fri, 14 Mar 2014 17:02:20 +0900
Subject: [PATCH 08/47] fix: use "cvAlloc" wrapper function for malloc.
---
modules/core/src/persistence.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/modules/core/src/persistence.cpp b/modules/core/src/persistence.cpp
index 4a6e0c9ec5..6847eec34d 100644
--- a/modules/core/src/persistence.cpp
+++ b/modules/core/src/persistence.cpp
@@ -4893,7 +4893,7 @@ cvUnregisterType( const char* type_name )
if( !CvType::first || !CvType::last )
CvType::first = CvType::last = 0;
- cvFree( info );
+ cvFree( &info );
}
}
From 5f94a205d14180d5fd9c0b451add6da35a2b3f9f Mon Sep 17 00:00:00 2001
From: berak
Date: Thu, 13 Mar 2014 08:39:15 +0100
Subject: [PATCH 09/47] fixed h / s ranges in histogram_calculation tutorial
literalinclude
literalinclude, dropped :lines:
---
.../histogram_comparison.rst | 90 ++-----------------
.../Histograms_Matching/compareHist_Demo.cpp | 6 +-
2 files changed, 9 insertions(+), 87 deletions(-)
diff --git a/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst b/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst
index 1a5c59de07..f5f636d08b 100644
--- a/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst
+++ b/doc/tutorials/imgproc/histograms/histogram_comparison/histogram_comparison.rst
@@ -84,88 +84,10 @@ Code
* **Code at glance:**
-.. code-block:: cpp
+.. literalinclude:: ../../../../../samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp
+ :language: cpp
+ :tab-width: 4
- #include "opencv2/highgui/highgui.hpp"
- #include "opencv2/imgproc/imgproc.hpp"
- #include
- #include
-
- using namespace std;
- using namespace cv;
-
- /** @function main */
- int main( int argc, char** argv )
- {
- Mat src_base, hsv_base;
- Mat src_test1, hsv_test1;
- Mat src_test2, hsv_test2;
- Mat hsv_half_down;
-
- /// Load three images with different environment settings
- if( argc < 4 )
- { printf("** Error. Usage: ./compareHist_Demo \n");
- return -1;
- }
-
- src_base = imread( argv[1], 1 );
- src_test1 = imread( argv[2], 1 );
- src_test2 = imread( argv[3], 1 );
-
- /// Convert to HSV
- cvtColor( src_base, hsv_base, CV_BGR2HSV );
- cvtColor( src_test1, hsv_test1, CV_BGR2HSV );
- cvtColor( src_test2, hsv_test2, CV_BGR2HSV );
-
- hsv_half_down = hsv_base( Range( hsv_base.rows/2, hsv_base.rows - 1 ), Range( 0, hsv_base.cols - 1 ) );
-
- /// Using 30 bins for hue and 32 for saturation
- int h_bins = 50; int s_bins = 60;
- int histSize[] = { h_bins, s_bins };
-
- // hue varies from 0 to 256, saturation from 0 to 180
- float h_ranges[] = { 0, 256 };
- float s_ranges[] = { 0, 180 };
-
- const float* ranges[] = { h_ranges, s_ranges };
-
- // Use the o-th and 1-st channels
- int channels[] = { 0, 1 };
-
- /// Histograms
- MatND hist_base;
- MatND hist_half_down;
- MatND hist_test1;
- MatND hist_test2;
-
- /// Calculate the histograms for the HSV images
- calcHist( &hsv_base, 1, channels, Mat(), hist_base, 2, histSize, ranges, true, false );
- normalize( hist_base, hist_base, 0, 1, NORM_MINMAX, -1, Mat() );
-
- calcHist( &hsv_half_down, 1, channels, Mat(), hist_half_down, 2, histSize, ranges, true, false );
- normalize( hist_half_down, hist_half_down, 0, 1, NORM_MINMAX, -1, Mat() );
-
- calcHist( &hsv_test1, 1, channels, Mat(), hist_test1, 2, histSize, ranges, true, false );
- normalize( hist_test1, hist_test1, 0, 1, NORM_MINMAX, -1, Mat() );
-
- calcHist( &hsv_test2, 1, channels, Mat(), hist_test2, 2, histSize, ranges, true, false );
- normalize( hist_test2, hist_test2, 0, 1, NORM_MINMAX, -1, Mat() );
-
- /// Apply the histogram comparison methods
- for( int i = 0; i < 4; i++ )
- { int compare_method = i;
- double base_base = compareHist( hist_base, hist_base, compare_method );
- double base_half = compareHist( hist_base, hist_half_down, compare_method );
- double base_test1 = compareHist( hist_base, hist_test1, compare_method );
- double base_test2 = compareHist( hist_base, hist_test2, compare_method );
-
- printf( " Method [%d] Perfect, Base-Half, Base-Test(1), Base-Test(2) : %f, %f, %f, %f \n", i, base_base, base_half , base_test1, base_test2 );
- }
-
- printf( "Done \n" );
-
- return 0;
- }
Explanation
@@ -211,11 +133,11 @@ Explanation
.. code-block:: cpp
- int h_bins = 50; int s_bins = 32;
+ int h_bins = 50; int s_bins = 60;
int histSize[] = { h_bins, s_bins };
- float h_ranges[] = { 0, 256 };
- float s_ranges[] = { 0, 180 };
+ float h_ranges[] = { 0, 180 };
+ float s_ranges[] = { 0, 256 };
const float* ranges[] = { h_ranges, s_ranges };
diff --git a/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp b/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp
index f4dd4e5e4e..424a38e93a 100644
--- a/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp
+++ b/samples/cpp/tutorial_code/Histograms_Matching/compareHist_Demo.cpp
@@ -40,13 +40,13 @@ int main( int argc, char** argv )
hsv_half_down = hsv_base( Range( hsv_base.rows/2, hsv_base.rows - 1 ), Range( 0, hsv_base.cols - 1 ) );
- /// Using 30 bins for hue and 32 for saturation
+ /// Using 50 bins for hue and 60 for saturation
int h_bins = 50; int s_bins = 60;
int histSize[] = { h_bins, s_bins };
- // hue varies from 0 to 256, saturation from 0 to 180
- float s_ranges[] = { 0, 256 };
+ // hue varies from 0 to 179, saturation from 0 to 255
float h_ranges[] = { 0, 180 };
+ float s_ranges[] = { 0, 256 };
const float* ranges[] = { h_ranges, s_ranges };
From 6b8de222d770174fc1fb31b1b64da8995e8a8ebf Mon Sep 17 00:00:00 2001
From: Alexander Smorkalov
Date: Wed, 5 Mar 2014 09:57:16 +0400
Subject: [PATCH 10/47] OpenCV_MODULES_SUFFIX variable added to
OpenCVConfig.cmake to enable custom module configurations.
---
cmake/templates/OpenCVConfig.cmake.in | 28 +++++++++++++++------------
1 file changed, 16 insertions(+), 12 deletions(-)
diff --git a/cmake/templates/OpenCVConfig.cmake.in b/cmake/templates/OpenCVConfig.cmake.in
index 6db61d2112..3222048282 100644
--- a/cmake/templates/OpenCVConfig.cmake.in
+++ b/cmake/templates/OpenCVConfig.cmake.in
@@ -18,8 +18,8 @@
# This file will define the following variables:
# - OpenCV_LIBS : The list of all imported targets for OpenCV modules.
# - OpenCV_INCLUDE_DIRS : The OpenCV include directories.
-# - OpenCV_COMPUTE_CAPABILITIES : The version of compute capability
-# - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API
+# - OpenCV_COMPUTE_CAPABILITIES : The version of compute capability.
+# - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API.
# - OpenCV_VERSION : The version of this OpenCV build: "@OPENCV_VERSION@"
# - OpenCV_VERSION_MAJOR : Major version part of OpenCV_VERSION: "@OPENCV_VERSION_MAJOR@"
# - OpenCV_VERSION_MINOR : Minor version part of OpenCV_VERSION: "@OPENCV_VERSION_MINOR@"
@@ -27,22 +27,26 @@
# - OpenCV_VERSION_TWEAK : Tweak version part of OpenCV_VERSION: "@OPENCV_VERSION_TWEAK@"
#
# Advanced variables:
-# - OpenCV_SHARED
-# - OpenCV_CONFIG_PATH
-# - OpenCV_INSTALL_PATH (not set on Windows)
-# - OpenCV_LIB_COMPONENTS
-# - OpenCV_USE_MANGLED_PATHS
-# - OpenCV_HAVE_ANDROID_CAMERA
+# - OpenCV_SHARED : Use OpenCV as shared library
+# - OpenCV_CONFIG_PATH : Path to this OpenCVConfig.cmake
+# - OpenCV_INSTALL_PATH : OpenCV location (not set on Windows)
+# - OpenCV_LIB_COMPONENTS : Present OpenCV modules list
+# - OpenCV_USE_MANGLED_PATHS : Mangled OpenCV path flag
+# - OpenCV_MODULES_SUFFIX : The suffix for OpenCVModules-XXX.cmake file
+# - OpenCV_HAVE_ANDROID_CAMERA : Presence of Android native camera wrappers
#
# ===================================================================================
-set(modules_file_suffix "")
-if(ANDROID)
- string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}")
+if(NOT DEFINED OpenCV_MODULES_SUFFIX)
+ if(ANDROID)
+ string(REPLACE - _ OpenCV_MODULES_SUFFIX "_${ANDROID_NDK_ABI_NAME}")
+ else()
+ set(OpenCV_MODULES_SUFFIX "")
+ endif()
endif()
if(NOT TARGET opencv_core)
- include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules${modules_file_suffix}.cmake)
+ include(${CMAKE_CURRENT_LIST_DIR}/OpenCVModules${OpenCV_MODULES_SUFFIX}.cmake)
endif()
# TODO All things below should be reviewed. What is about of moving this code into related modules (special vars/hooks/files)
From 6890aa0033d0cc092fd8c829ce2cedf4c1f0c079 Mon Sep 17 00:00:00 2001
From: vbystricky
Date: Mon, 17 Mar 2014 16:03:15 +0400
Subject: [PATCH 11/47] Fix problems on Intel HD graphics
---
modules/video/src/lkpyramid.cpp | 2 --
modules/video/src/opencl/pyrlk.cl | 46 ++-----------------------------
2 files changed, 2 insertions(+), 46 deletions(-)
diff --git a/modules/video/src/lkpyramid.cpp b/modules/video/src/lkpyramid.cpp
index cd57585658..a33e47664d 100644
--- a/modules/video/src/lkpyramid.cpp
+++ b/modules/video/src/lkpyramid.cpp
@@ -975,9 +975,7 @@ namespace cv
idxArg = kernel.set(idxArg, imageI); //image2d_t I
idxArg = kernel.set(idxArg, imageJ); //image2d_t J
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(prevPts)); // __global const float2* prevPts
- idxArg = kernel.set(idxArg, (int)prevPts.step); // int prevPtsStep
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(nextPts)); // __global const float2* nextPts
- idxArg = kernel.set(idxArg, (int)nextPts.step); // int nextPtsStep
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(status)); // __global uchar* status
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(err)); // __global float* err
idxArg = kernel.set(idxArg, (int)level); // const int level
diff --git a/modules/video/src/opencl/pyrlk.cl b/modules/video/src/opencl/pyrlk.cl
index c018554902..822e628f25 100644
--- a/modules/video/src/opencl/pyrlk.cl
+++ b/modules/video/src/opencl/pyrlk.cl
@@ -262,50 +262,9 @@ inline void GetError(image2d_t J, const float x, const float y, const float* Pch
*errval += fabs(diff);
}
-inline void SetPatch4(image2d_t I, const float x, const float y,
- float4* Pch, float4* Dx, float4* Dy,
- float* A11, float* A12, float* A22)
-{
- *Pch = read_imagef(I, sampler, (float2)(x, y));
-
- float4 dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) -
- (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)));
-
- float4 dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) -
- (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)));
-
-
- *Dx = dIdx;
- *Dy = dIdy;
- float4 sqIdx = dIdx * dIdx;
- *A11 += sqIdx.x + sqIdx.y + sqIdx.z;
- sqIdx = dIdx * dIdy;
- *A12 += sqIdx.x + sqIdx.y + sqIdx.z;
- sqIdx = dIdy * dIdy;
- *A22 += sqIdx.x + sqIdx.y + sqIdx.z;
-}
-
-inline void GetPatch4(image2d_t J, const float x, const float y,
- const float4* Pch, const float4* Dx, const float4* Dy,
- float* b1, float* b2)
-{
- float4 J_val = read_imagef(J, sampler, (float2)(x, y));
- float4 diff = (J_val - *Pch) * 32.0f;
- float4 xdiff = diff* *Dx;
- *b1 += xdiff.x + xdiff.y + xdiff.z;
- xdiff = diff* *Dy;
- *b2 += xdiff.x + xdiff.y + xdiff.z;
-}
-
-inline void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval)
-{
- float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch;
- *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z);
-}
-
#define GRIDSIZE 3
__kernel void lkSparse(image2d_t I, image2d_t J,
- __global const float2* prevPts, int prevPtsStep, __global float2* nextPts, int nextPtsStep, __global uchar* status, __global float* err,
+ __global const float2* prevPts, __global float2* nextPts, __global uchar* status, __global float* err,
const int level, const int rows, const int cols, int PATCH_X, int PATCH_Y, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr)
{
__local float smem1[BUFFER];
@@ -434,9 +393,8 @@ __kernel void lkSparse(image2d_t I, image2d_t J,
{
if (tid == 0 && level == 0)
status[gid] = 0;
- return;
+ break;
}
-
float b1 = 0;
float b2 = 0;
From 0c02e5de254d32e8f00bde4047f12624e77a7ef5 Mon Sep 17 00:00:00 2001
From: Anatoly Baksheev
Date: Mon, 17 Mar 2014 17:02:49 +0400
Subject: [PATCH 12/47] minor doc fix
---
modules/viz/doc/widget.rst | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/modules/viz/doc/widget.rst b/modules/viz/doc/widget.rst
index 7c1fc45751..0601ba1b72 100644
--- a/modules/viz/doc/widget.rst
+++ b/modules/viz/doc/widget.rst
@@ -1049,7 +1049,7 @@ viz::WWidgetMerger::WWidgetMerger
---------------------------------------
Constructs a WWidgetMerger.
-.. ocv:WWidgetMerger:: WWidgetMerger()
+.. ocv:function:: WWidgetMerger()
viz::WWidgetMerger::addCloud
-------------------------------
From 3940b6163b8370312d2a7fe66c96d43ab6d12efd Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Mon, 17 Mar 2014 18:52:28 +0400
Subject: [PATCH 13/47] remove intel guard since the code is 2 times faster on
AMD too
---
modules/ocl/perf/perf_filters.cpp | 6 +++---
modules/ocl/src/filtering.cpp | 6 +++---
modules/ocl/src/opencl/filtering_sep_filter_singlepass.cl | 2 ++
3 files changed, 8 insertions(+), 6 deletions(-)
diff --git a/modules/ocl/perf/perf_filters.cpp b/modules/ocl/perf/perf_filters.cpp
index b3ffc51b30..c542d647cf 100644
--- a/modules/ocl/perf/perf_filters.cpp
+++ b/modules/ocl/perf/perf_filters.cpp
@@ -262,13 +262,13 @@ OCL_PERF_TEST_P(SobelFixture, Sobel,
oclDst.download(dst);
- SANITY_CHECK(dst);
+ SANITY_CHECK(dst, 1e-3);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() cv::Sobel(src, dst, -1, dx, dy);
- SANITY_CHECK(dst);
+ SANITY_CHECK(dst, 1e-3);
}
else
OCL_PERF_ELSE
@@ -326,7 +326,7 @@ OCL_PERF_TEST_P(GaussianBlurFixture, GaussianBlur,
Mat src(srcSize, type), dst(srcSize, type);
declare.in(src, WARMUP_RNG).out(dst);
- const double eps = src.depth() == CV_8U ? 1 + DBL_EPSILON : 3e-4;
+ const double eps = src.depth() == CV_8U ? 1 + DBL_EPSILON : 5e-4;
if (RUN_OCL_IMPL)
{
diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp
index 35aa226de6..77052ffbf3 100644
--- a/modules/ocl/src/filtering.cpp
+++ b/modules/ocl/src/filtering.cpp
@@ -774,12 +774,12 @@ static void sepFilter2D_SinglePass(const oclMat &src, oclMat &dst,
option += " -D KERNEL_MATRIX_X=";
for(int i=0; i( &row_kernel.at(i) ) );
+ option += cv::format("DIG(0x%x)", *reinterpret_cast( &row_kernel.at(i) ) );
option += "0x0";
option += " -D KERNEL_MATRIX_Y=";
for(int i=0; i( &col_kernel.at(i) ) );
+ option += cv::format("DIG(0x%x)", *reinterpret_cast( &col_kernel.at(i) ) );
option += "0x0";
switch(src.type())
@@ -1410,7 +1410,7 @@ Ptr cv::ocl::createSeparableLinearFilter_GPU(int srcType, int
//if image size is non-degenerate and large enough
//and if filter support is reasonable to satisfy larger local memory requirements,
//then we can use single pass routine to avoid extra runtime calls overhead
- if( clCxt && clCxt->supportsFeature(FEATURE_CL_INTEL_DEVICE) &&
+ if( clCxt &&
rowKernel.rows <= 21 && columnKernel.rows <= 21 &&
(rowKernel.rows & 1) == 1 && (columnKernel.rows & 1) == 1 &&
imgSize.width > optimizedSepFilterLocalSize + (rowKernel.rows>>1) &&
diff --git a/modules/ocl/src/opencl/filtering_sep_filter_singlepass.cl b/modules/ocl/src/opencl/filtering_sep_filter_singlepass.cl
index c6555bff0f..c5f490284e 100644
--- a/modules/ocl/src/opencl/filtering_sep_filter_singlepass.cl
+++ b/modules/ocl/src/opencl/filtering_sep_filter_singlepass.cl
@@ -84,6 +84,8 @@
#define DST(_x,_y) (((global DSTTYPE*)(Dst+DstOffset+(_y)*DstPitch))[_x])
+#define DIG(a) a,
+
//horizontal and vertical filter kernels
//should be defined on host during compile time to avoid overhead
__constant uint mat_kernelX[] = {KERNEL_MATRIX_X};
From 82e6edfba28c91f80fbc20557a3f16906db95bc5 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Mon, 17 Mar 2014 19:59:35 +0400
Subject: [PATCH 14/47] optimized sep filter
---
modules/core/include/opencv2/core/ocl.hpp | 2 +-
modules/core/src/ocl.cpp | 4 +-
modules/imgproc/perf/opencl/perf_filters.cpp | 2 +-
modules/imgproc/src/filter.cpp | 99 +++++++---
.../src/opencl/filterSep_singlePass.cl | 177 ++++++++++++++++++
5 files changed, 251 insertions(+), 33 deletions(-)
create mode 100644 modules/imgproc/src/opencl/filterSep_singlePass.cl
diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp
index fb9ec24c56..fdb6f9a0aa 100644
--- a/modules/core/include/opencv2/core/ocl.hpp
+++ b/modules/core/include/opencv2/core/ocl.hpp
@@ -592,7 +592,7 @@ protected:
CV_EXPORTS const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf);
CV_EXPORTS const char* typeToStr(int t);
CV_EXPORTS const char* memopTypeToStr(int t);
-CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1);
+CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1, const char * name = NULL);
CV_EXPORTS void getPlatfomsInfo(std::vector& platform_info);
CV_EXPORTS int predictOptimalVectorWidth(InputArray src1, InputArray src2 = noArray(), InputArray src3 = noArray(),
InputArray src4 = noArray(), InputArray src5 = noArray(), InputArray src6 = noArray(),
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index 7c4f8de9e0..b56f84c16e 100644
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -4306,7 +4306,7 @@ static std::string kerToStr(const Mat & k)
return stream.str();
}
-String kernelToStr(InputArray _kernel, int ddepth)
+String kernelToStr(InputArray _kernel, int ddepth, const char * name)
{
Mat kernel = _kernel.getMat().reshape(1, 1);
@@ -4323,7 +4323,7 @@ String kernelToStr(InputArray _kernel, int ddepth)
const func_t func = funcs[depth];
CV_Assert(func != 0);
- return cv::format(" -D COEFF=%s", func(kernel).c_str());
+ return cv::format(" -D %s=%s", name ? name : "COEFF", func(kernel).c_str());
}
#define PROCESS_SRC(src) \
diff --git a/modules/imgproc/perf/opencl/perf_filters.cpp b/modules/imgproc/perf/opencl/perf_filters.cpp
index 57b928c289..f7329e3194 100644
--- a/modules/imgproc/perf/opencl/perf_filters.cpp
+++ b/modules/imgproc/perf/opencl/perf_filters.cpp
@@ -211,7 +211,7 @@ OCL_PERF_TEST_P(SobelFixture, Sobel,
OCL_TEST_CYCLE() cv::Sobel(src, dst, -1, dx, dy);
- SANITY_CHECK(dst);
+ SANITY_CHECK(dst, 1e-6);
}
///////////// Scharr ////////////////////////
diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp
index ea0baf6b09..bb54471c07 100644
--- a/modules/imgproc/src/filter.cpp
+++ b/modules/imgproc/src/filter.cpp
@@ -3350,27 +3350,8 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor,
int radiusY = (int)((buf.rows - src.rows) >> 1);
bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
- const char* btype = NULL;
- switch (borderType & ~BORDER_ISOLATED)
- {
- case BORDER_CONSTANT:
- btype = "BORDER_CONSTANT";
- break;
- case BORDER_REPLICATE:
- btype = "BORDER_REPLICATE";
- break;
- case BORDER_REFLECT:
- btype = "BORDER_REFLECT";
- break;
- case BORDER_WRAP:
- btype = "BORDER_WRAP";
- break;
- case BORDER_REFLECT101:
- btype = "BORDER_REFLECT_101";
- break;
- default:
- return false;
- }
+ const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" },
+ * const btype = borderMap[borderType & ~BORDER_ISOLATED];
bool extra_extrapolation = src.rows < (int)((-radiusY + globalsize[1]) >> 1) + 1;
extra_extrapolation |= src.rows < radiusY;
@@ -3463,36 +3444,96 @@ static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anc
return kernelCol.run(2, globalsize, localsize, sync);
}
+const int optimizedSepFilterLocalSize = 16;
+
+static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
+ InputArray _row_kernel, InputArray _col_kernel,
+ int borderType, int ddepth)
+{
+ Size size = _src.size(), wholeSize;
+ Point origin;
+ int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype),
+ esz = CV_ELEM_SIZE(stype), wdepth = std::max(std::max(sdepth, ddepth), CV_32F),
+ dtype = CV_MAKE_TYPE(ddepth, cn);
+ size_t src_step = _src.step(), src_offset = _src.offset();
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+ if ((src_offset % src_step) % esz != 0 || (!doubleSupport && sdepth == CV_64F) ||
+ !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE ||
+ borderType == BORDER_REFLECT || borderType == BORDER_WRAP ||
+ borderType == BORDER_REFLECT_101))
+ return false;
+
+ size_t lt2[2] = { optimizedSepFilterLocalSize, optimizedSepFilterLocalSize };
+ size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1] * (1 + (size.height - 1) / lt2[1]) };
+
+ char cvt[2][40];
+ const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP",
+ "BORDER_REFLECT_101" };
+
+ String opts = cv::format("-D BLK_X=%d -D BLK_Y=%d -D RADIUSX=%d -D RADIUSY=%d%s%s"
+ " -D srcT=%s -D convertToWT=%s -D WT=%s -D dstT=%s -D convertToDstT=%s"
+ " -D %s", (int)lt2[0], (int)lt2[1], _row_kernel.size().height / 2, _col_kernel.size().height / 2,
+ ocl::kernelToStr(_row_kernel, CV_32F, "KERNEL_MATRIX_X").c_str(),
+ ocl::kernelToStr(_col_kernel, CV_32F, "KERNEL_MATRIX_Y").c_str(),
+ ocl::typeToStr(stype), ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
+ ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), ocl::typeToStr(dtype),
+ ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType]);
+
+ ocl::Kernel k("sep_filter", ocl::imgproc::filterSep_singlePass_oclsrc, opts);
+ if (k.empty())
+ return false;
+
+ UMat src = _src.getUMat();
+ _dst.create(size, dtype);
+ UMat dst = _dst.getUMat();
+
+ int src_offset_x = static_cast((src_offset % src_step) / esz);
+ int src_offset_y = static_cast(src_offset / src_step);
+
+ src.locateROI(wholeSize, origin);
+
+ k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y,
+ wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst));
+
+ return k.run(2, gt2, lt2, false);
+}
+
static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
InputArray _kernelX, InputArray _kernelY, Point anchor,
double delta, int borderType )
{
+ Size imgSize = _src.size();
+
if (abs(delta)> FLT_MIN)
return false;
- int type = _src.type();
+ int type = _src.type(), cn = CV_MAT_CN(type);
if ( !( (type == CV_8UC1 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC4) &&
(ddepth == CV_32F || ddepth == CV_16S || ddepth == CV_8U || ddepth < 0) ) )
return false;
- int cn = CV_MAT_CN(type);
-
Mat kernelX = _kernelX.getMat().reshape(1, 1);
- if (1 != (kernelX.cols % 2))
+ if (kernelX.cols % 2 != 1)
return false;
Mat kernelY = _kernelY.getMat().reshape(1, 1);
- if (1 != (kernelY.cols % 2))
+ if (kernelY.cols % 2 != 1)
return false;
int sdepth = CV_MAT_DEPTH(type);
- if( anchor.x < 0 )
+ if (anchor.x < 0)
anchor.x = kernelX.cols >> 1;
- if( anchor.y < 0 )
+ if (anchor.y < 0)
anchor.y = kernelY.cols >> 1;
- if( ddepth < 0 )
+ if (ddepth < 0)
ddepth = sdepth;
+ CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 &&
+ imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) &&
+ imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1),
+ ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true)
+
UMat src = _src.getUMat();
Size srcWholeSize; Point srcOffset;
src.locateROI(srcWholeSize, srcOffset);
diff --git a/modules/imgproc/src/opencl/filterSep_singlePass.cl b/modules/imgproc/src/opencl/filterSep_singlePass.cl
new file mode 100644
index 0000000000..7284da0cbc
--- /dev/null
+++ b/modules/imgproc/src/opencl/filterSep_singlePass.cl
@@ -0,0 +1,177 @@
+/*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) 2014, Intel Corporation, 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*/
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+/////////////////////////////////Macro for border type////////////////////////////////////////////
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+#ifdef BORDER_CONSTANT
+// CCCCCC|abcdefgh|CCCCCCC
+#define EXTRAPOLATE(x, maxV)
+#elif defined BORDER_REPLICATE
+// aaaaaa|abcdefgh|hhhhhhh
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ (x) = max(min((x), (maxV) - 1), 0); \
+ }
+#elif defined BORDER_WRAP
+// cdefgh|abcdefgh|abcdefg
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ (x) = ( (x) + (maxV) ) % (maxV); \
+ }
+#elif defined BORDER_REFLECT
+// fedcba|abcdefgh|hgfedcb
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \
+ }
+#elif defined BORDER_REFLECT_101 || defined BORDER_REFLECT101
+// gfedcb|abcdefgh|gfedcba
+#define EXTRAPOLATE(x, maxV) \
+ { \
+ (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \
+ }
+#else
+#error No extrapolation method
+#endif
+
+#define SRC(_x,_y) convertToWT(((global srcT*)(Src+(_y)*src_step))[_x])
+
+#ifdef BORDER_CONSTANT
+// CCCCCC|abcdefgh|CCCCCCC
+#define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y))
+#else
+#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y))
+#endif
+
+#define DST(_x,_y) (((global dstT*)(Dst+dst_offset+(_y)*dst_step))[_x])
+
+#define noconvert
+
+// horizontal and vertical filter kernels
+// should be defined on host during compile time to avoid overhead
+#define DIG(a) a,
+__constant float mat_kernelX[] = { KERNEL_MATRIX_X };
+__constant float mat_kernelY[] = { KERNEL_MATRIX_Y };
+
+__kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width,
+ __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+{
+ // RADIUSX, RADIUSY are filter dimensions
+ // BLK_X, BLK_Y are local wrogroup sizes
+ // all these should be defined on host during compile time
+ // first lsmem array for source pixels used in first pass,
+ // second lsmemDy for storing first pass results
+ __local WT lsmem[BLK_Y+2*RADIUSY][BLK_X+2*RADIUSX];
+ __local WT lsmemDy[BLK_Y][BLK_X+2*RADIUSX];
+
+ // get local and global ids - used as image and local memory array indexes
+ int lix = get_local_id(0);
+ int liy = get_local_id(1);
+
+ int x = (int)get_global_id(0);
+ int y = (int)get_global_id(1);
+
+ // calculate pixel position in source image taking image offset into account
+ int srcX = x + srcOffsetX - RADIUSX;
+ int srcY = y + srcOffsetY - RADIUSY;
+ int xb = srcX;
+ int yb = srcY;
+
+ // extrapolate coordinates, if needed
+ // and read my own source pixel into local memory
+ // with account for extra border pixels, which will be read by starting workitems
+ int clocY = liy;
+ int cSrcY = srcY;
+ do
+ {
+ int yb = cSrcY;
+ EXTRAPOLATE(yb, (height));
+
+ int clocX = lix;
+ int cSrcX = srcX;
+ do
+ {
+ int xb = cSrcX;
+ EXTRAPOLATE(xb,(width));
+ lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 );
+
+ clocX += BLK_X;
+ cSrcX += BLK_X;
+ }
+ while(clocX < BLK_X+(RADIUSX*2));
+
+ clocY += BLK_Y;
+ cSrcY += BLK_Y;
+ }
+ while (clocY < BLK_Y+(RADIUSY*2));
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ // do vertical filter pass
+ // and store intermediate results to second local memory array
+ int i, clocX = lix;
+ WT sum = 0.0f;
+ do
+ {
+ sum = 0.0f;
+ for (i=0; i<=2*RADIUSY; i++)
+ sum = mad(lsmem[liy+i][clocX], mat_kernelY[i], sum);
+ lsmemDy[liy][clocX] = sum;
+ clocX += BLK_X;
+ }
+ while(clocX < BLK_X+(RADIUSX*2));
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ // if this pixel happened to be out of image borders because of global size rounding,
+ // then just return
+ if( x >= dst_cols || y >=dst_rows )
+ return;
+
+ // do second horizontal filter pass
+ // and calculate final result
+ sum = 0.0f;
+ for (i=0; i<=2*RADIUSX; i++)
+ sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum);
+
+ //store result into destination image
+ DST(x,y) = convertToDstT(sum);
+}
From b9cdde69911289b45c3686d2eee223f814380c97 Mon Sep 17 00:00:00 2001
From: yash
Date: Tue, 18 Mar 2014 08:44:33 +0530
Subject: [PATCH 15/47] edited sample code for mean/cam sihft and fixed an
error
---
doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst b/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst
index a111311af3..87ece69350 100644
--- a/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst
+++ b/doc/py_tutorials/py_video/py_meanshift/py_meanshift.rst
@@ -52,7 +52,7 @@ To use meanshift in OpenCV, first we need to setup the target, find its histogra
# set up the ROI for tracking
roi = frame[r:r+h, c:c+w]
- hsv_roi = cv2.cvtColor(frame, cv2.COLOR_BGR2HSV)
+ hsv_roi = cv2.cvtColor(roi, cv2.COLOR_BGR2HSV)
mask = cv2.inRange(hsv_roi, np.array((0., 60.,32.)), np.array((180.,255.,255.)))
roi_hist = cv2.calcHist([hsv_roi],[0],mask,[180],[0,180])
cv2.normalize(roi_hist,roi_hist,0,255,cv2.NORM_MINMAX)
@@ -127,7 +127,7 @@ It is almost same as meanshift, but it returns a rotated rectangle (that is our
# set up the ROI for tracking
roi = frame[r:r+h, c:c+w]
- hsv_roi = cv2.cvtColor(frame, cv2.COLOR_BGR2HSV)
+ hsv_roi = cv2.cvtColor(roi, cv2.COLOR_BGR2HSV)
mask = cv2.inRange(hsv_roi, np.array((0., 60.,32.)), np.array((180.,255.,255.)))
roi_hist = cv2.calcHist([hsv_roi],[0],mask,[180],[0,180])
cv2.normalize(roi_hist,roi_hist,0,255,cv2.NORM_MINMAX)
From 16869225ffb6029cea0f694b8ee9ffab2eb835bd Mon Sep 17 00:00:00 2001
From: RJ2
Date: Tue, 18 Mar 2014 08:01:18 +0100
Subject: [PATCH 16/47] It's will be better
---
doc/tutorials/core/adding_images/adding_images.rst | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/doc/tutorials/core/adding_images/adding_images.rst b/doc/tutorials/core/adding_images/adding_images.rst
index e3135693de..601dbc07ef 100644
--- a/doc/tutorials/core/adding_images/adding_images.rst
+++ b/doc/tutorials/core/adding_images/adding_images.rst
@@ -6,12 +6,12 @@ Adding (blending) two images using OpenCV
Goal
=====
-In this tutorial you will learn how to:
+In this tutorial you will learn:
.. container:: enumeratevisibleitemswithsquare
- * What is *linear blending* and why it is useful.
- * Add two images using :add_weighted:`addWeighted <>`
+ * what is *linear blending* and why it is useful;
+ * how to add two images using :add_weighted:`addWeighted <>`
Theory
=======
From 0470bb0e2958ae9725465fe6ab2767f19c5009fa Mon Sep 17 00:00:00 2001
From: RJ2
Date: Tue, 18 Mar 2014 08:59:53 +0100
Subject: [PATCH 17/47] I have changed one sentence in tutorial, making it more
understandable
---
doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst b/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst
index ef0f8640ca..b6a18fee88 100644
--- a/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst
+++ b/doc/tutorials/core/how_to_scan_images/how_to_scan_images.rst
@@ -18,7 +18,7 @@ We'll seek answers for the following questions:
Our test case
=============
-Let us consider a simple color reduction method. Using the unsigned char C and C++ type for matrix item storing a channel of pixel may have up to 256 different values. For a three channel image this can allow the formation of way too many colors (16 million to be exact). Working with so many color shades may give a heavy blow to our algorithm performance. However, sometimes it is enough to work with a lot less of them to get the same final result.
+Let us consider a simple color reduction method. By using the unsigned char C and C++ type for matrix item storing, a channel of pixel may have up to 256 different values. For a three channel image this can allow the formation of way too many colors (16 million to be exact). Working with so many color shades may give a heavy blow to our algorithm performance. However, sometimes it is enough to work with a lot less of them to get the same final result.
In this cases it's common that we make a *color space reduction*. This means that we divide the color space current value with a new input value to end up with fewer colors. For instance every value between zero and nine takes the new value zero, every value between ten and nineteen the value ten and so on.
From b4e4f13f9efdfe2d1f0909d78556d8c865d5f371 Mon Sep 17 00:00:00 2001
From: Alexander Smorkalov
Date: Wed, 5 Mar 2014 12:31:04 +0400
Subject: [PATCH 18/47] Superres module enabled for Android. GPU samples build
fixed for Android.
---
cmake/OpenCVModule.cmake | 66 +++++++++-----------
modules/superres/CMakeLists.txt | 9 ++-
modules/superres/src/btv_l1_gpu.cpp | 2 +-
modules/superres/src/frame_source.cpp | 2 +-
modules/superres/src/input_array_utility.cpp | 2 +-
modules/superres/src/optical_flow.cpp | 2 +-
modules/superres/src/precomp.hpp | 2 +-
modules/ts/CMakeLists.txt | 2 +-
samples/gpu/CMakeLists.txt | 2 +-
samples/gpu/brox_optical_flow.cpp | 1 +
samples/gpu/opticalflow_nvidia_api.cpp | 1 +
samples/gpu/super_resolution.cpp | 2 +
12 files changed, 49 insertions(+), 44 deletions(-)
diff --git a/cmake/OpenCVModule.cmake b/cmake/OpenCVModule.cmake
index c9c351113c..e11f5e671f 100644
--- a/cmake/OpenCVModule.cmake
+++ b/cmake/OpenCVModule.cmake
@@ -27,7 +27,8 @@
# The verbose template for OpenCV module:
#
# ocv_add_module(modname )
-# ocv_glob_module_sources() or glob them manually and ocv_set_module_sources(...)
+# ocv_glob_module_sources(([EXCLUDE_CUDA] )
+# or glob them manually and ocv_set_module_sources(...)
# ocv_module_include_directories()
# ocv_create_module()
#
@@ -478,14 +479,20 @@ endmacro()
# finds and sets headers and sources for the standard OpenCV module
# Usage:
-# ocv_glob_module_sources()
-macro(ocv_glob_module_sources EXCLUDE_CUDA EXCLUDE_OPENCL)
+# ocv_glob_module_sources([EXCLUDE_CUDA] )
+macro(ocv_glob_module_sources)
+ set(_argn ${ARGN})
+ list(FIND _argn "EXCLUDE_CUDA" exclude_cuda)
+ if(NOT exclude_cuda EQUAL -1)
+ list(REMOVE_AT _argn ${exclude_cuda})
+ endif()
+
file(GLOB_RECURSE lib_srcs "src/*.cpp")
file(GLOB_RECURSE lib_int_hdrs "src/*.hpp" "src/*.h")
file(GLOB lib_hdrs "include/opencv2/${name}/*.hpp" "include/opencv2/${name}/*.h")
file(GLOB lib_hdrs_detail "include/opencv2/${name}/detail/*.hpp" "include/opencv2/${name}/detail/*.h")
- if (NOT ${EXCLUDE_CUDA})
+ if (exclude_cuda EQUAL -1)
file(GLOB lib_cuda_srcs "src/cuda/*.cu")
set(cuda_objs "")
set(lib_cuda_hdrs "")
@@ -504,26 +511,22 @@ macro(ocv_glob_module_sources EXCLUDE_CUDA EXCLUDE_OPENCL)
source_group("Src" FILES ${lib_srcs} ${lib_int_hdrs})
- if (NOT ${EXCLUDE_OPENCL})
- file(GLOB cl_kernels "src/opencl/*.cl")
- if(HAVE_opencv_ocl AND cl_kernels)
- ocv_include_directories(${OPENCL_INCLUDE_DIRS})
- add_custom_command(
- OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp"
- COMMAND ${CMAKE_COMMAND} -DCL_DIR="${CMAKE_CURRENT_SOURCE_DIR}/src/opencl" -DOUTPUT="${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" -P "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake"
- DEPENDS ${cl_kernels} "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake")
- source_group("OpenCL" FILES ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
- list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
- endif()
+ file(GLOB cl_kernels "src/opencl/*.cl")
+ if(HAVE_opencv_ocl AND cl_kernels)
+ ocv_include_directories(${OPENCL_INCLUDE_DIRS})
+ add_custom_command(
+ OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp"
+ COMMAND ${CMAKE_COMMAND} -DCL_DIR="${CMAKE_CURRENT_SOURCE_DIR}/src/opencl" -DOUTPUT="${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" -P "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake"
+ DEPENDS ${cl_kernels} "${OpenCV_SOURCE_DIR}/cmake/cl2cpp.cmake")
+ source_group("OpenCL" FILES ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
+ list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
endif()
source_group("Include" FILES ${lib_hdrs})
source_group("Include\\detail" FILES ${lib_hdrs_detail})
- message(":${EXCLUDE_CUDA}: ${lib_cuda_srcs}")
-
- ocv_set_module_sources(${ARGN} HEADERS ${lib_hdrs} ${lib_hdrs_detail}
- SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs})
+ ocv_set_module_sources(${_argn} HEADERS ${lib_hdrs} ${lib_hdrs_detail}
+ SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs})
endmacro()
# creates OpenCV module in current folder
@@ -622,27 +625,20 @@ endmacro()
# short command for adding simple OpenCV module
# see ocv_add_module for argument details
# Usage:
-# ocv_define_module(module_name [INTERNAL] [REQUIRED] [] [OPTIONAL ])
+# ocv_define_module(module_name [INTERNAL] [EXCLUDE_CUDA] [REQUIRED] [] [OPTIONAL ])
macro(ocv_define_module module_name)
- set(_tmp_argn ${ARGN})
- set(exclude_cuda 0)
- set(exclude_opencl 0)
- set(argv0 ${ARGV1})
- set(argv1 ${ARGV2})
- set(argv2 ${ARGV3})
- foreach(i RANGE 0 2)
- if("${argv${i}}" STREQUAL "EXCLUDE_CUDA")
- set(exclude_cuda 1)
- list(REMOVE_AT _tmp_argn ${i})
- elseif ("${argv${i}}" STREQUAL "EXCLUDE_OPENCL")
- set(exclude_opencl 1)
- list(REMOVE_AT _tmp_argn ${i})
+ set(_argn ${ARGN})
+ set(exclude_cuda "")
+ foreach(arg ${_argn})
+ if("${arg}" STREQUAL "EXCLUDE_CUDA")
+ set(exclude_cuda "${arg}")
+ list(REMOVE_ITEM _argn ${arg})
endif()
endforeach()
- ocv_add_module(${module_name} ${_tmp_argn})
+ ocv_add_module(${module_name} ${_argn})
ocv_module_include_directories()
- ocv_glob_module_sources(${exclude_cuda} ${exclude_opencl})
+ ocv_glob_module_sources(${exclude_cuda})
ocv_create_module()
ocv_add_precompiled_headers(${the_module})
diff --git a/modules/superres/CMakeLists.txt b/modules/superres/CMakeLists.txt
index 82c61cccb4..8e3d7b2f2d 100644
--- a/modules/superres/CMakeLists.txt
+++ b/modules/superres/CMakeLists.txt
@@ -1,7 +1,12 @@
-if(ANDROID OR IOS)
+if(IOS)
ocv_module_disable(superres)
endif()
set(the_description "Super Resolution")
ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 -Wundef -Wshadow)
-ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_gpu opencv_highgui opencv_ocl ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
+if(ENABLE_DYNAMIC_CUDA)
+ add_definitions(-DDYNAMIC_CUDA_SUPPORT)
+ ocv_define_module(superres EXCLUDE_CUDA opencv_imgproc opencv_video OPTIONAL opencv_highgui opencv_ocl)
+else()
+ ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_gpu opencv_highgui opencv_ocl ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
+endif()
diff --git a/modules/superres/src/btv_l1_gpu.cpp b/modules/superres/src/btv_l1_gpu.cpp
index b93bcfd577..2f40fa3771 100644
--- a/modules/superres/src/btv_l1_gpu.cpp
+++ b/modules/superres/src/btv_l1_gpu.cpp
@@ -51,7 +51,7 @@ using namespace cv::gpu;
using namespace cv::superres;
using namespace cv::superres::detail;
-#if !defined(HAVE_CUDA) || !defined(HAVE_OPENCV_GPU)
+#if !defined(HAVE_CUDA) || !defined(HAVE_OPENCV_GPU) || defined(DYNAMIC_CUDA_SUPPORT)
Ptr cv::superres::createSuperResolution_BTVL1_GPU()
{
diff --git a/modules/superres/src/frame_source.cpp b/modules/superres/src/frame_source.cpp
index 20e45d9518..5f59a98b80 100644
--- a/modules/superres/src/frame_source.cpp
+++ b/modules/superres/src/frame_source.cpp
@@ -200,7 +200,7 @@ Ptr cv::superres::createFrameSource_Camera(int deviceId)
//////////////////////////////////////////////////////
// VideoFrameSource_GPU
-#ifndef HAVE_OPENCV_GPU
+#if !defined(HAVE_OPENCV_GPU) || defined(DYNAMIC_CUDA_SUPPORT)
Ptr cv::superres::createFrameSource_Video_GPU(const string& fileName)
{
diff --git a/modules/superres/src/input_array_utility.cpp b/modules/superres/src/input_array_utility.cpp
index 075cf95144..10fc1c96ed 100644
--- a/modules/superres/src/input_array_utility.cpp
+++ b/modules/superres/src/input_array_utility.cpp
@@ -207,7 +207,7 @@ namespace
switch (src.kind())
{
case _InputArray::GPU_MAT:
- #ifdef HAVE_OPENCV_GPU
+ #if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
gpu::cvtColor(src.getGpuMat(), dst.getGpuMatRef(), code, cn);
#else
CV_Error(CV_StsNotImplemented, "The called functionality is disabled for current build or platform");
diff --git a/modules/superres/src/optical_flow.cpp b/modules/superres/src/optical_flow.cpp
index e1e8a10275..617f83a5e1 100644
--- a/modules/superres/src/optical_flow.cpp
+++ b/modules/superres/src/optical_flow.cpp
@@ -344,7 +344,7 @@ Ptr cv::superres::createOptFlow_DualTVL1()
///////////////////////////////////////////////////////////////////
// GpuOpticalFlow
-#ifndef HAVE_OPENCV_GPU
+#if !defined(HAVE_OPENCV_GPU) || defined(DYNAMIC_CUDA_SUPPORT)
Ptr cv::superres::createOptFlow_Farneback_GPU()
{
diff --git a/modules/superres/src/precomp.hpp b/modules/superres/src/precomp.hpp
index 4cf49411b6..65fc7c8502 100644
--- a/modules/superres/src/precomp.hpp
+++ b/modules/superres/src/precomp.hpp
@@ -56,7 +56,7 @@
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/video/tracking.hpp"
-#ifdef HAVE_OPENCV_GPU
+#if defined(HAVE_OPENCV_GPU) && !defined(DYNAMIC_CUDA_SUPPORT)
#include "opencv2/gpu/gpu.hpp"
#ifdef HAVE_CUDA
#include "opencv2/gpu/stream_accessor.hpp"
diff --git a/modules/ts/CMakeLists.txt b/modules/ts/CMakeLists.txt
index dcd3e1563b..bb56da2d98 100644
--- a/modules/ts/CMakeLists.txt
+++ b/modules/ts/CMakeLists.txt
@@ -11,7 +11,7 @@ ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
ocv_add_module(ts opencv_core opencv_features2d)
-ocv_glob_module_sources(0 0)
+ocv_glob_module_sources()
ocv_module_include_directories()
ocv_create_module()
diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt
index 8fa539473b..d25c3a6b5d 100644
--- a/samples/gpu/CMakeLists.txt
+++ b/samples/gpu/CMakeLists.txt
@@ -41,7 +41,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})
- if(HAVE_CUDA)
+ if(HAVE_CUDA AND NOT ANDROID)
target_link_libraries(${the_target} ${CUDA_CUDA_LIBRARY})
endif()
diff --git a/samples/gpu/brox_optical_flow.cpp b/samples/gpu/brox_optical_flow.cpp
index 722e19f02f..7cd5089b41 100644
--- a/samples/gpu/brox_optical_flow.cpp
+++ b/samples/gpu/brox_optical_flow.cpp
@@ -1,6 +1,7 @@
#include
#include
#include
+#include
#include "cvconfig.h"
#include "opencv2/core/core.hpp"
diff --git a/samples/gpu/opticalflow_nvidia_api.cpp b/samples/gpu/opticalflow_nvidia_api.cpp
index 05a37ef69d..31ee569788 100644
--- a/samples/gpu/opticalflow_nvidia_api.cpp
+++ b/samples/gpu/opticalflow_nvidia_api.cpp
@@ -7,6 +7,7 @@
#include
#include
#include
+#include
#include "cvconfig.h"
#include
diff --git a/samples/gpu/super_resolution.cpp b/samples/gpu/super_resolution.cpp
index 6efd24144c..85cb6cff1c 100644
--- a/samples/gpu/super_resolution.cpp
+++ b/samples/gpu/super_resolution.cpp
@@ -1,6 +1,8 @@
#include
#include
#include
+#include
+
#include "opencv2/core/core.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/imgproc/imgproc.hpp"
From 0a16d93e1d2b8c23df77c769589cfc7982c68af1 Mon Sep 17 00:00:00 2001
From: Firat Kalaycilar
Date: Tue, 18 Mar 2014 17:26:24 +0200
Subject: [PATCH 19/47] Fixed an issue with weight assignment causing the
resulting GMM weights to be unsorted in the CUDA and OCL versions of
BackgroundSubtractorMOG2
---
modules/gpu/src/cuda/bgfg_mog.cu | 5 +++--
modules/ocl/src/opencl/bgfg_mog.cl | 5 +++--
2 files changed, 6 insertions(+), 4 deletions(-)
diff --git a/modules/gpu/src/cuda/bgfg_mog.cu b/modules/gpu/src/cuda/bgfg_mog.cu
index 89ad5ff0b5..055c9e041e 100644
--- a/modules/gpu/src/cuda/bgfg_mog.cu
+++ b/modules/gpu/src/cuda/bgfg_mog.cu
@@ -489,7 +489,7 @@ namespace cv { namespace gpu { namespace device
{
//need only weight if fit is found
float weight = alpha1 * gmm_weight(mode * frame.rows + y, x) + prune;
-
+ int swap_count = 0;
//fit not found yet
if (!fitsPDF)
{
@@ -540,6 +540,7 @@ namespace cv { namespace gpu { namespace device
if (weight < gmm_weight((i - 1) * frame.rows + y, x))
break;
+ swap_count++;
//swap one up
swap(gmm_weight, x, y, i - 1, frame.rows);
swap(gmm_variance, x, y, i - 1, frame.rows);
@@ -557,7 +558,7 @@ namespace cv { namespace gpu { namespace device
nmodes--;
}
- gmm_weight(mode * frame.rows + y, x) = weight; //update weight by the calculated value
+ gmm_weight((mode - swap_count) * frame.rows + y, x) = weight; //update weight by the calculated value
totalWeight += weight;
}
diff --git a/modules/ocl/src/opencl/bgfg_mog.cl b/modules/ocl/src/opencl/bgfg_mog.cl
index 6a95316f0f..a7479b929c 100644
--- a/modules/ocl/src/opencl/bgfg_mog.cl
+++ b/modules/ocl/src/opencl/bgfg_mog.cl
@@ -376,7 +376,7 @@ __kernel void mog2_kernel(__global T_FRAME * frame, __global int* fgmask, __glob
for (int mode = 0; mode < nmodes; ++mode)
{
float _weight = alpha1 * weight[(mode * frame_row + y) * weight_step + x] + prune;
-
+ int swap_count = 0;
if (!fitsPDF)
{
float var = variance[(mode * frame_row + y) * var_step + x];
@@ -404,6 +404,7 @@ __kernel void mog2_kernel(__global T_FRAME * frame, __global int* fgmask, __glob
{
if (_weight < weight[((i - 1) * frame_row + y) * weight_step + x])
break;
+ swap_count++;
swap(weight, x, y, i - 1, frame_row, weight_step);
swap(variance, x, y, i - 1, frame_row, var_step);
#if defined (CN1)
@@ -421,7 +422,7 @@ __kernel void mog2_kernel(__global T_FRAME * frame, __global int* fgmask, __glob
nmodes--;
}
- weight[(mode * frame_row + y) * weight_step + x] = _weight; //update weight by the calculated value
+ weight[((mode - swap_count) * frame_row + y) * weight_step + x] = _weight; //update weight by the calculated value
totalWeight += _weight;
}
From 8d97d0d6313d6e1e60ff91ad088e4a69a52bd1a5 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Tue, 18 Mar 2014 19:31:37 +0400
Subject: [PATCH 20/47] added 3-channels support to cv::flip
---
modules/core/src/copy.cpp | 7 +++--
modules/core/src/opencl/flip.cl | 55 ++++++++++++++++-----------------
2 files changed, 31 insertions(+), 31 deletions(-)
diff --git a/modules/core/src/copy.cpp b/modules/core/src/copy.cpp
index e3e959c950..5ac5f22c58 100644
--- a/modules/core/src/copy.cpp
+++ b/modules/core/src/copy.cpp
@@ -482,9 +482,9 @@ enum { FLIP_COLS = 1 << 0, FLIP_ROWS = 1 << 1, FLIP_BOTH = FLIP_ROWS | FLIP_COLS
static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
{
CV_Assert(flipCode >= - 1 && flipCode <= 1);
- int type = _src.type(), cn = CV_MAT_CN(type), flipType;
+ int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), flipType;
- if (cn > 4 || cn == 3)
+ if (cn > 4)
return false;
const char * kernelName;
@@ -506,7 +506,8 @@ static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
}
ocl::Kernel k(kernelName, ocl::core::flip_oclsrc,
- format( "-D type=%s", ocl::memopTypeToStr(type)));
+ format( "-D T=%s -D T1=%s -D cn=%d", ocl::memopTypeToStr(type),
+ ocl::memopTypeToStr(depth), cn));
if (k.empty())
return false;
diff --git a/modules/core/src/opencl/flip.cl b/modules/core/src/opencl/flip.cl
index 0c874dbe6f..bacfe7adfb 100644
--- a/modules/core/src/opencl/flip.cl
+++ b/modules/core/src/opencl/flip.cl
@@ -39,10 +39,18 @@
//
//M*/
-#define sizeoftype ((int)sizeof(type))
+#if cn != 3
+#define loadpix(addr) *(__global const T *)(addr)
+#define storepix(val, addr) *(__global T *)(addr) = val
+#define TSIZE (int)sizeof(T)
+#else
+#define loadpix(addr) vload3(0, (__global const T1 *)(addr))
+#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
+#define TSIZE ((int)sizeof(T1)*3)
+#endif
-__kernel void arithm_flip_rows(__global const uchar* srcptr, int srcstep, int srcoffset,
- __global uchar* dstptr, int dststep, int dstoffset,
+__kernel void arithm_flip_rows(__global const uchar * srcptr, int src_step, int src_offset,
+ __global uchar * dstptr, int dst_step, int dst_offset,
int rows, int cols, int thread_rows, int thread_cols)
{
int x = get_global_id(0);
@@ -50,19 +58,16 @@ __kernel void arithm_flip_rows(__global const uchar* srcptr, int srcstep, int sr
if (x < cols && y < thread_rows)
{
- __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x, sizeoftype, srcoffset)));
- __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, mad24(x, sizeoftype, srcoffset)));
+ T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
+ T src1 = loadpix(srcptr + mad24(rows - y - 1, src_step, mad24(x, TSIZE, src_offset)));
- __global type* dst0 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset)));
- __global type* dst1 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, mad24(x, sizeoftype, dstoffset)));
-
- dst0[0] = src1[0];
- dst1[0] = src0[0];
+ storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
+ storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x, TSIZE, dst_offset)));
}
}
-__kernel void arithm_flip_rows_cols(__global const uchar* srcptr, int srcstep, int srcoffset,
- __global uchar* dstptr, int dststep, int dstoffset,
+__kernel void arithm_flip_rows_cols(__global const uchar * srcptr, int src_step, int src_offset,
+ __global uchar * dstptr, int dst_step, int dst_offset,
int rows, int cols, int thread_rows, int thread_cols)
{
int x = get_global_id(0);
@@ -71,19 +76,16 @@ __kernel void arithm_flip_rows_cols(__global const uchar* srcptr, int srcstep, i
if (x < cols && y < thread_rows)
{
int x1 = cols - x - 1;
- __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x, sizeoftype, srcoffset)));
- __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, mad24(x1, sizeoftype, srcoffset)));
+ T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
+ T src1 = loadpix(srcptr + mad24(rows - y - 1, src_step, mad24(x1, TSIZE, src_offset)));
- __global type* dst0 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, mad24(x1, sizeoftype, dstoffset)));
- __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset)));
-
- dst0[0] = src0[0];
- dst1[0] = src1[0];
+ storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x1, TSIZE, dst_offset)));
+ storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
}
}
-__kernel void arithm_flip_cols(__global const uchar* srcptr, int srcstep, int srcoffset,
- __global uchar* dstptr, int dststep, int dstoffset,
+__kernel void arithm_flip_cols(__global const uchar * srcptr, int src_step, int src_offset,
+ __global uchar * dstptr, int dst_step, int dst_offset,
int rows, int cols, int thread_rows, int thread_cols)
{
int x = get_global_id(0);
@@ -92,13 +94,10 @@ __kernel void arithm_flip_cols(__global const uchar* srcptr, int srcstep, int sr
if (x < thread_cols && y < rows)
{
int x1 = cols - x - 1;
- __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x, sizeoftype, srcoffset)));
- __global const type* src1 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x1, sizeoftype, srcoffset)));
+ T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
+ T src1 = loadpix(srcptr + mad24(y, src_step, mad24(x1, TSIZE, src_offset)));
- __global type* dst0 = (__global type*)(dstptr + mad24(y, dststep, mad24(x1, sizeoftype, dstoffset)));
- __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset)));
-
- dst1[0] = src1[0];
- dst0[0] = src0[0];
+ storepix(src0, dstptr + mad24(y, dst_step, mad24(x1, TSIZE, dst_offset)));
+ storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
}
}
From d1cfcfcafd41d81522f8c4d3b64e62d3e7ecb9dc Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Tue, 18 Mar 2014 20:02:04 +0400
Subject: [PATCH 21/47] added 3-channels support to morphology operations
---
modules/imgproc/src/morph.cpp | 54 ++++------
modules/imgproc/src/opencl/morph.cl | 119 +++++++++++-----------
modules/imgproc/test/ocl/test_filters.cpp | 16 +--
3 files changed, 87 insertions(+), 102 deletions(-)
diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp
index ac958fc691..b4011eecee 100644
--- a/modules/imgproc/src/morph.cpp
+++ b/modules/imgproc/src/morph.cpp
@@ -42,7 +42,6 @@
#include "precomp.hpp"
#include
-#include
#include "opencl_kernels.hpp"
/****************************************************************************************\
@@ -1291,9 +1290,10 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
{
CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE);
+ int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
- if (_src.depth() == CV_64F && !doubleSupport)
+ if (depth == CV_64F && !doubleSupport)
return false;
UMat kernel8U;
@@ -1324,13 +1324,14 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
return false;
static const char * const op2str[] = { "ERODE", "DILATE" };
- String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s -D GENTYPE=%s -D DEPTH_%d",
- anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op],
+ String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s"
+ " -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y,
+ (int)localThreads[0], (int)localThreads[1], op2str[op],
doubleSupport ? " -D DOUBLE_SUPPORT" : "", rectKernel ? " -D RECTKERNEL" : "",
- ocl::typeToStr(_src.type()), _src.depth() );
+ ocl::typeToStr(_src.type()), _src.depth(), cn, ocl::typeToStr(depth));
std::vector kernels;
- for (int i = 0; i= (r_edge) ? (elem1) : (elem2)
-__kernel void morph(__global const uchar * restrict srcptr, int src_step, int src_offset,
+// BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
+#define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
+
+__kernel void morph(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset,
- int src_offset_x, int src_offset_y,
- int cols, int rows,
- __constant uchar * mat_kernel,
- int src_whole_cols, int src_whole_rows)
+ int src_offset_x, int src_offset_y, int cols, int rows,
+ __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows)
{
- int l_x = get_local_id(0);
- int l_y = get_local_id(1);
- int x = get_group_id(0)*LSIZE0;
- int y = get_group_id(1)*LSIZE1;
- int start_x = x+src_offset_x-RADIUSX;
- int end_x = x + src_offset_x+LSIZE0+RADIUSX;
- int width = end_x -(x+src_offset_x-RADIUSX)+1;
- int start_y = y+src_offset_y-RADIUSY;
- int point1 = mad24(l_y,LSIZE0,l_x);
- int point2 = point1 + LSIZE0*LSIZE1;
- int tl_x = point1 % width;
- int tl_y = point1 / width;
- int tl_x2 = point2 % width;
- int tl_y2 = point2 / width;
- int cur_x = start_x + tl_x;
- int cur_y = start_y + tl_y;
- int cur_x2 = start_x + tl_x2;
- int cur_y2 = start_y + tl_y2;
- int start_addr = mad24(cur_y,src_step, cur_x*(int)sizeof(GENTYPE));
- int start_addr2 = mad24(cur_y2,src_step, cur_x2*(int)sizeof(GENTYPE));
- GENTYPE temp0,temp1;
- __local GENTYPE LDS_DAT[2*LSIZE1*LSIZE0];
+ int gidx = get_global_id(0), gidy = get_global_id(1);
+ int l_x = get_local_id(0), l_y = get_local_id(1);
+ int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1;
+ int start_x = x + src_offset_x - RADIUSX;
+ int end_x = x + src_offset_x + LSIZE0 + RADIUSX;
+ int width = end_x - (x + src_offset_x - RADIUSX) + 1;
+ int start_y = y + src_offset_y - RADIUSY;
+ int point1 = mad24(l_y, LSIZE0, l_x);
+ int point2 = point1 + LSIZE0 * LSIZE1;
+ int tl_x = point1 % width, tl_y = point1 / width;
+ int tl_x2 = point2 % width, tl_y2 = point2 / width;
+ int cur_x = start_x + tl_x, cur_y = start_y + tl_y;
+ int cur_x2 = start_x + tl_x2, cur_y2 = start_y + tl_y2;
+ int start_addr = mad24(cur_y, src_step, cur_x * TSIZE);
+ int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE);
- int end_addr = mad24(src_whole_rows - 1,src_step,src_whole_cols*(int)sizeof(GENTYPE));
- //read pixels from src
- start_addr = ((start_addr < end_addr) && (start_addr > 0)) ? start_addr : 0;
- start_addr2 = ((start_addr2 < end_addr) && (start_addr2 > 0)) ? start_addr2 : 0;
- __global const GENTYPE * src;
- src = (__global const GENTYPE *)(srcptr+start_addr);
- temp0 = src[0];
- src = (__global const GENTYPE *)(srcptr+start_addr2);
- temp1 = src[0];
- //judge if read out of boundary
- temp0= ELEM(cur_x,0,src_whole_cols,(GENTYPE)VAL,temp0);
- temp0= ELEM(cur_y,0,src_whole_rows,(GENTYPE)VAL,temp0);
+ __local T LDS_DAT[2*LSIZE1*LSIZE0];
- temp1= ELEM(cur_x2,0,src_whole_cols,(GENTYPE)VAL,temp1);
- temp1= ELEM(cur_y2,0,src_whole_rows,(GENTYPE)VAL,temp1);
+ // read pixels from src
+ int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE);
+ start_addr = start_addr < end_addr && start_addr > 0 ? start_addr : 0;
+ start_addr2 = start_addr2 < end_addr && start_addr2 > 0 ? start_addr2 : 0;
+
+ T temp0 = loadpix(srcptr + start_addr);
+ T temp1 = loadpix(srcptr + start_addr2);
+
+ // judge if read out of boundary
+ temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL),temp0);
+ temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL),temp0);
+
+ temp1 = ELEM(cur_x2, 0, src_whole_cols, (T)(VAL), temp1);
+ temp1 = ELEM(cur_y2, 0, src_whole_rows, (T)(VAL), temp1);
LDS_DAT[point1] = temp0;
LDS_DAT[point2] = temp1;
barrier(CLK_LOCAL_MEM_FENCE);
- GENTYPE res = (GENTYPE)VAL;
- for(int i=0; i<2*RADIUSY+1; i++)
- for(int j=0; j<2*RADIUSX+1; j++)
+
+ T res = (T)(VAL);
+ for (int i = 0, sizey = 2 * RADIUSY + 1; i < sizey; i++)
+ for (int j = 0, sizex = 2 * RADIUSX + 1; j < sizex; j++)
{
res =
#ifndef RECTKERNEL
mat_kernel[i*(2*RADIUSX+1)+j] ?
#endif
- MORPH_OP(res,LDS_DAT[mad24(l_y+i,width,l_x+j)])
+ MORPH_OP(res, LDS_DAT[mad24(l_y + i, width, l_x + j)])
#ifndef RECTKERNEL
- :res
+ : res
#endif
;
}
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
- if(gidx
Date: Tue, 18 Mar 2014 19:42:04 +0400
Subject: [PATCH 22/47] added 3-channels support to cv::setIdentity
---
modules/core/src/matrix.cpp | 10 +++++-----
modules/core/src/opencl/set_identity.cl | 19 +++++++++++++++----
2 files changed, 20 insertions(+), 9 deletions(-)
diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp
index db1ce760f3..82d177eda4 100644
--- a/modules/core/src/matrix.cpp
+++ b/modules/core/src/matrix.cpp
@@ -2679,17 +2679,17 @@ namespace cv {
static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s )
{
- int type = _m.type(), cn = CV_MAT_CN(type);
- if (cn == 3)
- return false;
+ int type = _m.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
+ sctype = CV_MAKE_TYPE(depth, cn == 3 ? 4 : cn);
ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc,
- format("-D T=%s", ocl::memopTypeToStr(type)));
+ format("-D T=%s -D T1=%s -D cn=%d -D ST=%s", ocl::memopTypeToStr(type),
+ ocl::memopTypeToStr(depth), cn, ocl::memopTypeToStr(sctype)));
if (k.empty())
return false;
UMat m = _m.getUMat();
- k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, type, s)));
+ k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, sctype, s)));
size_t globalsize[2] = { m.cols, m.rows };
return k.run(2, globalsize, NULL, false);
diff --git a/modules/core/src/opencl/set_identity.cl b/modules/core/src/opencl/set_identity.cl
index d63ce793db..0e8f1424fb 100644
--- a/modules/core/src/opencl/set_identity.cl
+++ b/modules/core/src/opencl/set_identity.cl
@@ -43,17 +43,28 @@
//
//M*/
+#if cn != 3
+#define loadpix(addr) *(__global const T *)(addr)
+#define storepix(val, addr) *(__global T *)(addr) = val
+#define TSIZE (int)sizeof(T)
+#define scalar scalar_
+#else
+#define loadpix(addr) vload3(0, (__global const T1 *)(addr))
+#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
+#define TSIZE ((int)sizeof(T1)*3)
+#define scalar (T)(scalar_.x, scalar_.y, scalar_.z)
+#endif
+
__kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols,
- T scalar)
+ ST scalar_)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
- int src_index = mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset));
- __global T * src = (__global T *)(srcptr + src_index);
+ int src_index = mad24(y, src_step, mad24(x, TSIZE, src_offset));
- src[0] = x == y ? scalar : (T)(0);
+ storepix(x == y ? scalar : (T)(0), srcptr + src_index);
}
}
From b449b0bf71cf7fb69d8dd7a61ce43ebd5d239f91 Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Wed, 19 Mar 2014 15:59:00 +0400
Subject: [PATCH 23/47] simplified cv::sepFilter2D OpenCL part
---
modules/imgproc/src/filter.cpp | 133 +++---
modules/imgproc/src/opencl/filterSepCol.cl | 62 +--
modules/imgproc/src/opencl/filterSepRow.cl | 377 ++++--------------
.../src/opencl/filterSep_singlePass.cl | 11 +-
modules/imgproc/test/ocl/test_sepfilter2D.cpp | 4 +-
5 files changed, 167 insertions(+), 420 deletions(-)
diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp
index bb54471c07..ba2e347af0 100644
--- a/modules/imgproc/src/filter.cpp
+++ b/modules/imgproc/src/filter.cpp
@@ -41,6 +41,7 @@
//M*/
#include "precomp.hpp"
+#define CV_OPENCL_RUN_ASSERT
#include "opencl_kernels.hpp"
#include
@@ -3317,11 +3318,9 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
return kernel.run(2, globalsize, localsize, true);
}
-static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType, bool sync)
+static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType)
{
- int type = src.type();
- int cn = CV_MAT_CN(type);
- int sdepth = CV_MAT_DEPTH(type);
+ int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type);
Size bufSize = buf.size();
#ifdef ANDROID
@@ -3329,27 +3328,14 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor,
#else
size_t localsize[2] = {16, 16};
#endif
+
size_t globalsize[2] = {DIVUP(bufSize.width, localsize[0]) * localsize[0], DIVUP(bufSize.height, localsize[1]) * localsize[1]};
- if (CV_8U == sdepth)
- {
- switch (cn)
- {
- case 1:
- globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0];
- break;
- case 2:
- globalsize[0] = DIVUP((bufSize.width + 1) >> 1, localsize[0]) * localsize[0];
- break;
- case 4:
- globalsize[0] = DIVUP(bufSize.width, localsize[0]) * localsize[0];
- break;
- }
- }
+ if (type == CV_8UC1)
+ globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0];
- int radiusX = anchor;
- int radiusY = (int)((buf.rows - src.rows) >> 1);
+ int radiusX = anchor, radiusY = (buf.rows - src.rows) >> 1;
- bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
+ bool isolated = (borderType & BORDER_ISOLATED) != 0;
const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" },
* const btype = borderMap[borderType & ~BORDER_ISOLATED];
@@ -3358,49 +3344,38 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor,
extra_extrapolation |= src.cols < (int)((-radiusX + globalsize[0] + 8 * localsize[0] + 3) >> 1) + 1;
extra_extrapolation |= src.cols < radiusX;
- cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s",
- radiusX, (int)localsize[0], (int)localsize[1], cn,
- btype,
- extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
- isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED");
+ char cvt[40];
+ cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s"
+ " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s",
+ radiusX, (int)localsize[0], (int)localsize[1], cn, btype,
+ extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
+ isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED",
+ ocl::typeToStr(type), ocl::typeToStr(CV_32FC(cn)),
+ ocl::convertTypeStr(sdepth, CV_32F, cn, cvt),
+ ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F));
build_options += ocl::kernelToStr(kernelX, CV_32F);
Size srcWholeSize; Point srcOffset;
src.locateROI(srcWholeSize, srcOffset);
- std::stringstream strKernel;
- strKernel << "row_filter";
- if (-1 != cn)
- strKernel << "_C" << cn;
- if (-1 != sdepth)
- strKernel << "_D" << sdepth;
+ String kernelName("row_filter");
+ if (type == CV_8UC1)
+ kernelName += "_C1_D0";
- ocl::Kernel kernelRow;
- if (!kernelRow.create(strKernel.str().c_str(), cv::ocl::imgproc::filterSepRow_oclsrc,
- build_options))
+ ocl::Kernel k(kernelName.c_str(), cv::ocl::imgproc::filterSepRow_oclsrc,
+ build_options);
+ if (k.empty())
return false;
- int idxArg = 0;
- idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
- idxArg = kernelRow.set(idxArg, (int)(src.step / src.elemSize()));
+ k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x,
+ srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height,
+ ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()),
+ buf.cols, buf.rows, radiusY);
- idxArg = kernelRow.set(idxArg, srcOffset.x);
- idxArg = kernelRow.set(idxArg, srcOffset.y);
- idxArg = kernelRow.set(idxArg, src.cols);
- idxArg = kernelRow.set(idxArg, src.rows);
- idxArg = kernelRow.set(idxArg, srcWholeSize.width);
- idxArg = kernelRow.set(idxArg, srcWholeSize.height);
-
- idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrWriteOnly(buf));
- idxArg = kernelRow.set(idxArg, (int)(buf.step / buf.elemSize()));
- idxArg = kernelRow.set(idxArg, buf.cols);
- idxArg = kernelRow.set(idxArg, buf.rows);
- idxArg = kernelRow.set(idxArg, radiusY);
-
- return kernelRow.run(2, globalsize, localsize, sync);
+ return k.run(2, globalsize, localsize, false);
}
-static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor, bool sync)
+static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor)
{
#ifdef ANDROID
size_t localsize[2] = {16, 10};
@@ -3420,28 +3395,23 @@ static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anc
globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0];
char cvt[40];
- cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D GENTYPE_SRC=%s -D GENTYPE_DST=%s -D convert_to_DST=%s",
- anchor, (int)localsize[0], (int)localsize[1], cn, ocl::typeToStr(buf.type()),
- ocl::typeToStr(dtype), ocl::convertTypeStr(CV_32F, ddepth, cn, cvt));
+ cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d"
+ " -D srcT=%s -D dstT=%s -D convertToDstT=%s",
+ anchor, (int)localsize[0], (int)localsize[1], cn,
+ ocl::typeToStr(buf.type()), ocl::typeToStr(dtype),
+ ocl::convertTypeStr(CV_32F, ddepth, cn, cvt));
build_options += ocl::kernelToStr(kernelY, CV_32F);
- ocl::Kernel kernelCol;
- if (!kernelCol.create("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, build_options))
+ ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc,
+ build_options);
+ if (k.empty())
return false;
- int idxArg = 0;
- idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrReadOnly(buf));
- idxArg = kernelCol.set(idxArg, (int)(buf.step / buf.elemSize()));
- idxArg = kernelCol.set(idxArg, buf.cols);
- idxArg = kernelCol.set(idxArg, buf.rows);
+ k.args(ocl::KernelArg::PtrReadOnly(buf), (int)(buf.step / buf.elemSize()), buf.cols,
+ buf.rows, ocl::KernelArg::PtrWriteOnly(dst), (int)(dst.offset / dst.elemSize()),
+ (int)(dst.step / dst.elemSize()), dst.cols, dst.rows);
- idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
- idxArg = kernelCol.set(idxArg, (int)(dst.offset / dst.elemSize()));
- idxArg = kernelCol.set(idxArg, (int)(dst.step / dst.elemSize()));
- idxArg = kernelCol.set(idxArg, dst.cols);
- idxArg = kernelCol.set(idxArg, dst.rows);
-
- return kernelCol.run(2, globalsize, localsize, sync);
+ return k.run(2, globalsize, localsize, false);
}
const int optimizedSepFilterLocalSize = 16;
@@ -3473,12 +3443,14 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
String opts = cv::format("-D BLK_X=%d -D BLK_Y=%d -D RADIUSX=%d -D RADIUSY=%d%s%s"
" -D srcT=%s -D convertToWT=%s -D WT=%s -D dstT=%s -D convertToDstT=%s"
- " -D %s", (int)lt2[0], (int)lt2[1], _row_kernel.size().height / 2, _col_kernel.size().height / 2,
+ " -D %s -D srcT1=%s -D dstT1=%s -D cn=%d", (int)lt2[0], (int)lt2[1],
+ _row_kernel.size().height / 2, _col_kernel.size().height / 2,
ocl::kernelToStr(_row_kernel, CV_32F, "KERNEL_MATRIX_X").c_str(),
ocl::kernelToStr(_col_kernel, CV_32F, "KERNEL_MATRIX_Y").c_str(),
ocl::typeToStr(stype), ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), ocl::typeToStr(dtype),
- ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType]);
+ ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType],
+ ocl::typeToStr(sdepth), ocl::typeToStr(ddepth), cn);
ocl::Kernel k("sep_filter", ocl::imgproc::filterSep_singlePass_oclsrc, opts);
if (k.empty())
@@ -3529,10 +3501,13 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
if (ddepth < 0)
ddepth = sdepth;
- CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 &&
- imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) &&
- imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1),
- ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true)
+// printf("%d %d\n", imgSize.width, optimizedSepFilterLocalSize + (kernelX.rows >> 1));
+// printf("%d %d\n", imgSize.height, optimizedSepFilterLocalSize + (kernelY.rows >> 1));
+
+// CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 &&
+// imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) &&
+// imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1),
+// ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true)
UMat src = _src.getUMat();
Size srcWholeSize; Point srcOffset;
@@ -3546,12 +3521,12 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
Size srcSize = src.size();
Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1);
UMat buf; buf.create(bufSize, CV_MAKETYPE(CV_32F, cn));
- if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, false))
+ if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType))
return false;
_dst.create(srcSize, CV_MAKETYPE(ddepth, cn));
UMat dst = _dst.getUMat();
- return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y, false);
+ return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y);
}
#endif
diff --git a/modules/imgproc/src/opencl/filterSepCol.cl b/modules/imgproc/src/opencl/filterSepCol.cl
index 30a2221cf1..05717c6ad2 100644
--- a/modules/imgproc/src/opencl/filterSepCol.cl
+++ b/modules/imgproc/src/opencl/filterSepCol.cl
@@ -36,16 +36,6 @@
#define READ_TIMES_COL ((2*(RADIUSY+LSIZE1)-1)/LSIZE1)
#define RADIUS 1
-#if CN ==1
-#define ALIGN (((RADIUS)+3)>>2<<2)
-#elif CN==2
-#define ALIGN (((RADIUS)+1)>>1<<1)
-#elif CN==3
-#define ALIGN (((RADIUS)+3)>>2<<2)
-#elif CN==4
-#define ALIGN (RADIUS)
-#define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0)
-#endif
#define noconvert
@@ -65,16 +55,8 @@ The info above maybe obsolete.
#define DIG(a) a,
__constant float mat_kernel[] = { COEFF };
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
- (__global const GENTYPE_SRC * restrict src,
- const int src_step_in_pixel,
- const int src_whole_cols,
- const int src_whole_rows,
- __global GENTYPE_DST * dst,
- const int dst_offset_in_pixel,
- const int dst_step_in_pixel,
- const int dst_cols,
- const int dst_rows)
+__kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int src_whole_cols, int src_whole_rows,
+ __global dstT * dst, int dst_offset_in_pixel, int dst_step_in_pixel, int dst_cols, int dst_rows)
{
int x = get_global_id(0);
int y = get_global_id(1);
@@ -85,35 +67,35 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
int start_addr = mad24(y, src_step_in_pixel, x);
int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
- int i;
- GENTYPE_SRC sum, temp[READ_TIMES_COL];
- __local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
+ srcT sum, temp[READ_TIMES_COL];
+ __local srcT LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
- //read pixels from src
- for(i = 0;i>2<<2)
-#elif CN==2
-#define ALIGN (((RADIUS)+1)>>1<<1)
-#elif CN==3
-#define ALIGN (((RADIUS)+3)>>2<<2)
-#elif CN==4
-#define ALIGN (RADIUS)
-#endif
#ifdef BORDER_REPLICATE
-//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
+// BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr))
#endif
#ifdef BORDER_REFLECT
-//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
+// BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
#endif
#ifdef BORDER_REFLECT_101
-//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
+// BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
#endif
-//blur function does not support BORDER_WRAP
#ifdef BORDER_WRAP
-//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
+// BORDER_WRAP: cdefgh|abcdefgh|abcdefg
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
#endif
@@ -127,65 +115,56 @@
#endif //BORDER_CONSTANT
#endif //EXTRA_EXTRAPOLATION
-/**********************************************************************************
-These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur.
-Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle
-kernel must be in the center. ROI is not supported either.
-For channels =1,2,4, each kernels read 4 elements(not 4 pixels), and for channels =3,
-the kernel read 4 pixels, save them to LDS and read the data needed from LDS to
-calculate the result.
-The length of the convovle kernel supported is related to the LSIZE0 and the MAX size
-of LDS, which is HW related.
-For channels = 1,3 the RADIUS is no more than LSIZE0*2
-For channels = 2, the RADIUS is no more than LSIZE0
-For channels = 4, arbitary RADIUS is supported unless the LDS is not enough
-Niko
-6/29/2011
-The info above maybe obsolete.
-***********************************************************************************/
+#define noconvert
+
+#if cn != 3
+#define loadpix(addr) *(__global const srcT *)(addr)
+#define storepix(val, addr) *(__global dstT *)(addr) = val
+#define SRCSIZE ((int)sizeof(srcT))
+#define DSTSIZE ((int)sizeof(dstT))
+#else
+#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
+#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
+#define SRCSIZE ((int)sizeof(srcT1)*3)
+#define DSTSIZE ((int)sizeof(dstT1)*3)
+#endif
#define DIG(a) a,
__constant float mat_kernel[] = { COEFF };
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D0
- (__global uchar * restrict src,
- int src_step_in_pixel,
- int src_offset_x, int src_offset_y,
- int src_cols, int src_rows,
- int src_whole_cols, int src_whole_rows,
- __global float * dst,
- int dst_step_in_pixel,
- int dst_cols, int dst_rows,
- int radiusy)
+__kernel void row_filter_C1_D0(__global const uchar * src, int src_step_in_pixel, int src_offset_x, int src_offset_y,
+ int src_cols, int src_rows, int src_whole_cols, int src_whole_rows,
+ __global float * dst, int dst_step_in_pixel, int dst_cols, int dst_rows,
+ int radiusy)
{
int x = get_global_id(0)<<2;
int y = get_global_id(1);
int l_x = get_local_id(0);
int l_y = get_local_id(1);
- int start_x = x+src_offset_x - RADIUSX & 0xfffffffc;
+ int start_x = x + src_offset_x - RADIUSX & 0xfffffffc;
int offset = src_offset_x - RADIUSX & 3;
int start_y = y + src_offset_y - radiusy;
int start_addr = mad24(start_y, src_step_in_pixel, start_x);
- int i;
+
float4 sum;
uchar4 temp[READ_TIMES_ROW];
- __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1];
+ __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW * LSIZE0 + 1];
#ifdef BORDER_CONSTANT
int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
// read pixels from src
- for (i = 0; i < READ_TIMES_ROW; i++)
+ for (int i = 0; i < READ_TIMES_ROW; ++i)
{
- int current_addr = start_addr+i*LSIZE0*4;
- current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
- temp[i] = *(__global uchar4*)&src[current_addr];
+ int current_addr = mad24(i, LSIZE0 << 2, start_addr);
+ current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0;
+ temp[i] = *(__global const uchar4 *)&src[current_addr];
}
// judge if read out of boundary
#ifdef BORDER_ISOLATED
- for (i = 0; isrc_whole_cols)| (start_y<0) | (start_y >= src_whole_rows);
#endif
- int4 index[READ_TIMES_ROW];
- int4 addr;
+ int4 index[READ_TIMES_ROW], addr;
int s_y;
if (not_all_in_range)
{
// judge if read out of boundary
- for (i = 0; i < READ_TIMES_ROW; i++)
+ for (int i = 0; i < READ_TIMES_ROW; ++i)
{
- index[i] = (int4)(start_x+i*LSIZE0*4) + (int4)(0, 1, 2, 3);
+ index[i] = (int4)(mad24(i, LSIZE0 << 2, start_x)) + (int4)(0, 1, 2, 3);
#ifdef BORDER_ISOLATED
EXTRAPOLATE(index[i].x, src_offset_x, src_offset_x + src_cols);
EXTRAPOLATE(index[i].y, src_offset_x, src_offset_x + src_cols);
@@ -231,6 +209,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
EXTRAPOLATE(index[i].w, 0, src_whole_cols);
#endif
}
+
s_y = start_y;
#ifdef BORDER_ISOLATED
EXTRAPOLATE(s_y, src_offset_y, src_offset_y + src_rows);
@@ -239,9 +218,9 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
#endif
// read pixels from src
- for (i = 0; i 0)) ? current_addr : 0;
+ int current_addr = mad24(i, LSIZE0, start_addr);
+ current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0;
temp[i] = src[current_addr];
}
- //judge if read out of boundary
+ // judge if read out of boundary
#ifdef BORDER_ISOLATED
- for (i = 0; i 0)) ? current_addr : 0;
- temp[i] = src[current_addr];
- }
-
- // judge if read out of boundary
-#ifdef BORDER_ISOLATED
- for (i = 0; i 0)) ? current_addr : 0;
- temp[i] = src[current_addr];
- }
-
- // judge if read out of boundary
-#ifdef BORDER_ISOLATED
- for (i = 0; i
Date: Wed, 19 Mar 2014 17:30:13 +0400
Subject: [PATCH 24/47] Enabled Intel-specific optimizations for HOG detector.
---
modules/objdetect/src/hog.cpp | 16 +++++++++-------
modules/objdetect/src/opencl/objdetect_hog.cl | 18 +++++++++++++-----
2 files changed, 22 insertions(+), 12 deletions(-)
diff --git a/modules/objdetect/src/hog.cpp b/modules/objdetect/src/hog.cpp
index 18bb7afc22..0f4456ad51 100644
--- a/modules/objdetect/src/hog.cpp
+++ b/modules/objdetect/src/hog.cpp
@@ -1085,8 +1085,8 @@ static bool ocl_compute_gradients_8UC1(int height, int width, InputArray _img, f
size_t globalThreads[3] = { width, height, 1 };
char correctGamma = (correct_gamma) ? 1 : 0;
int grad_quadstep = (int)grad.step >> 3;
- int qangle_step_shift = 0;
- int qangle_step = (int)qangle.step >> (1 + qangle_step_shift);
+ int qangle_elem_size = CV_ELEM_SIZE1(qangle.type());
+ int qangle_step = (int)qangle.step / (2 * qangle_elem_size);
int idx = 0;
idx = k.set(idx, height);
@@ -1137,9 +1137,9 @@ static bool ocl_compute_hists(int nbins, int block_stride_x, int block_stride_y,
int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y)/block_stride_y;
int blocks_total = img_block_width * img_block_height;
- int qangle_step_shift = 0;
+ int qangle_elem_size = CV_ELEM_SIZE1(qangle.type());
int grad_quadstep = (int)grad.step >> 2;
- int qangle_step = (int)qangle.step >> qangle_step_shift;
+ int qangle_step = (int)qangle.step / qangle_elem_size;
int blocks_in_group = 4;
size_t localThreads[3] = { blocks_in_group * 24, 2, 1 };
@@ -1316,11 +1316,12 @@ static bool ocl_extract_descrs_by_cols(int win_height, int win_width, int block_
static bool ocl_compute(InputArray _img, Size win_stride, std::vector& _descriptors, int descr_format, Size blockSize,
Size cellSize, int nbins, Size blockStride, Size winSize, float sigma, bool gammaCorrection, double L2HysThreshold)
{
- Size imgSize = _img.size();
+ Size imgSize = _img.size();
Size effect_size = imgSize;
UMat grad(imgSize, CV_32FC2);
- UMat qangle(imgSize, CV_8UC2);
+ int qangle_type = ocl::Device::getDefault().isIntel() ? CV_32SC2 : CV_8UC2;
+ UMat qangle(imgSize, qangle_type);
const size_t block_hist_size = getBlockHistogramSize(blockSize, cellSize, nbins);
const Size blocks_per_img = numPartsWithin(imgSize, blockSize, blockStride);
@@ -1720,7 +1721,8 @@ static bool ocl_detect(InputArray img, std::vector &hits, double hit_thre
Size imgSize = img.size();
Size effect_size = imgSize;
UMat grad(imgSize, CV_32FC2);
- UMat qangle(imgSize, CV_8UC2);
+ int qangle_type = ocl::Device::getDefault().isIntel() ? CV_32SC2 : CV_8UC2;
+ UMat qangle(imgSize, qangle_type);
const size_t block_hist_size = getBlockHistogramSize(blockSize, cellSize, nbins);
const Size blocks_per_img = numPartsWithin(imgSize, blockSize, blockStride);
diff --git a/modules/objdetect/src/opencl/objdetect_hog.cl b/modules/objdetect/src/opencl/objdetect_hog.cl
index 082f9ab7fb..5c71aa1b45 100644
--- a/modules/objdetect/src/opencl/objdetect_hog.cl
+++ b/modules/objdetect/src/opencl/objdetect_hog.cl
@@ -50,6 +50,14 @@
#define NTHREADS 256
#define CV_PI_F 3.1415926535897932384626433832795f
+#ifdef INTEL_DEVICE
+#define QANGLE_TYPE int
+#define QANGLE_TYPE2 int2
+#else
+#define QANGLE_TYPE uchar
+#define QANGLE_TYPE2 uchar2
+#endif
+
//----------------------------------------------------------------------------
// Histogram computation
// 12 threads for a cell, 12x4 threads per block
@@ -59,7 +67,7 @@ __kernel void compute_hists_lut_kernel(
const int cnbins, const int cblock_hist_size, const int img_block_width,
const int blocks_in_group, const int blocks_total,
const int grad_quadstep, const int qangle_step,
- __global const float* grad, __global const uchar* qangle,
+ __global const float* grad, __global const QANGLE_TYPE* qangle,
__global const float* gauss_w_lut,
__global float* block_hists, __local float* smem)
{
@@ -86,7 +94,7 @@ __kernel void compute_hists_lut_kernel(
__global const float* grad_ptr = (gid < blocks_total) ?
grad + offset_y * grad_quadstep + (offset_x << 1) : grad;
- __global const uchar* qangle_ptr = (gid < blocks_total) ?
+ __global const QANGLE_TYPE* qangle_ptr = (gid < blocks_total) ?
qangle + offset_y * qangle_step + (offset_x << 1) : qangle;
__local float* hist = hists + 12 * (cell_y * CELLS_PER_BLOCK_Y + cell_x) +
@@ -101,7 +109,7 @@ __kernel void compute_hists_lut_kernel(
for (int dist_y = dist_y_begin; dist_y < dist_y_begin + 12; ++dist_y)
{
float2 vote = (float2) (grad_ptr[0], grad_ptr[1]);
- uchar2 bin = (uchar2) (qangle_ptr[0], qangle_ptr[1]);
+ QANGLE_TYPE2 bin = (QANGLE_TYPE2) (qangle_ptr[0], qangle_ptr[1]);
grad_ptr += grad_quadstep;
qangle_ptr += qangle_step;
@@ -558,7 +566,7 @@ __kernel void extract_descrs_by_cols_kernel(
__kernel void compute_gradients_8UC4_kernel(
const int height, const int width,
const int img_step, const int grad_quadstep, const int qangle_step,
- const __global uchar4 * img, __global float * grad, __global uchar * qangle,
+ const __global uchar4 * img, __global float * grad, __global QANGLE_TYPE * qangle,
const float angle_scale, const char correct_gamma, const int cnbins)
{
const int x = get_global_id(0);
@@ -660,7 +668,7 @@ __kernel void compute_gradients_8UC4_kernel(
__kernel void compute_gradients_8UC1_kernel(
const int height, const int width,
const int img_step, const int grad_quadstep, const int qangle_step,
- __global const uchar * img, __global float * grad, __global uchar * qangle,
+ __global const uchar * img, __global float * grad, __global QANGLE_TYPE * qangle,
const float angle_scale, const char correct_gamma, const int cnbins)
{
const int x = get_global_id(0);
From 284b2fc1e735886c569685333d558f624ce58719 Mon Sep 17 00:00:00 2001
From: Alexander Smorkalov
Date: Fri, 28 Feb 2014 18:18:20 +0400
Subject: [PATCH 25/47] Cut path to CUDA libraries to prevent generation of
OpenCVModules.cmake with abs path.
---
CMakeLists.txt | 11 ---------
cmake/OpenCVDetectAndroidSDK.cmake | 6 ++---
cmake/OpenCVDetectCUDA.cmake | 39 ++++++++++++++++++++++++++++++
3 files changed, 42 insertions(+), 14 deletions(-)
diff --git a/CMakeLists.txt b/CMakeLists.txt
index fb49497412..747d207e36 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -467,7 +467,6 @@ include(cmake/OpenCVFindLibsGUI.cmake)
include(cmake/OpenCVFindLibsVideo.cmake)
include(cmake/OpenCVFindLibsPerf.cmake)
-
# ----------------------------------------------------------------------------
# Detect other 3rd-party libraries/tools
# ----------------------------------------------------------------------------
@@ -513,16 +512,6 @@ if(NOT HAVE_CUDA)
set(ENABLE_DYNAMIC_CUDA OFF)
endif()
-if(HAVE_CUDA AND NOT ENABLE_DYNAMIC_CUDA)
- set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
- if(HAVE_CUBLAS)
- set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} ${CUDA_cublas_LIBRARY})
- endif()
- if(HAVE_CUFFT)
- set(OPENCV_LINKER_LIBS ${OPENCV_LINKER_LIBS} ${CUDA_cufft_LIBRARY})
- endif()
-endif()
-
# ----------------------------------------------------------------------------
# Solution folders:
# ----------------------------------------------------------------------------
diff --git a/cmake/OpenCVDetectAndroidSDK.cmake b/cmake/OpenCVDetectAndroidSDK.cmake
index af7427194e..273758967c 100644
--- a/cmake/OpenCVDetectAndroidSDK.cmake
+++ b/cmake/OpenCVDetectAndroidSDK.cmake
@@ -326,12 +326,12 @@ macro(add_android_project target path)
# copy all needed CUDA libs to project if EMBED_CUDA flag is present
if(android_proj_EMBED_CUDA)
- set(android_proj_culibs ${CUDA_npp_LIBRARY} ${CUDA_LIBRARIES})
+ set(android_proj_culibs ${CUDA_npp_LIBRARY_ABS} ${CUDA_LIBRARIES_ABS})
if(HAVE_CUFFT)
- list(INSERT android_proj_culibs 0 ${CUDA_cufft_LIBRARY})
+ list(INSERT android_proj_culibs 0 ${CUDA_cufft_LIBRARY_ABS})
endif()
if(HAVE_CUBLAS)
- list(INSERT android_proj_culibs 0 ${CUDA_cublas_LIBRARY})
+ list(INSERT android_proj_culibs 0 ${CUDA_cublas_LIBRARY_ABS})
endif()
foreach(lib ${android_proj_culibs})
get_filename_component(f "${lib}" NAME)
diff --git a/cmake/OpenCVDetectCUDA.cmake b/cmake/OpenCVDetectCUDA.cmake
index 56b142970e..24fbb03ce3 100644
--- a/cmake/OpenCVDetectCUDA.cmake
+++ b/cmake/OpenCVDetectCUDA.cmake
@@ -219,3 +219,42 @@ else()
unset(CUDA_ARCH_BIN CACHE)
unset(CUDA_ARCH_PTX CACHE)
endif()
+
+if(HAVE_CUDA)
+ set(CUDA_LIBS_PATH "")
+ foreach(p ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
+ get_filename_component(_tmp ${p} PATH)
+ list(APPEND CUDA_LIBS_PATH ${_tmp})
+ endforeach()
+
+ if(HAVE_CUBLAS)
+ foreach(p ${CUDA_cublas_LIBRARY})
+ get_filename_component(_tmp ${p} PATH)
+ list(APPEND CUDA_LIBS_PATH ${_tmp})
+ endforeach()
+ endif()
+
+ if(HAVE_CUFFT)
+ foreach(p ${CUDA_cufft_LIBRARY})
+ get_filename_component(_tmp ${p} PATH)
+ list(APPEND CUDA_LIBS_PATH ${_tmp})
+ endforeach()
+ endif()
+
+ list(REMOVE_DUPLICATES CUDA_LIBS_PATH)
+ link_directories(${CUDA_LIBS_PATH})
+
+ set(CUDA_LIBRARIES_ABS ${CUDA_LIBRARIES})
+ ocv_convert_to_lib_name(CUDA_LIBRARIES ${CUDA_LIBRARIES})
+ set(CUDA_npp_LIBRARY_ABS ${CUDA_npp_LIBRARY})
+ ocv_convert_to_lib_name(CUDA_npp_LIBRARY ${CUDA_npp_LIBRARY})
+ if(HAVE_CUBLAS)
+ set(CUDA_cublas_LIBRARY_ABS ${CUDA_cublas_LIBRARY})
+ ocv_convert_to_lib_name(CUDA_cublas_LIBRARY ${CUDA_cublas_LIBRARY})
+ endif()
+
+ if(HAVE_CUFFT)
+ set(CUDA_cufft_LIBRARY_ABS ${CUDA_cufft_LIBRARY})
+ ocv_convert_to_lib_name(CUDA_cufft_LIBRARY ${CUDA_cufft_LIBRARY})
+ endif()
+endif()
\ No newline at end of file
From 291458a8599d513eff1729ef8c0fcab89187621d Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Wed, 19 Mar 2014 18:49:33 +0400
Subject: [PATCH 26/47] generalized OpenCL version of cv::sepFilter2D; removed
some restrictions and added 3-channels support
---
modules/core/src/ocl.cpp | 4 +-
modules/imgproc/src/filter.cpp | 86 +++++++++++--------
modules/imgproc/src/opencl/filterSepCol.cl | 47 +++++-----
modules/imgproc/src/opencl/filterSepRow.cl | 45 ++++++----
modules/imgproc/test/ocl/test_filters.cpp | 2 +-
modules/imgproc/test/ocl/test_sepfilter2D.cpp | 17 ++--
6 files changed, 112 insertions(+), 89 deletions(-)
diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index b56f84c16e..3a7c718d4f 100644
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -4317,8 +4317,8 @@ String kernelToStr(InputArray _kernel, int ddepth, const char * name)
if (ddepth != depth)
kernel.convertTo(kernel, ddepth);
- typedef std::string (*func_t)(const Mat &);
- static const func_t funcs[] = { kerToStr, kerToStr, kerToStr,kerToStr,
+ typedef std::string (* func_t)(const Mat &);
+ static const func_t funcs[] = { kerToStr, kerToStr, kerToStr, kerToStr,
kerToStr, kerToStr, kerToStr, 0 };
const func_t func = funcs[depth];
CV_Assert(func != 0);
diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp
index ba2e347af0..c013a9b16c 100644
--- a/modules/imgproc/src/filter.cpp
+++ b/modules/imgproc/src/filter.cpp
@@ -41,7 +41,6 @@
//M*/
#include "precomp.hpp"
-#define CV_OPENCL_RUN_ASSERT
#include "opencl_kernels.hpp"
#include
@@ -3135,7 +3134,7 @@ template struct Filter2D : public BaseFi
// b e h b e h 0 0
// c f i c f i 0 0
template
-static int _prepareKernelFilter2D(std::vector& data, const Mat &kernel)
+static int _prepareKernelFilter2D(std::vector & data, const Mat & kernel)
{
Mat _kernel; kernel.convertTo(_kernel, DataDepth::value);
int size_y_aligned = ROUNDUP(kernel.rows * 2, 4);
@@ -3318,11 +3317,16 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
return kernel.run(2, globalsize, localsize, true);
}
-static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType)
+static bool ocl_sepRowFilter2D(const UMat & src, UMat & buf, const Mat & kernelX, int anchor,
+ int borderType, int ddepth, bool fast8uc1)
{
int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type);
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
Size bufSize = buf.size();
+ if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F))
+ return false;
+
#ifdef ANDROID
size_t localsize[2] = {16, 10};
#else
@@ -3330,7 +3334,7 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor,
#endif
size_t globalsize[2] = {DIVUP(bufSize.width, localsize[0]) * localsize[0], DIVUP(bufSize.height, localsize[1]) * localsize[1]};
- if (type == CV_8UC1)
+ if (fast8uc1)
globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0];
int radiusX = anchor, radiusY = (buf.rows - src.rows) >> 1;
@@ -3346,20 +3350,21 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor,
char cvt[40];
cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s"
- " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s",
+ " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s%s",
radiusX, (int)localsize[0], (int)localsize[1], cn, btype,
extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED",
ocl::typeToStr(type), ocl::typeToStr(CV_32FC(cn)),
ocl::convertTypeStr(sdepth, CV_32F, cn, cvt),
- ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F));
+ ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F),
+ doubleSupport ? " -D DOUBLE_SUPPORT" : "");
build_options += ocl::kernelToStr(kernelX, CV_32F);
Size srcWholeSize; Point srcOffset;
src.locateROI(srcWholeSize, srcOffset);
String kernelName("row_filter");
- if (type == CV_8UC1)
+ if (fast8uc1)
kernelName += "_C1_D0";
ocl::Kernel k(kernelName.c_str(), cv::ocl::imgproc::filterSepRow_oclsrc,
@@ -3367,39 +3372,47 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor,
if (k.empty())
return false;
- k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x,
- srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height,
- ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()),
- buf.cols, buf.rows, radiusY);
+ if (fast8uc1)
+ k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x,
+ srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height,
+ ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()),
+ buf.cols, buf.rows, radiusY);
+ else
+ k.args(ocl::KernelArg::PtrReadOnly(src), (int)src.step, srcOffset.x,
+ srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height,
+ ocl::KernelArg::PtrWriteOnly(buf), (int)buf.step, buf.cols, buf.rows, radiusY);
return k.run(2, globalsize, localsize, false);
}
-static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor)
+static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, int anchor)
{
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+ if (dst.depth() == CV_64F && !doubleSupport)
+ return false;
+
#ifdef ANDROID
- size_t localsize[2] = {16, 10};
+ size_t localsize[2] = { 16, 10 };
#else
- size_t localsize[2] = {16, 16};
+ size_t localsize[2] = { 16, 16 };
#endif
- size_t globalsize[2] = {0, 0};
+ size_t globalsize[2] = { 0, 0 };
int dtype = dst.type(), cn = CV_MAT_CN(dtype), ddepth = CV_MAT_DEPTH(dtype);
Size sz = dst.size();
globalsize[1] = DIVUP(sz.height, localsize[1]) * localsize[1];
-
- if (dtype == CV_8UC2)
- globalsize[0] = DIVUP((sz.width + 1) / 2, localsize[0]) * localsize[0];
- else
- globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0];
+ globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0];
char cvt[40];
cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d"
- " -D srcT=%s -D dstT=%s -D convertToDstT=%s",
+ " -D srcT=%s -D dstT=%s -D convertToDstT=%s"
+ " -D srcT1=%s -D dstT1=%s%s",
anchor, (int)localsize[0], (int)localsize[1], cn,
ocl::typeToStr(buf.type()), ocl::typeToStr(dtype),
- ocl::convertTypeStr(CV_32F, ddepth, cn, cvt));
+ ocl::convertTypeStr(CV_32F, ddepth, cn, cvt),
+ ocl::typeToStr(CV_32F), ocl::typeToStr(ddepth),
+ doubleSupport ? " -D DOUBLE_SUPPORT" : "");
build_options += ocl::kernelToStr(kernelY, CV_32F);
ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc,
@@ -3407,13 +3420,13 @@ static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anc
if (k.empty())
return false;
- k.args(ocl::KernelArg::PtrReadOnly(buf), (int)(buf.step / buf.elemSize()), buf.cols,
- buf.rows, ocl::KernelArg::PtrWriteOnly(dst), (int)(dst.offset / dst.elemSize()),
- (int)(dst.step / dst.elemSize()), dst.cols, dst.rows);
+ k.args(ocl::KernelArg::ReadOnly(buf), ocl::KernelArg::WriteOnly(dst));
return k.run(2, globalsize, localsize, false);
}
+#if 0
+
const int optimizedSepFilterLocalSize = 16;
static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
@@ -3471,18 +3484,19 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
return k.run(2, gt2, lt2, false);
}
+#endif
+
static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
InputArray _kernelX, InputArray _kernelY, Point anchor,
double delta, int borderType )
{
- Size imgSize = _src.size();
+// Size imgSize = _src.size();
if (abs(delta)> FLT_MIN)
return false;
int type = _src.type(), cn = CV_MAT_CN(type);
- if ( !( (type == CV_8UC1 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC4) &&
- (ddepth == CV_32F || ddepth == CV_16S || ddepth == CV_8U || ddepth < 0) ) )
+ if (cn > 4)
return false;
Mat kernelX = _kernelX.getMat().reshape(1, 1);
@@ -3501,9 +3515,6 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
if (ddepth < 0)
ddepth = sdepth;
-// printf("%d %d\n", imgSize.width, optimizedSepFilterLocalSize + (kernelX.rows >> 1));
-// printf("%d %d\n", imgSize.height, optimizedSepFilterLocalSize + (kernelY.rows >> 1));
-
// CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 &&
// imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) &&
// imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1),
@@ -3512,20 +3523,19 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
UMat src = _src.getUMat();
Size srcWholeSize; Point srcOffset;
src.locateROI(srcWholeSize, srcOffset);
- if ( (0 != (srcOffset.x % 4)) ||
- (0 != (src.cols % 4)) ||
- (0 != ((src.step / src.elemSize()) % 4))
- )
- return false;
+
+ bool fast8uc1 = type == CV_8UC1 && srcOffset.x % 4 == 0 &&
+ src.cols % 4 == 0 && src.step % 4 == 0;
Size srcSize = src.size();
Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1);
- UMat buf; buf.create(bufSize, CV_MAKETYPE(CV_32F, cn));
- if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType))
+ UMat buf(bufSize, CV_32FC(cn));
+ if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, ddepth, fast8uc1))
return false;
_dst.create(srcSize, CV_MAKETYPE(ddepth, cn));
UMat dst = _dst.getUMat();
+
return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y);
}
diff --git a/modules/imgproc/src/opencl/filterSepCol.cl b/modules/imgproc/src/opencl/filterSepCol.cl
index 05717c6ad2..f5d270cf42 100644
--- a/modules/imgproc/src/opencl/filterSepCol.cl
+++ b/modules/imgproc/src/opencl/filterSepCol.cl
@@ -34,29 +34,36 @@
//
//
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
#define READ_TIMES_COL ((2*(RADIUSY+LSIZE1)-1)/LSIZE1)
#define RADIUS 1
#define noconvert
-/**********************************************************************************
-These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur.
-Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle
-kernel must be in the center. ROI is not supported either.
-Each kernels read 4 elements(not 4 pixels), save them to LDS and read the data needed
-from LDS to calculate the result.
-The length of the convovle kernel supported is only related to the MAX size of LDS,
-which is HW related.
-Niko
-6/29/2011
-The info above maybe obsolete.
-***********************************************************************************/
+#if CN != 3
+#define loadpix(addr) *(__global const srcT *)(addr)
+#define storepix(val, addr) *(__global dstT *)(addr) = val
+#define SRCSIZE (int)sizeof(srcT)
+#define DSTSIZE (int)sizeof(dstT)
+#else
+#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
+#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
+#define SRCSIZE (int)sizeof(srcT1)*3
+#define DSTSIZE (int)sizeof(dstT1)*3
+#endif
#define DIG(a) a,
__constant float mat_kernel[] = { COEFF };
-__kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int src_whole_cols, int src_whole_rows,
- __global dstT * dst, int dst_offset_in_pixel, int dst_step_in_pixel, int dst_cols, int dst_rows)
+__kernel void col_filter(__global const uchar * src, int src_step, int src_offset, int src_whole_rows, int src_whole_cols,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
int x = get_global_id(0);
int y = get_global_id(1);
@@ -64,8 +71,8 @@ __kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int s
int l_x = get_local_id(0);
int l_y = get_local_id(1);
- int start_addr = mad24(y, src_step_in_pixel, x);
- int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
+ int start_addr = mad24(y, src_step, x * SRCSIZE);
+ int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * SRCSIZE);
srcT sum, temp[READ_TIMES_COL];
__local srcT LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
@@ -73,9 +80,9 @@ __kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int s
// read pixels from src
for (int i = 0; i < READ_TIMES_COL; ++i)
{
- int current_addr = mad24(i, LSIZE1 * src_step_in_pixel, start_addr);
+ int current_addr = mad24(i, LSIZE1 * src_step, start_addr);
current_addr = current_addr < end_addr ? current_addr : 0;
- temp[i] = src[current_addr];
+ temp[i] = loadpix(src + current_addr);
}
// save pixels to lds
@@ -95,7 +102,7 @@ __kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int s
// write the result to dst
if (x < dst_cols && y < dst_rows)
{
- start_addr = mad24(y, dst_step_in_pixel, x + dst_offset_in_pixel);
- dst[start_addr] = convertToDstT(sum);
+ start_addr = mad24(y, dst_step, mad24(DSTSIZE, x, dst_offset));
+ storepix(convertToDstT(sum), dst + start_addr);
}
}
diff --git a/modules/imgproc/src/opencl/filterSepRow.cl b/modules/imgproc/src/opencl/filterSepRow.cl
index 83968dfc1e..726de448e4 100644
--- a/modules/imgproc/src/opencl/filterSepRow.cl
+++ b/modules/imgproc/src/opencl/filterSepRow.cl
@@ -34,6 +34,14 @@
//
//
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
#define READ_TIMES_ROW ((2*(RADIUSX+LSIZE0)-1)/LSIZE0) //for c4 only
#define RADIUS 1
@@ -117,16 +125,16 @@
#define noconvert
-#if cn != 3
+#if CN != 3
#define loadpix(addr) *(__global const srcT *)(addr)
#define storepix(val, addr) *(__global dstT *)(addr) = val
-#define SRCSIZE ((int)sizeof(srcT))
-#define DSTSIZE ((int)sizeof(dstT))
+#define SRCSIZE (int)sizeof(srcT)
+#define DSTSIZE (int)sizeof(dstT)
#else
#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
-#define SRCSIZE ((int)sizeof(srcT1)*3)
-#define DSTSIZE ((int)sizeof(dstT1)*3)
+#define SRCSIZE (int)sizeof(srcT1)*3
+#define DSTSIZE (int)sizeof(dstT1)*3
#endif
#define DIG(a) a,
@@ -269,32 +277,33 @@ __kernel void row_filter_C1_D0(__global const uchar * src, int src_step_in_pixel
dst[start_addr] = sum.x;
}
-__kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int src_offset_x, int src_offset_y,
+__kernel void row_filter(__global const uchar * src, int src_step, int src_offset_x, int src_offset_y,
int src_cols, int src_rows, int src_whole_cols, int src_whole_rows,
- __global dstT * dst, int dst_step_in_pixel, int dst_cols, int dst_rows,
+ __global uchar * dst, int dst_step, int dst_cols, int dst_rows,
int radiusy)
{
int x = get_global_id(0);
int y = get_global_id(1);
int l_x = get_local_id(0);
int l_y = get_local_id(1);
+
int start_x = x + src_offset_x - RADIUSX;
int start_y = y + src_offset_y - radiusy;
- int start_addr = mad24(start_y, src_step_in_pixel, start_x);
+ int start_addr = mad24(start_y, src_step, start_x * SRCSIZE);
dstT sum;
srcT temp[READ_TIMES_ROW];
__local srcT LDS_DAT[LSIZE1][READ_TIMES_ROW * LSIZE0 + 1];
#ifdef BORDER_CONSTANT
- int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
+ int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * SRCSIZE);
// read pixels from src
for (int i = 0; i < READ_TIMES_ROW; i++)
{
- int current_addr = mad24(i, LSIZE0, start_addr);
- current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0;
- temp[i] = src[current_addr];
+ int current_addr = mad24(i, LSIZE0 * SRCSIZE, start_addr);
+ current_addr = current_addr < end_addr && current_addr >= 0 ? current_addr : 0;
+ temp[i] = loadpix(src + current_addr);
}
// judge if read out of boundary
@@ -312,8 +321,7 @@ __kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int s
}
#endif
#else
- int index[READ_TIMES_ROW];
- int s_x, s_y;
+ int index[READ_TIMES_ROW], s_x, s_y;
// judge if read out of boundary
for (int i = 0; i < READ_TIMES_ROW; ++i)
@@ -328,12 +336,12 @@ __kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int s
EXTRAPOLATE(s_x, 0, src_whole_cols);
EXTRAPOLATE(s_y, 0, src_whole_rows);
#endif
- index[i] = mad24(s_y, src_step_in_pixel, s_x);
+ index[i] = mad24(s_y, src_step, s_x * SRCSIZE);
}
// read pixels from src
for (int i = 0; i < READ_TIMES_ROW; ++i)
- temp[i] = src[index[i]];
+ temp[i] = loadpix(src + index[i]);
#endif // BORDER_CONSTANT
// save pixels to lds
@@ -349,10 +357,11 @@ __kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int s
temp[1] = LDS_DAT[l_y][l_x + RADIUSX + i];
sum += mad(convertToDstT(temp[0]), mat_kernel[RADIUSX - i], convertToDstT(temp[1]) * mat_kernel[RADIUSX + i]);
}
+
// write the result to dst
if (x < dst_cols && y < dst_rows)
{
- start_addr = mad24(y, dst_step_in_pixel, x);
- dst[start_addr] = sum;
+ start_addr = mad24(y, dst_step, x * DSTSIZE);
+ storepix(sum, dst + start_addr);
}
}
diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp
index fe16fe81d5..04b330527f 100644
--- a/modules/imgproc/test/ocl/test_filters.cpp
+++ b/modules/imgproc/test/ocl/test_filters.cpp
@@ -312,7 +312,7 @@ OCL_TEST_P(MorphologyEx, Mat)
(int)BORDER_REFLECT|BORDER_ISOLATED, (int)BORDER_WRAP|BORDER_ISOLATED, \
(int)BORDER_REFLECT_101|BORDER_ISOLATED*/) // WRAP and ISOLATED are not supported by cv:: version
-#define FILTER_TYPES Values(CV_8UC1, CV_8UC2, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4)
+#define FILTER_TYPES Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4)
OCL_INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine(
Values((MatType)CV_8UC1),
diff --git a/modules/imgproc/test/ocl/test_sepfilter2D.cpp b/modules/imgproc/test/ocl/test_sepfilter2D.cpp
index 09d01d157a..05d46cc2b3 100644
--- a/modules/imgproc/test/ocl/test_sepfilter2D.cpp
+++ b/modules/imgproc/test/ocl/test_sepfilter2D.cpp
@@ -75,9 +75,9 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool)
void random_roi()
{
Size ksize = randomSize(kernelMinSize, kernelMaxSize);
- if (1 != (ksize.width % 2))
+ if (1 != ksize.width % 2)
ksize.width++;
- if (1 != (ksize.height % 2))
+ if (1 != ksize.height % 2)
ksize.height++;
Mat temp = randomMat(Size(ksize.width, 1), CV_MAKE_TYPE(CV_32F, 1), -MAX_VALUE, MAX_VALUE);
@@ -86,24 +86,22 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool)
cv::normalize(temp, kernelY, 1.0, 0.0, NORM_L1);
Size roiSize = randomSize(ksize.width + 16, MAX_VALUE, ksize.height + 20, MAX_VALUE);
- std::cout << roiSize << std::endl;
int rest = roiSize.width % 4;
- if (0 != rest)
+ if (rest != 0)
roiSize.width += (4 - rest);
Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
rest = srcBorder.lef % 4;
- if (0 != rest)
+ if (rest != 0)
srcBorder.lef += (4 - rest);
rest = srcBorder.rig % 4;
- if (0 != rest)
+ if (rest != 0)
srcBorder.rig += (4 - rest);
randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
- anchor.x = -1;
- anchor.y = -1;
+ anchor.x = anchor.y = -1;
UMAT_UPLOAD_INPUT_PARAMETER(src)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
@@ -128,11 +126,10 @@ OCL_TEST_P(SepFilter2D, Mat)
}
}
-
OCL_INSTANTIATE_TEST_CASE_P(ImageProc, SepFilter2D,
Combine(
Values(CV_8U, CV_32F),
- Values(1, 4),
+ OCL_ALL_CHANNELS,
Values(
(BorderType)BORDER_CONSTANT,
(BorderType)BORDER_REPLICATE,
From b14c314fc3a18bd204364c7a4c5dec8fe218bdf0 Mon Sep 17 00:00:00 2001
From: Alexander Karsakov
Date: Wed, 19 Mar 2014 17:33:13 +0400
Subject: [PATCH 27/47] Fixed incorrect thread synchronizations
---
modules/objdetect/src/opencl/objdetect_hog.cl | 3 +--
modules/objdetect/test/opencl/test_hogdetector.cpp | 2 +-
2 files changed, 2 insertions(+), 3 deletions(-)
diff --git a/modules/objdetect/src/opencl/objdetect_hog.cl b/modules/objdetect/src/opencl/objdetect_hog.cl
index 5c71aa1b45..704dec4447 100644
--- a/modules/objdetect/src/opencl/objdetect_hog.cl
+++ b/modules/objdetect/src/opencl/objdetect_hog.cl
@@ -141,9 +141,8 @@ __kernel void compute_hists_lut_kernel(
final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] =
hist_[0] + hist_[1] + hist_[2];
}
-#ifdef CPU
+
barrier(CLK_LOCAL_MEM_FENCE);
-#endif
int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x;
if ((tid < cblock_hist_size) && (gid < blocks_total))
diff --git a/modules/objdetect/test/opencl/test_hogdetector.cpp b/modules/objdetect/test/opencl/test_hogdetector.cpp
index 8568352b69..b3ef6b48fb 100644
--- a/modules/objdetect/test/opencl/test_hogdetector.cpp
+++ b/modules/objdetect/test/opencl/test_hogdetector.cpp
@@ -110,7 +110,7 @@ OCL_TEST_P(HOG, Detect)
OCL_OFF(hog.detectMultiScale(img, cpu_found, 0, Size(8, 8), Size(0, 0), 1.05, 6));
OCL_ON(hog.detectMultiScale(uimg, gpu_found, 0, Size(8, 8), Size(0, 0), 1.05, 6));
- EXPECT_LT(checkRectSimilarity(img.size(), cpu_found, gpu_found), 1.0);
+ EXPECT_LT(checkRectSimilarity(img.size(), cpu_found, gpu_found), 0.05);
}
INSTANTIATE_TEST_CASE_P(OCL_ObjDetect, HOG, testing::Combine(
From 63d8a61b9b18db161e0ff94fc7bac0166f0b8e3b Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov
Date: Wed, 19 Mar 2014 19:12:37 +0400
Subject: [PATCH 28/47] enabled 3-channels support for
cv::createSuperResolution_BTVL1
---
modules/imgproc/src/filter.cpp | 15 ++++++---------
modules/imgproc/test/ocl/test_sepfilter2D.cpp | 4 ++--
modules/superres/src/btv_l1.cpp | 4 +---
3 files changed, 9 insertions(+), 14 deletions(-)
diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp
index c013a9b16c..2d7c7404f6 100644
--- a/modules/imgproc/src/filter.cpp
+++ b/modules/imgproc/src/filter.cpp
@@ -3425,8 +3425,6 @@ static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY
return k.run(2, globalsize, localsize, false);
}
-#if 0
-
const int optimizedSepFilterLocalSize = 16;
static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
@@ -3484,13 +3482,11 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
return k.run(2, gt2, lt2, false);
}
-#endif
-
static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
InputArray _kernelX, InputArray _kernelY, Point anchor,
double delta, int borderType )
{
-// Size imgSize = _src.size();
+ Size imgSize = _src.size();
if (abs(delta)> FLT_MIN)
return false;
@@ -3515,10 +3511,11 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
if (ddepth < 0)
ddepth = sdepth;
-// CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 &&
-// imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) &&
-// imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1),
-// ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true)
+ CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 &&
+ imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) &&
+ imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1) &&
+ (borderType & BORDER_ISOLATED) != 0,
+ ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true)
UMat src = _src.getUMat();
Size srcWholeSize; Point srcOffset;
diff --git a/modules/imgproc/test/ocl/test_sepfilter2D.cpp b/modules/imgproc/test/ocl/test_sepfilter2D.cpp
index 05d46cc2b3..b724641f45 100644
--- a/modules/imgproc/test/ocl/test_sepfilter2D.cpp
+++ b/modules/imgproc/test/ocl/test_sepfilter2D.cpp
@@ -85,7 +85,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool)
temp = randomMat(Size(1, ksize.height), CV_MAKE_TYPE(CV_32F, 1), -MAX_VALUE, MAX_VALUE);
cv::normalize(temp, kernelY, 1.0, 0.0, NORM_L1);
- Size roiSize = randomSize(ksize.width + 16, MAX_VALUE, ksize.height + 20, MAX_VALUE);
+ Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE);
int rest = roiSize.width % 4;
if (rest != 0)
roiSize.width += (4 - rest);
@@ -115,7 +115,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool)
OCL_TEST_P(SepFilter2D, Mat)
{
- for (int j = 0; j < test_loop_times; j++)
+ for (int j = 0; j < test_loop_times + 1; j++)
{
random_roi();
diff --git a/modules/superres/src/btv_l1.cpp b/modules/superres/src/btv_l1.cpp
index 1e4aa48a7d..d54b4b398a 100644
--- a/modules/superres/src/btv_l1.cpp
+++ b/modules/superres/src/btv_l1.cpp
@@ -1014,10 +1014,8 @@ namespace
return;
#ifdef HAVE_OPENCL
- if (isUmat_ && curFrame_.channels() == 1)
+ if (isUmat_)
curFrame_.copyTo(ucurFrame_);
- else
- isUmat_ = false;
#endif
++storePos_;
From 80a40ae3d7fd2f932addda446f0790e250f29de5 Mon Sep 17 00:00:00 2001
From: mlyashko
Date: Thu, 20 Mar 2014 16:15:43 +0400
Subject: [PATCH 29/47] changed epsilon for test pass on Win32
---
modules/video/test/test_tvl1optflow.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/modules/video/test/test_tvl1optflow.cpp b/modules/video/test/test_tvl1optflow.cpp
index 804eae8b62..274c13e65d 100644
--- a/modules/video/test/test_tvl1optflow.cpp
+++ b/modules/video/test/test_tvl1optflow.cpp
@@ -133,14 +133,13 @@ namespace
}
}
}
-
return sqrt(sum / (1e-9 + counter));
}
}
TEST(Video_calcOpticalFlowDual_TVL1, Regression)
{
- const double MAX_RMSE = 0.02;
+ const double MAX_RMSE = 0.03;
const string frame1_path = TS::ptr()->get_data_path() + "optflow/RubberWhale1.png";
const string frame2_path = TS::ptr()->get_data_path() + "optflow/RubberWhale2.png";
From d060d30fa0ee978a3356f619911fdc078e501d45 Mon Sep 17 00:00:00 2001
From: Andrey Pavlenko
Date: Thu, 20 Mar 2014 21:57:34 +0400
Subject: [PATCH 30/47] enabling OCL LBP branch for all devices
---
modules/objdetect/src/cascadedetect.cpp | 7 ++-----
1 file changed, 2 insertions(+), 5 deletions(-)
diff --git a/modules/objdetect/src/cascadedetect.cpp b/modules/objdetect/src/cascadedetect.cpp
index 3f0e6e38ce..2d5c0795dc 100644
--- a/modules/objdetect/src/cascadedetect.cpp
+++ b/modules/objdetect/src/cascadedetect.cpp
@@ -765,11 +765,8 @@ bool LBPEvaluator::read( const FileNode& node, Size _origWinSize )
nchannels = 1;
localSize = lbufSize = Size(0, 0);
if (ocl::haveOpenCL())
- {
- const ocl::Device& device = ocl::Device::getDefault();
- if (device.isAMD() && !device.hostUnifiedMemory())
- localSize = Size(8, 8);
- }
+ localSize = Size(8, 8);
+
return true;
}
From 640e180efe645c32863af549b7a856d91ef0276f Mon Sep 17 00:00:00 2001
From: Andrey Pavlenko
Date: Thu, 20 Mar 2014 22:22:55 +0400
Subject: [PATCH 31/47] switching to CV_HAAR_SCALE_IMAGE mode, enabling test
---
modules/ocl/perf/perf_haar.cpp | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
diff --git a/modules/ocl/perf/perf_haar.cpp b/modules/ocl/perf/perf_haar.cpp
index e70641d061..6b5100b6f2 100644
--- a/modules/ocl/perf/perf_haar.cpp
+++ b/modules/ocl/perf/perf_haar.cpp
@@ -92,12 +92,12 @@ PERF_TEST(HaarFixture, Haar)
typedef std::tr1::tuple Cascade_Image_MinSize_t;
typedef perf::TestBaseWithParam Cascade_Image_MinSize;
-OCL_PERF_TEST_P(Cascade_Image_MinSize, DISABLED_CascadeClassifier,
+OCL_PERF_TEST_P(Cascade_Image_MinSize, CascadeClassifier,
testing::Combine(testing::Values( string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt.xml"),
string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt2.xml") ),
- testing::Values(string("cv/shared/lena.png"),
- string("cv/cascadeandhog/images/bttf301.png")/*,
- string("cv/cascadeandhog/images/class57.png")*/ ),
+ testing::Values( string("cv/shared/lena.png"),
+ string("cv/cascadeandhog/images/bttf301.png"),
+ string("cv/cascadeandhog/images/class57.png") ),
testing::Values(30, 64, 90)))
{
const string cascasePath = get<0>(GetParam());
@@ -121,7 +121,7 @@ OCL_PERF_TEST_P(Cascade_Image_MinSize, DISABLED_CascadeClassifier,
faces.clear();
startTimer();
- cc.detectMultiScale(img, faces, 1.1, 3, 0, minSize);
+ cc.detectMultiScale(img, faces, 1.1, 3, CV_HAAR_SCALE_IMAGE, minSize);
stopTimer();
}
}
@@ -137,7 +137,7 @@ OCL_PERF_TEST_P(Cascade_Image_MinSize, DISABLED_CascadeClassifier,
ocl::finish();
startTimer();
- cc.detectMultiScale(uimg, faces, 1.1, 3, 0, minSize);
+ cc.detectMultiScale(uimg, faces, 1.1, 3, CV_HAAR_SCALE_IMAGE, minSize);
stopTimer();
}
}
From b7198ccf1c6f076586b74218a59727a968d55b26 Mon Sep 17 00:00:00 2001
From: Andrey Pavlenko
Date: Thu, 20 Mar 2014 22:30:16 +0400
Subject: [PATCH 32/47] dropping legacy modes testing
---
modules/objdetect/perf/opencl/perf_cascades.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/modules/objdetect/perf/opencl/perf_cascades.cpp b/modules/objdetect/perf/opencl/perf_cascades.cpp
index b660f59111..dd61cdb668 100644
--- a/modules/objdetect/perf/opencl/perf_cascades.cpp
+++ b/modules/objdetect/perf/opencl/perf_cascades.cpp
@@ -18,8 +18,6 @@ OCL_PERF_TEST_P(Cascade_Image_MinSize, CascadeClassifier,
testing::Combine(
testing::Values( string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt.xml"),
string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt2.xml"),
- string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt_old.xml"),
- string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt2_old.xml"),
string("cv/cascadeandhog/cascades/lbpcascade_frontalface.xml") ),
testing::Values( string("cv/shared/lena.png"),
string("cv/cascadeandhog/images/bttf301.png"),
From 0bd4fd3a87ddee808419e1e4d656c2cbe7dc94ca Mon Sep 17 00:00:00 2001
From: Alexander Karsakov
Date: Mon, 17 Mar 2014 12:18:55 +0400
Subject: [PATCH 33/47] Workaround for Intel platform: replace min() with
ternary operator
---
modules/imgproc/src/opencl/morph.cl | 5 +++++
1 file changed, 5 insertions(+)
diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl
index cb6e733ed4..35c0a27ff6 100644
--- a/modules/imgproc/src/opencl/morph.cl
+++ b/modules/imgproc/src/opencl/morph.cl
@@ -69,8 +69,13 @@
#endif
#ifdef ERODE
+#ifdef INTEL_DEVICE
+// workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older)
+#define MORPH_OP(A,B) ((A) < (B) ? (A) : (B))
+#else
#define MORPH_OP(A,B) min((A),(B))
#endif
+#endif
#ifdef DILATE
#define MORPH_OP(A,B) max((A),(B))
#endif
From b0ad84cfa2f4ecb884de542c9706177f22dade6d Mon Sep 17 00:00:00 2001
From: Alexander Smorkalov
Date: Fri, 21 Mar 2014 12:48:38 +0400
Subject: [PATCH 34/47] Libraries filter update after abs path cut.
---
cmake/OpenCVGenAndroidMK.cmake | 7 +++++--
1 file changed, 5 insertions(+), 2 deletions(-)
diff --git a/cmake/OpenCVGenAndroidMK.cmake b/cmake/OpenCVGenAndroidMK.cmake
index ee52fa6886..2622d2aaed 100644
--- a/cmake/OpenCVGenAndroidMK.cmake
+++ b/cmake/OpenCVGenAndroidMK.cmake
@@ -56,8 +56,11 @@ if(ANDROID)
# remove CUDA runtime and NPP from regular deps
# it can be added separately if needed.
- ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "libcu")
- ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "libnpp")
+ ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "cusparse")
+ ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "cufft")
+ ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "cublas")
+ ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "npp")
+ ocv_list_filterout(OPENCV_EXTRA_COMPONENTS_CONFIGMAKE "cudart")
if(HAVE_CUDA)
# CUDA runtime libraries and are required always
From 846266fde4c82fe392ccf122ef64ea30b6da4f01 Mon Sep 17 00:00:00 2001
From: Alexander Smorkalov
Date: Wed, 19 Feb 2014 16:20:40 +0400
Subject: [PATCH 35/47] Native camera fix for some deivices with Qualcomm SoC
like Samsung Galaxy S4.
---
.../armeabi-v7a/libnative_camera_r2.2.0.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r2.3.3.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r3.0.1.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r4.0.0.so | Bin 263644 -> 259548 bytes
.../armeabi-v7a/libnative_camera_r4.0.3.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r4.1.1.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r4.2.0.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r4.3.0.so | Bin 275932 -> 271836 bytes
.../armeabi-v7a/libnative_camera_r4.4.0.so | Bin 275932 -> 271836 bytes
.../lib/armeabi/libnative_camera_r2.2.0.so | Bin 288212 -> 284116 bytes
.../lib/armeabi/libnative_camera_r2.3.3.so | Bin 288212 -> 284116 bytes
.../lib/armeabi/libnative_camera_r3.0.1.so | Bin 284124 -> 280028 bytes
.../lib/armeabi/libnative_camera_r4.0.0.so | Bin 271832 -> 267736 bytes
.../lib/armeabi/libnative_camera_r4.0.3.so | Bin 284120 -> 280024 bytes
.../lib/armeabi/libnative_camera_r4.1.1.so | Bin 288216 -> 284120 bytes
.../lib/armeabi/libnative_camera_r4.2.0.so | Bin 288216 -> 284120 bytes
.../lib/armeabi/libnative_camera_r4.3.0.so | Bin 288220 -> 284124 bytes
.../lib/armeabi/libnative_camera_r4.4.0.so | Bin 284124 -> 280028 bytes
3rdparty/lib/mips/libnative_camera_r4.0.3.so | Bin 546012 -> 545972 bytes
3rdparty/lib/mips/libnative_camera_r4.1.1.so | Bin 546104 -> 546008 bytes
3rdparty/lib/mips/libnative_camera_r4.2.0.so | Bin 546108 -> 546012 bytes
3rdparty/lib/mips/libnative_camera_r4.3.0.so | Bin 546104 -> 546008 bytes
3rdparty/lib/mips/libnative_camera_r4.4.0.so | Bin 550340 -> 550308 bytes
3rdparty/lib/x86/libnative_camera_r2.3.3.so | Bin 427444 -> 423348 bytes
3rdparty/lib/x86/libnative_camera_r3.0.1.so | Bin 427444 -> 423348 bytes
3rdparty/lib/x86/libnative_camera_r4.0.3.so | Bin 427444 -> 423348 bytes
3rdparty/lib/x86/libnative_camera_r4.1.1.so | Bin 431540 -> 427444 bytes
3rdparty/lib/x86/libnative_camera_r4.2.0.so | Bin 447940 -> 443844 bytes
3rdparty/lib/x86/libnative_camera_r4.3.0.so | Bin 447940 -> 443844 bytes
3rdparty/lib/x86/libnative_camera_r4.4.0.so | Bin 456132 -> 456132 bytes
.../camera_wrapper/CMakeLists.txt | 2 +-
.../camera_wrapper/camera_wrapper.cpp | 171 ++++++++++--------
.../src/java/android+NativeCameraView.java | 1 -
33 files changed, 101 insertions(+), 73 deletions(-)
diff --git a/3rdparty/lib/armeabi-v7a/libnative_camera_r2.2.0.so b/3rdparty/lib/armeabi-v7a/libnative_camera_r2.2.0.so
index 5b618a87459123ff5f70d6f0647ca5d1785140cb..9b8352e03fe48531bab41f5b28dccdc403c7898d 100755
GIT binary patch
delta 80209
zcmZ_14}6XF|NnnoXB)$4V;F{EY8WQes;SW&^RFn?NK8>v)RbyOXNuk`)kGZ?#XI#@
zsYVnL1?00`&=XrKsozM4o-EOZvp7-bT`Fg#sKdnu>YtecCvaLi^n)w2qFndvj-Oe>FV^I@VJJ$;t
zi+Z?v+G&f%x!QRSJYrkrntGEghECG_tT28YKBRLfUB3)JGsAY~RJz&(e+Cbru5c;5
z^I_Y%RN0PTrA@W15wO&+A7Oc{gP*yal_C|ok?`4d5msxha25Q_s0eG8@*sHo)wXp;
zc{CiHWQHfghsQe`%b9_P;EmVVRthoc@YS&VHzGV#AR(6onSpP^5-P(uPBx)V(qSLC
zf2M7%#+DAIhQ*(RPZ7`8_;&cV{0QoIDD5lBR0^!u;*qR2~?2bA{8D9OIRPqN8mo>Cux3ceGHwX!C_(iC_IXM
zb`q!kxBeOTSik<$VKrkQ@4VJ_D#$9j1U>-EiW0sOJ_ySgkySn>jAz1C#AVPq+Z+x~
zvPM6HyJQ9Hd!nu1{^5wWnl-SjYwHrCt-fJ=3w-#q(4c)3o=yHH8kY`VfP<5C_z|4e
zJhTIU0}mXiJFtfezu=gFLpFgl(6pgt-Snt!6~MB@ec{=}gG+ojJm;a%pj`lWcH0(5
z7^nWyFkTJ!Bfdl9-@(DjBXfA%a0oYOWLbMSgh+!j13lprupFj@Z-N_4V3V+?3Eu-R
zVH?XPk}>o&eCi3?86y&33a{WeAY(vyJ$y&D&CIg@%M$x>1Sgq+8h9QF>DUr)6JuGM
zVcC?z1L5E#`NP9_6nuCvGpM`SLt*jNbQ4YIy)W8<_k7+wMQ&<<~d
z7aL;>KZZy28qV_UYRPSRfAFrE$v`*&vG-(d;c
z4QB0HcJp9!fkAss7|#ylEn)0x@z46{EkgY}d%^uAtQccke)43KyaK;;YiKW63onDE
z9f@xjLnm3nUGT`^fpfORe}wamL3shLy3MvS$(1pa)sp4Ef{Q6R0E%NuOUt^&;KlG2
z9D0NC4T(&V~`bf;3P9V#}npwJB$y6vDM15irA0Y
zS)2w_;UrEjLC%DOlWem5Fn%qJJ-fpk&RIrVkOq3mSw|<~31Pf2j6VxS5O32oRHmkpb?j)
z923U3!wqhVu<|(nIIHMB90PL#t3=kwqj3M*BAkM7DQKWy@+7#;w9QHH8rcHL?I3!}7m?7v07Emn^*`{EFkryCR%Jsqi^?p0NoUCUIyb
z-c|G4!o{6Jht^K;nrU|3?Q<5q5ti!S5t+%kA))WmuOK01x_zZ?%FVKI)%7NLZ<
zVf+42d^?RJ9XNU#^}EB-M*S>!xW{N
zI(aN>s?os9VHMWGbBy>-*j!^i_yr?g1Dk`cb+To>MLc+JxDswVJ#bkrV{VKG$2=Sb
zR$vg$geN>0y6j#CXS2%W{v?i_=>zyPBfb}2!V(4JW$*%nYv3IQH|@-wk#QC5Nyp)1
z)CU_F3HRq-EmvpgL3o7GzydgDSA=s37H2iG-h`8k8QcnI6PF7TiSLES8uiQIXyUTh
z5cgOYa4g0l$8o99>|)C*G%BRPiNxhvQsVvLRYv~Ja36#3f%DyVcW#HE?6@?SiM%J@EEfp}n9SE--eZdMRvE;=uz^2iVNdfd4Y;
zkAdU3wL=2u|NF77P>0On3OI)j0;{lzwFZ8Jc(8+Q@B}zZXZQ#B;FEUUvA!HOJ8am6
z-PRZ*$?$IO1?S8ChqD{txQRn?a4(o93Gf;1U@`nzVJQ9%9J@O-2)Dzja9|KOvi8A6
zbQn#2*=vr%Im9C!Ze*Q>{fvP~`gaCJW6mlCMnVU8gu$1?6Ac~;S2_&@>W_mnjrfCb
zq`^)87LIdG9<~a`*o%*xkmskT_MmYQ-T*Nbe=}!jzmua&9F~>=cqlxuB
z2}}PMN;n4lSY^TZ-|zxskTvba=~oXla(~|$o<;m3+Ur0E>99G&cfdLHC+7;;1T#E1
z@^Qqf<2m>Ma~#~}Z^Gs_+zgwKR=CK@X-k|MuhK=?;
zmK#TKYKk!eZp_K#TWpzuNpSb4Y-^(O)9^z0j4Kd-9X`C-w&rU5Q+Pc~n529PUT!R)
zs}GyFj=BGBiDN;Xqp8&uUS;qQc%{L2!K)3P3l|x@0xmH4Q+Teyd*N9@mh<0lIHm_3
zu+`VH<{R7sp7#t-!|lKlcY(XJ#Orhkv*0Ip+tw;wBlF>%-`Lhx&Hol|uoKpJ++`ft
zzOt>Q#GU*9TX2lt#*r%`(7~(lK2{-nma_?t!uI>2tJoU&z_WacrSY`O+2%!oPr+o9
z-3A~1%2jtG20c$2{&z*`O83GX$y6du8DEf*kiT=)}S
zMZB5AjjV=O=pgp!HtC4N+{S%ja|wsT{oe>&V#yPZyWri%^ML}m*Wa$XBk9Xw9fqYnPZJ!{;CUQ-NML@PgpNFk*k|nL*TDTb&;$?FFTleMJ_(ODxc^m_
zHOXKvJk8*f@a)w*e+f2tO$MXbNLU2FVDNExg~1v97!w9Bgf|*|4Bl?=RabLa|DJ8#
ztP6ZAY`@R>Ke$RBz+oeF(&@-a<;07#V
z@Ot29c(cJx2Jk4C15t2{T>?-0$mM*l=p3j9c-ZD2h8&AXcxyw5cfySgu7z(i8cevB
zBOM1OInYRhm%<~B_z?K85x)bDZW~(TvpqO=v=2GF@WM+Yti5)i!Ow&*kFb7JJ_!#r
zI*>&_Xmrp5zQTz2fDah)LGbc+p-t+!9mfvCF&m!xfy?5i*IA-h;pXdIb)Vbqgu5FZ
z9D?^59sCXF(_R|&Wy~}g$TcC%mUZ}IIE@9?G2^j@;^_BRXdB%FAK>IZQD@)Eh5Bxy?_^6l9|No6+8n2>AP`2UE@MyS&
za_d2C!+H^QH>bD2!Aa_jhZ`JmIZGt2r^79kkHe{LLJhdGxq{gjinoQ$BU?Y%JaCPM%>&ni@MhhF&iVgo9EWkZ
z=unnmCA`Mi#vj86jrjlIH;i~0e87lXgE?>+@uu)TBYp||OduXO|7YOXVVEP<4F>NgJ=sp*;s|&!wZdg)JP7sMm!OI$(W%lg_}jx&Cm#VyDbQ4Q(D#V$v6*_-gli*24gU`Tb{q^t+;=wgq2G9B06?~=BWBG9`Bq2D*
z&AC87Xmr>W-fVO*2sZO4!h4Lk7q;o(QC-3};iE?TU%-tkf(yj?e?N}6HKA4VS6GFn
zw{S8co}>+Cz`KoAJ{%rMJh%xSfQK9T%iw)R{(JD$V7@&6-;3ix94XqtMLBFkDr{43
z2QMSuLit8`JsezP_riza;5L0SEWQ%{i+D$^|1E6m3_0ijpK-KW>*6RK7}b%Z*$d!c
zg&y!(uJM9Pcr`rS$e#$`2B&F#5Bv-qoS~P);_KlR#CLK0adxA>aGaup7;T`z81DVJ
z2+dW#1pdq5f$&Z1Lh(uPdV}Y|=0$5UJd%spV0%BoYaG_|e>ILc?gN4eZExl51P3od
zv*6LEL$_8FVDrAg3peJ*BAEXvY~ITK1e-S=abvk;ll59CH9
z*uh%Zyx02V)=F&bzzo|B2eneZBe
zAB1-ryb7LYw6_&LY{VG;cUTF``;LY%8BR)~`jrgvNLd74;2#C5j6orC)ssk;PsnB2cG_L^If4$k^}EE;?KYhjQC3U
zFC+egurY&`u&4Q_p$aCY)fz$H2Kp1JmIVKL&2ar2cboA7cjIhWApxrMABv
zjyCH1JUG&5psgm@llWZEXrLoJiw1)6YvFjKfr)UeQGXu1$yk6l;cTP+SMU-e?kU4z
z9%}2|`OiabGHf1dvtaX3I}tVytFN3GS
z{SBTEk2H7%JYg5-{}^px4UVZK1bGWQ!{FWUJcECM7aIH*yv*Rp`#6vp+zeh12WPk&
zyxn6s2IAOf@NMu>gCBoFWBB4VlUhnuBQw1AGkT}Qr-Z!f`iA29k3@UDfIIA
z9vrg{1-|tt4KnFk#Fbol8e!P!+bga
z8y&=vkYiM61?L$Zbbu!s9i+hdMt(1NtkJ=haFNk|f4I=dADqwl%P~3_Nx~?jgB*B?
zQDFi+%;?~5c!SZwG5iIM*XoMAM$
z8cv60o69Cy3zr%BAH#i&_C25Di0%;D#yeoYQQfZSkp)_I;aPyhbmZ2;6$T?7I2QyK_Z-OfatbTA5zuIr$YH3pt<u7`BWK9)+ua4sGMPaIulU0QM6P?v{(;5~KZ>;jNnQk^OoFjx$sU?zgMpN~44K
z;QdAiAHwBE2cN=zql0a5nUVjEu+jd1;Uj_eJ8|~%5`EYO(jDz!x{5xUK6r+QCam+V5m<|^h6&{fWj1C@y
zCmJ2hhl`B-A~@G*e+j(Yh`;s#=l`)rgKv_MV>GxL4o*!m*1;Y2bF0Q8$u-~iaAn^}
zi|3czm^8Bv!AlH20jKu~#s7pyTo&S*2iX5IjSBUrvjj$krtm7GLR+}Zh^N5jXT6zl
zHu*A!WQNAUGfG|75}n~fxR_VFxDO%T)Oyu}V~)>d&F9!FOY|vhewh3{Jh0qlT}=gv
zH+_(cP&h+*IQ$tLyce7c?}TLrB)=HmZSY}u58OoKp7@7&EM_DOfbCPEXScV*<_8Fm
z!M%+5zhU$9fzRPxha=>BiY#Fx>kwRQa5cQ2cTDt9?y=$?WvaWRd}+&U&B)k_QTVHEYJThp20Q>I^Y}NJq8!T
z`wV^$K49>Va9Z!;xfED1ZkV6<&bNPd3-W
zoBs$sIsFMfU~uFt?vxC^6ps5dl%ETCcbNTO2F(H-gN=ms@FatOh36O?H=BDsgZsf7
z4ZaiZ#Y-%SI;IxG!wvo%o-8cqf0=;_9P^BX#7B7qWAG4ov%wF+8NY-cTo%DO2EPp#
z7`zQmF!*OU4d(r~z$mZ5F$zbJWzNzKz6zdh@EG_>K3WY9x;gMO2Csq_8vGsng28{o
z>zi@@4-Sf$!he2P+z~!XgfKi2f)kPL|C7*%H(~(+u#XpBdl+gXTZTp7NQU~
z^Iw4{dHS#l_2jS-huOeBcu$3yZ%y
zmr?kaQQ;dBuE+{iI02jS`tw)=*`eQPx)NSt@Km^FNGQGtjv5-`jqsAcLZ1ih_29^7
ziRGD&tg_$XVk4p96I>)3@y_rneH1I3Bpsee{zc9vYGe(CON{t9c)k&z2AiAA^E?i7
z6MO&%Cs}PDhw+v${xXce4dd^__~(BzzyI@3hjjv$-*9-A&gIDXCw!JSon;NlZf5_B
zRZM)F#v8&1`OPSqBtHQ@3eR)}xI-BCfKL#A$cc0Qzx*E#eiQLv&{b7nv;Q|K&JNlQ2dTR)_Hh*!(*FcK8%u&!1uk8vGyp
z7d%n9G>lKc=I;Ug0goOPVI|X_^ZZ}Fh!=FoDr^i-AR&0@Y!$|x;n~C&=nVG`<4oB6
z*{)%5;Oq8*8NLagc5{Sv8SkGvt7Kfr;qX21Ten15JW+M{fiQjqUPJs&jXxg7Ps1CD
zkJI?#FkV{6?EkXFD{-`9kOfa3>){Tt-0Mkvdl>J7dl3&FSdPG>4gM3JYOw1`#*D#n
zursLX-x<~IaZDH;`gVIC_-QyrXJ7~%{LKcLp&Q{fd>ek0j`Cc1-#rmlrsm%Zn?K3$
z2;6*1gvIH?S)jS_;QKiL@6`sL#c?z@aHNvm>q~m2f`D^>ixeg(I^~~#zt7T>k{O{U%@SupMnp<1C(Ed+w=PX!A-Ck?gg_v
z&JunLe*<%E>hMw6p2V2Y1wIF#%;o&wQ5_AR=Gbn`aYy*H!B@il488&GZg3tv+~6l+
zbBUM1<}W0C2zT)qKbi2AG+^+rVO#@GGve{=esi$jT6#S0C_rRB$
zgJ?RAsJkPq5PNE1qAX{(moFzZWcW3d?a6eSSa8~g@9DtMv_8Sf8N!(g
zsUws8?4L}0&TX9_muqWhYD+vqW0&9`p!s_OeR!_&m0(?Fb79)j;r
zDko@6ZtrE1uZCPsod>mzk5!jYkMG!9uPRHfU)$nMaA&$gJ_}r+&Vb9Z9wnGa;fpxV
z^g8yxwSxRMVH0eArqXhwtI%~?zZ~9={a?z}%pBj7wZ0S^|6KGqdLHdgY!x=oT;yJi
z-xUG!C|hxSOqq#oQx@qA@Pk{{E$|2``cb|q=uDixtkJMchk`zG_~xJWfchmC&E^eG
z8;G}{-f-ewu{~uNdM3A8Q1PHv{2!VMcOqG)$*MnNN73L8Vy|F7Ngh{+mJN^5oTu<{
ziE2HK9-=HJ#*JFtbPv87e1qh>;5ERNB>zSDN}O_Vl<6jlkC;q+(a!n=do}il=ua>=
z%_E7gXP((QAo~g_n|f82vZC
zVfem*pNF5)+zgb*P!=D$JJWU8yd`2?3)NyrQa+-5h;Jlq+$wp}|C3rl&O7M@gOhx6
zz7}|s_|51zjkOFWppD3pX_&SlcB1BuK<^>`9$nYfwq1Gc9cqXUWD`8(6~
zXgS;v{~zevl$DfXicH*$S$iqZ5dR$Bf&Dq<2Y82*!2-d25wC7?9(Y41>O6dU6X;LZ
zKLH*cSV>f@Z;aM0Mk6J}zoo>nM1SJv0^Q=#qBH#mKd+fN({%iI1x)Af!!VDit>2)0
z+P_=#4!{Fs`;W&tOe^$IweVHzstvk+hq?E4rVc@;YA+h?K_|=6FUi}0y+}Kflo6EM
z;RD3Dh_mLv&Cvxi8oxFY7sLN0=_*Qa>Ve%6m`Lnn>{m3_jCwA$Yl+`agR%JT(e~2B
zN1YhT_mtJd-=wUia1OT0M8D@x2}Pz7lD1(FK*v+Eu-D;VO@lm!uu3(LA6K!yBK|47
zfx+9{mC>4EQ|ObSlZj?S=KURuxKRY{(vIW&u|)L5BwVX4Z59jr`BDF-cDOr68jbZ5p*5g6s=F$
zf^Q?`b$o@i=i!BG=e2#lOjkGkLSmFM-(k0Yg@e->xE?dM65m7`ilzjo?HFmw7m>FZ
zevYzC>ldkV+hQe<_Zk;Rn4#$;2`^Id44qHVi|T10t`5}eMeeT6y}lJ|`IvalD>=6{LZge_A6
zI0S(F=b%E-E18+#wlAU*>Jz>VJux
zCe-^F&L@W#;+*@pEAT&0`3v5P%9N_@AB~I&IOZewMlPCjuC-?%Ia%apXpSVbV;*EW
zi8jSQk+@99(36y3@JFz{f5u)*{R!yRTE8B)OvT!U8~??s`*8e*ZXkIl_7{|7
z_y`=FmcRx0E$`tA5!Ws%I+wDj;9Tya%5w>JD*JLzyBB{ikLpT>+q>s*#CUH^-^ge$
zA34u$5Ilfql5*{$^z-p{qeTnPG_x<2#P1fxT!^-BU)23Vn?=jdM90a4D04dOq~1r>
zOX5B<1K5L{*f#mwp(HH!AA_I4{#52NF2-1of8f1{x-IQrV9U}=hdz;d}0ov+)V3U
zw2e4)dR?_MvEXzeXrNQo|DF2e!*eE=HC|?e6&~3+E8wd8n}^k#!l&0(1NFWsD=1DL
zWKUO*NH&Yu$?ix?&04(+&(hA~uqUbgf@~b`%81C=&Ob^WW6IH#pz8~agEZNuFTcW9
zH?30pRrM@W-CLItFj{%vwe6Oj{|D&~=J>a0_3MoEyMoJZKjNL{vfGTFN$_hDSDJ9fd1o)in4-YuzHN!d@ePPgH|bWA!aW
zkEt!OrfTn!MTqNeL?qhByD-9T*0~L{Hk?p9TrG7rB|Jd>Fl{5|kp?{SNrE|LYIZmD
zdn3@*=;#K`dB*Tb>UjLQ8ar;({o1hi7pCLM
z;;zw-#9JWtqJ^iLwH3M$R_6=lN1S!Tt{mfe*+{uT>M-3#xe>S
z@z2F>fc}Iwrqsuth~AINl!X3@_N6?D{WAJEn>fo
zs2C5M=AHwNIbsa^cO5s
zGQMj9rgb@W-&6kw+RhKMQ+WPobalPaid39R$<ll-w-e1?TP=0%5*&?$FV|xE+cjm
zb#5SUCH8i$FBuPGmr=gM?;`#(_8Zz(L_j^(1f79klYWxg5d5C<4)z2B%PAF_kb-{@
zbuLq0MNVt%bo`HLTl?W!SfynUZHJYP=6rG4ayOt(H{stM&ew`%|>EBxKjM&-l_SG
zh>yb8iR;#Pv2!R*;NTRc_3NweoYp(7{uStEkHY(CjM`VAnIylbd>aj9!#5g@OD8^!
ziyg~M*vc#6TD9BYpGLjWa3OjTEYm7VXY2--@aHLlbG#j5>}I_ZG#-sUOwI=SU7<5B
z_4d=oEu_!ad8u%!#T65JTwU3O(*vhMuC_gwf}IJyM9R0=)#w^jrfJ%n*f&z;y7EyX(AJ>Ma
zqT}$%bT9r+2YN6^}
z)%mImREty>t1eMps=8ctrKpU*Rq9x+TCBQGb%W|A)vc=ARClQEQY}&4tGZvcRMn??
z1ogPppO)pgC-|D=Nr>e!^ZRdt){4%J<%C8~Q>_p6qw
z`c#jomZ|zxt5q#eysk>sXw?|iSk-vd1l2^JruEsw-7jsjgGqpt?zQtLiS*64kw`
z`$c8{_o?HEYME-eYNhHK)oN9Fb4Aufq-wNktZKY!f@-2_l4>gI?Eh)%=%bphnyH$l
znyorYHAi)J-&{)mf^As&iZD{y$$G3sj3#m#QvTU8%ZCb+u}->IT(Ks#{fe
zsg|hjRo&l0_y1CL_*9RmmZ_dm^{ZB^$|sz%31U=ZRpV6?RFhRxR8v*cJZfa9W~yeX
zW~<77`jC!BsphE4fAA20uIfb9DXRIZ(^U&p3spUH)tIj;|D!}2EK*&pD*www{7Y4r
ztFBaCrMg;mo$3bFO{!Z}cLda9?NUdHYN@JE^@wVj>KRqPYPG7xwXw`#jB2cEylR4~
z{7)3gOLmm~FGU@FRMS;6R5MlOzr9GsVX7lk$ExP4PE?hDtReYRRP$A5sTPXL{y$e8
z^Ht>^WJtxus`5Mb!b??Gs;*L9ty-+QPIZIoCe^K~+f;X;&i=nk9VM!JRrjlws`^xq
zs8*`_RposinbAnqSk(m8WYrYa)KUclYPxEMYNl$IYPRYyRry74>3Ec?{F1-$
zSk*k$>8b^)vs&r?KUW>|RTrohsV-JsqPkRdx#~*Q)vCp+>r^+WZc^QMAanKWYtvFG}S(;>8cs3nW|Z;
z*{Z`-N2%thjt!{C%2me{)qK_Iss*Zrs&iH6t1eJ2QeCRLQgyZJI@JxTn;d2T+p3Oj
zsykG7sg|hjRrRTssaC4`RjXAkZtk7EK{Z-6UNuov_Wxvc$WIANg*4SZs_Cj3s+p=;
zsv}f$RC86QsOGCqS1nMTg*yBHTy@M>U7%W|x>$9o>T=bUs;g92s}`%SQ{AAtNp-90
zwl=!|?@-4s)e_ZGRiEk+)iTv`)k@Vfs(#gKRe4p{*%YcVs|HrE%K{Zh|Sv5s9
zRW(huk7|Z$rfQaIw(2m|5vrq9b3AH{Rn1kMsG6rbMKxb_mTIBuT-Eui3sj3#7ppE&
zU8?F?uEt8$RjR91i&fXDZcyE%x>a?X>JHUiswJv>Rrjlw2GnEu)Nw?$OtoCKQuU0g
zU$t6QUdwW}t!j*FtZKY!f@-2_lB06|S4XO9nyUOhK)
z>NeFKs(V%UtCp(zRF9~Zsg|o&s-98xt5&O8?R5W_U*3}?j!}(Ol^@g-e}ZbFYO-pI
zYN~3QY9G~f)eO~4)hyNQcDnx$Q^yF^QK~trV^!rR24#jOs^+OqS1nMTrCO*uS9QMX
z0@Wgq8jDqzs4i7qsk%yawQ8~II@JxTn^d=|?ousL-K)A^)l;g5PxXjunQEo#8CAb(
zWc$FHh*phJja7|TO;Al#O;Sw`D9``Y(ML61HA6L1HA^*Hb(rc1)lsTBsuNZ7RHvxs
zt4?>6{jWeBvs4RJ=c>+EU7%W|xT=bUs;g92s}`%SQ{5mc`~N0&Y*pQ+x=Xc0
zb+77v)lyZTYPo8q>KRqPYPG7x`;Rh+qfux7k5NagYP@QKYNBeAYO-poYMN>v)eO~4
z)hyL))nTfml63#iQO8)-iK=<3Q&jU+r>o9VEmWPWxXFO{7RjXAaJLoo2jZuwNjaN-nO;Sx(
zO;t@(?W3BmnxX0~YhovOqc5^kyy+KV<%l(3tW{#|7t1eJN-S1(ELMS7OT^kCR+(4{
zadv_|#oH^+ZrMP74M~RAEN}lfyS-Rwg(X(I8L2Gr|~v%7ez)hV9_N-$D2Mpb@)Nc{1t@^eDMiK@w}^0PtW
zPgU)sDnBnI{tVSj)vSOvwT7ufeild)M|ls#**(3nO<99!P3@#si#4!RRep{~5?6XB
zG_~7%=ab_tCcD)U4V0;tt5$mVG__lHkoWeaUZiTYYP@QK_cUn@k`+_Db~83omSVQ*
zFx4E@v8uVMQ&jU+r>hpI&QhJPTBN#Ib&2ZIfL>}XSI0`#RjR91H>hq>-RfN^^SxiO
zRJBa?jH+L?T2($&b=Hn*v}%lMtZKY!f@-2_l4^25n_4OANL5Yqo^EEhw6nZ+yxqRR
zDD~!eTgBU#beW=jR*H8*bGCnmVy1Umb9+#Oe8mFqf#&ulBNrs3#LWV3KayuJEHC8o2HQBqfg+0E3{E)4rjrI0#
z$r8*`Ec8xjXWQ-i>5&fIZxbev(?-Y2KXHcC20CozR-v3pB0BJEt}4V5QJDb4c`r^iqL=cLw`d>Uwh-k#KcDF3ys{vg&inSF
zSMo}VsJ!UOoHxINW3uQ(F1$o}t0zX3mnmXJ@8V^fcrkeWHbHbUrytR~Eh|Zs*MCw)
z7qd*FFYuC)D6j9Ni+VZzh%Vvv6j9#j$P(oR#JPOj=z5C|`Q%7+y6SxK^X^)#@Y`}N
z>=?WWAc1!{9}1(hRExyFni&*E3soaoQ?B<}GEv?(i4k2x!=h_xN0j%G5=7UrV4~|;
zFwqa$(L_BTu|#5QV2MO|H!(x>W0pwt6P8GH6EEC|ZsuXU=%<`dMYr&lkLXqo5~80m
zLu}IKpEEu6JNUdr*33a>0B*tS4aumdMHHh;9F17jqP*XkCVCMwCK}6(i8f`%
zMB|w;(dNvUXnST%v?DVnn#`9vMLRQtq8BrRqAARvXg6j=G?hEb9$X8z=+591M-K*{
zXiq+#6z$J86U}B5iQd2_5#_C%LQ&qToGUtxSARsuvk65f@Ij8~?HrdxALnIE(YaiQ
ziSjzga?vNa(CERnaEpKOI)XUn^Tj;TC;4W9=u=#fi9XF|<)XYQvPtwAo}Y^H4%#-+
z=lCv=Xc2ph=tAC56@8w)MsyK-jVLd3mx?kNeWDD)BOWmrY-OShu5!`s>_(!0@nVf=
zC2uQ={>iIKqVf`7PqdtEE9z&Pi|%8aiymbViXLMSiXLJRivGYL6y<%_BvCg%d6F!~
zAB+LfGRA;t31dL?cgBDyul1yh{>T^*-OCsd-OU&f{V!uc^e4uEC~rcK5Iw>e5apGv
z9MMw7fGF>-<%$mAYo-&$;FX>{QC`5DBFa}(@N8uqM^?fGST2!5Gj*b@zVtYMEb>|utTW!M>p-N~4LjMfJ%QDgXawR7
zJI1gh4ZHemX!V{k>~h0CV%VjIz1OgJc?{TQ*qaP{onfyw?3ISS)UX#D_5#D6YuK~Y
z_Q+tFZUm+n_C&)TYuKX!`@`r>kNCfVXrjorG~xOuooEiT*ICvwnz8Z=|*6R
zVNW#dv4%a$u!k9TmSJZYb|1q|HEd6^0f~kkZ`d)09ckFrRmLVT>~h0CV%VjIy*FTc
zoTJq)Be2b|HyQRi!(MIJD-C<8VJ|l91%^G>uxC|yLQ62+2uv~TiH1Gautyp8FvHF=
z>`+-hN6TZ&$TP8vNn^Bvw|-@6r@gBCXe$
z)8~BmT&VqR*559xQFhGW==B9zGqP%{T33wN*dV*%;D$qlWtKOWMQ3)Z9KWf4c0rb_
zvR|#W#myI7W~cD|A9Nb};U6x`|CG!?^srf3Nw>ILSGK6BH@NEN`7bZKPP_~FT@Y{!&@=Pu5Q$@Jjorq
zxtiKby=w;BZCoEzdz%inlSkCMAc55cmekb_za4B}X+P#|FvRW`lXucp
z%#1wDZygtSZx~{yrVTrpl2^2`?vw3?)Jy8%o|)CKcyn1k
z3-;2-wN(R7dfyvjUwU=o$tYiAjn`JiJ9UD24}WZJnSR=3t!D1_(^CV!BG#4SrVOP#
zT;Odt)K0Vi_Fg~K?vNZ&*>p?n`gpoL_FMXzj5Qsw%6_xb61Lb$v72Kz%ffm;9BNN=
zy;b8)9A*z`xU441oyKrkR_)CnW+ydwRbEl+^4Y5rDi@x$U+I)vdz^*z{yxlZexfKfAs_f}^6>lzgJ-xZCa!$%U_g+`Q(*8{O
zma>Dgma>#`FQw~)F6-`x_$PVT{`FDrChnGQD|y|3`aj;}PqwaLVJl7~t+D(3b^Q9u
z+B83l>)-56yUtE<$8GFPwDiPH{^*NptA0N5dfv5bTC;6#@W;d?u%${*Se;(aeLb)D
zKfEQ5H+{3U<=N>j>+FNxXRotcdpb{YC4X}&{jO!H6VP(>!UL`(E0I<3+XMWMP;?hM1Fohgr8hXa
zp4zH&Wea^HSlSUotQTZE)mFV&Hs5z0p6i_DFZ5~EDH4)v6EBY44$Je7!Y4$#`t$A}z+M2FIPvl(btevKh
zl{NESqE)YAShOzlK04BF!Fq`)i*ok3%FQus=k;vor}$w78AOj#rc-k6cUc8!Hu^GJ
z%+Y5ZWdr3U_vRa-yx7T#dT_pJO~j#cTz8|5{2%Wm1BZfA`t>*VaHLC=M%{oaRf
zu-is`EtR~U8|>!xChz+<*lj!=H?*u=_h~-gWlI^_g=2sB6P5kt=pT8yPjuwz#QKq^
zqw7VSj&Wu5cI8H%9v1b*#I1QFHZmM5qx|g#w;K|XBqQfS)pftSd>5)V{C@g)ChKx%
zZ90FytWD$ZB;kYF+T)V`J}MRNK^y*`Doua4
zD|2g#PPi-|D{9Am9P-)Sf8*zRHhkxO_C`A)(ysiV*8BF2_QkHL)!zT!XlJ{&Rd^Fe
z**(%ytL?P81FZdt8>2a_MpRzyk4=p}kmR=Cs;wGaZNHK?A!R~BR!VL}aTk`MYUCXYFUl-N4jO{zNu|JI3lU_kZaw${Z>(6rH!rV|3`wwCBykb
zcgdOyRl6!$`XqPHh1%mOlQDKR4NKUu_B}&((6pP@KP6>^IrGaye6U
z)~@{aeEu=6M!YL-wlAsw$0@G2sw%wS-E6nI^x(OOHp>P?90*)fRbH5oxOzb9mIkc`
zx0*A+wSUEclr7P2Iq#fgM`^gGW98;e-qyF+ee8YS>u<4t@LYGQgS%rw3NxP0iQ=0r
zeE(^E6Zga+QSOOsg_&8sx$e4S*oCU%@<}~rWm%Q;3R{=DdSsR~Dt+h!yGLf_e+tX;
z?6+KRU#Kc8=U*M%o#u<|(WT_-pQ4DpRCv?Svb@7N7pjhxcl71;m|4=K^v)hrOYS1}
zbm8?9^AKEGe%;XM4lcJZ=K`D8KI$1{?_;P=5M6)*Y;g1XNuE(=wPDL!8=o2iscp@zqwPpL^+er^_|4#96M7Dvr&J_O%}z
zmDJujudpMX+)`mZGP1m7qew=@8NQIc;BNl&-VyH~qwTIuZTLra@PiZ|<(K=sJ;vDW
z5)&r7te4>0s_IYWz}SCkuQf?&|DNiqtXJ#3@TV)uJAVv&|Fajo#bfMI_8M>0t#+r@
zdux)|bw@A)`-A)zDhI^2xK-mMMq518%i5Tv@f=!dvZx7MJ56jlb(*4!aj>Yw!Gu
z|4{Lj_q|)~)~@%@d%s4zuQ)HQM6Z85uhF%S=UzLx)d0K1uDozAnrr9oC#tS(?2hD$
z>yUqbVMJvkZ`@eBt>?|^H1`z@`||2+{{B?mVraTM#~%?JQCYtxWnw#@{hnQU?pzml
zM|MWpg}An?y!d=$UoUqO`Fp4xQMt7`#U0t9w(28hG2wiqPa+?(yZ*w(d4%lvk5BRO
zuN9f;s`YzL+x4!U9O+KqVqH9EVmt46$-e+&KQ`*EhD@RTsQSN2_gu)nBMLU+FwyLpr&u#YL?k&!B%dVMAa&@C_EVKMsGR_*xMuDtA=T^Z>eG|s-me%|}eIJG(!H-4R6*uPhsIV@Wpi(&gN>
ziS_q%%5yn0lEjQO^}aLS?%v>;^P}BSMHi};cn^=ad)Z^X&2sI&*{}8SEawNSt1lnf2a)Mn%_GS`iWXSM5n5^R!=102SWfbM4Na
z+s;-xXT~4=fkS3oP2ddqo!_c_?n2{NBI}kR`|Q%=cJD?nJUt-hg|!D7IyVomo|OhS
z`VVqejd`ij%e7V4oZoz0eDC}Jcig^e$?-IHsQ=vMk_EedOYtFpCN~UGiw^lSE}NBA
z|M}LQw;#!>R}^tNp?;A&Auqc3jsD2QPX;u1M{Muv{(DbCUdOlF%y{0NlG|V8x>Dp_e8Q*PY9P%%3
z&Urhj+cb`D%MSUMbmkZ1UYMB`waoQGy=C3yx;7pjJo
z)pw8Bc*vjMydJk;`JLr#TvN63ck2@GFBhr?l%>1-vw5;SWrgnfYYzEGHIE=|gr*h!
z-uQG%gKYQKK`HLRr2SCai<9<9tW@%g$WL|hQ#8MT{L}`$+@Ttu)+SIx>NcfDF*W%0
zPkOMlMlLndsNu0YhAOSB&8EWQ+N-EK+Feui_wSFpM|{xu^oWKd*rE$Au*%}4k!{H{
zvRb_&%acJfeWF{vlu7mx>~xwO;O54!ceZ=K&io9g)ycJzv;X%I^zeFZLwXQ85*`o633`nzR+PiF>i`27#A7UkGnpKF{hSeMpDy}}|zKbqU<
zgF}8#OnfBW$2gtTRDE9&yE^Kn2CpX3)Qi8Tu@lx*E#X44rt0P2lSBO;
zPVzNXYb#C=I#CyTzT)^$zY}@0;uw)*b&)43j%Z{_#VnnO8d0%NBNHq364_f9$*TCTMsBVsAyQHo
zxvb(_jSQ*SO=Ndnq-(`D8p){GNn~dr;$g)kR(z$Qw2ChYeOXs2wqlz`I#zs6jS}L;fRcUG9xTt>p7rk7PwHTIg=&x8M2r
zG(U+@^6}}^i|mpi?lZ^N4ro^5`SPSICe0o7((PO`$8N2u`ncjCTEHdgX!q@Ou(rZV
ziQY2SEtTz(>)gK_w^F|0Jw!iqVZYip_=
ztJp}rs1h$Xyk6(RB+~aHSBEb;b|-g@r|Q*H%a1o(eKzk?XK0P92n@PhmP>}(?G-n#
zSwQ6lP6dYEsrXejJf@5CJxMc9I^I8zH(1i&ZFP!YR8y5#5$TrkURIO5vB8q44mDL1
zE24dcv`{#t6Bnei)82XN^zqr$o=t6!IL4i7utXXiv1TTTGqtN5PNmb(qVsPbeTd|T
z)SG=O!zV|Hc{RNqom11((b+ZK9G$`WBa;~~J5_sJ_VqP2?b+9@%Aac5_+-YOt7*x#
z1;>_C%|ty_f1J9=C#$OBRD=`wjf&F1F<2_hrNI<*c#Ty#{JdD%HR4#&g$_fXKoijd^tbUYD-V5hyvJo-i!lmEB6=m-AN_tD
z4WiGYjnFvC-Po~c0@ouWuw$?ju`k1Jf?df?Q6hFz>?Hj18@L`8*pFg+JMLiE7Io{1CO}NkAHQcXGbD!(|PkJT#W2zLOb6
zKSHlWKSYz!L+IbPx~!w6vQH-nbt$aIx=-`J#Xmd>5*%PqECu9l2uZL+jiQbtpzQQoFhQkt=_gD7|0)x*UxSlC)S13+;-T#n{KpEMyC@m|4uneY0j}
z&0^-nu>2WULiu<77h=t@vRJcNEN=e`@xJ$V|NNfkeJ|qk{q=afdY;#x^Z)Prob!I)
zp$1iJFLWiFoYN%cyl+r78SoI`sqTHm9f`VJ!!l=(pf*Qqu=_RD?}#7r_llJj&mdFC
zBKFEmq9GHQZYJqNUSqX0$s5#A#QGhwmb^~>
zW%oK!f6|b`_Rb>x+P`NhsK!P~{7Yp8)y{+~#;~SYSSAED7js(}e`pqHTTYKAth8isF)JE)a
zvybPH-W^LSg@%U)1Oe=3{97;d=rMHYxKYDu*pWH7^L7S1q1d%Qb@Vb$EZ%TG_<4*<
z&3EmEhBeC5&RAxgLjp$lO-76dUwHGsF~Nxzc~d2fa{ch|?B-{YXM-D`E8O+yc4lRqWHsG*)^eM(les6l7FnSf0(HgG=Kp`m-JKoAQytxi#>D7M;p
zwmE-y&r{tabiwAj3(a(O&*}Cw)7^%yd2#a96))&?*g(#ven#4MOtCXh%G@T;Q;g!-
zx1Zr_d$CYAk2~Ty{><)Y^|znnSAdT?NY(CdIH}!-e;(5*QM{)W7Q~q4D|M}7g{~=Z
z`C&Rq%x#6Ic(WW-C#wB&TMT#c3&BK^;XBMvZI;98gobfu;go+q75XAm(GH5y4Nf6l
zL?`oUDwOM{eNK9jjyK0YpJ^{VH@lxmWV;uU9^R2Vd3v1GFW@YSmhs)i46Go{a=RLK
ze-R0K=`}RkB0N0rIg0@tpVj)>X0y2-w8>X(t|v%$v5@q3DYVI6v>mlA{m%;J@TW
z>RZF+|Cjh_+SLRkUUzwLwFT;whbZ_ZpKytNt5NzSAf8KHYX~G+XE0(T^c3KSpq7pX
z?}b#m{U9
z(4?QAFWL2YrgT3s5^~pr(TT$o&9Yl18?}V=Z)=vL(eZ5ZRy^~#mlT0L6SKU%f~`#@alJRy_ww>nFYYwQ2P+}Ni6QaRRfXNv5eS9WOC9|tV|y&w
z4@X(N=n8FCIhu9&lDy&JZApaSXWJuv!y$Ax9Va&bOER*374|)B@~6;jfF_F*yZR-0
zv;9#;@j4VS5RT!*`YeNDg`#*7ii(3&FmYl(aK!{gamV^dp9bL;PV6BRBSI9#K`6W+
z9N{1#01sdh;#{f1O&B^id6?_Lm;9RKLG^sF=#=OckDa7bD-K|KIQV0ZkT@M#>|4*$
zz9K8=qdMlboJ7;YIyPZB@uwH+*kaD_sbhJ|NeEq5$7+_7K(d&%{hD;4i|Sa!*Cd?I
zu44&blZEXD)D8A=QMV`*{DTvFcWjc~>vRoY6F)*zYgoH)$lMn9Z9?$vqT~~6*tg%1
zYI2(`NWZa5k9T~NJQvG}s1wEev2fbY
z>e9%NHnXe!W;V$)sx$FFv6>C~mUudK#Q^akd-q%NDJf^ae@pt2?X1O0(z(s?ruM^I
z`8#DSh=Fkxi(E;fUi`tptH`hEkzP18X&44DjG%qll9i+fEycO;N+Qwq)hzrw%-cb1
z$#-O+CacO>Ed<9E4p9psFBiqnXp()as%LyS<=-h=r}UcnRqo80Z&iKa8jT`dt607a
zOYdVmnJuek&J3D5b?Q|^lyn%+1g#6}XPy#xE0tiW&3)|%&&H$5>Y3Z8{yUX-hr#Op
z&^A`y*Y^W9RQ{sx=LAPhB^#1X#v^;Zz9RwbWI74!bCOS2sNL}slZTBN+06YW6;=
zxxxpkRw-vE5(9a4)evKmyE=JTU)FUM#+4v85I`Tj!6vLC{x9CfC%tXmX$fyD?@n}u
zop9;`dv_HHpod1XL#wdTNn;&XldfJ(a$SXe1mTC>W4vjh7gASV!N#n{N+gQ?cQsnb
z(UEM^YVsjHRLPvbCqX1l*YkVwuIOUoheRj**tcu3vPoj6)*>a7*xzeOkASiDE=oJ~
zPK@*MiuaEXSDWN_>dCshI0R2VQG!WgHS5kucjC|9W~3)+#g;HKu3NeFX`eO8>9x=H
z*(SM-73(+F#K>2a*0a9@a%h4~h=
zH%{#J^<-fNMa@K#M{1L99!TnDWx3g9`A;tU{gs1m&3+uay`1kqA
zpi&1Tl`hzg1ngVC-OeI
z$^Mf`I`{tFXqLaT@sKYlRnX1&KyN|}^I@3^s_BD+gdZ`DJ1Ca@oJn5mo>|#@dduuj
zrUy>`)jMG_mMH2Lm}Z*f(WZdueCG2`o-;jSdXubUtv8e2eNFQ13cu+>z}a!dINb}F
z{!ho3CcE$cwc5qyh)>VS!RlF)G2)ndz-utPPM+=$uc=LJ-e!`hGRdneSlMQh*2aXn
zE;^<^wuMcyx`Hjz5-+xT3yGrPP3+26|mTH?b}wh>%%
zVgGC!r`KwDW%#oBJ4qOwX=B@WlG#qK
zwY{g0VMRMg`*uex%_B0&U2JSn79t9&c^r{njXhhXS;QmwHrm(fmvJqiH+CI*p{<2t
zX_zwIXT=b&WIwD`Ncb3UBM!Zo*DhrKWy`Z^T4%Lqe#$OnUR!NX*_B;L*+46!*(AAp
zNp<(>3E4BBYp5o9h52zZr%w-=Zju|=kJ;pNw?|bxmA%zQ`F@o>mHl^lw{aTUEQ>1f@
zn{X6ul4luD=nM2q^ctrBM!!v;sXwQ$$G(ZH6=zmjOV~yV`7-yWT
z|4DxdgU#v?JfU{OZV|{6WDP_^02P7`gVuvqfaZcGg7~4!?uQp<`KZ6h5dvJ)%hfNk
zQ@P{?l~KM>!LHQQ65<#1Ya+TRs5*oE{2=?iiZ0#+}vM(xC`LsJ|p2i1GiC5uW-iHFKfuNc_&{Ul=GKe*Hxj-~rB98EUK#kPC;
zCR~H4)4OX;@(YS=hOJ4yXfdM1ve8`<1JHHw5!-E?*oz17+!V*&JU}{hG|YeYxe2G_
z@b7*?V*ll8?Pt%c?>f}Gxmx-k-lg_MbcY{Kcv6Q|Fs#b)Y|T>^
z0IMq83ynI#5{{EV&Ex%(wW&|{OI9{>uW068*vvh@nR|9K_lM2glbX5Tf!ltEDJpk!
zVhiuC4I+`ci|ohaSln5w7TZRky^;=w-u^
zk%lT3b&>>)_Od9!-HR8Dazq9j;D3(Ig)h0m)hASFvN@X+2o1o=ISz>Hp)AGnmPG8~{
za{*h$ogc$l!=3F}VIJ{tdKY{$tIQ)k#b{d$>v$TDV4DXU&K>SH{4xbMO*>6IoSUq|
zDODodewu`ex2+}Y>S?0)H)8c+k8L}U=Wq*<=Wr8FppEj626podi3r2^W-!UoR=aMY
z;c>b@g^Q<~Sxj8M*S9
zML2bUb<2lri)A?*#~mvyTJ~i=l%GPG!gBLTpOH^@9*ocIJjnmE^I(uAlw-8Ta8L+v
zQVaK7UD~K$H>h6ih$|tyCuUZ#h;t;olR-|aFv%S({H9x46K^?h*&gV}V>DaK{(FwR
zo>$NA!;)(2ng1^&l&06SxL=@`>RB?!Wcd393#T~P{{p+H
z`n3fR2iG${j$ZXFnxktyi!VTMhB~&i0A^)%Y+nJqwX9|Lfu2=0pUm*h`eBH3;X9qY
zFmFYuWx1?|sedJr)D