Merge branch 'master' of https://github.com/Itseez/opencv into back_proj_fix

CMakeLists.txt

@@ -438,7 +438,6 @@ include(cmake/OpenCVFindLibsGUI.cmake)
 include(cmake/OpenCVFindLibsVideo.cmake)
 include(cmake/OpenCVFindLibsPerf.cmake)
 
-
 # ----------------------------------------------------------------------------
 #  Detect other 3rd-party libraries/tools
 # ----------------------------------------------------------------------------

cmake/OpenCVDetectCUDA.cmake

@@ -217,3 +217,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()

cmake/OpenCVModule.cmake

@@ -27,7 +27,8 @@
 # The verbose template for OpenCV module:
 #
 #   ocv_add_module(modname <dependencies>)
-#   ocv_glob_module_sources() or glob them manually and ocv_set_module_sources(...)
+#   ocv_glob_module_sources(([EXCLUDE_CUDA] <extra sources&headers>)
+#                          or glob them manually and ocv_set_module_sources(...)
 #   ocv_module_include_directories(<extra include directories>)
 #   ocv_create_module()
 #   <add extra link dependencies, compiler options, etc>
@@ -478,9 +479,15 @@ endmacro()
 
 # finds and sets headers and sources for the standard OpenCV module
 # Usage:
-# ocv_glob_module_sources(<extra sources&headers in the same format as used in ocv_set_module_sources>)
+# ocv_glob_module_sources([EXCLUDE_CUDA] <extra sources&headers in the same format as used in ocv_set_module_sources>)
 macro(ocv_glob_module_sources)
-  file(GLOB_RECURSE lib_srcs     "src/*.cpp")
+  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/*.hpp" "include/opencv2/${name}/*.hpp" "include/opencv2/${name}/*.h")
   file(GLOB lib_hdrs_detail "include/opencv2/${name}/detail/*.hpp" "include/opencv2/${name}/detail/*.h")
@@ -492,15 +499,21 @@ macro(ocv_glob_module_sources)
   ocv_source_group("Src" DIRBASE "${CMAKE_CURRENT_SOURCE_DIR}/src" FILES ${lib_srcs} ${lib_int_hdrs})
   ocv_source_group("Include" DIRBASE "${CMAKE_CURRENT_SOURCE_DIR}/include" FILES ${lib_hdrs} ${lib_hdrs_detail})
 
-  file(GLOB lib_cuda_srcs "src/cuda/*.cu")
+  if (exclude_cuda EQUAL -1)
-  set(cuda_objs "")
+    file(GLOB lib_cuda_srcs "src/cuda/*.cu")
-  set(lib_cuda_hdrs "")
+    set(cuda_objs "")
-  if(HAVE_CUDA AND lib_cuda_srcs)
+    set(lib_cuda_hdrs "")
-    ocv_include_directories(${CUDA_INCLUDE_DIRS})
+    if(HAVE_CUDA)
-    file(GLOB lib_cuda_hdrs "src/cuda/*.hpp")
+      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})
+      ocv_cuda_compile(cuda_objs ${lib_cuda_srcs} ${lib_cuda_hdrs})
-    source_group("Src\\Cuda" FILES ${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()
 
   file(GLOB cl_kernels "src/opencl/*.cl")
@@ -516,8 +529,8 @@ macro(ocv_glob_module_sources)
     list(APPEND lib_srcs ${cl_kernels} "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.cpp" "${CMAKE_CURRENT_BINARY_DIR}/opencl_kernels.hpp")
   endif()
 
-  ocv_set_module_sources(${ARGN} HEADERS ${lib_hdrs} ${lib_hdrs_detail}
+  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})
+                         SOURCES ${lib_srcs} ${lib_int_hdrs} ${cuda_objs} ${lib_cuda_srcs} ${lib_cuda_hdrs})
 endmacro()
 
 # creates OpenCV module in current folder
@@ -622,11 +635,20 @@ endmacro()
 # short command for adding simple OpenCV module
 # see ocv_add_module for argument details
 # Usage:
-# ocv_define_module(module_name  [INTERNAL] [REQUIRED] [<list of dependencies>] [OPTIONAL <list of optional dependencies>])
+# ocv_define_module(module_name  [INTERNAL] [EXCLUDE_CUDA] [REQUIRED] [<list of dependencies>] [OPTIONAL <list of optional dependencies>])
 macro(ocv_define_module module_name)
-  ocv_add_module(${module_name} ${ARGN})
+  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} ${_argn})
   ocv_module_include_directories()
-  ocv_glob_module_sources()
+  ocv_glob_module_sources(${exclude_cuda})
   ocv_create_module()
   ocv_add_precompiled_headers(${the_module})
 

cmake/templates/OpenCVConfig.cmake.in

@@ -19,8 +19,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_COMPUTE_CAPABILITIES     : The version of compute capability.
-#      - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API
+#      - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API.
 #      - OpenCV_VERSION                  : The version of this OpenCV build: "@OPENCV_VERSION_PLAIN@"
 #      - OpenCV_VERSION_MAJOR            : Major version part of OpenCV_VERSION: "@OPENCV_VERSION_MAJOR@"
 #      - OpenCV_VERSION_MINOR            : Minor version part of OpenCV_VERSION: "@OPENCV_VERSION_MINOR@"
@@ -28,25 +28,29 @@
 #      - OpenCV_VERSION_STATUS           : Development status of this build: "@OPENCV_VERSION_STATUS@"
 #
 #    Advanced variables:
-#      - OpenCV_SHARED
+#      - OpenCV_SHARED                   : Use OpenCV as shared library
-#      - OpenCV_CONFIG_PATH
+#      - OpenCV_CONFIG_PATH              : Path to this OpenCVConfig.cmake
-#      - OpenCV_INSTALL_PATH  (not set on Windows)
+#      - OpenCV_INSTALL_PATH             : OpenCV location (not set on Windows)
-#      - OpenCV_LIB_COMPONENTS
+#      - OpenCV_LIB_COMPONENTS           : Present OpenCV modules list
-#      - OpenCV_USE_MANGLED_PATHS
+#      - OpenCV_USE_MANGLED_PATHS        : Mangled OpenCV path flag
-#      - OpenCV_HAVE_ANDROID_CAMERA
+#      - OpenCV_MODULES_SUFFIX           : The suffix for OpenCVModules-XXX.cmake file
+#      - OpenCV_HAVE_ANDROID_CAMERA      : Presence of Android native camera wrappers
 #
 #    Deprecated variables:
 #      - OpenCV_VERSION_TWEAK            : Always "0"
 #
 # ===================================================================================
 
-set(modules_file_suffix "")
+if(NOT DEFINED OpenCV_MODULES_SUFFIX)
-if(ANDROID)
+  if(ANDROID)
-  string(REPLACE - _ modules_file_suffix "_${ANDROID_NDK_ABI_NAME}")
+    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)

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)

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)

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.
+   * what is *linear blending* and why it is useful;
-   * Add two images using :add_weighted:`addWeighted <>`
+   * how to add two images using :add_weighted:`addWeighted <>`
 
 Theory
 =======

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.
 

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.hpp"
-   #include "opencv2/imgproc.hpp"
-   #include <iostream>
-   #include <stdio.h>
-
-   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 <image_settings0> <image_setting1> <image_settings2>\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 h_ranges[] = { 0, 180 };
-     float s_ranges[] = { 0, 180 };
+     float s_ranges[] = { 0, 256 };
 
      const float* ranges[] = { h_ranges, s_ranges };
 

modules/androidcamera/camera_wrapper/CMakeLists.txt

@@ -58,7 +58,7 @@ SET_TARGET_PROPERTIES(${the_target} PROPERTIES
                       RUNTIME_OUTPUT_DIRECTORY ${EXECUTABLE_OUTPUT_PATH}
                       )
 
-if (NOT (CMAKE_BUILD_TYPE MATCHES "debug"))
+if (NOT (CMAKE_BUILD_TYPE MATCHES "Debug"))
     ADD_CUSTOM_COMMAND( TARGET ${the_target} POST_BUILD COMMAND ${CMAKE_STRIP} --strip-unneeded "${LIBRARY_OUTPUT_PATH}/lib${the_target}.so" )
 endif()
 

modules/androidcamera/camera_wrapper/camera_wrapper.cpp

@@ -61,6 +61,12 @@
 
 using namespace android;
 
+// non-public camera related classes are not binary compatible
+// objects of these classes have different sizeof on different platforms
+// additional memory tail to all system objects to overcome sizeof issue
+#define MAGIC_TAIL 4096
+
+
 void debugShowFPS();
 
 #if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
@@ -90,6 +96,7 @@ public:
 };
 #endif
 
+
 std::string getProcessName()
 {
     std::string result;
@@ -142,12 +149,22 @@ class CameraHandler: public CameraListener
 protected:
     int cameraId;
     sp<Camera> camera;
-    CameraParameters params;
+#if defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
+    sp<SurfaceTexture> surface;
+#endif
+#if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0) || defined(ANDROID_r4_4_0)
+    sp<BufferQueue> queue;
+    sp<ConsumerListenerStub> listener;
+#endif
+    CameraParameters* params;
     CameraCallback cameraCallback;
     void* userData;
 
     int emptyCameraCallbackReported;
 
+    int width;
+    int height;
+
     static const char* flashModesNames[ANDROID_CAMERA_FLASH_MODES_NUM];
     static const char* focusModesNames[ANDROID_CAMERA_FOCUS_MODES_NUM];
     static const char* whiteBalanceModesNames[ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM];
@@ -258,7 +275,7 @@ protected:
 
     int is_supported(const char* supp_modes_key, const char* mode)
     {
-        const char* supported_modes = params.get(supp_modes_key);
+        const char* supported_modes = params->get(supp_modes_key);
         return (supported_modes && mode && (strstr(supported_modes, mode) > 0));
     }
 
@@ -268,7 +285,7 @@ protected:
         if (focus_distance_type >= 0 && focus_distance_type < 3)
     {
             float focus_distances[3];
-            const char* output = params.get(CameraParameters::KEY_FOCUS_DISTANCES);
+            const char* output = params->get(CameraParameters::KEY_FOCUS_DISTANCES);
             int val_num = CameraHandler::split_float(output, focus_distances, ',', 3);
             if(val_num == 3)
         {
@@ -300,10 +317,15 @@ public:
         emptyCameraCallbackReported(0)
     {
         LOGD("Instantiated new CameraHandler (%p, %p)", callback, _userData);
+        void* params_buffer = operator new(sizeof(CameraParameters) + MAGIC_TAIL);
+        params = new(params_buffer) CameraParameters();
     }
 
     virtual ~CameraHandler()
     {
+        if (params)
+            params->~CameraParameters();
+            operator delete(params);
         LOGD("CameraHandler destructor is called");
     }
 
@@ -371,10 +393,18 @@ const char* CameraHandler::focusModesNames[ANDROID_CAMERA_FOCUS_MODES_NUM] =
     CameraParameters::FOCUS_MODE_AUTO,
 #if !defined(ANDROID_r2_2_0)
     CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO,
+#else
+    CameraParameters::FOCUS_MODE_AUTO,
 #endif
     CameraParameters::FOCUS_MODE_EDOF,
     CameraParameters::FOCUS_MODE_FIXED,
-    CameraParameters::FOCUS_MODE_INFINITY
+    CameraParameters::FOCUS_MODE_INFINITY,
+    CameraParameters::FOCUS_MODE_MACRO,
+#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1)
+    CameraParameters::FOCUS_MODE_CONTINUOUS_PICTURE
+#else
+    CameraParameters::FOCUS_MODE_AUTO
+#endif
 };
 
 const char* CameraHandler::whiteBalanceModesNames[ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM] =
@@ -534,39 +564,39 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
     {
         LOGI("initCameraConnect: Setting paramers from previous camera handler");
         camera->setParameters(prevCameraParameters->flatten());
-        handler->params.unflatten(prevCameraParameters->flatten());
+        handler->params->unflatten(prevCameraParameters->flatten());
     }
     else
     {
         android::String8 params_str = camera->getParameters();
         LOGI("initCameraConnect: [%s]", params_str.string());
 
-        handler->params.unflatten(params_str);
+        handler->params->unflatten(params_str);
 
-        LOGD("Supported Cameras: %s", handler->params.get("camera-indexes"));
+        LOGD("Supported Cameras: %s", handler->params->get("camera-indexes"));
-        LOGD("Supported Picture Sizes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PICTURE_SIZES));
+        LOGD("Supported Picture Sizes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PICTURE_SIZES));
-        LOGD("Supported Picture Formats: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PICTURE_FORMATS));
+        LOGD("Supported Picture Formats: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PICTURE_FORMATS));
-        LOGD("Supported Preview Sizes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES));
+        LOGD("Supported Preview Sizes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES));
-        LOGD("Supported Preview Formats: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS));
+        LOGD("Supported Preview Formats: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS));
-        LOGD("Supported Preview Frame Rates: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_FRAME_RATES));
+        LOGD("Supported Preview Frame Rates: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_FRAME_RATES));
-        LOGD("Supported Thumbnail Sizes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_JPEG_THUMBNAIL_SIZES));
+        LOGD("Supported Thumbnail Sizes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_JPEG_THUMBNAIL_SIZES));
-        LOGD("Supported Whitebalance Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_WHITE_BALANCE));
+        LOGD("Supported Whitebalance Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_WHITE_BALANCE));
-        LOGD("Supported Effects: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_EFFECTS));
+        LOGD("Supported Effects: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_EFFECTS));
-        LOGD("Supported Scene Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_SCENE_MODES));
+        LOGD("Supported Scene Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_SCENE_MODES));
-        LOGD("Supported Focus Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES));
+        LOGD("Supported Focus Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES));
-        LOGD("Supported Antibanding Options: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_ANTIBANDING));
+        LOGD("Supported Antibanding Options: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_ANTIBANDING));
-        LOGD("Supported Flash Modes: %s", handler->params.get(CameraParameters::KEY_SUPPORTED_FLASH_MODES));
+        LOGD("Supported Flash Modes: %s", handler->params->get(CameraParameters::KEY_SUPPORTED_FLASH_MODES));
 
 #if !defined(ANDROID_r2_2_0)
         // Set focus mode to continuous-video if supported
-        const char* available_focus_modes = handler->params.get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES);
+        const char* available_focus_modes = handler->params->get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES);
         if (available_focus_modes != 0)
         {
             if (strstr(available_focus_modes, "continuous-video") != NULL)
             {
-                handler->params.set(CameraParameters::KEY_FOCUS_MODE, CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO);
+                handler->params->set(CameraParameters::KEY_FOCUS_MODE, CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO);
 
-                status_t resParams = handler->camera->setParameters(handler->params.flatten());
+                status_t resParams = handler->camera->setParameters(handler->params->flatten());
 
                 if (resParams != 0)
                 {
@@ -581,7 +611,7 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
 #endif
 
         //check if yuv420sp format available. Set this format as preview format.
-        const char* available_formats = handler->params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS);
+        const char* available_formats = handler->params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_FORMATS);
         if (available_formats != 0)
         {
             const char* format_to_set = 0;
@@ -607,9 +637,9 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
 
             if (0 != format_to_set)
             {
-                handler->params.setPreviewFormat(format_to_set);
+                handler->params->setPreviewFormat(format_to_set);
 
-                status_t resParams = handler->camera->setParameters(handler->params.flatten());
+                status_t resParams = handler->camera->setParameters(handler->params->flatten());
 
                 if (resParams != 0)
                     LOGE("initCameraConnect: failed to set preview format to %s", format_to_set);
@@ -617,6 +647,13 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
                     LOGD("initCameraConnect: preview format is set to %s", format_to_set);
             }
         }
+
+        handler->params->setPreviewSize(640, 480);
+        status_t resParams = handler->camera->setParameters(handler->params->flatten());
+        if (resParams != 0)
+            LOGE("initCameraConnect: failed to set preview resolution to 640x480");
+        else
+            LOGD("initCameraConnect: preview format is set to 640x480");
     }
 
     status_t bufferStatus;
@@ -627,22 +664,27 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
 #elif defined(ANDROID_r2_3_3)
     /* Do nothing in case of 2.3 for now */
 #elif defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
-    sp<SurfaceTexture> surfaceTexture = new SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID);
+    void* surface_texture_obj = operator new(sizeof(SurfaceTexture) + MAGIC_TAIL);
-    bufferStatus = camera->setPreviewTexture(surfaceTexture);
+    handler->surface = new(surface_texture_obj) SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID);
+    bufferStatus = camera->setPreviewTexture(handler->surface);
     if (bufferStatus != 0)
         LOGE("initCameraConnect: failed setPreviewTexture call (status %d); camera might not work correctly", bufferStatus);
 #elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
-    sp<BufferQueue> bufferQueue = new BufferQueue();
+    void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL);
-    sp<BufferQueue::ConsumerListener> queueListener = new ConsumerListenerStub();
+    handler->queue = new(buffer_queue_obj) BufferQueue();
-    bufferQueue->consumerConnect(queueListener);
+    void* consumer_listener_obj = operator new(sizeof(ConsumerListenerStub) + MAGIC_TAIL);
-    bufferStatus = camera->setPreviewTexture(bufferQueue);
+    handler->listener = new(consumer_listener_obj) ConsumerListenerStub();
+    handler->queue->consumerConnect(handler->listener);
+    bufferStatus = camera->setPreviewTexture(handler->queue);
     if (bufferStatus != 0)
         LOGE("initCameraConnect: failed setPreviewTexture call; camera might not work correctly");
 # elif defined(ANDROID_r4_4_0)
-    sp<BufferQueue> bufferQueue = new BufferQueue();
+    void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL);
-    sp<IConsumerListener> queueListener = new ConsumerListenerStub();
+    handler->queue = new(buffer_queue_obj) BufferQueue();
-    bufferQueue->consumerConnect(queueListener, true);
+    void* consumer_listener_obj = operator new(sizeof(ConsumerListenerStub) + MAGIC_TAIL);
-    bufferStatus = handler->camera->setPreviewTarget(bufferQueue);
+    handler->listener = new(consumer_listener_obj) ConsumerListenerStub();
+    handler->queue->consumerConnect(handler->listener, true);
+    bufferStatus = handler->camera->setPreviewTarget(handler->queue);
     if (bufferStatus != 0)
         LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly");
 # endif
@@ -723,18 +765,18 @@ double CameraHandler::getProperty(int propIdx)
     case ANDROID_CAMERA_PROPERTY_FRAMEWIDTH:
     {
         int w,h;
-        params.getPreviewSize(&w, &h);
+        params->getPreviewSize(&w, &h);
         return w;
     }
     case ANDROID_CAMERA_PROPERTY_FRAMEHEIGHT:
     {
         int w,h;
-        params.getPreviewSize(&w, &h);
+        params->getPreviewSize(&w, &h);
         return h;
     }
     case ANDROID_CAMERA_PROPERTY_SUPPORTED_PREVIEW_SIZES_STRING:
     {
-        cameraPropertySupportedPreviewSizesString = params.get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES);
+        cameraPropertySupportedPreviewSizesString = params->get(CameraParameters::KEY_SUPPORTED_PREVIEW_SIZES);
         union {const char* str;double res;} u;
         memset(&u.res, 0, sizeof(u.res));
         u.str = cameraPropertySupportedPreviewSizesString.c_str();
@@ -742,7 +784,7 @@ double CameraHandler::getProperty(int propIdx)
     }
     case ANDROID_CAMERA_PROPERTY_PREVIEW_FORMAT_STRING:
     {
-        const char* fmt = params.get(CameraParameters::KEY_PREVIEW_FORMAT);
+        const char* fmt = params->get(CameraParameters::KEY_PREVIEW_FORMAT);
         if (fmt == CameraParameters::PIXEL_FORMAT_YUV422SP)
             fmt = "yuv422sp";
         else if (fmt == CameraParameters::PIXEL_FORMAT_YUV420SP)
@@ -762,44 +804,44 @@ double CameraHandler::getProperty(int propIdx)
     }
     case ANDROID_CAMERA_PROPERTY_EXPOSURE:
     {
-        int exposure = params.getInt(CameraParameters::KEY_EXPOSURE_COMPENSATION);
+        int exposure = params->getInt(CameraParameters::KEY_EXPOSURE_COMPENSATION);
         return exposure;
     }
     case ANDROID_CAMERA_PROPERTY_FPS:
     {
-        return params.getPreviewFrameRate();
+        return params->getPreviewFrameRate();
     }
     case ANDROID_CAMERA_PROPERTY_FLASH_MODE:
     {
         int flash_mode = getModeNum(CameraHandler::flashModesNames,
                                     ANDROID_CAMERA_FLASH_MODES_NUM,
-                                    params.get(CameraParameters::KEY_FLASH_MODE));
+                                    params->get(CameraParameters::KEY_FLASH_MODE));
         return flash_mode;
     }
     case ANDROID_CAMERA_PROPERTY_FOCUS_MODE:
     {
         int focus_mode = getModeNum(CameraHandler::focusModesNames,
                                     ANDROID_CAMERA_FOCUS_MODES_NUM,
-                                    params.get(CameraParameters::KEY_FOCUS_MODE));
+                                    params->get(CameraParameters::KEY_FOCUS_MODE));
         return focus_mode;
     }
     case ANDROID_CAMERA_PROPERTY_WHITE_BALANCE:
     {
         int white_balance = getModeNum(CameraHandler::whiteBalanceModesNames,
                                        ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM,
-                                       params.get(CameraParameters::KEY_WHITE_BALANCE));
+                                       params->get(CameraParameters::KEY_WHITE_BALANCE));
         return white_balance;
     }
     case ANDROID_CAMERA_PROPERTY_ANTIBANDING:
     {
         int antibanding = getModeNum(CameraHandler::antibandingModesNames,
                                      ANDROID_CAMERA_ANTIBANDING_MODES_NUM,
-                                     params.get(CameraParameters::KEY_ANTIBANDING));
+                                     params->get(CameraParameters::KEY_ANTIBANDING));
         return antibanding;
     }
     case ANDROID_CAMERA_PROPERTY_FOCAL_LENGTH:
     {
-        float focal_length = params.getFloat(CameraParameters::KEY_FOCAL_LENGTH);
+        float focal_length = params->getFloat(CameraParameters::KEY_FOCAL_LENGTH);
         return focal_length;
     }
     case ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_NEAR:
@@ -814,6 +856,24 @@ double CameraHandler::getProperty(int propIdx)
     {
         return getFocusDistance(ANDROID_CAMERA_FOCUS_DISTANCE_FAR_INDEX);
     }
+#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1)
+    case ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK:
+    {
+        const char* status = params->get(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK);
+        if (status == CameraParameters::TRUE)
+            return 1.;
+        else
+            return 0.;
+    }
+    case ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK:
+    {
+        const char* status = params->get(CameraParameters::KEY_AUTO_EXPOSURE_LOCK);
+        if (status == CameraParameters::TRUE)
+            return 1.;
+        else
+            return 0.;
+    }
+#endif
     default:
         LOGW("CameraHandler::getProperty - Unsupported property.");
     };
@@ -824,99 +884,151 @@ void CameraHandler::setProperty(int propIdx, double value)
 {
     LOGD("CameraHandler::setProperty(%d, %f)", propIdx, value);
 
+    android::String8 params_str;
+    params_str = camera->getParameters();
+    LOGI("Params before set: [%s]", params_str.string());
+
     switch (propIdx)
     {
     case ANDROID_CAMERA_PROPERTY_FRAMEWIDTH:
     {
         int w,h;
-        params.getPreviewSize(&w, &h);
+        params->getPreviewSize(&w, &h);
-        w = (int)value;
+        width = (int)value;
-        params.setPreviewSize(w, h);
     }
     break;
     case ANDROID_CAMERA_PROPERTY_FRAMEHEIGHT:
     {
         int w,h;
-        params.getPreviewSize(&w, &h);
+        params->getPreviewSize(&w, &h);
-        h = (int)value;
+        height = (int)value;
-        params.setPreviewSize(w, h);
     }
     break;
     case ANDROID_CAMERA_PROPERTY_EXPOSURE:
     {
-        int max_exposure = params.getInt("max-exposure-compensation");
+        int max_exposure = params->getInt("max-exposure-compensation");
-        int min_exposure = params.getInt("min-exposure-compensation");
+        int min_exposure = params->getInt("min-exposure-compensation");
-        if(max_exposure && min_exposure){
+        if(max_exposure && min_exposure)
+        {
             int exposure = (int)value;
-            if(exposure >= min_exposure && exposure <= max_exposure){
+            if(exposure >= min_exposure && exposure <= max_exposure)
-                params.set("exposure-compensation", exposure);
+                params->set("exposure-compensation", exposure);
-            } else {
+            else
                 LOGE("Exposure compensation not in valid range (%i,%i).", min_exposure, max_exposure);
-            }
+        } else
-        } else {
             LOGE("Exposure compensation adjust is not supported.");
-        }
+
+        camera->setParameters(params->flatten());
     }
     break;
     case ANDROID_CAMERA_PROPERTY_FLASH_MODE:
     {
         int new_val = (int)value;
-        if(new_val >= 0 && new_val < ANDROID_CAMERA_FLASH_MODES_NUM){
+        if(new_val >= 0 && new_val < ANDROID_CAMERA_FLASH_MODES_NUM)
+        {
             const char* mode_name = flashModesNames[new_val];
             if(is_supported(CameraParameters::KEY_SUPPORTED_FLASH_MODES, mode_name))
-                params.set(CameraParameters::KEY_FLASH_MODE, mode_name);
+                params->set(CameraParameters::KEY_FLASH_MODE, mode_name);
             else
                 LOGE("Flash mode %s is not supported.", mode_name);
-        } else {
-            LOGE("Flash mode value not in valid range.");
         }
+        else
+            LOGE("Flash mode value not in valid range.");
+
+        camera->setParameters(params->flatten());
     }
     break;
     case ANDROID_CAMERA_PROPERTY_FOCUS_MODE:
     {
         int new_val = (int)value;
-        if(new_val >= 0 && new_val < ANDROID_CAMERA_FOCUS_MODES_NUM){
+        if(new_val >= 0 && new_val < ANDROID_CAMERA_FOCUS_MODES_NUM)
+        {
             const char* mode_name = focusModesNames[new_val];
             if(is_supported(CameraParameters::KEY_SUPPORTED_FOCUS_MODES, mode_name))
-                params.set(CameraParameters::KEY_FOCUS_MODE, mode_name);
+                params->set(CameraParameters::KEY_FOCUS_MODE, mode_name);
             else
                 LOGE("Focus mode %s is not supported.", mode_name);
-        } else {
-            LOGE("Focus mode value not in valid range.");
         }
+        else
+            LOGE("Focus mode value not in valid range.");
+
+        camera->setParameters(params->flatten());
     }
     break;
     case ANDROID_CAMERA_PROPERTY_WHITE_BALANCE:
     {
         int new_val = (int)value;
-        if(new_val >= 0 && new_val < ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM){
+        if(new_val >= 0 && new_val < ANDROID_CAMERA_WHITE_BALANCE_MODES_NUM)
+        {
             const char* mode_name = whiteBalanceModesNames[new_val];
             if(is_supported(CameraParameters::KEY_SUPPORTED_WHITE_BALANCE, mode_name))
-                params.set(CameraParameters::KEY_WHITE_BALANCE, mode_name);
+                params->set(CameraParameters::KEY_WHITE_BALANCE, mode_name);
             else
                 LOGE("White balance mode %s is not supported.", mode_name);
-        } else {
-            LOGE("White balance mode value not in valid range.");
         }
+        else
+            LOGE("White balance mode value not in valid range.");
+
+        camera->setParameters(params->flatten());
     }
     break;
     case ANDROID_CAMERA_PROPERTY_ANTIBANDING:
     {
         int new_val = (int)value;
-        if(new_val >= 0 && new_val < ANDROID_CAMERA_ANTIBANDING_MODES_NUM){
+        if(new_val >= 0 && new_val < ANDROID_CAMERA_ANTIBANDING_MODES_NUM)
+        {
             const char* mode_name = antibandingModesNames[new_val];
             if(is_supported(CameraParameters::KEY_SUPPORTED_ANTIBANDING, mode_name))
-                params.set(CameraParameters::KEY_ANTIBANDING, mode_name);
+                params->set(CameraParameters::KEY_ANTIBANDING, mode_name);
             else
                 LOGE("Antibanding mode %s is not supported.", mode_name);
-        } else {
-            LOGE("Antibanding mode value not in valid range.");
         }
+        else
+            LOGE("Antibanding mode value not in valid range.");
+
+        camera->setParameters(params->flatten());
     }
     break;
+#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1)
+    case ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK:
+    {
+        if (is_supported(CameraParameters::KEY_AUTO_EXPOSURE_LOCK_SUPPORTED, "true"))
+        {
+            if (value != 0)
+                params->set(CameraParameters::KEY_AUTO_EXPOSURE_LOCK, CameraParameters::TRUE);
+            else
+                params->set(CameraParameters::KEY_AUTO_EXPOSURE_LOCK, CameraParameters::FALSE);
+            LOGE("Expose lock is set");
+        }
+        else
+            LOGE("Expose lock is not supported");
+
+        camera->setParameters(params->flatten());
+    }
+    break;
+    case ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK:
+    {
+        if (is_supported(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK_SUPPORTED, "true"))
+        {
+            if (value != 0)
+                params->set(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK, CameraParameters::TRUE);
+            else
+                params->set(CameraParameters::KEY_AUTO_WHITEBALANCE_LOCK, CameraParameters::FALSE);
+            LOGE("White balance lock is set");
+        }
+        else
+            LOGE("White balance lock is not supported");
+
+        camera->setParameters(params->flatten());
+    }
+    break;
+#endif
     default:
         LOGW("CameraHandler::setProperty - Unsupported property.");
     };
+
+    params_str = camera->getParameters();
+    LOGI("Params after set: [%s]", params_str.string());
 }
 
 void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
@@ -935,7 +1047,10 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
         return;
     }
 
-    CameraParameters curCameraParameters((*ppcameraHandler)->params.flatten());
+    // delayed resolution setup to exclude errors during other parameres setup on the fly
+    // without camera restart
+    if (((*ppcameraHandler)->width != 0) && ((*ppcameraHandler)->height != 0))
+        (*ppcameraHandler)->params->setPreviewSize((*ppcameraHandler)->width, (*ppcameraHandler)->height);
 
 #if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) \
  || defined(ANDROID_r4_3_0) || defined(ANDROID_r4_4_0)
@@ -951,27 +1066,27 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
         return;
     }
 
-    handler->camera->setParameters(curCameraParameters.flatten());
+    handler->camera->setParameters((*ppcameraHandler)->params->flatten());
-    handler->params.unflatten(curCameraParameters.flatten());
 
     status_t bufferStatus;
 # if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
-    sp<SurfaceTexture> surfaceTexture = new SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID);
+    void* surface_texture_obj = operator new(sizeof(SurfaceTexture) + MAGIC_TAIL);
-    bufferStatus = handler->camera->setPreviewTexture(surfaceTexture);
+    handler->surface = new(surface_texture_obj) SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID);
+    bufferStatus = handler->camera->setPreviewTexture(handler->surface);
     if (bufferStatus != 0)
         LOGE("applyProperties: failed setPreviewTexture call (status %d); camera might not work correctly", bufferStatus);
 # elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
-    sp<BufferQueue> bufferQueue = new BufferQueue();
+    void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL);
-    sp<BufferQueue::ConsumerListener> queueListener = new ConsumerListenerStub();
+    handler->queue = new(buffer_queue_obj) BufferQueue();
-    bufferQueue->consumerConnect(queueListener);
+    handler->queue->consumerConnect(handler->listener);
-    bufferStatus = handler->camera->setPreviewTexture(bufferQueue);
+    bufferStatus = handler->camera->setPreviewTexture(handler->queue);
     if (bufferStatus != 0)
         LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly");
 # elif defined(ANDROID_r4_4_0)
-    sp<BufferQueue> bufferQueue = new BufferQueue();
+    void* buffer_queue_obj = operator new(sizeof(BufferQueue) + MAGIC_TAIL);
-    sp<IConsumerListener> queueListener = new ConsumerListenerStub();
+    handler->queue = new(buffer_queue_obj) BufferQueue();
-    bufferQueue->consumerConnect(queueListener, true);
+    handler->queue->consumerConnect(handler->listener, true);
-    bufferStatus = handler->camera->setPreviewTarget(bufferQueue);
+    bufferStatus = handler->camera->setPreviewTarget(handler->queue);
     if (bufferStatus != 0)
         LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly");
 # endif
@@ -1002,7 +1117,7 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
     LOGD("CameraHandler::applyProperties(): after previousCameraHandler->closeCameraConnect");
 
     LOGD("CameraHandler::applyProperties(): before initCameraConnect");
-    CameraHandler* handler=initCameraConnect(cameraCallback, cameraId, userData, &curCameraParameters);
+    CameraHandler* handler=initCameraConnect(cameraCallback, cameraId, userData, (*ppcameraHandler)->params);
     LOGD("CameraHandler::applyProperties(): after initCameraConnect, handler=0x%x", (int)handler);
     if (handler == NULL) {
         LOGE("ERROR in applyProperties --- cannot reinit camera");

modules/androidcamera/include/camera_properties.h

@@ -15,7 +15,9 @@ enum {
     ANDROID_CAMERA_PROPERTY_FOCAL_LENGTH = 105,
     ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_NEAR = 106,
     ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_OPTIMAL = 107,
-    ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_FAR = 108
+    ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_FAR = 108,
+    ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK = 109,
+    ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK = 110
 };
 
 
@@ -30,12 +32,12 @@ enum {
 
 enum {
     ANDROID_CAMERA_FOCUS_MODE_AUTO = 0,
-    ANDROID_CAMERA_FOCUS_MODE_CONTINUOUS_PICTURE,
     ANDROID_CAMERA_FOCUS_MODE_CONTINUOUS_VIDEO,
     ANDROID_CAMERA_FOCUS_MODE_EDOF,
     ANDROID_CAMERA_FOCUS_MODE_FIXED,
     ANDROID_CAMERA_FOCUS_MODE_INFINITY,
     ANDROID_CAMERA_FOCUS_MODE_MACRO,
+    ANDROID_CAMERA_FOCUS_MODE_CONTINUOUS_PICTURE,
     ANDROID_CAMERA_FOCUS_MODES_NUM
 };
 

modules/core/include/opencv2/core/mat.hpp

@@ -118,6 +118,8 @@ public:
 
     virtual int kind() const;
     virtual int dims(int i=-1) const;
+    virtual int cols(int i=-1) const;
+    virtual int rows(int i=-1) const;
     virtual Size size(int i=-1) const;
     virtual int sizend(int* sz, int i=-1) const;
     virtual bool sameSize(const _InputArray& arr) const;

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<PlatformInfo>& platform_info);
 CV_EXPORTS int predictOptimalVectorWidth(InputArray src1, InputArray src2 = noArray(), InputArray src3 = noArray(),
                                          InputArray src4 = noArray(), InputArray src5 = noArray(), InputArray src6 = noArray(),

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;
 

modules/core/src/matrix.cpp

@@ -1416,6 +1416,16 @@ int _InputArray::kind() const
     return flags & KIND_MASK;
 }
 
+int _InputArray::rows(int i) const
+{
+    return size(i).height;
+}
+
+int _InputArray::cols(int i) const
+{
+    return size(i).width;
+}
+
 Size _InputArray::size(int i) const
 {
     int k = kind();
@@ -2078,45 +2088,45 @@ void _OutputArray::create(Size _sz, int mtype, int i, bool allowTransposed, int
     create(2, sizes, mtype, i, allowTransposed, fixedDepthMask);
 }
 
-void _OutputArray::create(int rows, int cols, int mtype, int i, bool allowTransposed, int fixedDepthMask) const
+void _OutputArray::create(int _rows, int _cols, int mtype, int i, bool allowTransposed, int fixedDepthMask) const
 {
     int k = kind();
     if( k == MAT && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
-        CV_Assert(!fixedSize() || ((Mat*)obj)->size.operator()() == Size(cols, rows));
+        CV_Assert(!fixedSize() || ((Mat*)obj)->size.operator()() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((Mat*)obj)->type() == mtype);
-        ((Mat*)obj)->create(rows, cols, mtype);
+        ((Mat*)obj)->create(_rows, _cols, mtype);
         return;
     }
     if( k == UMAT && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
-        CV_Assert(!fixedSize() || ((UMat*)obj)->size.operator()() == Size(cols, rows));
+        CV_Assert(!fixedSize() || ((UMat*)obj)->size.operator()() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((UMat*)obj)->type() == mtype);
-        ((UMat*)obj)->create(rows, cols, mtype);
+        ((UMat*)obj)->create(_rows, _cols, mtype);
         return;
     }
     if( k == GPU_MAT && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
-        CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(cols, rows));
+        CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((cuda::GpuMat*)obj)->type() == mtype);
-        ((cuda::GpuMat*)obj)->create(rows, cols, mtype);
+        ((cuda::GpuMat*)obj)->create(_rows, _cols, mtype);
         return;
     }
     if( k == OPENGL_BUFFER && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
-        CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(cols, rows));
+        CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((ogl::Buffer*)obj)->type() == mtype);
-        ((ogl::Buffer*)obj)->create(rows, cols, mtype);
+        ((ogl::Buffer*)obj)->create(_rows, _cols, mtype);
         return;
     }
     if( k == CUDA_MEM && i < 0 && !allowTransposed && fixedDepthMask == 0 )
     {
-        CV_Assert(!fixedSize() || ((cuda::CudaMem*)obj)->size() == Size(cols, rows));
+        CV_Assert(!fixedSize() || ((cuda::CudaMem*)obj)->size() == Size(_cols, _rows));
         CV_Assert(!fixedType() || ((cuda::CudaMem*)obj)->type() == mtype);
-        ((cuda::CudaMem*)obj)->create(rows, cols, mtype);
+        ((cuda::CudaMem*)obj)->create(_rows, _cols, mtype);
         return;
     }
-    int sizes[] = {rows, cols};
+    int sizes[] = {_rows, _cols};
     create(2, sizes, mtype, i, allowTransposed, fixedDepthMask);
 }
 
@@ -2679,17 +2689,17 @@ namespace cv {
 
 static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s )
 {
-    int type = _m.type(), cn = CV_MAT_CN(type);
+    int type = _m.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
-    if (cn == 3)
+            sctype = CV_MAKE_TYPE(depth, cn == 3 ? 4 : cn);
-        return false;
 
     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);

modules/core/src/ocl.cpp

@@ -2179,7 +2179,6 @@ static cl_device_id selectOpenCLDevice()
             goto not_found;
         }
     }
-
     if (deviceTypes.size() == 0)
     {
         if (!isID)
@@ -2193,13 +2192,16 @@ static cl_device_id selectOpenCLDevice()
     for (size_t t = 0; t < deviceTypes.size(); t++)
     {
         int deviceType = 0;
-        if (deviceTypes[t] == "GPU")
+        std::string tempStrDeviceType = deviceTypes[t];
+        std::transform( tempStrDeviceType.begin(), tempStrDeviceType.end(), tempStrDeviceType.begin(), tolower );
+
+        if (tempStrDeviceType == "gpu" || tempStrDeviceType == "dgpu" || tempStrDeviceType == "igpu")
             deviceType = Device::TYPE_GPU;
-        else if (deviceTypes[t] == "CPU")
+        else if (tempStrDeviceType == "cpu")
             deviceType = Device::TYPE_CPU;
-        else if (deviceTypes[t] == "ACCELERATOR")
+        else if (tempStrDeviceType == "accelerator")
             deviceType = Device::TYPE_ACCELERATOR;
-        else if (deviceTypes[t] == "ALL")
+        else if (tempStrDeviceType == "all")
             deviceType = Device::TYPE_ALL;
         else
         {
@@ -2229,7 +2231,14 @@ static cl_device_id selectOpenCLDevice()
         {
             std::string name;
             CV_OclDbgAssert(getStringInfo(clGetDeviceInfo, devices[i], CL_DEVICE_NAME, name) == CL_SUCCESS);
-            if (isID || name.find(deviceName) != std::string::npos)
+            cl_bool useGPU = true;
+            if(tempStrDeviceType == "dgpu" || tempStrDeviceType == "igpu")
+            {
+                cl_bool isIGPU = CL_FALSE;
+                clGetDeviceInfo(devices[i], CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(isIGPU), &isIGPU, NULL);
+                useGPU = tempStrDeviceType == "dgpu" ? !isIGPU : isIGPU;
+            }
+            if ( (isID || name.find(deviceName) != std::string::npos) && useGPU)
             {
                 // TODO check for OpenCL 1.1
                 return devices[i];
@@ -4307,7 +4316,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);
 
@@ -4318,13 +4327,13 @@ String kernelToStr(InputArray _kernel, int ddepth)
     if (ddepth != depth)
         kernel.convertTo(kernel, ddepth);
 
-    typedef std::string (*func_t)(const Mat &);
+    typedef std::string (* func_t)(const Mat &);
-    static const func_t funcs[] = { kerToStr<uchar>, kerToStr<char>, kerToStr<ushort>,kerToStr<short>,
+    static const func_t funcs[] = { kerToStr<uchar>, kerToStr<char>, kerToStr<ushort>, kerToStr<short>,
                                     kerToStr<int>, kerToStr<float>, kerToStr<double>, 0 };
     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) \

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,
+__kernel void arithm_flip_rows(__global const uchar * srcptr, int src_step, int src_offset,
-                               __global uchar* dstptr, int dststep, int dstoffset,
+                               __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)));
+        T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
-        __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, mad24(x, sizeoftype, srcoffset)));
+        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)));
+        storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
-        __global type* dst1 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, mad24(x, sizeoftype, dstoffset)));
+        storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x, TSIZE, dst_offset)));
-
-        dst0[0] = src1[0];
-        dst1[0] = src0[0];
     }
 }
 
-__kernel void arithm_flip_rows_cols(__global const uchar* srcptr, int srcstep, int srcoffset,
+__kernel void arithm_flip_rows_cols(__global const uchar * srcptr, int src_step, int src_offset,
-                                    __global uchar* dstptr, int dststep, int dstoffset,
+                                    __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)));
+        T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
-        __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, mad24(x1, sizeoftype, srcoffset)));
+        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)));
+        storepix(src0, dstptr + mad24(rows - y - 1, dst_step, mad24(x1, TSIZE, dst_offset)));
-        __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset)));
+        storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
-
-        dst0[0] = src0[0];
-        dst1[0] = src1[0];
     }
 }
 
-__kernel void arithm_flip_cols(__global const uchar* srcptr, int srcstep, int srcoffset,
+__kernel void arithm_flip_cols(__global const uchar * srcptr, int src_step, int src_offset,
-                               __global uchar* dstptr, int dststep, int dstoffset,
+                               __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)));
+        T src0 = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset)));
-        __global const type* src1 = (__global const type*)(srcptr + mad24(y, srcstep, mad24(x1, sizeoftype, srcoffset)));
+        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)));
+        storepix(src0, dstptr + mad24(y, dst_step, mad24(x1, TSIZE, dst_offset)));
-        __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, mad24(x, sizeoftype, dstoffset)));
+        storepix(src1, dstptr + mad24(y, dst_step, mad24(x, TSIZE, dst_offset)));
-
-        dst1[0] = src1[0];
-        dst0[0] = src0[0];
     }
 }

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));
+        int src_index = mad24(y, src_step, mad24(x, TSIZE, src_offset));
-        __global T * src = (__global T *)(srcptr + src_index);
 
-        src[0] = x == y ? scalar : (T)(0);
+        storepix(x == y ? scalar : (T)(0), srcptr + src_index);
     }
 }

modules/core/src/persistence.cpp

@@ -4824,7 +4824,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);
@@ -4862,7 +4862,7 @@ cvUnregisterType( const char* type_name )
         if( !CvType::first || !CvType::last )
             CvType::first = CvType::last = 0;
 
-        free( info );
+        cvFree( &info );
     }
 }
 

modules/cudabgsegm/src/cuda/mog2.cu

@@ -163,7 +163,7 @@ namespace cv { namespace cuda { 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)
                 {
@@ -214,6 +214,7 @@ namespace cv { namespace cuda { 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);
@@ -231,7 +232,7 @@ namespace cv { namespace cuda { 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;
             }
 

modules/flann/include/opencv2/flann/kmeans_index.h

@@ -758,10 +758,13 @@ private:
 
                     for (int k=0; k<indices_length; ++k) {
                         if (belongs_to[k]==j) {
-                            belongs_to[k] = i;
+                            // for cluster j, we move the furthest element from the center to the empty cluster i
-                            count[j]--;
+                            if ( distance_(dataset_[indices[k]], dcenters[j], veclen_) == radiuses[j] ) {
-                            count[i]++;
+                                belongs_to[k] = i;
-                            break;
+                                count[j]--;
+                                count[i]++;
+                                break;
+                            }
                         }
                     }
                     converged = false;

modules/highgui/doc/reading_and_writing_images_and_video.rst

@@ -483,7 +483,7 @@ VideoWriter constructors
 
     :param filename: Name of the output video file.
 
-    :param fourcc: 4-character code of codec used to compress the frames. For example, ``CV_FOURCC('P','I','M,'1')``  is a MPEG-1 codec, ``CV_FOURCC('M','J','P','G')``  is a motion-jpeg codec etc. List of codes can be obtained at `Video Codecs by FOURCC <http://www.fourcc.org/codecs.php>`_ page.
+    :param fourcc: 4-character code of codec used to compress the frames. For example, ``CV_FOURCC('P','I','M','1')``  is a MPEG-1 codec, ``CV_FOURCC('M','J','P','G')``  is a motion-jpeg codec etc. List of codes can be obtained at `Video Codecs by FOURCC <http://www.fourcc.org/codecs.php>`_ page.
 
     :param fps: Framerate of the created video stream.
 

modules/highgui/include/opencv2/highgui/highgui_c.h

@@ -463,6 +463,8 @@ enum
     CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_NEAR = 8006,
     CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_OPTIMAL = 8007,
     CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_FAR = 8008,
+    CV_CAP_PROP_ANDROID_EXPOSE_LOCK = 8009,
+    CV_CAP_PROP_ANDROID_WHITEBALANCE_LOCK = 8010,
 
     // Properties of cameras available through AVFOUNDATION interface
     CV_CAP_PROP_IOS_DEVICE_FOCUS = 9001,
@@ -543,6 +545,7 @@ enum
 enum
 {
     CV_CAP_ANDROID_FOCUS_MODE_AUTO = 0,
+    CV_CAP_ANDROID_FOCUS_MODE_CONTINUOUS_PICTURE,
     CV_CAP_ANDROID_FOCUS_MODE_CONTINUOUS_VIDEO,
     CV_CAP_ANDROID_FOCUS_MODE_EDOF,
     CV_CAP_ANDROID_FOCUS_MODE_FIXED,

modules/highgui/src/cap_android.cpp

@@ -289,6 +289,10 @@ double CvCapture_Android::getProperty( int propIdx )
         return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_OPTIMAL);
     case CV_CAP_PROP_ANDROID_FOCUS_DISTANCE_FAR:
         return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_FOCUS_DISTANCE_FAR);
+    case CV_CAP_PROP_ANDROID_EXPOSE_LOCK:
+        return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK);
+    case CV_CAP_PROP_ANDROID_WHITEBALANCE_LOCK:
+        return (double)m_activity->getProperty(ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK);
     default:
         CV_Error( CV_StsOutOfRange, "Failed attempt to GET unsupported camera property." );
         break;
@@ -327,14 +331,23 @@ bool CvCapture_Android::setProperty( int propIdx, double propValue )
         case CV_CAP_PROP_ANDROID_ANTIBANDING:
             m_activity->setProperty(ANDROID_CAMERA_PROPERTY_ANTIBANDING, propValue);
             break;
+        case CV_CAP_PROP_ANDROID_EXPOSE_LOCK:
+            m_activity->setProperty(ANDROID_CAMERA_PROPERTY_EXPOSE_LOCK, propValue);
+            break;
+        case CV_CAP_PROP_ANDROID_WHITEBALANCE_LOCK:
+            m_activity->setProperty(ANDROID_CAMERA_PROPERTY_WHITEBALANCE_LOCK, propValue);
+            break;
         default:
             CV_Error( CV_StsOutOfRange, "Failed attempt to SET unsupported camera property." );
             return false;
         }
 
-        if (propIdx != CV_CAP_PROP_AUTOGRAB) {// property for highgui class CvCapture_Android only
+        // Only changes in frame size require camera restart
+        if ((propIdx == CV_CAP_PROP_FRAME_WIDTH) || (propIdx == CV_CAP_PROP_FRAME_HEIGHT))
+        {   // property for highgui class CvCapture_Android only
             m_CameraParamsChanged = true;
         }
+
         res = true;
     }
 

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 ////////////////////////

modules/imgproc/src/canny.cpp

@@ -100,19 +100,29 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float
         low_thresh = std::min(32767.0f, low_thresh);
         high_thresh = std::min(32767.0f, high_thresh);
 
-        if (low_thresh > 0) low_thresh *= low_thresh;
+        if (low_thresh > 0)
-        if (high_thresh > 0) high_thresh *= high_thresh;
+            low_thresh *= low_thresh;
+        if (high_thresh > 0)
+            high_thresh *= high_thresh;
     }
     int low = cvFloor(low_thresh), high = cvFloor(high_thresh);
     Size esize(size.width + 2, size.height + 2);
 
     UMat mag;
-    size_t globalsize[2] = { size.width * cn, size.height }, localsize[2] = { 16, 16 };
+    size_t globalsize[2] = { size.width, size.height }, localsize[2] = { 16, 16 };
 
     if (aperture_size == 3 && !_src.isSubmatrix())
     {
         // Sobel calculation
-        ocl::Kernel calcSobelRowPassKernel("calcSobelRowPass", ocl::imgproc::canny_oclsrc);
+        char cvt[2][40];
+        ocl::Kernel calcSobelRowPassKernel("calcSobelRowPass", ocl::imgproc::canny_oclsrc,
+                                           format("-D OP_SOBEL -D cn=%d -D shortT=%s -D ucharT=%s"
+                                                  " -D convertToIntT=%s -D intT=%s -D convertToShortT=%s", cn,
+                                                  ocl::typeToStr(CV_16SC(cn)),
+                                                  ocl::typeToStr(CV_8UC(cn)),
+                                                  ocl::convertTypeStr(CV_8U, CV_32S, cn, cvt[0]),
+                                                  ocl::typeToStr(CV_32SC(cn)),
+                                                  ocl::convertTypeStr(CV_32S, CV_16S, cn, cvt[1])));
         if (calcSobelRowPassKernel.empty())
             return false;
 
@@ -126,58 +136,62 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float
 
         // magnitude calculation
         ocl::Kernel magnitudeKernel("calcMagnitude_buf", ocl::imgproc::canny_oclsrc,
-                                    L2gradient ? " -D L2GRAD" : "");
+                                    format("-D cn=%d%s -D OP_MAG_BUF -D shortT=%s -D convertToIntT=%s -D intT=%s",
+                                           cn, L2gradient ? " -D L2GRAD" : "",
+                                           ocl::typeToStr(CV_16SC(cn)),
+                                           ocl::convertTypeStr(CV_16S, CV_32S, cn, cvt[0]),
+                                           ocl::typeToStr(CV_32SC(cn))));
         if (magnitudeKernel.empty())
             return false;
 
-        mag = UMat(esize, CV_32SC(cn), Scalar::all(0));
+        mag = UMat(esize, CV_32SC1, Scalar::all(0));
         dx.create(size, CV_16SC(cn));
         dy.create(size, CV_16SC(cn));
 
         magnitudeKernel.args(ocl::KernelArg::ReadOnlyNoSize(dxBuf), ocl::KernelArg::ReadOnlyNoSize(dyBuf),
                              ocl::KernelArg::WriteOnlyNoSize(dx), ocl::KernelArg::WriteOnlyNoSize(dy),
-                             ocl::KernelArg::WriteOnlyNoSize(mag, cn), size.height, size.width);
+                             ocl::KernelArg::WriteOnlyNoSize(mag), size.height, size.width);
 
         if (!magnitudeKernel.run(2, globalsize, localsize, false))
             return false;
     }
     else
     {
-        dx.create(size, CV_16SC(cn));
+        Sobel(_src, dx, CV_16S, 1, 0, aperture_size, 1, 0, BORDER_REPLICATE);
-        dy.create(size, CV_16SC(cn));
+        Sobel(_src, dy, CV_16S, 0, 1, aperture_size, 1, 0, BORDER_REPLICATE);
-
-        Sobel(_src, dx, CV_16SC1, 1, 0, aperture_size, 1, 0, BORDER_REPLICATE);
-        Sobel(_src, dy, CV_16SC1, 0, 1, aperture_size, 1, 0, BORDER_REPLICATE);
 
         // magnitude calculation
         ocl::Kernel magnitudeKernel("calcMagnitude", ocl::imgproc::canny_oclsrc,
-                                    L2gradient ? " -D L2GRAD" : "");
+                                    format("-D OP_MAG -D cn=%d%s -D intT=int -D shortT=short -D convertToIntT=convert_int_sat",
+                                           cn, L2gradient ? " -D L2GRAD" : ""));
         if (magnitudeKernel.empty())
             return false;
 
-        mag = UMat(esize, CV_32SC(cn), Scalar::all(0));
+        mag = UMat(esize, CV_32SC1, Scalar::all(0));
         magnitudeKernel.args(ocl::KernelArg::ReadOnlyNoSize(dx), ocl::KernelArg::ReadOnlyNoSize(dy),
-                             ocl::KernelArg::WriteOnlyNoSize(mag, cn), size.height, size.width);
+                             ocl::KernelArg::WriteOnlyNoSize(mag), size.height, size.width);
 
         if (!magnitudeKernel.run(2, globalsize, NULL, false))
             return false;
     }
 
     // map calculation
-    ocl::Kernel calcMapKernel("calcMap", ocl::imgproc::canny_oclsrc);
+    ocl::Kernel calcMapKernel("calcMap", ocl::imgproc::canny_oclsrc,
+                              format("-D OP_MAP -D cn=%d", cn));
     if (calcMapKernel.empty())
         return false;
 
-    UMat map(esize, CV_32SC(cn));
+    UMat map(esize, CV_32SC1);
     calcMapKernel.args(ocl::KernelArg::ReadOnlyNoSize(dx), ocl::KernelArg::ReadOnlyNoSize(dy),
-                       ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::WriteOnlyNoSize(map, cn),
+                       ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::WriteOnlyNoSize(map),
                        size.height, size.width, low, high);
 
     if (!calcMapKernel.run(2, globalsize, localsize, false))
         return false;
 
     // local hysteresis thresholding
-    ocl::Kernel edgesHysteresisLocalKernel("edgesHysteresisLocal", ocl::imgproc::canny_oclsrc);
+    ocl::Kernel edgesHysteresisLocalKernel("edgesHysteresisLocal", ocl::imgproc::canny_oclsrc,
+                                           "-D OP_HYST_LOCAL");
     if (edgesHysteresisLocalKernel.empty())
         return false;
 
@@ -193,7 +207,8 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float
 
     for ( ; ; )
     {
-        ocl::Kernel edgesHysteresisGlobalKernel("edgesHysteresisGlobal", ocl::imgproc::canny_oclsrc);
+        ocl::Kernel edgesHysteresisGlobalKernel("edgesHysteresisGlobal", ocl::imgproc::canny_oclsrc,
+                                                "-D OP_HYST_GLOBAL");
         if (edgesHysteresisGlobalKernel.empty())
             return false;
 
@@ -221,14 +236,15 @@ static bool ocl_Canny(InputArray _src, OutputArray _dst, float low_thresh, float
     }
 
     // get edges
-    ocl::Kernel getEdgesKernel("getEdges", ocl::imgproc::canny_oclsrc);
+    ocl::Kernel getEdgesKernel("getEdges", ocl::imgproc::canny_oclsrc, "-D OP_EDGES");
     if (getEdgesKernel.empty())
         return false;
 
-    _dst.create(size, CV_8UC(cn));
+    _dst.create(size, CV_8UC1);
     UMat dst = _dst.getUMat();
 
     getEdgesKernel.args(ocl::KernelArg::ReadOnlyNoSize(map), ocl::KernelArg::WriteOnly(dst));
+
     return getEdgesKernel.run(2, globalsize, NULL, false);
 }
 
@@ -254,12 +270,12 @@ void cv::Canny( InputArray _src, OutputArray _dst,
     }
 
     if ((aperture_size & 1) == 0 || (aperture_size != -1 && (aperture_size < 3 || aperture_size > 7)))
-        CV_Error(CV_StsBadFlag, "");
+        CV_Error(CV_StsBadFlag, "Aperture size should be odd");
 
     if (low_thresh > high_thresh)
         std::swap(low_thresh, high_thresh);
 
-    CV_OCL_RUN(_dst.isUMat() && cn == 1,
+    CV_OCL_RUN(_dst.isUMat() && (cn == 1 || cn == 3),
                ocl_Canny(_src, _dst, (float)low_thresh, (float)high_thresh, aperture_size, L2gradient, cn, size))
 
     Mat src = _src.getMat(), dst = _dst.getMat();

modules/imgproc/src/filter.cpp

@@ -3134,7 +3134,7 @@ template<typename ST, class CastOp, class VecOp> struct Filter2D : public BaseFi
 // b e h b e h 0 0
 // c f i c f i 0 0
 template <typename T>
-static int _prepareKernelFilter2D(std::vector<T>& data, const Mat &kernel)
+static int _prepareKernelFilter2D(std::vector<T> & data, const Mat & kernel)
 {
     Mat _kernel; kernel.convertTo(_kernel, DataDepth<T>::value);
     int size_y_aligned = ROUNDUP(kernel.rows * 2, 4);
@@ -3317,200 +3317,224 @@ 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(const UMat & src, UMat & buf, const Mat & kernelX, int anchor,
+                               int borderType, int ddepth, bool fast8uc1)
 {
-    int type = src.type();
+    int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type);
-    int cn = CV_MAT_CN(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
-    int sdepth = CV_MAT_DEPTH(type);
     Size bufSize = buf.size();
 
+    if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F))
+        return false;
+
 #ifdef ANDROID
     size_t localsize[2] = {16, 10};
 #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)
+    if (fast8uc1)
-    {
+        globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0];
-        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;
-        }
-    }
 
-    int radiusX = anchor;
+    int radiusX = anchor, radiusY = (buf.rows - src.rows) >> 1;
-    int radiusY = (int)((buf.rows - src.rows) >> 1);
 
-    bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
+    bool isolated = (borderType & BORDER_ISOLATED) != 0;
-    const char* btype = NULL;
+    const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" },
-    switch (borderType & ~BORDER_ISOLATED)
+        * const btype = borderMap[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;
-    }
 
     bool extra_extrapolation = src.rows < (int)((-radiusY + globalsize[1]) >> 1) + 1;
     extra_extrapolation |= src.rows < radiusY;
     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",
+    char cvt[40];
-        radiusX, (int)localsize[0], (int)localsize[1], cn,
+    cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s"
-        btype,
+                                          " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s%s",
-        extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
+                                          radiusX, (int)localsize[0], (int)localsize[1], cn, btype,
-        isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED");
+                                          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),
+                                          doubleSupport ? " -D DOUBLE_SUPPORT" : "");
     build_options += ocl::kernelToStr(kernelX, CV_32F);
 
     Size srcWholeSize; Point srcOffset;
     src.locateROI(srcWholeSize, srcOffset);
 
-    std::stringstream strKernel;
+    String kernelName("row_filter");
-    strKernel << "row_filter";
+    if (fast8uc1)
-    if (-1 != cn)
+        kernelName += "_C1_D0";
-        strKernel << "_C" << cn;
-    if (-1 != sdepth)
-        strKernel << "_D" << sdepth;
 
-    ocl::Kernel kernelRow;
+    ocl::Kernel k(kernelName.c_str(), cv::ocl::imgproc::filterSepRow_oclsrc,
-    if (!kernelRow.create(strKernel.str().c_str(), cv::ocl::imgproc::filterSepRow_oclsrc,
+                  build_options);
-                          build_options))
+    if (k.empty())
         return false;
 
-    int idxArg = 0;
+    if (fast8uc1)
-    idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
+        k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x,
-    idxArg = kernelRow.set(idxArg, (int)(src.step / src.elemSize()));
+               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);
 
-    idxArg = kernelRow.set(idxArg, srcOffset.x);
+    return k.run(2, globalsize, localsize, false);
-    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);
 }
 
-static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor, bool sync)
+static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, double delta, 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];
-
+    globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0];
-    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];
 
     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",
+    cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d"
-                    anchor, (int)localsize[0], (int)localsize[1], cn, ocl::typeToStr(buf.type()),
+                                          " -D srcT=%s -D dstT=%s -D convertToDstT=%s"
-                                          ocl::typeToStr(dtype), ocl::convertTypeStr(CV_32F, ddepth, cn, cvt));
+                                          " -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::typeToStr(CV_32F), ocl::typeToStr(ddepth),
+                                          doubleSupport ? " -D DOUBLE_SUPPORT" : "");
     build_options += ocl::kernelToStr(kernelY, CV_32F);
 
-    ocl::Kernel kernelCol;
+    ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc,
-    if (!kernelCol.create("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, build_options))
+                  build_options);
+    if (k.empty())
         return false;
 
-    int idxArg = 0;
+    k.args(ocl::KernelArg::ReadOnly(buf), ocl::KernelArg::WriteOnly(dst),
-    idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrReadOnly(buf));
+           static_cast<float>(delta));
-    idxArg = kernelCol.set(idxArg, (int)(buf.step / buf.elemSize()));
-    idxArg = kernelCol.set(idxArg, buf.cols);
-    idxArg = kernelCol.set(idxArg, buf.rows);
 
-    idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
+    return k.run(2, globalsize, localsize, false);
-    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);
+const int optimizedSepFilterLocalSize = 16;
+
+static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst,
+                                       Mat row_kernel, Mat col_kernel,
+                                       double delta, 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 || ddepth == 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 -D srcT1=%s -D dstT1=%s -D CN=%d", (int)lt2[0], (int)lt2[1],
+                             row_kernel.cols / 2, col_kernel.cols / 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::typeToStr(sdepth), ocl::typeToStr(ddepth), cn);
+
+    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<int>((src_offset % src_step) / esz);
+    int src_offset_y = static_cast<int>(src_offset / src_step);
+
+    src.locateROI(wholeSize, origin);
+
+    k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y,
+           wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst),
+           static_cast<float>(delta));
+
+    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 )
 {
-    if (abs(delta)> FLT_MIN)
+    const ocl::Device & d = ocl::Device::getDefault();
-        return false;
+    Size imgSize = _src.size();
 
-    int type = _src.type();
+    int type = _src.type(), sdepth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
-    if ( !( (type == CV_8UC1 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC4) &&
+    if (cn > 4)
-            (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 (ddepth < 0)
-    if( anchor.x < 0 )
-        anchor.x = kernelX.cols >> 1;
-    if( anchor.y < 0 )
-        anchor.y = kernelY.cols >> 1;
-
-    if( ddepth < 0 )
         ddepth = sdepth;
 
+    CV_OCL_RUN_(kernelY.cols <= 21 && kernelX.cols <= 21 &&
+                imgSize.width > optimizedSepFilterLocalSize + (kernelX.cols >> 1) &&
+                imgSize.height > optimizedSepFilterLocalSize + (kernelY.cols >> 1) &&
+                (!(borderType & BORDER_ISOLATED) || _src.offset() == 0) && anchor == Point(-1, -1) &&
+                (d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())),
+                ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, delta,
+                                           borderType & ~BORDER_ISOLATED, ddepth), true)
+
+    if (anchor.x < 0)
+        anchor.x = kernelX.cols >> 1;
+    if (anchor.y < 0)
+        anchor.y = kernelY.cols >> 1;
+
     UMat src = _src.getUMat();
     Size srcWholeSize; Point srcOffset;
     src.locateROI(srcWholeSize, srcOffset);
-    if ( (0 != (srcOffset.x % 4))   ||
+
-         (0 != (src.cols % 4))      ||
+    bool fast8uc1 = type == CV_8UC1 && srcOffset.x % 4 == 0 &&
-         (0 != ((src.step / src.elemSize()) % 4))
+            src.cols % 4 == 0 && src.step % 4 == 0;
-       )
-        return false;
 
     Size srcSize = src.size();
     Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1);
-    UMat buf; buf.create(bufSize, CV_MAKETYPE(CV_32F, cn));
+    UMat buf(bufSize, CV_32FC(cn));
-    if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, false))
+    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, false);
+
+    return ocl_sepColFilter2D(buf, dst, kernelY, delta, anchor.y);
 }
 
 #endif

modules/imgproc/src/morph.cpp

@@ -42,7 +42,6 @@
 
 #include "precomp.hpp"
 #include <limits.h>
-#include <stdio.h>
 #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",
+    String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s"
-                                 anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op],
+                                 " -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<ocl::Kernel> kernels;
-    for (int i = 0; i<iterations; i++)
+    for (int i = 0; i < iterations; i++)
     {
         ocl::Kernel k("morph", ocl::imgproc::morph_oclsrc, buildOptions);
         if (k.empty())
@@ -1341,38 +1342,35 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
     _dst.create(src.size(), src.type());
     UMat dst = _dst.getUMat();
 
-    if( iterations== 1 && src.u != dst.u)
+    if (iterations == 1 && src.u != dst.u)
     {
         Size wholesize;
         Point ofs;
         src.locateROI(wholesize, ofs);
         int wholecols = wholesize.width, wholerows = wholesize.height;
 
-        int idxArg = 0;
+        kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst),
-        idxArg = kernels[0].set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src));
+                        ofs.x, ofs.y, src.cols, src.rows, ocl::KernelArg::PtrReadOnly(kernel8U),
-        idxArg = kernels[0].set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst));
+                        wholecols, wholerows);
-        idxArg = kernels[0].set(idxArg, ofs.x);
-        idxArg = kernels[0].set(idxArg, ofs.y);
-        idxArg = kernels[0].set(idxArg, src.cols);
-        idxArg = kernels[0].set(idxArg, src.rows);
-        idxArg = kernels[0].set(idxArg, ocl::KernelArg::PtrReadOnly(kernel8U));
-        idxArg = kernels[0].set(idxArg, wholecols);
-        idxArg = kernels[0].set(idxArg, wholerows);
 
         return kernels[0].run(2, globalThreads, localThreads, false);
     }
 
-    for(int i = 0; i< iterations; i++)
+    for (int i = 0; i < iterations; i++)
     {
         UMat source;
         Size wholesize;
         Point ofs;
-        if( i == 0)
+
+        if (i == 0)
         {
             int cols =  src.cols, rows = src.rows;
             src.locateROI(wholesize,ofs);
             src.adjustROI(ofs.y, wholesize.height - rows - ofs.y, ofs.x, wholesize.width - cols - ofs.x);
-            src.copyTo(source);
+            if(src.u != dst.u)
+                source = src;
+            else
+                src.copyTo(source);
             src.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
             source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
         }
@@ -1385,20 +1383,11 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
             dst.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
             source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
         }
-
         source.locateROI(wholesize, ofs);
-        int wholecols = wholesize.width, wholerows = wholesize.height;
 
-        int idxArg = 0;
+        kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst),
-        idxArg = kernels[i].set(idxArg, ocl::KernelArg::ReadOnlyNoSize(source));
+                        ofs.x, ofs.y, source.cols, source.rows, ocl::KernelArg::PtrReadOnly(kernel8U),
-        idxArg = kernels[i].set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst));
+                        wholesize.width, wholesize.height);
-        idxArg = kernels[i].set(idxArg, ofs.x);
-        idxArg = kernels[i].set(idxArg, ofs.y);
-        idxArg = kernels[i].set(idxArg, source.cols);
-        idxArg = kernels[i].set(idxArg, source.rows);
-        idxArg = kernels[i].set(idxArg, ocl::KernelArg::PtrReadOnly(kernel8U));
-        idxArg = kernels[i].set(idxArg, wholecols);
-        idxArg = kernels[i].set(idxArg, wholerows);
 
         if (!kernels[i].run(2, globalThreads, localThreads, false))
             return false;
@@ -1414,7 +1403,7 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
                      int borderType, const Scalar& borderValue )
 {
 #ifdef HAVE_OPENCL
-    int src_type = _src.type(), dst_type = _dst.type(),
+    int src_type = _src.type(),
         src_cn = CV_MAT_CN(src_type), src_depth = CV_MAT_DEPTH(src_type);
 #endif
 
@@ -1427,13 +1416,13 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
         return;
 #endif
 
-    if( iterations == 0 || kernel.rows*kernel.cols == 1 )
+    if (iterations == 0 || kernel.rows*kernel.cols == 1)
     {
         _src.copyTo(_dst);
         return;
     }
 
-    if( !kernel.data )
+    if (!kernel.data)
     {
         kernel = getStructuringElement(MORPH_RECT, Size(1+iterations*2,1+iterations*2));
         anchor = Point(iterations, iterations);
@@ -1449,8 +1438,7 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
         iterations = 1;
     }
 
-    CV_OCL_RUN(_dst.isUMat() && _src.size() == _dst.size() && src_type == dst_type &&
+    CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && src_cn <= 4 &&
-               _src.dims() <= 2 && (src_cn == 1 || src_cn == 4) &&
                (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) &&
                borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() &&
                (op == MORPH_ERODE || op == MORPH_DILATE),

modules/imgproc/src/opencl/canny.cl

@@ -43,6 +43,18 @@
 //
 //M*/
 
+#ifdef OP_SOBEL
+
+#if cn != 3
+#define loadpix(addr) convertToIntT(*(__global const ucharT *)(addr))
+#define storepix(val, addr) *(__global shortT *)(addr) = convertToShortT(val)
+#define shortSize (int)sizeof(shortT)
+#else
+#define loadpix(addr) convertToIntT(vload3(0, (__global const uchar *)(addr)))
+#define storepix(val, addr) vstore3(convertToShortT(val), 0, (__global short *)(addr))
+#define shortSize (int)sizeof(short) * cn
+#endif
+
 // Smoothing perpendicular to the derivative direction with a triangle filter
 // only support 3x3 Sobel kernel
 // h (-1) =  1, h (0) =  2, h (1) =  1
@@ -54,11 +66,9 @@
 // dx_buf	output dx buffer
 // dy_buf	output dy buffer
 
-__kernel void __attribute__((reqd_work_group_size(16, 16, 1)))
+__kernel void calcSobelRowPass(__global const uchar * src, int src_step, int src_offset, int rows, int cols,
-calcSobelRowPass
+                               __global uchar * dx_buf, int dx_buf_step, int dx_buf_offset,
-    (__global const uchar * src, int src_step, int src_offset, int rows, int cols,
+                               __global uchar * dy_buf, int dy_buf_step, int dy_buf_offset)
-     __global uchar * dx_buf, int dx_buf_step, int dx_buf_offset,
-     __global uchar * dy_buf, int dy_buf_step, int dy_buf_offset)
 {
     int gidx = get_global_id(0);
     int gidy = get_global_id(1);
@@ -66,34 +76,39 @@ calcSobelRowPass
     int lidx = get_local_id(0);
     int lidy = get_local_id(1);
 
-    __local int smem[16][18];
+    __local intT smem[16][18];
 
-    smem[lidy][lidx + 1] = src[mad24(src_step, min(gidy, rows - 1), gidx + src_offset)];
+    smem[lidy][lidx + 1] = loadpix(src + mad24(src_step, min(gidy, rows - 1), mad24(gidx, cn, src_offset)));
     if (lidx == 0)
     {
-        smem[lidy][0]  = src[mad24(src_step, min(gidy, rows - 1), max(gidx - 1,  0)        + src_offset)];
+        smem[lidy][0]  = loadpix(src + mad24(src_step, min(gidy, rows - 1), mad24(max(gidx - 1,  0), cn, src_offset)));
-        smem[lidy][17] = src[mad24(src_step, min(gidy, rows - 1), min(gidx + 16, cols - 1) + src_offset)];
+        smem[lidy][17] = loadpix(src + mad24(src_step, min(gidy, rows - 1), mad24(min(gidx + 16, cols - 1), cn, src_offset)));
     }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (gidy < rows && gidx < cols)
     {
-        *(__global short *)(dx_buf + mad24(gidy, dx_buf_step, gidx * (int)sizeof(short) + dx_buf_offset)) =
+        storepix(smem[lidy][lidx + 2] - smem[lidy][lidx],
-            smem[lidy][lidx + 2] - smem[lidy][lidx];
+                 dx_buf + mad24(gidy, dx_buf_step, mad24(gidx, shortSize, dx_buf_offset)));
-        *(__global short *)(dy_buf + mad24(gidy, dy_buf_step, gidx * (int)sizeof(short) + dy_buf_offset)) =
+        storepix(mad24(2, smem[lidy][lidx + 1], smem[lidy][lidx] + smem[lidy][lidx + 2]),
-            smem[lidy][lidx] + 2 * smem[lidy][lidx + 1] + smem[lidy][lidx + 2];
+                 dy_buf + mad24(gidy, dy_buf_step, mad24(gidx, shortSize, dy_buf_offset)));
     }
 }
 
-inline int calc(short x, short y)
+#elif defined OP_MAG_BUF || defined OP_MAG
+
+inline intT calc(shortT x, shortT y)
 {
 #ifdef L2GRAD
-    return x * x + y * y;
+    intT intx = convertToIntT(x), inty = convertToIntT(y);
+    return intx * intx + inty * inty;
 #else
-    return (x >= 0 ? x : -x) + (y >= 0 ? y : -y);
+    return convertToIntT( (x >= (shortT)(0) ? x : -x) + (y >= (shortT)(0) ? y : -y) );
 #endif
 }
 
+#ifdef OP_MAG
+
 // calculate the magnitude of the filter pass combining both x and y directions
 // This is the non-buffered version(non-3x3 sobel)
 //
@@ -112,18 +127,43 @@ __kernel void calcMagnitude(__global const uchar * dxptr, int dx_step, int dx_of
 
     if (y < rows && x < cols)
     {
-        int dx_index = mad24(dx_step, y, x * (int)sizeof(short) + dx_offset);
+        int dx_index = mad24(dx_step, y, mad24(x, (int)sizeof(short) * cn, dx_offset));
-        int dy_index = mad24(dy_step, y, x * (int)sizeof(short) + dy_offset);
+        int dy_index = mad24(dy_step, y, mad24(x, (int)sizeof(short) * cn, dy_offset));
-        int mag_index = mad24(mag_step, y + 1, (x + 1) * (int)sizeof(int) + mag_offset);
+        int mag_index = mad24(mag_step, y + 1, mad24(x + 1, (int)sizeof(int), mag_offset));
 
-        __global const short * dx = (__global const short *)(dxptr + dx_index);
+        __global short * dx = (__global short *)(dxptr + dx_index);
-        __global const short * dy = (__global const short *)(dyptr + dy_index);
+        __global short * dy = (__global short *)(dyptr + dy_index);
         __global int * mag = (__global int *)(magptr + mag_index);
 
-        mag[0] = calc(dx[0], dy[0]);
+        int cmag = calc(dx[0], dy[0]);
+#if cn > 1
+        short cx = dx[0], cy = dy[0];
+        int pmag;
+
+        #pragma unroll
+        for (int i = 1; i < cn; ++i)
+        {
+            pmag = calc(dx[i], dy[i]);
+            if (pmag > cmag)
+                cmag = pmag, cx = dx[i], cy = dy[i];
+        }
+
+        dx[0] = cx, dy[0] = cy;
+#endif
+        mag[0] = cmag;
     }
 }
 
+#elif defined OP_MAG_BUF
+
+#if cn != 3
+#define loadpix(addr) *(__global const shortT *)(addr)
+#define shortSize (int)sizeof(shortT)
+#else
+#define loadpix(addr) vload3(0, (__global const short *)(addr))
+#define shortSize (int)sizeof(short)*cn
+#endif
+
 // calculate the magnitude of the filter pass combining both x and y directions
 // This is the buffered version(3x3 sobel)
 //
@@ -132,59 +172,64 @@ __kernel void calcMagnitude(__global const uchar * dxptr, int dx_step, int dx_of
 // dx			direvitive in x direction output
 // dy			direvitive in y direction output
 // mag			magnitude direvitive of xy output
-__kernel void __attribute__((reqd_work_group_size(16, 16, 1)))
+__kernel void calcMagnitude_buf(__global const uchar * dx_buf, int dx_buf_step, int dx_buf_offset,
-calcMagnitude_buf
+                                __global const uchar * dy_buf, int dy_buf_step, int dy_buf_offset,
-    (__global const short * dx_buf, int dx_buf_step, int dx_buf_offset,
+                                __global uchar * dx, int dx_step, int dx_offset,
-     __global const short * dy_buf, int dy_buf_step, int dy_buf_offset,
+                                __global uchar * dy, int dy_step, int dy_offset,
-     __global short * dx, int dx_step, int dx_offset,
+                                __global uchar * mag, int mag_step, int mag_offset, int rows, int cols)
-     __global short * dy, int dy_step, int dy_offset,
-     __global int * mag, int mag_step, int mag_offset,
-     int rows, int cols)
 {
-    dx_buf_step    /= sizeof(*dx_buf);
-    dx_buf_offset  /= sizeof(*dx_buf);
-    dy_buf_step    /= sizeof(*dy_buf);
-    dy_buf_offset  /= sizeof(*dy_buf);
-    dx_step    /= sizeof(*dx);
-    dx_offset  /= sizeof(*dx);
-    dy_step    /= sizeof(*dy);
-    dy_offset  /= sizeof(*dy);
-    mag_step   /= sizeof(*mag);
-    mag_offset /= sizeof(*mag);
-
     int gidx = get_global_id(0);
     int gidy = get_global_id(1);
 
     int lidx = get_local_id(0);
     int lidy = get_local_id(1);
 
-    __local short sdx[18][16];
+    __local shortT sdx[18][16];
-    __local short sdy[18][16];
+    __local shortT sdy[18][16];
 
-    sdx[lidy + 1][lidx] = dx_buf[gidx + min(gidy, rows - 1) * dx_buf_step + dx_buf_offset];
+    sdx[lidy + 1][lidx] = loadpix(dx_buf + mad24(min(gidy, rows - 1), dx_buf_step, mad24(gidx, shortSize, dx_buf_offset)));
-    sdy[lidy + 1][lidx] = dy_buf[gidx + min(gidy, rows - 1) * dy_buf_step + dy_buf_offset];
+    sdy[lidy + 1][lidx] = loadpix(dy_buf + mad24(min(gidy, rows - 1), dy_buf_step, mad24(gidx, shortSize, dy_buf_offset)));
     if (lidy == 0)
     {
-        sdx[0][lidx]  = dx_buf[gidx + min(max(gidy - 1, 0), rows - 1) * dx_buf_step + dx_buf_offset];
+        sdx[0][lidx]  = loadpix(dx_buf + mad24(clamp(gidy - 1, 0, rows - 1), dx_buf_step, mad24(gidx, shortSize, dx_buf_offset)));
-        sdx[17][lidx] = dx_buf[gidx + min(gidy + 16, rows - 1)        * dx_buf_step + dx_buf_offset];
+        sdx[17][lidx] = loadpix(dx_buf + mad24(min(gidy + 16, rows - 1), dx_buf_step, mad24(gidx, shortSize, dx_buf_offset)));
 
-        sdy[0][lidx]  = dy_buf[gidx + min(max(gidy - 1, 0), rows - 1) * dy_buf_step + dy_buf_offset];
+        sdy[0][lidx]  = loadpix(dy_buf + mad24(clamp(gidy - 1, 0, rows - 1), dy_buf_step, mad24(gidx, shortSize, dy_buf_offset)));
-        sdy[17][lidx] = dy_buf[gidx + min(gidy + 16, rows - 1)        * dy_buf_step + dy_buf_offset];
+        sdy[17][lidx] = loadpix(dy_buf + mad24(min(gidy + 16, rows - 1), dy_buf_step, mad24(gidx, shortSize, dy_buf_offset)));
     }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (gidx < cols && gidy < rows)
     {
-        short x =  sdx[lidy][lidx] + 2 * sdx[lidy + 1][lidx] + sdx[lidy + 2][lidx];
+        shortT x = sdx[lidy + 1][lidx] * (shortT)(2) + sdx[lidy][lidx] + sdx[lidy + 2][lidx];
-        short y = -sdy[lidy][lidx] + sdy[lidy + 2][lidx];
+        shortT y = -sdy[lidy][lidx] + sdy[lidy + 2][lidx];
 
-        dx[gidx + gidy * dx_step + dx_offset] = x;
+#if cn == 1
-        dy[gidx + gidy * dy_step + dy_offset] = y;
+        *(__global short *)(dx + mad24(gidy, dx_step, mad24(gidx, shortSize, dx_offset))) = x;
+        *(__global short *)(dy + mad24(gidy, dy_step, mad24(gidx, shortSize, dy_offset))) = y;
 
-        mag[(gidx + 1) + (gidy + 1) * mag_step + mag_offset] = calc(x, y);
+        *(__global int *)(mag + mad24(gidy + 1, mag_step, mad24(gidx + 1, (int)sizeof(int), mag_offset))) = calc(x, y);
+#elif cn == 3
+        intT magv = calc(x, y);
+        short cx = x.x, cy = y.x;
+        int cmag = magv.x;
+
+        if (cmag < magv.y)
+            cx = x.y, cy = y.y, cmag = magv.y;
+        if (cmag < magv.z)
+            cx = x.z, cy = y.z, cmag = magv.z;
+
+        *(__global short *)(dx + mad24(gidy, dx_step, mad24(gidx, shortSize, dx_offset))) = cx;
+        *(__global short *)(dy + mad24(gidy, dy_step, mad24(gidx, shortSize, dy_offset))) = cy;
+
+        *(__global int *)(mag + mad24(gidy + 1, mag_step, mad24(gidx + 1, (int)sizeof(int), mag_offset))) = cmag;
+#endif
     }
 }
 
+#endif
+
+#elif defined OP_MAP
 
 //////////////////////////////////////////////////////////////////////////////////////////
 // 0.4142135623730950488016887242097 is tan(22.5)
@@ -208,13 +253,11 @@ calcMagnitude_buf
 // mag			magnitudes calculated from calcMagnitude function
 // map			output containing raw edge types
 
-__kernel void __attribute__((reqd_work_group_size(16,16,1)))
+__kernel void calcMap(__global const uchar * dx, int dx_step, int dx_offset,
-calcMap(
+                      __global const uchar * dy, int dy_step, int dy_offset,
-    __global const uchar * dx, int dx_step, int dx_offset,
+                      __global const uchar * mag, int mag_step, int mag_offset,
-    __global const uchar * dy, int dy_step, int dy_offset,
+                      __global uchar * map, int map_step, int map_offset,
-    __global const uchar * mag, int mag_step, int mag_offset,
+                      int rows, int cols, int low_thresh, int high_thresh)
-    __global uchar * map, int map_step, int map_offset,
-    int rows, int cols, int low_thresh, int high_thresh)
 {
     __local int smem[18][18];
 
@@ -227,7 +270,7 @@ calcMap(
     int grp_idx = get_global_id(0) & 0xFFFFF0;
     int grp_idy = get_global_id(1) & 0xFFFFF0;
 
-    int tid = lidx + lidy * 16;
+    int tid = mad24(lidy, 16, lidx);
     int lx = tid % 18;
     int ly = tid / 18;
 
@@ -250,8 +293,8 @@ calcMap(
 
         if (m > low_thresh)
         {
-            short xs = *(__global const short *)(dx + mad24(gidy, dx_step, dx_offset + (int)sizeof(short) * gidx));
+            short xs = *(__global const short *)(dx + mad24(gidy, dx_step, mad24(gidx, (int)sizeof(short) * cn, dx_offset)));
-            short ys = *(__global const short *)(dy + mad24(gidy, dy_step, dy_offset + (int)sizeof(short) * gidx));
+            short ys = *(__global const short *)(dy + mad24(gidy, dy_step, mad24(gidx, (int)sizeof(short) * cn, dy_offset)));
             int x = abs(xs), y = abs(ys);
 
             int tg22x = x * TG22;
@@ -278,13 +321,15 @@ calcMap(
                 }
             }
         }
-        *(__global int *)(map + mad24(map_step, gidy + 1, (gidx + 1) * (int)sizeof(int) + map_offset)) = edge_type;
+        *(__global int *)(map + mad24(map_step, gidy + 1, mad24(gidx + 1, (int)sizeof(int), + map_offset))) = edge_type;
     }
 }
 
 #undef CANNY_SHIFT
 #undef TG22
 
+#elif defined OP_HYST_LOCAL
+
 struct PtrStepSz
 {
     __global uchar * ptr;
@@ -312,11 +357,9 @@ inline void set(struct PtrStepSz data, int y, int x, int value)
 // stack	the potiential edge points found in this kernel call
 // counter	the number of potiential edge points
 
-__kernel void __attribute__((reqd_work_group_size(16,16,1)))
+__kernel void edgesHysteresisLocal(__global uchar * map_ptr, int map_step, int map_offset,
-edgesHysteresisLocal
+                                   __global ushort2 * st, __global unsigned int * counter,
-    (__global uchar * map_ptr, int map_step, int map_offset,
+                                   int rows, int cols)
-     __global ushort2 * st, __global unsigned int * counter,
-    int rows, int cols)
 {
     struct PtrStepSz map = { map_ptr + map_offset, map_step, rows + 1, cols + 1 };
 
@@ -402,6 +445,8 @@ edgesHysteresisLocal
     }
 }
 
+#elif defined OP_HYST_GLOBAL
+
 __constant int c_dx[8] = {-1,  0,  1, -1, 1, -1, 0, 1};
 __constant int c_dy[8] = {-1, -1, -1,  0, 0,  1, 1, 1};
 
@@ -409,10 +454,9 @@ __constant int c_dy[8] = {-1, -1, -1,  0, 0,  1, 1, 1};
 #define stack_size 512
 #define map_index mad24(map_step, pos.y, pos.x * (int)sizeof(int))
 
-__kernel void __attribute__((reqd_work_group_size(128, 1, 1)))
+__kernel void edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset,
-edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset,
+                                    __global ushort2 * st1, __global ushort2 * st2, __global int * counter,
-    __global ushort2 * st1, __global ushort2 * st2, __global int * counter,
+                                    int rows, int cols, int count)
-    int rows, int cols, int count)
 {
     map += map_offset;
 
@@ -492,6 +536,8 @@ edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset,
 #undef map_index
 #undef stack_size
 
+#elif defined OP_EDGES
+
 // Get the edge result. egde type of value 2 will be marked as an edge point and set to 255. Otherwise 0.
 // map		edge type mappings
 // dst		edge output
@@ -504,7 +550,7 @@ __kernel void getEdges(__global const uchar * mapptr, int map_step, int map_offs
 
     if (y < rows && x < cols)
     {
-        int map_index = mad24(map_step, y + 1, (x + 1) * (int)sizeof(int) + map_offset);
+        int map_index = mad24(map_step, y + 1, mad24(x + 1, (int)sizeof(int), map_offset));
         int dst_index = mad24(dst_step, y, x + dst_offset);
 
         __global const int * map = (__global const int *)(mapptr + map_index);
@@ -512,3 +558,5 @@ __kernel void getEdges(__global const uchar * mapptr, int map_step, int map_offs
         dst[dst_index] = (uchar)(-(map[0] >> 1));
     }
 }
+
+#endif

modules/imgproc/src/opencl/filterSepCol.cl

@@ -34,47 +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
-#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
 
-/**********************************************************************************
+#if CN != 3
-These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur.
+#define loadpix(addr) *(__global const srcT *)(addr)
-Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle
+#define storepix(val, addr)  *(__global dstT *)(addr) = val
-kernel must be in the center. ROI is not supported either.
+#define SRCSIZE (int)sizeof(srcT)
-Each kernels read 4 elements(not 4 pixels), save them to LDS and read the data needed
+#define DSTSIZE (int)sizeof(dstT)
-from LDS to calculate the result.
+#else
-The length of the convovle kernel supported is only related to the MAX size of LDS,
+#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
-which is HW related.
+#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
-Niko
+#define SRCSIZE (int)sizeof(srcT1)*3
-6/29/2011
+#define DSTSIZE (int)sizeof(dstT1)*3
-The info above maybe obsolete.
+#endif
-***********************************************************************************/
 
 #define DIG(a) a,
 __constant float mat_kernel[] = { COEFF };
 
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
+__kernel void col_filter(__global const uchar * src, int src_step, int src_offset, int src_whole_rows, int src_whole_cols,
-                        (__global const GENTYPE_SRC * restrict src,
+                         __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, float delta)
-                         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)
 {
     int x = get_global_id(0);
     int y = get_global_id(1);
@@ -82,38 +71,38 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
     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 start_addr = mad24(y, src_step, x * SRCSIZE);
-    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);
 
-    int i;
+    srcT sum, temp[READ_TIMES_COL];
-    GENTYPE_SRC sum, temp[READ_TIMES_COL];
+    __local srcT LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
-    __local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
 
-    //read pixels from src
+    // read pixels from src
-    for(i = 0;i<READ_TIMES_COL;i++)
+    for (int i = 0; i < READ_TIMES_COL; ++i)
     {
-        int current_addr = start_addr+i*LSIZE1*src_step_in_pixel;
+        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
-    for(i = 0;i<READ_TIMES_COL;i++)
-    {
-        LDS_DAT[l_y+i*LSIZE1][l_x] = temp[i];
     }
+
+    // save pixels to lds
+    for (int i = 0; i < READ_TIMES_COL; ++i)
+        LDS_DAT[mad24(i, LSIZE1, l_y)][l_x] = temp[i];
     barrier(CLK_LOCAL_MEM_FENCE);
-    //read pixels from lds and calculate the result
+
-    sum = LDS_DAT[l_y+RADIUSY][l_x]*mat_kernel[RADIUSY];
+    // read pixels from lds and calculate the result
-    for(i=1;i<=RADIUSY;i++)
+    sum = LDS_DAT[l_y + RADIUSY][l_x] * mat_kernel[RADIUSY];
+    for (int i = 1; i <= RADIUSY; ++i)
     {
-        temp[0]=LDS_DAT[l_y+RADIUSY-i][l_x];
+        temp[0] = LDS_DAT[l_y + RADIUSY - i][l_x];
-        temp[1]=LDS_DAT[l_y+RADIUSY+i][l_x];
+        temp[1] = LDS_DAT[l_y + RADIUSY + i][l_x];
-        sum += temp[0] * mat_kernel[RADIUSY-i]+temp[1] * mat_kernel[RADIUSY+i];
+        sum += mad(temp[0], mat_kernel[RADIUSY - i], temp[1] * mat_kernel[RADIUSY + i]);
     }
-    //write the result to dst
+
-    if((x<dst_cols) & (y<dst_rows))
+    // 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);
+        start_addr = mad24(y, dst_step, mad24(DSTSIZE, x, dst_offset));
-        dst[start_addr] = convert_to_DST(sum);
+        storepix(convertToDstT(sum + (srcT)(delta)), dst + start_addr);
     }
 }

modules/imgproc/src/opencl/filterSepRow.cl

@@ -34,41 +34,37 @@
 //
 //
 
-#define READ_TIMES_ROW ((2*(RADIUSX+LSIZE0)-1)/LSIZE0) //for c4 only
+#ifdef DOUBLE_SUPPORT
-#define READ_TIMES_COL ((2*(RADIUSY+LSIZE1)-1)/LSIZE1)
+#ifdef cl_amd_fp64
-//#pragma OPENCL EXTENSION cl_amd_printf : enable
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
-#define RADIUS 1
+#elif defined (cl_khr_fp64)
-#if CN ==1
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
-#define ALIGN (((RADIUS)+3)>>2<<2)
+#endif
-#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
 
+#define READ_TIMES_ROW ((2*(RADIUSX+LSIZE0)-1)/LSIZE0) //for c4 only
+#define RADIUS 1
+
 #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 +123,56 @@
     #endif //BORDER_CONSTANT
 #endif //EXTRA_EXTRAPOLATION
 
-/**********************************************************************************
+#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
+#if CN != 3
-kernel must be in the center. ROI is not supported either.
+#define loadpix(addr) *(__global const srcT *)(addr)
-For channels =1,2,4, each kernels read 4 elements(not 4 pixels), and for channels =3,
+#define storepix(val, addr)  *(__global dstT *)(addr) = val
-the kernel read 4 pixels, save them to LDS and read the data needed from LDS to
+#define SRCSIZE (int)sizeof(srcT)
-calculate the result.
+#define DSTSIZE (int)sizeof(dstT)
-The length of the convovle kernel supported is related to the LSIZE0 and the MAX size
+#else
-of LDS, which is HW related.
+#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
-For channels = 1,3 the RADIUS is no more than LSIZE0*2
+#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
-For channels = 2, the RADIUS is no more than LSIZE0
+#define SRCSIZE (int)sizeof(srcT1)*3
-For channels = 4, arbitary RADIUS is supported unless the LDS is not enough
+#define DSTSIZE (int)sizeof(dstT1)*3
-Niko
+#endif
-6/29/2011
-The info above maybe obsolete.
-***********************************************************************************/
 
 #define DIG(a) a,
 __constant float mat_kernel[] = { COEFF };
 
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D0
+__kernel void row_filter_C1_D0(__global const uchar * src, int src_step_in_pixel, int src_offset_x, int src_offset_y,
-    (__global uchar * restrict src,
+                               int src_cols, int src_rows, int src_whole_cols, int src_whole_rows,
-     int src_step_in_pixel,
+                               __global float * dst, int dst_step_in_pixel, int dst_cols, int dst_rows,
-     int src_offset_x, int src_offset_y,
+                               int radiusy)
-     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;
+        int current_addr = mad24(i, LSIZE0 << 2, start_addr);
-        current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
+        current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0;
-        temp[i] = *(__global uchar4*)&src[current_addr];
+        temp[i] = *(__global const uchar4 *)&src[current_addr];
     }
 
     // judge if read out of boundary
 #ifdef BORDER_ISOLATED
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
     {
         temp[i].x = ELEM(start_x+i*LSIZE0*4,   src_offset_x, src_offset_x + src_cols, 0,         temp[i].x);
         temp[i].y = ELEM(start_x+i*LSIZE0*4+1, src_offset_x, src_offset_x + src_cols, 0,         temp[i].y);
@@ -194,7 +181,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
         temp[i]   = ELEM(start_y,              src_offset_y, src_offset_y + src_rows, (uchar4)0, temp[i]);
     }
 #else
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
     {
         temp[i].x = ELEM(start_x+i*LSIZE0*4,   0, src_whole_cols, 0,         temp[i].x);
         temp[i].y = ELEM(start_x+i*LSIZE0*4+1, 0, src_whole_cols, 0,         temp[i].y);
@@ -209,16 +196,15 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
 #else
     int not_all_in_range = (start_x<0) | (start_x + READ_TIMES_ROW*LSIZE0*4+4>src_whole_cols)| (start_y<0) | (start_y >= src_whole_rows);
 #endif
-    int4 index[READ_TIMES_ROW];
+    int4 index[READ_TIMES_ROW], addr;
-    int4 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 +217,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 +226,9 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
 #endif
 
         // read pixels from src
-        for (i = 0; i<READ_TIMES_ROW; i++)
+        for (int i = 0; i < READ_TIMES_ROW; ++i)
         {
-            addr = mad24((int4)s_y,(int4)src_step_in_pixel,index[i]);
+            addr = mad24((int4)s_y, (int4)src_step_in_pixel, index[i]);
             temp[i].x = src[addr.x];
             temp[i].y = src[addr.y];
             temp[i].z = src[addr.z];
@@ -251,26 +238,26 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
     else
     {
         // read pixels from src
-        for (i = 0; i<READ_TIMES_ROW; i++)
+        for (int i = 0; i < READ_TIMES_ROW; ++i)
-            temp[i] = *(__global uchar4*)&src[start_addr+i*LSIZE0*4];
+            temp[i] = *(__global uchar4*)&src[mad24(i, LSIZE0 << 2, start_addr)];
     }
 #endif //BORDER_CONSTANT
 
     // save pixels to lds
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
-        LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
+        LDS_DAT[l_y][mad24(i, LSIZE0, l_x)] = temp[i];
     barrier(CLK_LOCAL_MEM_FENCE);
 
     // read pixels from lds and calculate the result
-    sum =convert_float4(vload4(0,(__local uchar*)&LDS_DAT[l_y][l_x]+RADIUSX+offset))*mat_kernel[RADIUSX];
+    sum = convert_float4(vload4(0,(__local uchar *)&LDS_DAT[l_y][l_x]+RADIUSX+offset)) * mat_kernel[RADIUSX];
-    for (i=1; i<=RADIUSX; i++)
+    for (int i = 1; i <= RADIUSX; ++i)
     {
         temp[0] = vload4(0, (__local uchar*)&LDS_DAT[l_y][l_x] + RADIUSX + offset - i);
         temp[1] = vload4(0, (__local uchar*)&LDS_DAT[l_y][l_x] + RADIUSX + offset + i);
-        sum += convert_float4(temp[0]) * mat_kernel[RADIUSX-i] + convert_float4(temp[1]) * mat_kernel[RADIUSX+i];
+        sum += mad(convert_float4(temp[0]), mat_kernel[RADIUSX-i], convert_float4(temp[1]) * mat_kernel[RADIUSX + i]);
     }
 
-    start_addr = mad24(y,dst_step_in_pixel,x);
+    start_addr = mad24(y, dst_step_in_pixel, x);
 
     // write the result to dst
     if ((x+3<dst_cols) & (y<dst_rows))
@@ -290,63 +277,58 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
         dst[start_addr] = sum.x;
 }
 
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C4_D0
+__kernel void row_filter(__global const uchar * src, int src_step, int src_offset_x, int src_offset_y,
-    (__global uchar4 * restrict src,
+                         int src_cols, int src_rows, int src_whole_cols, int src_whole_rows,
-     int src_step_in_pixel,
+                         __global uchar * dst, int dst_step, int dst_cols, int dst_rows,
-     int src_offset_x, int src_offset_y,
+                         int radiusy)
-     int src_cols, int src_rows,
-     int src_whole_cols, int src_whole_rows,
-     __global float4 * dst,
-     int dst_step_in_pixel,
-     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 i;
-    float4 sum;
-    uchar4 temp[READ_TIMES_ROW];
 
-    __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+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, 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 (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; i++)
     {
-        int current_addr = start_addr+i*LSIZE0;
+        int current_addr = mad24(i, LSIZE0 * SRCSIZE, start_addr);
-        current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
+        current_addr = current_addr < end_addr && current_addr >= 0 ? current_addr : 0;
-        temp[i] = src[current_addr];
+        temp[i] = loadpix(src + current_addr);
     }
 
-    //judge if read out of boundary
+    // judge if read out of boundary
 #ifdef BORDER_ISOLATED
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
     {
-        temp[i]= ELEM(start_x+i*LSIZE0, src_offset_x, src_offset_x + src_cols, (uchar4)0, temp[i]);
+        temp[i] = ELEM(mad24(i, LSIZE0, start_x), src_offset_x, src_offset_x + src_cols, (srcT)(0), temp[i]);
-        temp[i]= ELEM(start_y,          src_offset_y, src_offset_y + src_rows, (uchar4)0, temp[i]);
+        temp[i] = ELEM(start_y,                   src_offset_y, src_offset_y + src_rows, (srcT)(0), temp[i]);
     }
 #else
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
     {
-        temp[i]= ELEM(start_x+i*LSIZE0, 0, src_whole_cols, (uchar4)0, temp[i]);
+        temp[i] = ELEM(mad24(i, LSIZE0, start_x), 0, src_whole_cols, (srcT)(0), temp[i]);
-        temp[i]= ELEM(start_y,          0, src_whole_rows, (uchar4)0, temp[i]);
+        temp[i] = ELEM(start_y,                   0, src_whole_rows, (srcT)(0), temp[i]);
     }
 #endif
 #else
-    int index[READ_TIMES_ROW];
+    int index[READ_TIMES_ROW], s_x, s_y;
-    int s_x,s_y;
 
     // judge if read out of boundary
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
     {
-        s_x = start_x+i*LSIZE0;
+        s_x = mad24(i, LSIZE0, start_x);
         s_y = start_y;
+
 #ifdef BORDER_ISOLATED
         EXTRAPOLATE(s_x, src_offset_x, src_offset_x + src_cols);
         EXTRAPOLATE(s_y, src_offset_y, src_offset_y + src_rows);
@@ -354,216 +336,32 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_
         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 (i = 0; i<READ_TIMES_ROW; i++)
-        temp[i] = src[index[i]];
-#endif //BORDER_CONSTANT
-
-    //save pixels to lds
-    for (i = 0; i<READ_TIMES_ROW; i++)
-        LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    //read pixels from lds and calculate the result
-    sum =convert_float4(LDS_DAT[l_y][l_x+RADIUSX])*mat_kernel[RADIUSX];
-    for (i=1; i<=RADIUSX; i++)
-    {
-        temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
-        temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
-        sum += convert_float4(temp[0])*mat_kernel[RADIUSX-i]+convert_float4(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;
-    }
-}
-
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D5
-    (__global float * 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)
-{
-    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 i;
-    float sum;
-    float temp[READ_TIMES_ROW];
-
-    __local float 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)
-    {
+        temp[i] = loadpix(src + index[i]);
-        int current_addr = start_addr+i*LSIZE0;
+#endif // BORDER_CONSTANT
-        current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
-        temp[i] = src[current_addr];
-    }
-
-    // judge if read out of boundary
-#ifdef BORDER_ISOLATED
-    for (i = 0; i<READ_TIMES_ROW; i++)
-    {
-        temp[i]= ELEM(start_x+i*LSIZE0, src_offset_x, src_offset_x + src_cols, (float)0,temp[i]);
-        temp[i]= ELEM(start_y,          src_offset_y, src_offset_y + src_rows, (float)0,temp[i]);
-    }
-#else
-    for (i = 0; i<READ_TIMES_ROW; i++)
-    {
-        temp[i]= ELEM(start_x+i*LSIZE0, 0, src_whole_cols, (float)0,temp[i]);
-        temp[i]= ELEM(start_y,          0, src_whole_rows, (float)0,temp[i]);
-    }
-#endif
-#else // BORDER_CONSTANT
-    int index[READ_TIMES_ROW];
-    int s_x,s_y;
-    // judge if read out of boundary
-    for (i = 0; i<READ_TIMES_ROW; i++)
-    {
-        s_x = start_x + i*LSIZE0, s_y = start_y;
-#ifdef BORDER_ISOLATED
-        EXTRAPOLATE(s_x, src_offset_x, src_offset_x + src_cols);
-        EXTRAPOLATE(s_y, src_offset_y, src_offset_y + src_rows);
-#else
-        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);
-    }
-    // read pixels from src
-    for (i = 0; i<READ_TIMES_ROW; i++)
-        temp[i] = src[index[i]];
-#endif// BORDER_CONSTANT
-
-    //save pixels to lds
-    for (i = 0; i<READ_TIMES_ROW; i++)
-        LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    // read pixels from lds and calculate the result
-    sum =LDS_DAT[l_y][l_x+RADIUSX]*mat_kernel[RADIUSX];
-    for (i=1; i<=RADIUSX; i++)
-    {
-        temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
-        temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
-        sum += temp[0]*mat_kernel[RADIUSX-i]+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;
-    }
-}
-
-__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C4_D5
-    (__global float4 * 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 float4 * dst,
-     int dst_step_in_pixel,
-     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 i;
-    float4 sum;
-    float4 temp[READ_TIMES_ROW];
-
-    __local float4 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++)
-    {
-        int current_addr = start_addr+i*LSIZE0;
-        current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0;
-        temp[i] = src[current_addr];
-    }
-
-    // judge if read out of boundary
-#ifdef BORDER_ISOLATED
-    for (i = 0; i<READ_TIMES_ROW; i++)
-    {
-        temp[i]= ELEM(start_x+i*LSIZE0, src_offset_x, src_offset_x + src_cols, (float4)0,temp[i]);
-        temp[i]= ELEM(start_y,          src_offset_y, src_offset_y + src_rows, (float4)0,temp[i]);
-    }
-#else
-    for (i = 0; i<READ_TIMES_ROW; i++)
-    {
-        temp[i]= ELEM(start_x+i*LSIZE0, 0, src_whole_cols, (float4)0,temp[i]);
-        temp[i]= ELEM(start_y,          0, src_whole_rows, (float4)0,temp[i]);
-    }
-#endif
-#else
-    int index[READ_TIMES_ROW];
-    int s_x,s_y;
-
-    // judge if read out of boundary
-    for (i = 0; i<READ_TIMES_ROW; i++)
-    {
-        s_x = start_x + i*LSIZE0, s_y = start_y;
-#ifdef BORDER_ISOLATED
-        EXTRAPOLATE(s_x, src_offset_x, src_offset_x + src_cols);
-        EXTRAPOLATE(s_y, src_offset_y, src_offset_y + src_rows);
-#else
-        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);
-    }
-    // read pixels from src
-    for (i = 0; i<READ_TIMES_ROW; i++)
-        temp[i] = src[index[i]];
-#endif
 
     // save pixels to lds
-    for (i = 0; i<READ_TIMES_ROW; i++)
+    for (int i = 0; i < READ_TIMES_ROW; ++i)
-        LDS_DAT[l_y][l_x+i*LSIZE0]=temp[i];
+        LDS_DAT[l_y][mad24(i, LSIZE0, l_x)] = temp[i];
     barrier(CLK_LOCAL_MEM_FENCE);
 
     // read pixels from lds and calculate the result
-    sum =LDS_DAT[l_y][l_x+RADIUSX]*mat_kernel[RADIUSX];
+    sum = convertToDstT(LDS_DAT[l_y][l_x + RADIUSX]) * mat_kernel[RADIUSX];
-    for (i=1; i<=RADIUSX; i++)
+    for (int i = 1; i <= RADIUSX; ++i)
     {
-        temp[0]=LDS_DAT[l_y][l_x+RADIUSX-i];
+        temp[0] = LDS_DAT[l_y][l_x + RADIUSX - i];
-        temp[1]=LDS_DAT[l_y][l_x+RADIUSX+i];
+        temp[1] = LDS_DAT[l_y][l_x + RADIUSX + i];
-        sum += temp[0]*mat_kernel[RADIUSX-i]+temp[1]*mat_kernel[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)
+    if (x < dst_cols && y < dst_rows)
     {
-        start_addr = mad24(y,dst_step_in_pixel,x);
+        start_addr = mad24(y, dst_step, x * DSTSIZE);
-        dst[start_addr] = sum;
+        storepix(sum, dst + start_addr);
     }
 }

modules/imgproc/src/opencl/filterSep_singlePass.cl

@@ -0,0 +1,187 @@
+/*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
+
+#if CN != 3
+#define loadpix(addr) *(__global const srcT *)(addr)
+#define storepix(val, addr)  *(__global dstT *)(addr) = val
+#define SRCSIZE (int)sizeof(srcT)
+#define DSTSIZE (int)sizeof(dstT)
+#else
+#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
+#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
+#define SRCSIZE (int)sizeof(srcT1)*3
+#define DSTSIZE (int)sizeof(dstT1)*3
+#endif
+
+#define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x)))
+
+#ifdef BORDER_CONSTANT
+// CCCCCC|abcdefgh|CCCCCCC
+#define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y))
+#else
+#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y))
+#endif
+
+#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, float delta)
+{
+    // 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 = get_global_id(0);
+    int y = 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
+    storepix(convertToDstT(sum + (WT)(delta)), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset)));
+}

modules/imgproc/src/opencl/morph.cl

@@ -43,6 +43,16 @@
 #endif
 #endif
 
+#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
+
 #ifdef DEPTH_0
 #ifdef ERODE
 #define VAL 255
@@ -50,16 +60,14 @@
 #ifdef DILATE
 #define VAL 0
 #endif
-#endif
+#elif defined DEPTH_5
-#ifdef DEPTH_5
 #ifdef ERODE
 #define VAL FLT_MAX
 #endif
 #ifdef DILATE
 #define VAL -FLT_MAX
 #endif
-#endif
+#elif defined DEPTH_6
-#ifdef DEPTH_6
 #ifdef ERODE
 #define VAL DBL_MAX
 #endif
@@ -69,84 +77,80 @@
 #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
-//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 * 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 src_offset_x, int src_offset_y, int cols, int rows,
-                    int cols, int rows,
+                    __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows)
-                    __constant uchar * mat_kernel,
-                    int src_whole_cols, int src_whole_rows)
 {
-    int l_x = get_local_id(0);
+    int gidx = get_global_id(0), gidy = get_global_id(1);
-    int l_y = get_local_id(1);
+    int l_x = get_local_id(0), l_y = get_local_id(1);
-    int x = get_group_id(0)*LSIZE0;
+    int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1;
-    int y = get_group_id(1)*LSIZE1;
+    int start_x = x + src_offset_x - RADIUSX;
-    int start_x = x+src_offset_x-RADIUSX;
+    int end_x = x + src_offset_x + LSIZE0 + RADIUSX;
-    int end_x = x + src_offset_x+LSIZE0+RADIUSX;
+    int width = end_x - (x + src_offset_x - RADIUSX) + 1;
-    int width = end_x -(x+src_offset_x-RADIUSX)+1;
+    int start_y = y + src_offset_y - RADIUSY;
-    int start_y = y+src_offset_y-RADIUSY;
+    int point1 = mad24(l_y, LSIZE0, l_x);
-    int point1 = mad24(l_y,LSIZE0,l_x);
+    int point2 = point1 + LSIZE0 * LSIZE1;
-    int point2 = point1 + LSIZE0*LSIZE1;
+    int tl_x = point1 % width, tl_y = point1 / width;
-    int tl_x = point1 % width;
+    int tl_x2 = point2 % width, tl_y2 = point2 / width;
-    int tl_y = point1 / width;
+    int cur_x = start_x + tl_x, cur_y = start_y + tl_y;
-    int tl_x2 = point2 % width;
+    int cur_x2 = start_x + tl_x2, cur_y2 = start_y + tl_y2;
-    int tl_y2 = point2 / width;
+    int start_addr = mad24(cur_y, src_step, cur_x * TSIZE);
-    int cur_x = start_x + tl_x;
+    int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE);
-    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 end_addr = mad24(src_whole_rows - 1,src_step,src_whole_cols*(int)sizeof(GENTYPE));
+    __local T LDS_DAT[2*LSIZE1*LSIZE0];
-    //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);
 
-    temp1= ELEM(cur_x2,0,src_whole_cols,(GENTYPE)VAL,temp1);
+    // read pixels from src
-    temp1= ELEM(cur_y2,0,src_whole_rows,(GENTYPE)VAL,temp1);
+    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++)
+    T res = (T)(VAL);
-        for(int j=0; j<2*RADIUSX+1; j++)
+    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<cols && gidy<rows)
-    {
-        int dst_index = mad24(gidy, dst_step, dst_offset + gidx * (int)sizeof(GENTYPE));
-        __global GENTYPE * dst = (__global GENTYPE *)(dstptr + dst_index);
-        dst[0] = res;
-    }
 
+    if (gidx < cols && gidy < rows)
+    {
+        int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset));
+        storepix(res, dstptr + dst_index);
+    }
 }

modules/imgproc/test/ocl/test_canny.cpp

@@ -58,9 +58,9 @@ IMPLEMENT_PARAM_CLASS(AppertureSize, int)
 IMPLEMENT_PARAM_CLASS(L2gradient, bool)
 IMPLEMENT_PARAM_CLASS(UseRoi, bool)
 
-PARAM_TEST_CASE(Canny, AppertureSize, L2gradient, UseRoi)
+PARAM_TEST_CASE(Canny, Channels, AppertureSize, L2gradient, UseRoi)
 {
-    int apperture_size;
+    int cn, apperture_size;
     bool useL2gradient, use_roi;
 
     TEST_DECLARE_INPUT_PARAMETER(src);
@@ -68,19 +68,19 @@ PARAM_TEST_CASE(Canny, AppertureSize, L2gradient, UseRoi)
 
     virtual void SetUp()
     {
-        apperture_size = GET_PARAM(0);
+        cn = GET_PARAM(0);
-        useL2gradient = GET_PARAM(1);
+        apperture_size = GET_PARAM(1);
-        use_roi = GET_PARAM(2);
+        useL2gradient = GET_PARAM(2);
+        use_roi = GET_PARAM(3);
     }
 
     void generateTestData()
     {
-        Mat img = readImage("shared/fruits.png", IMREAD_GRAYSCALE);
+        Mat img = readImageType("shared/fruits.png", CV_8UC(cn));
         ASSERT_FALSE(img.empty()) << "cann't load shared/fruits.png";
 
         Size roiSize = img.size();
         int type = img.type();
-        ASSERT_EQ(CV_8UC1, type);
 
         Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
         randomSubMat(src, src_roi, roiSize, srcBorder, type, 2, 100);
@@ -108,6 +108,7 @@ OCL_TEST_P(Canny, Accuracy)
 }
 
 OCL_INSTANTIATE_TEST_CASE_P(ImgProc, Canny, testing::Combine(
+                                testing::Values(1, 3),
                                 testing::Values(AppertureSize(3), AppertureSize(5)),
                                 testing::Values(L2gradient(false), L2gradient(true)),
                                 testing::Values(UseRoi(false), UseRoi(true))));

modules/imgproc/test/ocl/test_filters.cpp

@@ -306,7 +306,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),
@@ -349,28 +349,28 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine(
                             Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
-                            Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
                             Values(3, 5, 7),
                             Values(Size(0,0)),//not used
                             Values((BorderType)BORDER_CONSTANT),//not used
                             Values(1.0, 2.0, 3.0),
-                            Bool() ) );
+                            Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(
-                            Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
                             Values(3, 5, 7),
                             Values(Size(0,0)),//not used
                             Values((BorderType)BORDER_CONSTANT),//not used
                             Values(1.0, 2.0, 3.0),
-                            Bool() ) );
+                            Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine(
-                            Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
                             Values(3, 5, 7),
-                            Values(Size(0,0), Size(0,1), Size(0,2), Size(0,3), Size(0,4), Size(0,5),Size(0,6)),//uses as generator of operations
+                            Values(Size(0, 0), Size(0, 1), Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations
-                            Values((BorderType)BORDER_CONSTANT),//not used
+                            Values((BorderType)BORDER_CONSTANT),// not used
                             Values(1.0, 2.0, 3.0),
-                            Bool() ) );
+                            Bool()));
 
 
 } } // namespace cvtest::ocl

modules/imgproc/test/ocl/test_sepfilter2D.cpp

@@ -61,6 +61,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool)
     int borderType;
     bool useRoi;
     Mat kernelX, kernelY;
+    double delta;
 
     TEST_DECLARE_INPUT_PARAMETER(src);
     TEST_DECLARE_OUTPUT_PARAMETER(dst);
@@ -75,33 +76,25 @@ 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);
         cv::normalize(temp, kernelX, 1.0, 0.0, NORM_L1);
         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, MAX_VALUE, ksize.height, MAX_VALUE);
-        int rest = roiSize.width % 4;
-        if (0 != rest)
-            roiSize.width += (4 - rest);
         Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
-        rest = srcBorder.lef % 4;
-        if (0 != rest)
-            srcBorder.lef += (4 - rest);
-        rest = srcBorder.rig % 4;
-        if (0 != rest)
-            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.x = anchor.y = -1;
-        anchor.y = -1;
+        delta = randomDouble(-100, 100);
 
         UMAT_UPLOAD_INPUT_PARAMETER(src);
         UMAT_UPLOAD_OUTPUT_PARAMETER(dst);
@@ -115,22 +108,21 @@ 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 + 3; j++)
     {
         random_roi();
 
-        OCL_OFF(cv::sepFilter2D(src_roi, dst_roi, -1, kernelX, kernelY, anchor, 0.0, borderType));
+        OCL_OFF(cv::sepFilter2D(src_roi, dst_roi, -1, kernelX, kernelY, anchor, delta, borderType));
-        OCL_ON(cv::sepFilter2D(usrc_roi, udst_roi, -1, kernelX, kernelY, anchor, 0.0, borderType));
+        OCL_ON(cv::sepFilter2D(usrc_roi, udst_roi, -1, kernelX, kernelY, anchor, delta, borderType));
 
         Near(1.0);
     }
 }
 
-
 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,

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

@@ -175,7 +175,6 @@ public class NativeCameraView extends CameraBridgeViewBase {
                 }
 
                 deliverAndDrawFrame(mFrame);
-
             } while (!mStopThread);
         }
     }

modules/nonfree/perf/perf_precomp.hpp

@@ -9,12 +9,13 @@
 #ifndef __OPENCV_PERF_PRECOMP_HPP__
 #define __OPENCV_PERF_PRECOMP_HPP__
 
+#include "cvconfig.h"
+
 #include "opencv2/ts.hpp"
 #include "opencv2/nonfree.hpp"
 #include "opencv2/highgui.hpp"
 
 #include "opencv2/opencv_modules.hpp"
-#include "cvconfig.h"
 
 #ifdef HAVE_OPENCV_OCL
 #  include "opencv2/nonfree/ocl.hpp"

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"),

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())
-    {
+        localSize = Size(8, 8);
-        const ocl::Device& device = ocl::Device::getDefault();
+
-        if (device.isAMD() && !device.hostUnifiedMemory())
-            localSize = Size(8, 8);
-    }
     return true;
 }
 

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_elem_size = CV_ELEM_SIZE1(qangle.type());
-    int qangle_step = (int)qangle.step >> (1 + qangle_step_shift);
+    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<float>& _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<Point> &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);

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;
@@ -133,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))
@@ -558,7 +565,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 +667,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);

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(

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<int>(logf(static_cast<float>(min(size.width, size.height))) / logf(2.0f));

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)

modules/superres/CMakeLists.txt

@@ -1,4 +1,4 @@
-if(ANDROID OR IOS)
+if(IOS)
   ocv_module_disable(superres)
 endif()
 

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_;
 

modules/ts/src/ts_perf.cpp

@@ -20,7 +20,7 @@ static std::vector<std::string> available_impls;
 static std::string  param_impl;
 
 static enum PERF_STRATEGY strategyForce = PERF_STRATEGY_DEFAULT;
-static enum PERF_STRATEGY strategyModule = PERF_STRATEGY_BASE;
+static enum PERF_STRATEGY strategyModule = PERF_STRATEGY_SIMPLE;
 
 static double       param_max_outliers;
 static double       param_max_deviation;

modules/video/src/bgfg_gaussmix2.cpp

@@ -578,7 +578,7 @@ public:
                 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 )
@@ -643,6 +643,7 @@ public:
                                 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++ )
@@ -660,7 +661,7 @@ public:
                         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

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

modules/video/src/opencl/bgfg_mog2.cl

@@ -102,7 +102,7 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame
         {
 
             float c_weight = alpha1 * _weight[(mode * frame_row + y) * weight_step + x] + prune;
-
+            int swap_count = 0;
             if (!fitsPDF)
             {
                 float c_var = _variance[(mode * frame_row + y) * var_step + x];
@@ -132,6 +132,7 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame
                     {
                         if (c_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 (CN==1)
@@ -149,7 +150,7 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame
                 nmodes--;
             }
 
-            _weight[(mode * frame_row + y) * weight_step + x] = c_weight; //update weight by the calculated value
+            _weight[((mode - swap_count) * frame_row + y) * weight_step + x] = c_weight; //update weight by the calculated value
             totalWeight += c_weight;
         }
 

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;
 

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";

modules/viz/doc/widget.rst

@@ -1050,7 +1050,7 @@ viz::WWidgetMerger::WWidgetMerger
 ---------------------------------------
 Constructs a WWidgetMerger.
 
-.. ocv:WWidgetMerger:: WWidgetMerger()
+.. ocv:function:: WWidgetMerger()
 
 viz::WWidgetMerger::addCloud
 -------------------------------

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
+    // hue varies from 0 to 179, saturation from 0 to 255
-    float s_ranges[] = { 0, 256 };
     float h_ranges[] = { 0, 180 };
+    float s_ranges[] = { 0, 256 };
 
     const float* ranges[] = { h_ranges, s_ranges };
 

samples/gpu/CMakeLists.txt

@@ -49,8 +49,8 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
 
     target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${OPENCV_CUDA_SAMPLES_REQUIRED_DEPS})
 
-    if(HAVE_CUDA)
+    if(HAVE_CUDA AND NOT ANDROID)
-        target_link_libraries(${the_target} ${CUDA_CUDA_LIBRARY})
+      target_link_libraries(${the_target} ${CUDA_CUDA_LIBRARY})
     endif()
 
     if(HAVE_opencv_nonfree)

samples/gpu/brox_optical_flow.cpp

@@ -1,7 +1,7 @@
 #include <iostream>
 #include <iomanip>
 #include <string>
-#include <cctype>
+#include <ctype.h>
 
 #include "opencv2/core.hpp"
 #include "opencv2/core/utility.hpp"

samples/gpu/opticalflow_nvidia_api.cpp

@@ -7,6 +7,7 @@
 #include <memory>
 #include <exception>
 #include <ctime>
+#include <ctype.h>
 
 #include "cvconfig.h"
 #include <iostream>

samples/gpu/super_resolution.cpp

@@ -1,6 +1,8 @@
 #include <iostream>
 #include <iomanip>
 #include <string>
+#include <ctype.h>
+
 #include "opencv2/core.hpp"
 #include "opencv2/core/utility.hpp"
 #include "opencv2/highgui.hpp"