Merge pull request #1664 from SpecLad:merge-2.4

modules/bioinspired/CMakeLists.txt

@@ -1,2 +1,3 @@
 set(the_description "Biologically inspired algorithms")
+ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef)
 ocv_define_module(bioinspired opencv_core OPTIONAL opencv_highgui opencv_ocl)

modules/contrib/CMakeLists.txt

@@ -1 +1 @@
-ocv_define_module(contrib opencv_imgproc opencv_calib3d opencv_ml opencv_video opencv_objdetect OPTIONAL opencv_highgui)
+ocv_define_module(contrib opencv_imgproc opencv_calib3d opencv_ml opencv_video opencv_objdetect OPTIONAL opencv_highgui opencv_nonfree)

modules/contrib/src/featuretracker.cpp

@@ -45,6 +45,12 @@
 #include "opencv2/calib3d.hpp"
 #include "opencv2/contrib/hybridtracker.hpp"
 
+#ifdef HAVE_OPENCV_NONFREE
+#include "opencv2/nonfree/nonfree.hpp"
+
+static bool makeUseOfNonfree = initModule_nonfree();
+#endif
+
 using namespace cv;
 
 CvFeatureTracker::CvFeatureTracker(CvFeatureTrackerParams _params) :

modules/highgui/src/ffmpeg_codecs.hpp

@@ -60,14 +60,6 @@ extern "C" {
 #include <errno.h>
 #endif
 
-#ifdef WIN32
-#  ifdef __OPENCV_BUILD
-#    define AVUTIL_COMMON_H
-#    define MKBETAG(a,b,c,d) ((d) | ((c) << 8) | ((b) << 16) | ((unsigned)(a) << 24))
-#  endif
-#  include <libavformat/avformat.h>
-#else
-
 // if the header path is not specified explicitly, let's deduce it
 #if !defined HAVE_FFMPEG_AVCODEC_H && !defined HAVE_LIBAVCODEC_AVCODEC_H
 
@@ -81,14 +73,12 @@ extern "C" {
   #include <ffmpeg/avformat.h>
 #endif
 
-#if defined(HAVE_LIBAVFORMAT_AVFORMAT_H)
+#if defined(HAVE_LIBAVFORMAT_AVFORMAT_H) || defined(WIN32)
   #include <libavformat/avformat.h>
 #endif
 
 #endif
 
-#endif
-
 #ifdef __cplusplus
 }
 #endif

modules/highgui/src/window_QT.cpp

@@ -1736,7 +1736,8 @@ void CvWindow::displayStatusBar(QString text, int delayms)
 
 void CvWindow::enablePropertiesButton()
 {
-    vect_QActions[9]->setDisabled(false);
+    if (!vect_QActions.empty())
+        vect_QActions[9]->setDisabled(false);
 }
 
 

modules/highgui/src/window_w32.cpp

@@ -1087,12 +1087,7 @@ cvShowImage( const char* name, const CvArr* arr )
     window = icvFindWindowByName(name);
     if(!window)
     {
-        #ifndef HAVE_OPENGL
-            cvNamedWindow(name, CV_WINDOW_AUTOSIZE);
-        #else
-            cvNamedWindow(name, CV_WINDOW_AUTOSIZE | CV_WINDOW_OPENGL);
-        #endif
-
+        cvNamedWindow(name, CV_WINDOW_AUTOSIZE);
         window = icvFindWindowByName(name);
     }
 

modules/highgui/test/test_ffmpeg.cpp

@@ -47,8 +47,6 @@ using namespace cv;
 
 #ifdef HAVE_FFMPEG
 
-#include "ffmpeg_codecs.hpp"
-
 using namespace std;
 
 class CV_FFmpegWriteBigVideoTest : public cvtest::BaseTest
@@ -61,32 +59,34 @@ public:
         const double fps0 = 15;
         const double time_sec = 1;
 
-        const size_t n = sizeof(codec_bmp_tags)/sizeof(codec_bmp_tags[0]);
+        const int tags[] = {
+            0,
+            //VideoWriter::fourcc('D', 'I', 'V', '3'),
+            //VideoWriter::fourcc('D', 'I', 'V', 'X'),
+            VideoWriter::fourcc('D', 'X', '5', '0'),
+            VideoWriter::fourcc('F', 'L', 'V', '1'),
+            VideoWriter::fourcc('H', '2', '6', '1'),
+            VideoWriter::fourcc('H', '2', '6', '3'),
+            VideoWriter::fourcc('I', '4', '2', '0'),
+            //VideoWriter::fourcc('j', 'p', 'e', 'g'),
+            VideoWriter::fourcc('M', 'J', 'P', 'G'),
+            VideoWriter::fourcc('m', 'p', '4', 'v'),
+            VideoWriter::fourcc('M', 'P', 'E', 'G'),
+            //VideoWriter::fourcc('W', 'M', 'V', '1'),
+            //VideoWriter::fourcc('W', 'M', 'V', '2'),
+            VideoWriter::fourcc('X', 'V', 'I', 'D'),
+            //VideoWriter::fourcc('Y', 'U', 'Y', '2'),
+        };
+
+        const size_t n = sizeof(tags)/sizeof(tags[0]);
 
         bool created = false;
 
         for (size_t j = 0; j < n; ++j)
         {
-        stringstream s; s << codec_bmp_tags[j].tag;
-        int tag = codec_bmp_tags[j].tag;
-
-        if( tag != MKTAG('H', '2', '6', '3') &&
-            tag != MKTAG('H', '2', '6', '1') &&
-            //tag != MKTAG('D', 'I', 'V', 'X') &&
-            tag != MKTAG('D', 'X', '5', '0') &&
-            tag != MKTAG('X', 'V', 'I', 'D') &&
-            tag != MKTAG('m', 'p', '4', 'v') &&
-            //tag != MKTAG('D', 'I', 'V', '3') &&
-            //tag != MKTAG('W', 'M', 'V', '1') &&
-            //tag != MKTAG('W', 'M', 'V', '2') &&
-            tag != MKTAG('M', 'P', 'E', 'G') &&
-            tag != MKTAG('M', 'J', 'P', 'G') &&
-            //tag != MKTAG('j', 'p', 'e', 'g') &&
-            tag != 0 &&
-            tag != MKTAG('I', '4', '2', '0') &&
-            //tag != MKTAG('Y', 'U', 'Y', '2') &&
-            tag != MKTAG('F', 'L', 'V', '1') )
-            continue;
+        int tag = tags[j];
+        stringstream s;
+        s << tag;
 
         const string filename = "output_"+s.str()+".avi";
 
@@ -104,7 +104,10 @@ public:
                 frame_s = Size(1920, 1080);*/
 
             if( tag == VideoWriter::fourcc('M', 'P', 'E', 'G') )
+            {
+                frame_s = Size(720, 576);
                 fps = 25;
+            }
 
             VideoWriter writer(filename, tag, fps, frame_s);
 

modules/highgui/test/test_video_io.cpp

@@ -480,18 +480,34 @@ void CV_HighGuiTest::SpecificVideoTest(const string& dir, const cvtest::VideoFor
 
     size_t FRAME_COUNT = (size_t)cap.get(CAP_PROP_FRAME_COUNT);
 
-    if (FRAME_COUNT != IMAGE_COUNT )
+    size_t allowed_extra_frames = 0;
+
+    // Hack! Newer FFmpeg versions in this combination produce a file
+    // whose reported duration is one frame longer than needed, and so
+    // the calculated frame count is also off by one. Ideally, we'd want
+    // to fix both writing (to produce the correct duration) and reading
+    // (to correctly report frame count for such files), but I don't know
+    // how to do either, so this is a workaround for now.
+    // See also the same hack in CV_PositioningTest::run.
+    if (fourcc == VideoWriter::fourcc('M', 'P', 'E', 'G') && ext == "mkv")
+        allowed_extra_frames = 1;
+
+    if (FRAME_COUNT < IMAGE_COUNT || FRAME_COUNT > IMAGE_COUNT + allowed_extra_frames)
     {
         ts->printf(ts->LOG, "\nFrame count checking for video_%s.%s...\n", fourcc_str.c_str(), ext.c_str());
         ts->printf(ts->LOG, "Video codec: %s\n", fourcc_str.c_str());
-        ts->printf(ts->LOG, "Required frame count: %d; Returned frame count: %d\n", IMAGE_COUNT, FRAME_COUNT);
+        if (allowed_extra_frames != 0)
+            ts->printf(ts->LOG, "Required frame count: %d-%d; Returned frame count: %d\n",
+                       IMAGE_COUNT, IMAGE_COUNT + allowed_extra_frames, FRAME_COUNT);
+        else
+            ts->printf(ts->LOG, "Required frame count: %d; Returned frame count: %d\n", IMAGE_COUNT, FRAME_COUNT);
         ts->printf(ts->LOG, "Error: Incorrect frame count in the video.\n");
         ts->printf(ts->LOG, "Continue checking...\n");
         ts->set_failed_test_info(ts->FAIL_BAD_ACCURACY);
         return;
     }
 
-    for (int i = 0; (size_t)i < FRAME_COUNT; i++)
+    for (int i = 0; (size_t)i < IMAGE_COUNT; i++)
     {
         Mat frame; cap >> frame;
         if (frame.empty())

modules/highgui/test/test_video_pos.cpp

@@ -114,16 +114,21 @@ public:
             cap.set(CAP_PROP_POS_FRAMES, 0);
             int N = (int)cap.get(CAP_PROP_FRAME_COUNT);
 
-            if (N != n_frames || N != N0)
+            // See the same hack in CV_HighGuiTest::SpecificVideoTest for explanation.
+            int allowed_extra_frames = 0;
+            if (fmt.fourcc == VideoWriter::fourcc('M', 'P', 'E', 'G') && fmt.ext == "mkv")
+                allowed_extra_frames = 1;
+
+            if (N < n_frames || N > n_frames + allowed_extra_frames || N != N0)
             {
                 ts->printf(ts->LOG, "\nError: returned frame count (N0=%d, N=%d) is different from the reference number %d\n", N0, N, n_frames);
                 ts->set_failed_test_info(ts->FAIL_INVALID_OUTPUT);
                 return;
             }
 
-            for (int k = 0; k < N; ++k)
+            for (int k = 0; k < n_frames; ++k)
             {
-                int idx = theRNG().uniform(0, N);
+                int idx = theRNG().uniform(0, n_frames);
 
                 if( !cap.set(CAP_PROP_POS_FRAMES, idx) )
                 {

modules/imgproc/include/opencv2/imgproc.hpp

@@ -1147,7 +1147,7 @@ CV_EXPORTS_W void HoughLines( InputArray image, OutputArray lines,
                               double rho, double theta, int threshold,
                               double srn = 0, double stn = 0 );
 
-//! finds line segments in the black-n-white image using probabalistic Hough transform
+//! finds line segments in the black-n-white image using probabilistic Hough transform
 CV_EXPORTS_W void HoughLinesP( InputArray image, OutputArray lines,
                                double rho, double theta, int threshold,
                                double minLineLength = 0, double maxLineGap = 0 );

modules/java/CMakeLists.txt

@@ -330,16 +330,16 @@ if(ENABLE_SOLUTION_FOLDERS)
 endif()
 
 if(ANDROID)
-  install(TARGETS ${the_module}
+  ocv_install_target(${the_module} EXPORT OpenCVModules
           LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
           ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main)
 else()
   if(NOT INSTALL_CREATE_DISTRIB)
-    install(TARGETS ${the_module}
+    ocv_install_target(${the_module} EXPORT OpenCVModules
             RUNTIME DESTINATION ${JAR_INSTALL_DIR} COMPONENT main
             LIBRARY DESTINATION ${JAR_INSTALL_DIR} COMPONENT main)
   else()
-    install(TARGETS ${the_module}
+    ocv_install_target(${the_module} EXPORT OpenCVModules
             RUNTIME DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT main
             LIBRARY DESTINATION ${JAR_INSTALL_DIR}/${OpenCV_ARCH} COMPONENT main)
   endif()

modules/java/generator/gen_java.py

@@ -806,12 +806,22 @@ public class %(jc)s {
             version_suffix =  ''.join( (epoch, major, minor) )
             self.classes[class_name].imports.add("java.lang.String")
             self.java_code[class_name]["j_code"].write("""
-    public static final String VERSION = "%(v)s", NATIVE_LIBRARY_NAME = "opencv_java%(vs)s";
-    public static final int VERSION_EPOCH = %(ep)s;
-    public static final int VERSION_MAJOR = %(ma)s;
-    public static final int VERSION_MINOR = %(mi)s;
-    public static final int VERSION_REVISION = %(re)s;
-    public static final String VERSION_STATUS = "%(st)s";
+    // these constants are wrapped inside functions to prevent inlining
+    private static String getVersion() { return "%(v)s"; }
+    private static String getNativeLibraryName() { return "opencv_java%(vs)s"; }
+    private static int getVersionEpoch() { return %(ep)s; }
+    private static int getVersionMajor() { return %(ma)s; }
+    private static int getVersionMinor() { return %(mi)s; }
+    private static int getVersionRevision() { return %(re)s; }
+    private static String getVersionStatus() { return "%(st)s"; }
+
+    public static final String VERSION = getVersion();
+    public static final String NATIVE_LIBRARY_NAME = getNativeLibraryName();
+    public static final int VERSION_EPOCH = getVersionEpoch();
+    public static final int VERSION_MAJOR = getVersionMajor();
+    public static final int VERSION_MINOR = getVersionMinor();
+    public static final int VERSION_REVISION = getVersionRevision();
+    public static final String VERSION_STATUS = getVersionStatus();
 """ % { 'v' : version_str, 'vs' : version_suffix, 'ep' : epoch, 'ma' : major, 'mi' : minor, 're' : revision, 'st': status } )
 
 

modules/java/test/CMakeLists.txt

@@ -10,6 +10,9 @@ set(opencv_test_java_bin_dir "${CMAKE_CURRENT_BINARY_DIR}/.build")
 set(android_source_dir "${CMAKE_CURRENT_SOURCE_DIR}/../android_test")
 set(java_source_dir ${CMAKE_CURRENT_SOURCE_DIR})
 
+# make sure the build directory exists
+file(MAKE_DIRECTORY "${opencv_test_java_bin_dir}")
+
 # get project sources
 file(GLOB_RECURSE opencv_test_java_files RELATIVE "${android_source_dir}" "${android_source_dir}/res/*" "${android_source_dir}/src/*.java")
 # These are the files that need to be updated for pure Java.
@@ -58,14 +61,6 @@ add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/build.xml"
                    COMMENT "Copying build.xml"
                   )
 
-# Create a script for running the Java tests and place it in build/bin.
-#if(WIN32)
-  #file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java.cmd"   "cd ${opencv_test_java_bin_dir}\nset PATH=${EXECUTABLE_OUTPUT_PATH}/Release;%PATH%\nant -DjavaLibraryPath=${EXECUTABLE_OUTPUT_PATH}/Release buildAndTest")
-  #file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java_D.cmd" "cd ${opencv_test_java_bin_dir}\nset PATH=${EXECUTABLE_OUTPUT_PATH}/Debug;%PATH%\nant   -DjavaLibraryPath=${EXECUTABLE_OUTPUT_PATH}/Debug   buildAndTest")
-#else()
-  #file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java.sh" "cd ${opencv_test_java_bin_dir};\nant -DjavaLibraryPath=${LIBRARY_OUTPUT_PATH} buildAndTest;\ncd -")
-#endif()
-
 add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/build/jar/opencv-test.jar"
                    COMMAND "${ANT_EXECUTABLE}" build
                    WORKING_DIRECTORY "${opencv_test_java_bin_dir}"
@@ -73,8 +68,19 @@ add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/build/jar/opencv-test.jar
                    COMMENT "Build Java tests"
                   )
 
-add_custom_target(${PROJECT_NAME} ALL SOURCES "${opencv_test_java_bin_dir}/build/jar/opencv-test.jar")
-add_dependencies(${PROJECT_NAME} ${the_module})
+# Not add_custom_command because generator expressions aren't supported in
+# OUTPUT file names, and we need to generate different files for different
+# configurations.
+add_custom_target(${PROJECT_NAME}_properties
+                  COMMAND "${CMAKE_COMMAND}" -E echo "opencv.lib.path = $<TARGET_FILE_DIR:${the_module}>"
+                    > "${opencv_test_java_bin_dir}/ant-$<CONFIGURATION>.properties"
+                  VERBATIM
+                 )
+
+add_custom_target(${PROJECT_NAME} ALL
+                  DEPENDS ${the_module} ${PROJECT_NAME}_properties
+                  SOURCES "${opencv_test_java_bin_dir}/build/jar/opencv-test.jar"
+                 )
 
 if(ENABLE_SOLUTION_FOLDERS)
   set_target_properties(${PROJECT_NAME} PROPERTIES FOLDER "tests accuracy")

modules/java/test/build.xml

@@ -1,4 +1,6 @@
 <project>
+  <property file="ant-${opencv.build.type}.properties"/>
+
   <path id="master-classpath">
     <fileset dir="lib">
       <include name="*.jar"/>
@@ -34,8 +36,8 @@
   <target name="test">
     <mkdir dir="testResults"/>
     <junit printsummary="true" haltonfailure="false" haltonerror="false" showoutput="false" logfailedtests="true" maxmemory="256m">
-      <sysproperty key="java.library.path" path="${javaLibraryPath}"/>
-      <env key="PATH" path="${javaLibraryPath}"/>
+      <sysproperty key="java.library.path" path="${opencv.lib.path}"/>
+      <env key="PATH" path="${opencv.lib.path}"/>
       <classpath refid="master-classpath"/>
       <classpath>
         <pathelement location="build/classes"/>
@@ -61,4 +63,4 @@
     <antcall target="jar"/>
     <antcall target="test"/>
   </target>
-</project>
+</project>

modules/java/test/src/org/opencv/test/OpenCVTestRunner.java

@@ -2,7 +2,6 @@ package org.opencv.test;
 
 import java.io.File;
 import java.io.IOException;
-import junit.framework.Assert;
 
 import org.opencv.core.Mat;
 

modules/ocl/CMakeLists.txt

@@ -5,4 +5,4 @@ endif()
 
 set(the_description "OpenCL-accelerated Computer Vision")
 ocv_define_module(ocl opencv_core opencv_imgproc opencv_features2d opencv_objdetect opencv_video opencv_calib3d opencv_ml "${OPENCL_LIBRARIES}")
-ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow)
+ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow -Wundef)

modules/ocl/doc/image_filtering.rst

@@ -701,3 +701,17 @@ Performs linear blending of two images.
     :param weights2: Weights for second image. Must have tha same size as ``img2`` . Supports only ``CV_32F`` type.
 
     :param result: Destination image.
+
+ocl::medianFilter
+--------------------
+Blurs an image using the median filter.
+
+.. ocv:function:: void ocl::medianFilter(const oclMat &src, oclMat &dst, int m)
+
+    :param src: input ```1-``` or ```4```-channel image; the image depth should be ```CV_8U```, ```CV_32F```.
+
+    :param dst: destination array of the same size and type as ```src```.
+
+    :param m: aperture linear size; it must be odd and greater than ```1```. Currently only ```3```, ```5``` are supported.
+
+The function smoothes an image using the median filter with the \texttt{m} \times \texttt{m} aperture. Each channel of a multi-channel image is processed independently. In-place operation is supported.

modules/ocl/include/opencv2/ocl.hpp

@@ -858,11 +858,8 @@ namespace cv
         //! Applies a generic geometrical transformation to an image.
 
         // Supports INTER_NEAREST, INTER_LINEAR.
-
         // Map1 supports CV_16SC2, CV_32FC2  types.
-
         // Src supports CV_8UC1, CV_8UC2, CV_8UC4.
-
         CV_EXPORTS void remap(const oclMat &src, oclMat &dst, oclMat &map1, oclMat &map2, int interpolation, int bordertype, const Scalar &value = Scalar());
 
         //! copies 2D array to a larger destination array and pads borders with user-specifiable constant
@@ -870,7 +867,7 @@ namespace cv
         CV_EXPORTS void copyMakeBorder(const oclMat &src, oclMat &dst, int top, int bottom, int left, int right, int boardtype, const Scalar &value = Scalar());
 
         //! Smoothes image using median filter
-        // The source 1- or 4-channel image. When m is 3 or 5, the image depth should be CV 8U or CV 32F.
+        // The source 1- or 4-channel image. m should be 3 or 5, the image depth should be CV_8U or CV_32F.
         CV_EXPORTS void medianFilter(const oclMat &src, oclMat &dst, int m);
 
         //! warps the image using affine transformation

modules/ocl/include/opencv2/ocl/cl_runtime/cl_runtime.hpp

@@ -3,14 +3,20 @@
 
 #ifdef HAVE_OPENCL
 
-#if defined(HAVE_OPENCL12)
-#include "cl_runtime_opencl12.hpp"
-#elif defined(HAVE_OPENCL11)
-#include "cl_runtime_opencl11.hpp"
+#if defined(HAVE_OPENCL_STATIC)
+
+#if defined __APPLE__
+#include <OpenCL/cl.h>
 #else
-#error Invalid OpenCL configuration
+#include <CL/cl.h>
 #endif
 
-#endif
+#else // HAVE_OPENCL_STATIC
+
+#include "cl_runtime_opencl.hpp"
+
+#endif // HAVE_OPENCL_STATIC
+
+#endif // HAVE_OPENCL
 
 #endif // __OPENCV_OCL_CL_RUNTIME_HPP__

modules/ocl/include/opencv2/ocl/cl_runtime/cl_runtime_opencl.hpp

@@ -4,12 +4,6 @@
 #ifndef __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__
 #define __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__
 
-#ifdef HAVE_OPENCL
-
-#if defined __APPLE__ && !defined(IOS)
-#include <OpenCL/cl.h>
-#else
-
 // generated by parser_cl.py
 #define clGetPlatformIDs clGetPlatformIDs_
 #define clGetPlatformInfo clGetPlatformInfo_
@@ -382,8 +376,4 @@ extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueBarrier)(cl_command_queue)
 extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clUnloadCompiler)();
 extern CL_RUNTIME_EXPORT void* (CL_API_CALL*clGetExtensionFunctionAddress)(const char*);
 
-#endif
-
-#endif
-
 #endif // __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__

modules/ocl/include/opencv2/ocl/cl_runtime/cl_runtime_opencl11.hpp

@@ -1,333 +0,0 @@
-//
-// AUTOGENERATED, DO NOT EDIT
-//
-#ifndef __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__
-#define __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__
-
-#ifdef HAVE_OPENCL
-
-#if defined __APPLE__ && !defined(IOS)
-#include <OpenCL/cl.h>
-#else
-
-// generated by parser_cl.py
-#define clGetPlatformIDs clGetPlatformIDs_
-#define clGetPlatformInfo clGetPlatformInfo_
-#define clGetDeviceIDs clGetDeviceIDs_
-#define clGetDeviceInfo clGetDeviceInfo_
-#define clCreateContext clCreateContext_
-#define clCreateContextFromType clCreateContextFromType_
-#define clRetainContext clRetainContext_
-#define clReleaseContext clReleaseContext_
-#define clGetContextInfo clGetContextInfo_
-#define clCreateCommandQueue clCreateCommandQueue_
-#define clRetainCommandQueue clRetainCommandQueue_
-#define clReleaseCommandQueue clReleaseCommandQueue_
-#define clGetCommandQueueInfo clGetCommandQueueInfo_
-#define clSetCommandQueueProperty clSetCommandQueueProperty_
-#define clCreateBuffer clCreateBuffer_
-#define clCreateSubBuffer clCreateSubBuffer_
-#define clCreateImage2D clCreateImage2D_
-#define clCreateImage3D clCreateImage3D_
-#define clRetainMemObject clRetainMemObject_
-#define clReleaseMemObject clReleaseMemObject_
-#define clGetSupportedImageFormats clGetSupportedImageFormats_
-#define clGetMemObjectInfo clGetMemObjectInfo_
-#define clGetImageInfo clGetImageInfo_
-#define clSetMemObjectDestructorCallback clSetMemObjectDestructorCallback_
-#define clCreateSampler clCreateSampler_
-#define clRetainSampler clRetainSampler_
-#define clReleaseSampler clReleaseSampler_
-#define clGetSamplerInfo clGetSamplerInfo_
-#define clCreateProgramWithSource clCreateProgramWithSource_
-#define clCreateProgramWithBinary clCreateProgramWithBinary_
-#define clRetainProgram clRetainProgram_
-#define clReleaseProgram clReleaseProgram_
-#define clBuildProgram clBuildProgram_
-#define clUnloadCompiler clUnloadCompiler_
-#define clGetProgramInfo clGetProgramInfo_
-#define clGetProgramBuildInfo clGetProgramBuildInfo_
-#define clCreateKernel clCreateKernel_
-#define clCreateKernelsInProgram clCreateKernelsInProgram_
-#define clRetainKernel clRetainKernel_
-#define clReleaseKernel clReleaseKernel_
-#define clSetKernelArg clSetKernelArg_
-#define clGetKernelInfo clGetKernelInfo_
-#define clGetKernelWorkGroupInfo clGetKernelWorkGroupInfo_
-#define clWaitForEvents clWaitForEvents_
-#define clGetEventInfo clGetEventInfo_
-#define clCreateUserEvent clCreateUserEvent_
-#define clRetainEvent clRetainEvent_
-#define clReleaseEvent clReleaseEvent_
-#define clSetUserEventStatus clSetUserEventStatus_
-#define clSetEventCallback clSetEventCallback_
-#define clGetEventProfilingInfo clGetEventProfilingInfo_
-#define clFlush clFlush_
-#define clFinish clFinish_
-#define clEnqueueReadBuffer clEnqueueReadBuffer_
-#define clEnqueueReadBufferRect clEnqueueReadBufferRect_
-#define clEnqueueWriteBuffer clEnqueueWriteBuffer_
-#define clEnqueueWriteBufferRect clEnqueueWriteBufferRect_
-#define clEnqueueCopyBuffer clEnqueueCopyBuffer_
-#define clEnqueueCopyBufferRect clEnqueueCopyBufferRect_
-#define clEnqueueReadImage clEnqueueReadImage_
-#define clEnqueueWriteImage clEnqueueWriteImage_
-#define clEnqueueCopyImage clEnqueueCopyImage_
-#define clEnqueueCopyImageToBuffer clEnqueueCopyImageToBuffer_
-#define clEnqueueCopyBufferToImage clEnqueueCopyBufferToImage_
-#define clEnqueueMapBuffer clEnqueueMapBuffer_
-#define clEnqueueMapImage clEnqueueMapImage_
-#define clEnqueueUnmapMemObject clEnqueueUnmapMemObject_
-#define clEnqueueNDRangeKernel clEnqueueNDRangeKernel_
-#define clEnqueueTask clEnqueueTask_
-#define clEnqueueNativeKernel clEnqueueNativeKernel_
-#define clEnqueueMarker clEnqueueMarker_
-#define clEnqueueWaitForEvents clEnqueueWaitForEvents_
-#define clEnqueueBarrier clEnqueueBarrier_
-#define clGetExtensionFunctionAddress clGetExtensionFunctionAddress_
-
-#if defined __APPLE__
-#include <OpenCL/cl.h>
-#else
-#include <CL/cl.h>
-#endif
-
-// generated by parser_cl.py
-#undef clGetPlatformIDs
-#define clGetPlatformIDs clGetPlatformIDs_pfn
-#undef clGetPlatformInfo
-#define clGetPlatformInfo clGetPlatformInfo_pfn
-#undef clGetDeviceIDs
-#define clGetDeviceIDs clGetDeviceIDs_pfn
-#undef clGetDeviceInfo
-#define clGetDeviceInfo clGetDeviceInfo_pfn
-#undef clCreateContext
-#define clCreateContext clCreateContext_pfn
-#undef clCreateContextFromType
-#define clCreateContextFromType clCreateContextFromType_pfn
-#undef clRetainContext
-#define clRetainContext clRetainContext_pfn
-#undef clReleaseContext
-#define clReleaseContext clReleaseContext_pfn
-#undef clGetContextInfo
-#define clGetContextInfo clGetContextInfo_pfn
-#undef clCreateCommandQueue
-#define clCreateCommandQueue clCreateCommandQueue_pfn
-#undef clRetainCommandQueue
-#define clRetainCommandQueue clRetainCommandQueue_pfn
-#undef clReleaseCommandQueue
-#define clReleaseCommandQueue clReleaseCommandQueue_pfn
-#undef clGetCommandQueueInfo
-#define clGetCommandQueueInfo clGetCommandQueueInfo_pfn
-#undef clSetCommandQueueProperty
-#define clSetCommandQueueProperty clSetCommandQueueProperty_pfn
-#undef clCreateBuffer
-#define clCreateBuffer clCreateBuffer_pfn
-#undef clCreateSubBuffer
-#define clCreateSubBuffer clCreateSubBuffer_pfn
-#undef clCreateImage2D
-#define clCreateImage2D clCreateImage2D_pfn
-#undef clCreateImage3D
-#define clCreateImage3D clCreateImage3D_pfn
-#undef clRetainMemObject
-#define clRetainMemObject clRetainMemObject_pfn
-#undef clReleaseMemObject
-#define clReleaseMemObject clReleaseMemObject_pfn
-#undef clGetSupportedImageFormats
-#define clGetSupportedImageFormats clGetSupportedImageFormats_pfn
-#undef clGetMemObjectInfo
-#define clGetMemObjectInfo clGetMemObjectInfo_pfn
-#undef clGetImageInfo
-#define clGetImageInfo clGetImageInfo_pfn
-#undef clSetMemObjectDestructorCallback
-#define clSetMemObjectDestructorCallback clSetMemObjectDestructorCallback_pfn
-#undef clCreateSampler
-#define clCreateSampler clCreateSampler_pfn
-#undef clRetainSampler
-#define clRetainSampler clRetainSampler_pfn
-#undef clReleaseSampler
-#define clReleaseSampler clReleaseSampler_pfn
-#undef clGetSamplerInfo
-#define clGetSamplerInfo clGetSamplerInfo_pfn
-#undef clCreateProgramWithSource
-#define clCreateProgramWithSource clCreateProgramWithSource_pfn
-#undef clCreateProgramWithBinary
-#define clCreateProgramWithBinary clCreateProgramWithBinary_pfn
-#undef clRetainProgram
-#define clRetainProgram clRetainProgram_pfn
-#undef clReleaseProgram
-#define clReleaseProgram clReleaseProgram_pfn
-#undef clBuildProgram
-#define clBuildProgram clBuildProgram_pfn
-#undef clUnloadCompiler
-#define clUnloadCompiler clUnloadCompiler_pfn
-#undef clGetProgramInfo
-#define clGetProgramInfo clGetProgramInfo_pfn
-#undef clGetProgramBuildInfo
-#define clGetProgramBuildInfo clGetProgramBuildInfo_pfn
-#undef clCreateKernel
-#define clCreateKernel clCreateKernel_pfn
-#undef clCreateKernelsInProgram
-#define clCreateKernelsInProgram clCreateKernelsInProgram_pfn
-#undef clRetainKernel
-#define clRetainKernel clRetainKernel_pfn
-#undef clReleaseKernel
-#define clReleaseKernel clReleaseKernel_pfn
-#undef clSetKernelArg
-#define clSetKernelArg clSetKernelArg_pfn
-#undef clGetKernelInfo
-#define clGetKernelInfo clGetKernelInfo_pfn
-#undef clGetKernelWorkGroupInfo
-#define clGetKernelWorkGroupInfo clGetKernelWorkGroupInfo_pfn
-#undef clWaitForEvents
-#define clWaitForEvents clWaitForEvents_pfn
-#undef clGetEventInfo
-#define clGetEventInfo clGetEventInfo_pfn
-#undef clCreateUserEvent
-#define clCreateUserEvent clCreateUserEvent_pfn
-#undef clRetainEvent
-#define clRetainEvent clRetainEvent_pfn
-#undef clReleaseEvent
-#define clReleaseEvent clReleaseEvent_pfn
-#undef clSetUserEventStatus
-#define clSetUserEventStatus clSetUserEventStatus_pfn
-#undef clSetEventCallback
-#define clSetEventCallback clSetEventCallback_pfn
-#undef clGetEventProfilingInfo
-#define clGetEventProfilingInfo clGetEventProfilingInfo_pfn
-#undef clFlush
-#define clFlush clFlush_pfn
-#undef clFinish
-#define clFinish clFinish_pfn
-#undef clEnqueueReadBuffer
-#define clEnqueueReadBuffer clEnqueueReadBuffer_pfn
-#undef clEnqueueReadBufferRect
-#define clEnqueueReadBufferRect clEnqueueReadBufferRect_pfn
-#undef clEnqueueWriteBuffer
-#define clEnqueueWriteBuffer clEnqueueWriteBuffer_pfn
-#undef clEnqueueWriteBufferRect
-#define clEnqueueWriteBufferRect clEnqueueWriteBufferRect_pfn
-#undef clEnqueueCopyBuffer
-#define clEnqueueCopyBuffer clEnqueueCopyBuffer_pfn
-#undef clEnqueueCopyBufferRect
-#define clEnqueueCopyBufferRect clEnqueueCopyBufferRect_pfn
-#undef clEnqueueReadImage
-#define clEnqueueReadImage clEnqueueReadImage_pfn
-#undef clEnqueueWriteImage
-#define clEnqueueWriteImage clEnqueueWriteImage_pfn
-#undef clEnqueueCopyImage
-#define clEnqueueCopyImage clEnqueueCopyImage_pfn
-#undef clEnqueueCopyImageToBuffer
-#define clEnqueueCopyImageToBuffer clEnqueueCopyImageToBuffer_pfn
-#undef clEnqueueCopyBufferToImage
-#define clEnqueueCopyBufferToImage clEnqueueCopyBufferToImage_pfn
-#undef clEnqueueMapBuffer
-#define clEnqueueMapBuffer clEnqueueMapBuffer_pfn
-#undef clEnqueueMapImage
-#define clEnqueueMapImage clEnqueueMapImage_pfn
-#undef clEnqueueUnmapMemObject
-#define clEnqueueUnmapMemObject clEnqueueUnmapMemObject_pfn
-#undef clEnqueueNDRangeKernel
-#define clEnqueueNDRangeKernel clEnqueueNDRangeKernel_pfn
-#undef clEnqueueTask
-#define clEnqueueTask clEnqueueTask_pfn
-#undef clEnqueueNativeKernel
-#define clEnqueueNativeKernel clEnqueueNativeKernel_pfn
-#undef clEnqueueMarker
-#define clEnqueueMarker clEnqueueMarker_pfn
-#undef clEnqueueWaitForEvents
-#define clEnqueueWaitForEvents clEnqueueWaitForEvents_pfn
-#undef clEnqueueBarrier
-#define clEnqueueBarrier clEnqueueBarrier_pfn
-#undef clGetExtensionFunctionAddress
-#define clGetExtensionFunctionAddress clGetExtensionFunctionAddress_pfn
-
-#ifndef CL_RUNTIME_EXPORT
-#if (defined(BUILD_SHARED_LIBS) || defined(OPENCV_OCL_SHARED)) && (defined WIN32 || defined _WIN32 || defined WINCE)
-#define CL_RUNTIME_EXPORT __declspec(dllimport)
-#else
-#define CL_RUNTIME_EXPORT
-#endif
-#endif
-
-// generated by parser_cl.py
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetPlatformIDs)(cl_uint, cl_platform_id*, cl_uint*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetPlatformInfo)(cl_platform_id, cl_platform_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetDeviceIDs)(cl_platform_id, cl_device_type, cl_uint, cl_device_id*, cl_uint*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetDeviceInfo)(cl_device_id, cl_device_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_context (CL_API_CALL*clCreateContext)(const cl_context_properties*, cl_uint, const cl_device_id*, void (CL_CALLBACK*) (const char*, const void*, size_t, void*), void*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_context (CL_API_CALL*clCreateContextFromType)(const cl_context_properties*, cl_device_type, void (CL_CALLBACK*) (const char*, const void*, size_t, void*), void*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainContext)(cl_context);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseContext)(cl_context);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetContextInfo)(cl_context, cl_context_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_command_queue (CL_API_CALL*clCreateCommandQueue)(cl_context, cl_device_id, cl_command_queue_properties, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainCommandQueue)(cl_command_queue);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseCommandQueue)(cl_command_queue);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetCommandQueueInfo)(cl_command_queue, cl_command_queue_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clSetCommandQueueProperty)(cl_command_queue, cl_command_queue_properties, cl_bool, cl_command_queue_properties*);
-extern CL_RUNTIME_EXPORT cl_mem (CL_API_CALL*clCreateBuffer)(cl_context, cl_mem_flags, size_t, void*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_mem (CL_API_CALL*clCreateSubBuffer)(cl_mem, cl_mem_flags, cl_buffer_create_type, const void*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_mem (CL_API_CALL*clCreateImage2D)(cl_context, cl_mem_flags, const cl_image_format*, size_t, size_t, size_t, void*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_mem (CL_API_CALL*clCreateImage3D)(cl_context, cl_mem_flags, const cl_image_format*, size_t, size_t, size_t, size_t, size_t, void*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainMemObject)(cl_mem);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseMemObject)(cl_mem);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetSupportedImageFormats)(cl_context, cl_mem_flags, cl_mem_object_type, cl_uint, cl_image_format*, cl_uint*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetMemObjectInfo)(cl_mem, cl_mem_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetImageInfo)(cl_mem, cl_image_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clSetMemObjectDestructorCallback)(cl_mem, void (CL_CALLBACK*) (cl_mem, void*), void*);
-extern CL_RUNTIME_EXPORT cl_sampler (CL_API_CALL*clCreateSampler)(cl_context, cl_bool, cl_addressing_mode, cl_filter_mode, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainSampler)(cl_sampler);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseSampler)(cl_sampler);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetSamplerInfo)(cl_sampler, cl_sampler_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_program (CL_API_CALL*clCreateProgramWithSource)(cl_context, cl_uint, const char**, const size_t*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_program (CL_API_CALL*clCreateProgramWithBinary)(cl_context, cl_uint, const cl_device_id*, const size_t*, const unsigned char**, cl_int*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainProgram)(cl_program);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseProgram)(cl_program);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clBuildProgram)(cl_program, cl_uint, const cl_device_id*, const char*, void (CL_CALLBACK*) (cl_program, void*), void*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clUnloadCompiler)();
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetProgramInfo)(cl_program, cl_program_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetProgramBuildInfo)(cl_program, cl_device_id, cl_program_build_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_kernel (CL_API_CALL*clCreateKernel)(cl_program, const char*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clCreateKernelsInProgram)(cl_program, cl_uint, cl_kernel*, cl_uint*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainKernel)(cl_kernel);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseKernel)(cl_kernel);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clSetKernelArg)(cl_kernel, cl_uint, size_t, const void*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetKernelInfo)(cl_kernel, cl_kernel_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetKernelWorkGroupInfo)(cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clWaitForEvents)(cl_uint, const cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetEventInfo)(cl_event, cl_event_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_event (CL_API_CALL*clCreateUserEvent)(cl_context, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clRetainEvent)(cl_event);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clReleaseEvent)(cl_event);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clSetUserEventStatus)(cl_event, cl_int);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clSetEventCallback)(cl_event, cl_int, void (CL_CALLBACK*) (cl_event, cl_int, void*), void*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clGetEventProfilingInfo)(cl_event, cl_profiling_info, size_t, void*, size_t*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clFlush)(cl_command_queue);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clFinish)(cl_command_queue);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueReadBuffer)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueReadBufferRect)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueWriteBuffer)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueWriteBufferRect)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueCopyBuffer)(cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueCopyBufferRect)(cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueReadImage)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueWriteImage)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueCopyImage)(cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueCopyImageToBuffer)(cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, size_t, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueCopyBufferToImage)(cl_command_queue, cl_mem, cl_mem, size_t, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT void* (CL_API_CALL*clEnqueueMapBuffer)(cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event*, cl_event*, cl_int*);
-extern CL_RUNTIME_EXPORT void* (CL_API_CALL*clEnqueueMapImage)(cl_command_queue, cl_mem, cl_bool, cl_map_flags, const size_t*, const size_t*, size_t*, size_t*, cl_uint, const cl_event*, cl_event*, cl_int*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueUnmapMemObject)(cl_command_queue, cl_mem, void*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueNDRangeKernel)(cl_command_queue, cl_kernel, cl_uint, const size_t*, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueTask)(cl_command_queue, cl_kernel, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueNativeKernel)(cl_command_queue, void (CL_CALLBACK* user_func) (void*), void*, size_t, cl_uint, const cl_mem*, const void**, cl_uint, const cl_event*, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueMarker)(cl_command_queue, cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueWaitForEvents)(cl_command_queue, cl_uint, const cl_event*);
-extern CL_RUNTIME_EXPORT cl_int (CL_API_CALL*clEnqueueBarrier)(cl_command_queue);
-extern CL_RUNTIME_EXPORT void* (CL_API_CALL*clGetExtensionFunctionAddress)(const char*);
-
-#endif
-
-#endif
-
-#endif // __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__

modules/ocl/include/opencv2/ocl/cl_runtime/cl_runtime_opencl11_wrappers.hpp

@@ -1,231 +0,0 @@
-//
-// AUTOGENERATED, DO NOT EDIT
-//
-#ifndef __OPENCV_OCL_CL_RUNTIME_OPENCL_WRAPPERS_HPP__
-#define __OPENCV_OCL_CL_RUNTIME_OPENCL_WRAPPERS_HPP__
-
-// generated by parser_cl.py
-#undef clGetPlatformIDs
-#define clGetPlatformIDs clGetPlatformIDs_fn
-inline cl_int clGetPlatformIDs(cl_uint p0, cl_platform_id* p1, cl_uint* p2) { return clGetPlatformIDs_pfn(p0, p1, p2); }
-#undef clGetPlatformInfo
-#define clGetPlatformInfo clGetPlatformInfo_fn
-inline cl_int clGetPlatformInfo(cl_platform_id p0, cl_platform_info p1, size_t p2, void* p3, size_t* p4) { return clGetPlatformInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clGetDeviceIDs
-#define clGetDeviceIDs clGetDeviceIDs_fn
-inline cl_int clGetDeviceIDs(cl_platform_id p0, cl_device_type p1, cl_uint p2, cl_device_id* p3, cl_uint* p4) { return clGetDeviceIDs_pfn(p0, p1, p2, p3, p4); }
-#undef clGetDeviceInfo
-#define clGetDeviceInfo clGetDeviceInfo_fn
-inline cl_int clGetDeviceInfo(cl_device_id p0, cl_device_info p1, size_t p2, void* p3, size_t* p4) { return clGetDeviceInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateContext
-#define clCreateContext clCreateContext_fn
-inline cl_context clCreateContext(const cl_context_properties* p0, cl_uint p1, const cl_device_id* p2, void (CL_CALLBACK*p3) (const char*, const void*, size_t, void*), void* p4, cl_int* p5) { return clCreateContext_pfn(p0, p1, p2, p3, p4, p5); }
-#undef clCreateContextFromType
-#define clCreateContextFromType clCreateContextFromType_fn
-inline cl_context clCreateContextFromType(const cl_context_properties* p0, cl_device_type p1, void (CL_CALLBACK*p2) (const char*, const void*, size_t, void*), void* p3, cl_int* p4) { return clCreateContextFromType_pfn(p0, p1, p2, p3, p4); }
-#undef clRetainContext
-#define clRetainContext clRetainContext_fn
-inline cl_int clRetainContext(cl_context p0) { return clRetainContext_pfn(p0); }
-#undef clReleaseContext
-#define clReleaseContext clReleaseContext_fn
-inline cl_int clReleaseContext(cl_context p0) { return clReleaseContext_pfn(p0); }
-#undef clGetContextInfo
-#define clGetContextInfo clGetContextInfo_fn
-inline cl_int clGetContextInfo(cl_context p0, cl_context_info p1, size_t p2, void* p3, size_t* p4) { return clGetContextInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateCommandQueue
-#define clCreateCommandQueue clCreateCommandQueue_fn
-inline cl_command_queue clCreateCommandQueue(cl_context p0, cl_device_id p1, cl_command_queue_properties p2, cl_int* p3) { return clCreateCommandQueue_pfn(p0, p1, p2, p3); }
-#undef clRetainCommandQueue
-#define clRetainCommandQueue clRetainCommandQueue_fn
-inline cl_int clRetainCommandQueue(cl_command_queue p0) { return clRetainCommandQueue_pfn(p0); }
-#undef clReleaseCommandQueue
-#define clReleaseCommandQueue clReleaseCommandQueue_fn
-inline cl_int clReleaseCommandQueue(cl_command_queue p0) { return clReleaseCommandQueue_pfn(p0); }
-#undef clGetCommandQueueInfo
-#define clGetCommandQueueInfo clGetCommandQueueInfo_fn
-inline cl_int clGetCommandQueueInfo(cl_command_queue p0, cl_command_queue_info p1, size_t p2, void* p3, size_t* p4) { return clGetCommandQueueInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clSetCommandQueueProperty
-#define clSetCommandQueueProperty clSetCommandQueueProperty_fn
-inline cl_int clSetCommandQueueProperty(cl_command_queue p0, cl_command_queue_properties p1, cl_bool p2, cl_command_queue_properties* p3) { return clSetCommandQueueProperty_pfn(p0, p1, p2, p3); }
-#undef clCreateBuffer
-#define clCreateBuffer clCreateBuffer_fn
-inline cl_mem clCreateBuffer(cl_context p0, cl_mem_flags p1, size_t p2, void* p3, cl_int* p4) { return clCreateBuffer_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateSubBuffer
-#define clCreateSubBuffer clCreateSubBuffer_fn
-inline cl_mem clCreateSubBuffer(cl_mem p0, cl_mem_flags p1, cl_buffer_create_type p2, const void* p3, cl_int* p4) { return clCreateSubBuffer_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateImage2D
-#define clCreateImage2D clCreateImage2D_fn
-inline cl_mem clCreateImage2D(cl_context p0, cl_mem_flags p1, const cl_image_format* p2, size_t p3, size_t p4, size_t p5, void* p6, cl_int* p7) { return clCreateImage2D_pfn(p0, p1, p2, p3, p4, p5, p6, p7); }
-#undef clCreateImage3D
-#define clCreateImage3D clCreateImage3D_fn
-inline cl_mem clCreateImage3D(cl_context p0, cl_mem_flags p1, const cl_image_format* p2, size_t p3, size_t p4, size_t p5, size_t p6, size_t p7, void* p8, cl_int* p9) { return clCreateImage3D_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9); }
-#undef clRetainMemObject
-#define clRetainMemObject clRetainMemObject_fn
-inline cl_int clRetainMemObject(cl_mem p0) { return clRetainMemObject_pfn(p0); }
-#undef clReleaseMemObject
-#define clReleaseMemObject clReleaseMemObject_fn
-inline cl_int clReleaseMemObject(cl_mem p0) { return clReleaseMemObject_pfn(p0); }
-#undef clGetSupportedImageFormats
-#define clGetSupportedImageFormats clGetSupportedImageFormats_fn
-inline cl_int clGetSupportedImageFormats(cl_context p0, cl_mem_flags p1, cl_mem_object_type p2, cl_uint p3, cl_image_format* p4, cl_uint* p5) { return clGetSupportedImageFormats_pfn(p0, p1, p2, p3, p4, p5); }
-#undef clGetMemObjectInfo
-#define clGetMemObjectInfo clGetMemObjectInfo_fn
-inline cl_int clGetMemObjectInfo(cl_mem p0, cl_mem_info p1, size_t p2, void* p3, size_t* p4) { return clGetMemObjectInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clGetImageInfo
-#define clGetImageInfo clGetImageInfo_fn
-inline cl_int clGetImageInfo(cl_mem p0, cl_image_info p1, size_t p2, void* p3, size_t* p4) { return clGetImageInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clSetMemObjectDestructorCallback
-#define clSetMemObjectDestructorCallback clSetMemObjectDestructorCallback_fn
-inline cl_int clSetMemObjectDestructorCallback(cl_mem p0, void (CL_CALLBACK*p1) (cl_mem, void*), void* p2) { return clSetMemObjectDestructorCallback_pfn(p0, p1, p2); }
-#undef clCreateSampler
-#define clCreateSampler clCreateSampler_fn
-inline cl_sampler clCreateSampler(cl_context p0, cl_bool p1, cl_addressing_mode p2, cl_filter_mode p3, cl_int* p4) { return clCreateSampler_pfn(p0, p1, p2, p3, p4); }
-#undef clRetainSampler
-#define clRetainSampler clRetainSampler_fn
-inline cl_int clRetainSampler(cl_sampler p0) { return clRetainSampler_pfn(p0); }
-#undef clReleaseSampler
-#define clReleaseSampler clReleaseSampler_fn
-inline cl_int clReleaseSampler(cl_sampler p0) { return clReleaseSampler_pfn(p0); }
-#undef clGetSamplerInfo
-#define clGetSamplerInfo clGetSamplerInfo_fn
-inline cl_int clGetSamplerInfo(cl_sampler p0, cl_sampler_info p1, size_t p2, void* p3, size_t* p4) { return clGetSamplerInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateProgramWithSource
-#define clCreateProgramWithSource clCreateProgramWithSource_fn
-inline cl_program clCreateProgramWithSource(cl_context p0, cl_uint p1, const char** p2, const size_t* p3, cl_int* p4) { return clCreateProgramWithSource_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateProgramWithBinary
-#define clCreateProgramWithBinary clCreateProgramWithBinary_fn
-inline cl_program clCreateProgramWithBinary(cl_context p0, cl_uint p1, const cl_device_id* p2, const size_t* p3, const unsigned char** p4, cl_int* p5, cl_int* p6) { return clCreateProgramWithBinary_pfn(p0, p1, p2, p3, p4, p5, p6); }
-#undef clRetainProgram
-#define clRetainProgram clRetainProgram_fn
-inline cl_int clRetainProgram(cl_program p0) { return clRetainProgram_pfn(p0); }
-#undef clReleaseProgram
-#define clReleaseProgram clReleaseProgram_fn
-inline cl_int clReleaseProgram(cl_program p0) { return clReleaseProgram_pfn(p0); }
-#undef clBuildProgram
-#define clBuildProgram clBuildProgram_fn
-inline cl_int clBuildProgram(cl_program p0, cl_uint p1, const cl_device_id* p2, const char* p3, void (CL_CALLBACK*p4) (cl_program, void*), void* p5) { return clBuildProgram_pfn(p0, p1, p2, p3, p4, p5); }
-#undef clUnloadCompiler
-#define clUnloadCompiler clUnloadCompiler_fn
-inline cl_int clUnloadCompiler() { return clUnloadCompiler_pfn(); }
-#undef clGetProgramInfo
-#define clGetProgramInfo clGetProgramInfo_fn
-inline cl_int clGetProgramInfo(cl_program p0, cl_program_info p1, size_t p2, void* p3, size_t* p4) { return clGetProgramInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clGetProgramBuildInfo
-#define clGetProgramBuildInfo clGetProgramBuildInfo_fn
-inline cl_int clGetProgramBuildInfo(cl_program p0, cl_device_id p1, cl_program_build_info p2, size_t p3, void* p4, size_t* p5) { return clGetProgramBuildInfo_pfn(p0, p1, p2, p3, p4, p5); }
-#undef clCreateKernel
-#define clCreateKernel clCreateKernel_fn
-inline cl_kernel clCreateKernel(cl_program p0, const char* p1, cl_int* p2) { return clCreateKernel_pfn(p0, p1, p2); }
-#undef clCreateKernelsInProgram
-#define clCreateKernelsInProgram clCreateKernelsInProgram_fn
-inline cl_int clCreateKernelsInProgram(cl_program p0, cl_uint p1, cl_kernel* p2, cl_uint* p3) { return clCreateKernelsInProgram_pfn(p0, p1, p2, p3); }
-#undef clRetainKernel
-#define clRetainKernel clRetainKernel_fn
-inline cl_int clRetainKernel(cl_kernel p0) { return clRetainKernel_pfn(p0); }
-#undef clReleaseKernel
-#define clReleaseKernel clReleaseKernel_fn
-inline cl_int clReleaseKernel(cl_kernel p0) { return clReleaseKernel_pfn(p0); }
-#undef clSetKernelArg
-#define clSetKernelArg clSetKernelArg_fn
-inline cl_int clSetKernelArg(cl_kernel p0, cl_uint p1, size_t p2, const void* p3) { return clSetKernelArg_pfn(p0, p1, p2, p3); }
-#undef clGetKernelInfo
-#define clGetKernelInfo clGetKernelInfo_fn
-inline cl_int clGetKernelInfo(cl_kernel p0, cl_kernel_info p1, size_t p2, void* p3, size_t* p4) { return clGetKernelInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clGetKernelWorkGroupInfo
-#define clGetKernelWorkGroupInfo clGetKernelWorkGroupInfo_fn
-inline cl_int clGetKernelWorkGroupInfo(cl_kernel p0, cl_device_id p1, cl_kernel_work_group_info p2, size_t p3, void* p4, size_t* p5) { return clGetKernelWorkGroupInfo_pfn(p0, p1, p2, p3, p4, p5); }
-#undef clWaitForEvents
-#define clWaitForEvents clWaitForEvents_fn
-inline cl_int clWaitForEvents(cl_uint p0, const cl_event* p1) { return clWaitForEvents_pfn(p0, p1); }
-#undef clGetEventInfo
-#define clGetEventInfo clGetEventInfo_fn
-inline cl_int clGetEventInfo(cl_event p0, cl_event_info p1, size_t p2, void* p3, size_t* p4) { return clGetEventInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clCreateUserEvent
-#define clCreateUserEvent clCreateUserEvent_fn
-inline cl_event clCreateUserEvent(cl_context p0, cl_int* p1) { return clCreateUserEvent_pfn(p0, p1); }
-#undef clRetainEvent
-#define clRetainEvent clRetainEvent_fn
-inline cl_int clRetainEvent(cl_event p0) { return clRetainEvent_pfn(p0); }
-#undef clReleaseEvent
-#define clReleaseEvent clReleaseEvent_fn
-inline cl_int clReleaseEvent(cl_event p0) { return clReleaseEvent_pfn(p0); }
-#undef clSetUserEventStatus
-#define clSetUserEventStatus clSetUserEventStatus_fn
-inline cl_int clSetUserEventStatus(cl_event p0, cl_int p1) { return clSetUserEventStatus_pfn(p0, p1); }
-#undef clSetEventCallback
-#define clSetEventCallback clSetEventCallback_fn
-inline cl_int clSetEventCallback(cl_event p0, cl_int p1, void (CL_CALLBACK*p2) (cl_event, cl_int, void*), void* p3) { return clSetEventCallback_pfn(p0, p1, p2, p3); }
-#undef clGetEventProfilingInfo
-#define clGetEventProfilingInfo clGetEventProfilingInfo_fn
-inline cl_int clGetEventProfilingInfo(cl_event p0, cl_profiling_info p1, size_t p2, void* p3, size_t* p4) { return clGetEventProfilingInfo_pfn(p0, p1, p2, p3, p4); }
-#undef clFlush
-#define clFlush clFlush_fn
-inline cl_int clFlush(cl_command_queue p0) { return clFlush_pfn(p0); }
-#undef clFinish
-#define clFinish clFinish_fn
-inline cl_int clFinish(cl_command_queue p0) { return clFinish_pfn(p0); }
-#undef clEnqueueReadBuffer
-#define clEnqueueReadBuffer clEnqueueReadBuffer_fn
-inline cl_int clEnqueueReadBuffer(cl_command_queue p0, cl_mem p1, cl_bool p2, size_t p3, size_t p4, void* p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueReadBuffer_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueReadBufferRect
-#define clEnqueueReadBufferRect clEnqueueReadBufferRect_fn
-inline cl_int clEnqueueReadBufferRect(cl_command_queue p0, cl_mem p1, cl_bool p2, const size_t* p3, const size_t* p4, const size_t* p5, size_t p6, size_t p7, size_t p8, size_t p9, void* p10, cl_uint p11, const cl_event* p12, cl_event* p13) { return clEnqueueReadBufferRect_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13); }
-#undef clEnqueueWriteBuffer
-#define clEnqueueWriteBuffer clEnqueueWriteBuffer_fn
-inline cl_int clEnqueueWriteBuffer(cl_command_queue p0, cl_mem p1, cl_bool p2, size_t p3, size_t p4, const void* p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueWriteBuffer_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueWriteBufferRect
-#define clEnqueueWriteBufferRect clEnqueueWriteBufferRect_fn
-inline cl_int clEnqueueWriteBufferRect(cl_command_queue p0, cl_mem p1, cl_bool p2, const size_t* p3, const size_t* p4, const size_t* p5, size_t p6, size_t p7, size_t p8, size_t p9, const void* p10, cl_uint p11, const cl_event* p12, cl_event* p13) { return clEnqueueWriteBufferRect_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13); }
-#undef clEnqueueCopyBuffer
-#define clEnqueueCopyBuffer clEnqueueCopyBuffer_fn
-inline cl_int clEnqueueCopyBuffer(cl_command_queue p0, cl_mem p1, cl_mem p2, size_t p3, size_t p4, size_t p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueCopyBuffer_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueCopyBufferRect
-#define clEnqueueCopyBufferRect clEnqueueCopyBufferRect_fn
-inline cl_int clEnqueueCopyBufferRect(cl_command_queue p0, cl_mem p1, cl_mem p2, const size_t* p3, const size_t* p4, const size_t* p5, size_t p6, size_t p7, size_t p8, size_t p9, cl_uint p10, const cl_event* p11, cl_event* p12) { return clEnqueueCopyBufferRect_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12); }
-#undef clEnqueueReadImage
-#define clEnqueueReadImage clEnqueueReadImage_fn
-inline cl_int clEnqueueReadImage(cl_command_queue p0, cl_mem p1, cl_bool p2, const size_t* p3, const size_t* p4, size_t p5, size_t p6, void* p7, cl_uint p8, const cl_event* p9, cl_event* p10) { return clEnqueueReadImage_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10); }
-#undef clEnqueueWriteImage
-#define clEnqueueWriteImage clEnqueueWriteImage_fn
-inline cl_int clEnqueueWriteImage(cl_command_queue p0, cl_mem p1, cl_bool p2, const size_t* p3, const size_t* p4, size_t p5, size_t p6, const void* p7, cl_uint p8, const cl_event* p9, cl_event* p10) { return clEnqueueWriteImage_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10); }
-#undef clEnqueueCopyImage
-#define clEnqueueCopyImage clEnqueueCopyImage_fn
-inline cl_int clEnqueueCopyImage(cl_command_queue p0, cl_mem p1, cl_mem p2, const size_t* p3, const size_t* p4, const size_t* p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueCopyImage_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueCopyImageToBuffer
-#define clEnqueueCopyImageToBuffer clEnqueueCopyImageToBuffer_fn
-inline cl_int clEnqueueCopyImageToBuffer(cl_command_queue p0, cl_mem p1, cl_mem p2, const size_t* p3, const size_t* p4, size_t p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueCopyImageToBuffer_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueCopyBufferToImage
-#define clEnqueueCopyBufferToImage clEnqueueCopyBufferToImage_fn
-inline cl_int clEnqueueCopyBufferToImage(cl_command_queue p0, cl_mem p1, cl_mem p2, size_t p3, const size_t* p4, const size_t* p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueCopyBufferToImage_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueMapBuffer
-#define clEnqueueMapBuffer clEnqueueMapBuffer_fn
-inline void* clEnqueueMapBuffer(cl_command_queue p0, cl_mem p1, cl_bool p2, cl_map_flags p3, size_t p4, size_t p5, cl_uint p6, const cl_event* p7, cl_event* p8, cl_int* p9) { return clEnqueueMapBuffer_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9); }
-#undef clEnqueueMapImage
-#define clEnqueueMapImage clEnqueueMapImage_fn
-inline void* clEnqueueMapImage(cl_command_queue p0, cl_mem p1, cl_bool p2, cl_map_flags p3, const size_t* p4, const size_t* p5, size_t* p6, size_t* p7, cl_uint p8, const cl_event* p9, cl_event* p10, cl_int* p11) { return clEnqueueMapImage_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11); }
-#undef clEnqueueUnmapMemObject
-#define clEnqueueUnmapMemObject clEnqueueUnmapMemObject_fn
-inline cl_int clEnqueueUnmapMemObject(cl_command_queue p0, cl_mem p1, void* p2, cl_uint p3, const cl_event* p4, cl_event* p5) { return clEnqueueUnmapMemObject_pfn(p0, p1, p2, p3, p4, p5); }
-#undef clEnqueueNDRangeKernel
-#define clEnqueueNDRangeKernel clEnqueueNDRangeKernel_fn
-inline cl_int clEnqueueNDRangeKernel(cl_command_queue p0, cl_kernel p1, cl_uint p2, const size_t* p3, const size_t* p4, const size_t* p5, cl_uint p6, const cl_event* p7, cl_event* p8) { return clEnqueueNDRangeKernel_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8); }
-#undef clEnqueueTask
-#define clEnqueueTask clEnqueueTask_fn
-inline cl_int clEnqueueTask(cl_command_queue p0, cl_kernel p1, cl_uint p2, const cl_event* p3, cl_event* p4) { return clEnqueueTask_pfn(p0, p1, p2, p3, p4); }
-#undef clEnqueueNativeKernel
-#define clEnqueueNativeKernel clEnqueueNativeKernel_fn
-inline cl_int clEnqueueNativeKernel(cl_command_queue p0, void (CL_CALLBACK*p1) (void*), void* p2, size_t p3, cl_uint p4, const cl_mem* p5, const void** p6, cl_uint p7, const cl_event* p8, cl_event* p9) { return clEnqueueNativeKernel_pfn(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9); }
-#undef clEnqueueMarker
-#define clEnqueueMarker clEnqueueMarker_fn
-inline cl_int clEnqueueMarker(cl_command_queue p0, cl_event* p1) { return clEnqueueMarker_pfn(p0, p1); }
-#undef clEnqueueWaitForEvents
-#define clEnqueueWaitForEvents clEnqueueWaitForEvents_fn
-inline cl_int clEnqueueWaitForEvents(cl_command_queue p0, cl_uint p1, const cl_event* p2) { return clEnqueueWaitForEvents_pfn(p0, p1, p2); }
-#undef clEnqueueBarrier
-#define clEnqueueBarrier clEnqueueBarrier_fn
-inline cl_int clEnqueueBarrier(cl_command_queue p0) { return clEnqueueBarrier_pfn(p0); }
-#undef clGetExtensionFunctionAddress
-#define clGetExtensionFunctionAddress clGetExtensionFunctionAddress_fn
-inline void* clGetExtensionFunctionAddress(const char* p0) { return clGetExtensionFunctionAddress_pfn(p0); }
-
-#endif // __OPENCV_OCL_CL_RUNTIME_OPENCL_WRAPPERS_HPP__

modules/ocl/perf/perf_arithm.cpp

@@ -99,7 +99,7 @@ PERF_TEST_P(ExpFixture, Exp, OCL_TYPICAL_MAT_SIZES)
 {
     // getting params
     const Size srcSize = GetParam();
-    const double eps = 3e-1;
+    const double eps = 1e-6;
 
     // creating src data
     Mat src(srcSize, CV_32FC1), dst(srcSize, CV_32FC1);
@@ -114,17 +114,15 @@ PERF_TEST_P(ExpFixture, Exp, OCL_TYPICAL_MAT_SIZES)
         OCL_TEST_CYCLE() cv::ocl::exp(oclSrc, oclDst);
 
         oclDst.download(dst);
-
-        SANITY_CHECK(dst, eps);
     }
     else if (RUN_PLAIN_IMPL)
     {
         TEST_CYCLE() cv::exp(src, dst);
-
-        SANITY_CHECK(dst, eps);
     }
     else
         OCL_PERF_ELSE
+
+    SANITY_CHECK(dst, eps, ERROR_RELATIVE);
 }
 
 ///////////// LOG ////////////////////////
@@ -135,7 +133,7 @@ PERF_TEST_P(LogFixture, Log, OCL_TYPICAL_MAT_SIZES)
 {
     // getting params
     const Size srcSize = GetParam();
-    const double eps = 1e-5;
+    const double eps = 1e-6;
 
     // creating src data
     Mat src(srcSize, CV_32F), dst(srcSize, src.type());
@@ -153,17 +151,15 @@ PERF_TEST_P(LogFixture, Log, OCL_TYPICAL_MAT_SIZES)
         OCL_TEST_CYCLE() cv::ocl::log(oclSrc, oclDst);
 
         oclDst.download(dst);
-
-        SANITY_CHECK(dst, eps);
     }
     else if (RUN_PLAIN_IMPL)
     {
         TEST_CYCLE() cv::log(src, dst);
-
-        SANITY_CHECK(dst, eps);
     }
     else
         OCL_PERF_ELSE
+
+    SANITY_CHECK(dst, eps, ERROR_RELATIVE);
 }
 
 ///////////// Add ////////////////////////
@@ -818,8 +814,9 @@ typedef TestBaseWithParam<Size> PowFixture;
 PERF_TEST_P(PowFixture, pow, OCL_TYPICAL_MAT_SIZES)
 {
     const Size srcSize = GetParam();
+    const double eps = 1e-6;
 
-   Mat src(srcSize, CV_32F), dst(srcSize, CV_32F);
+    Mat src(srcSize, CV_32F), dst(srcSize, CV_32F);
     declare.in(src, WARMUP_RNG).out(dst);
 
     if (RUN_OCL_IMPL)
@@ -829,17 +826,15 @@ PERF_TEST_P(PowFixture, pow, OCL_TYPICAL_MAT_SIZES)
         OCL_TEST_CYCLE() cv::ocl::pow(oclSrc, -2.0, oclDst);
 
         oclDst.download(dst);
-
-        SANITY_CHECK(dst, 5e-2);
     }
     else if (RUN_PLAIN_IMPL)
     {
         TEST_CYCLE() cv::pow(src, -2.0, dst);
-
-        SANITY_CHECK(dst, 5e-2);
     }
     else
         OCL_PERF_ELSE
+
+    SANITY_CHECK(dst, eps, ERROR_RELATIVE);
 }
 
 ///////////// AddWeighted////////////////////////

modules/ocl/src/cl_runtime/cl_runtime.cpp

@@ -1,6 +1,6 @@
 #include "precomp.hpp"
 
-#if defined(HAVE_OPENCL) && (!defined(__APPLE__) || defined(IOS))
+#if defined(HAVE_OPENCL) && !defined(HAVE_OPENCL_STATIC)
 
 #include "opencv2/ocl/cl_runtime/cl_runtime.hpp"
 
@@ -87,12 +87,6 @@ static void* opencl_check_fn(int ID)
     return func;
 }
 
-#if defined(HAVE_OPENCL12)
-#include "cl_runtime_opencl12_impl.hpp"
-#elif defined(HAVE_OPENCL11)
-#include "cl_runtime_opencl11_impl.hpp"
-#else
-#error Invalid OpenCL configuration
-#endif
+#include "cl_runtime_opencl_impl.hpp"
 
 #endif

modules/ocl/src/cl_runtime/cl_runtime_opencl11_impl.hpp

@@ -1,435 +0,0 @@
-//
-// AUTOGENERATED, DO NOT EDIT
-//
-// generated by parser_cl.py
-enum OPENCL_FN_ID {
-    OPENCL_FN_clGetPlatformIDs = 0,
-    OPENCL_FN_clGetPlatformInfo,
-    OPENCL_FN_clGetDeviceIDs,
-    OPENCL_FN_clGetDeviceInfo,
-    OPENCL_FN_clCreateContext,
-    OPENCL_FN_clCreateContextFromType,
-    OPENCL_FN_clRetainContext,
-    OPENCL_FN_clReleaseContext,
-    OPENCL_FN_clGetContextInfo,
-    OPENCL_FN_clCreateCommandQueue,
-    OPENCL_FN_clRetainCommandQueue,
-    OPENCL_FN_clReleaseCommandQueue,
-    OPENCL_FN_clGetCommandQueueInfo,
-    OPENCL_FN_clSetCommandQueueProperty,
-    OPENCL_FN_clCreateBuffer,
-    OPENCL_FN_clCreateSubBuffer,
-    OPENCL_FN_clCreateImage2D,
-    OPENCL_FN_clCreateImage3D,
-    OPENCL_FN_clRetainMemObject,
-    OPENCL_FN_clReleaseMemObject,
-    OPENCL_FN_clGetSupportedImageFormats,
-    OPENCL_FN_clGetMemObjectInfo,
-    OPENCL_FN_clGetImageInfo,
-    OPENCL_FN_clSetMemObjectDestructorCallback,
-    OPENCL_FN_clCreateSampler,
-    OPENCL_FN_clRetainSampler,
-    OPENCL_FN_clReleaseSampler,
-    OPENCL_FN_clGetSamplerInfo,
-    OPENCL_FN_clCreateProgramWithSource,
-    OPENCL_FN_clCreateProgramWithBinary,
-    OPENCL_FN_clRetainProgram,
-    OPENCL_FN_clReleaseProgram,
-    OPENCL_FN_clBuildProgram,
-    OPENCL_FN_clUnloadCompiler,
-    OPENCL_FN_clGetProgramInfo,
-    OPENCL_FN_clGetProgramBuildInfo,
-    OPENCL_FN_clCreateKernel,
-    OPENCL_FN_clCreateKernelsInProgram,
-    OPENCL_FN_clRetainKernel,
-    OPENCL_FN_clReleaseKernel,
-    OPENCL_FN_clSetKernelArg,
-    OPENCL_FN_clGetKernelInfo,
-    OPENCL_FN_clGetKernelWorkGroupInfo,
-    OPENCL_FN_clWaitForEvents,
-    OPENCL_FN_clGetEventInfo,
-    OPENCL_FN_clCreateUserEvent,
-    OPENCL_FN_clRetainEvent,
-    OPENCL_FN_clReleaseEvent,
-    OPENCL_FN_clSetUserEventStatus,
-    OPENCL_FN_clSetEventCallback,
-    OPENCL_FN_clGetEventProfilingInfo,
-    OPENCL_FN_clFlush,
-    OPENCL_FN_clFinish,
-    OPENCL_FN_clEnqueueReadBuffer,
-    OPENCL_FN_clEnqueueReadBufferRect,
-    OPENCL_FN_clEnqueueWriteBuffer,
-    OPENCL_FN_clEnqueueWriteBufferRect,
-    OPENCL_FN_clEnqueueCopyBuffer,
-    OPENCL_FN_clEnqueueCopyBufferRect,
-    OPENCL_FN_clEnqueueReadImage,
-    OPENCL_FN_clEnqueueWriteImage,
-    OPENCL_FN_clEnqueueCopyImage,
-    OPENCL_FN_clEnqueueCopyImageToBuffer,
-    OPENCL_FN_clEnqueueCopyBufferToImage,
-    OPENCL_FN_clEnqueueMapBuffer,
-    OPENCL_FN_clEnqueueMapImage,
-    OPENCL_FN_clEnqueueUnmapMemObject,
-    OPENCL_FN_clEnqueueNDRangeKernel,
-    OPENCL_FN_clEnqueueTask,
-    OPENCL_FN_clEnqueueNativeKernel,
-    OPENCL_FN_clEnqueueMarker,
-    OPENCL_FN_clEnqueueWaitForEvents,
-    OPENCL_FN_clEnqueueBarrier,
-    OPENCL_FN_clGetExtensionFunctionAddress,
-};
-// generated by parser_cl.py
-const char* opencl_fn_names[] = {
-    "clGetPlatformIDs",
-    "clGetPlatformInfo",
-    "clGetDeviceIDs",
-    "clGetDeviceInfo",
-    "clCreateContext",
-    "clCreateContextFromType",
-    "clRetainContext",
-    "clReleaseContext",
-    "clGetContextInfo",
-    "clCreateCommandQueue",
-    "clRetainCommandQueue",
-    "clReleaseCommandQueue",
-    "clGetCommandQueueInfo",
-    "clSetCommandQueueProperty",
-    "clCreateBuffer",
-    "clCreateSubBuffer",
-    "clCreateImage2D",
-    "clCreateImage3D",
-    "clRetainMemObject",
-    "clReleaseMemObject",
-    "clGetSupportedImageFormats",
-    "clGetMemObjectInfo",
-    "clGetImageInfo",
-    "clSetMemObjectDestructorCallback",
-    "clCreateSampler",
-    "clRetainSampler",
-    "clReleaseSampler",
-    "clGetSamplerInfo",
-    "clCreateProgramWithSource",
-    "clCreateProgramWithBinary",
-    "clRetainProgram",
-    "clReleaseProgram",
-    "clBuildProgram",
-    "clUnloadCompiler",
-    "clGetProgramInfo",
-    "clGetProgramBuildInfo",
-    "clCreateKernel",
-    "clCreateKernelsInProgram",
-    "clRetainKernel",
-    "clReleaseKernel",
-    "clSetKernelArg",
-    "clGetKernelInfo",
-    "clGetKernelWorkGroupInfo",
-    "clWaitForEvents",
-    "clGetEventInfo",
-    "clCreateUserEvent",
-    "clRetainEvent",
-    "clReleaseEvent",
-    "clSetUserEventStatus",
-    "clSetEventCallback",
-    "clGetEventProfilingInfo",
-    "clFlush",
-    "clFinish",
-    "clEnqueueReadBuffer",
-    "clEnqueueReadBufferRect",
-    "clEnqueueWriteBuffer",
-    "clEnqueueWriteBufferRect",
-    "clEnqueueCopyBuffer",
-    "clEnqueueCopyBufferRect",
-    "clEnqueueReadImage",
-    "clEnqueueWriteImage",
-    "clEnqueueCopyImage",
-    "clEnqueueCopyImageToBuffer",
-    "clEnqueueCopyBufferToImage",
-    "clEnqueueMapBuffer",
-    "clEnqueueMapImage",
-    "clEnqueueUnmapMemObject",
-    "clEnqueueNDRangeKernel",
-    "clEnqueueTask",
-    "clEnqueueNativeKernel",
-    "clEnqueueMarker",
-    "clEnqueueWaitForEvents",
-    "clEnqueueBarrier",
-    "clGetExtensionFunctionAddress",
-};
-
-namespace {
-// generated by parser_cl.py
-template <int ID, typename _R>
-struct opencl_fn0
-{
-    typedef _R (CL_API_CALL*FN)();
-    static _R CL_API_CALL switch_fn()
-    { return ((FN)opencl_check_fn(ID))(); }
-};
-
-template <int ID, typename _R, typename _T1>
-struct opencl_fn1
-{
-    typedef _R (CL_API_CALL*FN)(_T1);
-    static _R CL_API_CALL switch_fn(_T1 p1)
-    { return ((FN)opencl_check_fn(ID))(p1); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2>
-struct opencl_fn2
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2)
-    { return ((FN)opencl_check_fn(ID))(p1, p2); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3>
-struct opencl_fn3
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4>
-struct opencl_fn4
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5>
-struct opencl_fn5
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6>
-struct opencl_fn6
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7>
-struct opencl_fn7
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8>
-struct opencl_fn8
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8, typename _T9>
-struct opencl_fn9
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8, _T9);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8, _T9 p9)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8, p9); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8, typename _T9, typename _T10>
-struct opencl_fn10
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8, _T9, _T10);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8, _T9 p9, _T10 p10)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8, typename _T9, typename _T10, typename _T11>
-struct opencl_fn11
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8, _T9, _T10, _T11);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8, _T9 p9, _T10 p10, _T11 p11)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8, typename _T9, typename _T10, typename _T11, typename _T12>
-struct opencl_fn12
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8, _T9, _T10, _T11, _T12);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8, _T9 p9, _T10 p10, _T11 p11, _T12 p12)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8, typename _T9, typename _T10, typename _T11, typename _T12, typename _T13>
-struct opencl_fn13
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8, _T9, _T10, _T11, _T12, _T13);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8, _T9 p9, _T10 p10, _T11 p11, _T12 p12, _T13 p13)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13); }
-};
-
-template <int ID, typename _R, typename _T1, typename _T2, typename _T3, typename _T4, typename _T5, typename _T6, typename _T7, typename _T8, typename _T9, typename _T10, typename _T11, typename _T12, typename _T13, typename _T14>
-struct opencl_fn14
-{
-    typedef _R (CL_API_CALL*FN)(_T1, _T2, _T3, _T4, _T5, _T6, _T7, _T8, _T9, _T10, _T11, _T12, _T13, _T14);
-    static _R CL_API_CALL switch_fn(_T1 p1, _T2 p2, _T3 p3, _T4 p4, _T5 p5, _T6 p6, _T7 p7, _T8 p8, _T9 p9, _T10 p10, _T11 p11, _T12 p12, _T13 p13, _T14 p14)
-    { return ((FN)opencl_check_fn(ID))(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14); }
-};
-
-}
-
-// generated by parser_cl.py
-cl_int (CL_API_CALL*clGetPlatformIDs)(cl_uint, cl_platform_id*, cl_uint*) = opencl_fn3<OPENCL_FN_clGetPlatformIDs, cl_int, cl_uint, cl_platform_id*, cl_uint*>::switch_fn;
-cl_int (CL_API_CALL*clGetPlatformInfo)(cl_platform_id, cl_platform_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetPlatformInfo, cl_int, cl_platform_id, cl_platform_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clGetDeviceIDs)(cl_platform_id, cl_device_type, cl_uint, cl_device_id*, cl_uint*) = opencl_fn5<OPENCL_FN_clGetDeviceIDs, cl_int, cl_platform_id, cl_device_type, cl_uint, cl_device_id*, cl_uint*>::switch_fn;
-cl_int (CL_API_CALL*clGetDeviceInfo)(cl_device_id, cl_device_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetDeviceInfo, cl_int, cl_device_id, cl_device_info, size_t, void*, size_t*>::switch_fn;
-cl_context (CL_API_CALL*clCreateContext)(const cl_context_properties*, cl_uint, const cl_device_id*, void (CL_CALLBACK*) (const char*, const void*, size_t, void*), void*, cl_int*) = opencl_fn6<OPENCL_FN_clCreateContext, cl_context, const cl_context_properties*, cl_uint, const cl_device_id*, void (CL_CALLBACK*) (const char*, const void*, size_t, void*), void*, cl_int*>::switch_fn;
-cl_context (CL_API_CALL*clCreateContextFromType)(const cl_context_properties*, cl_device_type, void (CL_CALLBACK*) (const char*, const void*, size_t, void*), void*, cl_int*) = opencl_fn5<OPENCL_FN_clCreateContextFromType, cl_context, const cl_context_properties*, cl_device_type, void (CL_CALLBACK*) (const char*, const void*, size_t, void*), void*, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clRetainContext)(cl_context) = opencl_fn1<OPENCL_FN_clRetainContext, cl_int, cl_context>::switch_fn;
-cl_int (CL_API_CALL*clReleaseContext)(cl_context) = opencl_fn1<OPENCL_FN_clReleaseContext, cl_int, cl_context>::switch_fn;
-cl_int (CL_API_CALL*clGetContextInfo)(cl_context, cl_context_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetContextInfo, cl_int, cl_context, cl_context_info, size_t, void*, size_t*>::switch_fn;
-cl_command_queue (CL_API_CALL*clCreateCommandQueue)(cl_context, cl_device_id, cl_command_queue_properties, cl_int*) = opencl_fn4<OPENCL_FN_clCreateCommandQueue, cl_command_queue, cl_context, cl_device_id, cl_command_queue_properties, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clRetainCommandQueue)(cl_command_queue) = opencl_fn1<OPENCL_FN_clRetainCommandQueue, cl_int, cl_command_queue>::switch_fn;
-cl_int (CL_API_CALL*clReleaseCommandQueue)(cl_command_queue) = opencl_fn1<OPENCL_FN_clReleaseCommandQueue, cl_int, cl_command_queue>::switch_fn;
-cl_int (CL_API_CALL*clGetCommandQueueInfo)(cl_command_queue, cl_command_queue_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetCommandQueueInfo, cl_int, cl_command_queue, cl_command_queue_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clSetCommandQueueProperty)(cl_command_queue, cl_command_queue_properties, cl_bool, cl_command_queue_properties*) = opencl_fn4<OPENCL_FN_clSetCommandQueueProperty, cl_int, cl_command_queue, cl_command_queue_properties, cl_bool, cl_command_queue_properties*>::switch_fn;
-cl_mem (CL_API_CALL*clCreateBuffer)(cl_context, cl_mem_flags, size_t, void*, cl_int*) = opencl_fn5<OPENCL_FN_clCreateBuffer, cl_mem, cl_context, cl_mem_flags, size_t, void*, cl_int*>::switch_fn;
-cl_mem (CL_API_CALL*clCreateSubBuffer)(cl_mem, cl_mem_flags, cl_buffer_create_type, const void*, cl_int*) = opencl_fn5<OPENCL_FN_clCreateSubBuffer, cl_mem, cl_mem, cl_mem_flags, cl_buffer_create_type, const void*, cl_int*>::switch_fn;
-cl_mem (CL_API_CALL*clCreateImage2D)(cl_context, cl_mem_flags, const cl_image_format*, size_t, size_t, size_t, void*, cl_int*) = opencl_fn8<OPENCL_FN_clCreateImage2D, cl_mem, cl_context, cl_mem_flags, const cl_image_format*, size_t, size_t, size_t, void*, cl_int*>::switch_fn;
-cl_mem (CL_API_CALL*clCreateImage3D)(cl_context, cl_mem_flags, const cl_image_format*, size_t, size_t, size_t, size_t, size_t, void*, cl_int*) = opencl_fn10<OPENCL_FN_clCreateImage3D, cl_mem, cl_context, cl_mem_flags, const cl_image_format*, size_t, size_t, size_t, size_t, size_t, void*, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clRetainMemObject)(cl_mem) = opencl_fn1<OPENCL_FN_clRetainMemObject, cl_int, cl_mem>::switch_fn;
-cl_int (CL_API_CALL*clReleaseMemObject)(cl_mem) = opencl_fn1<OPENCL_FN_clReleaseMemObject, cl_int, cl_mem>::switch_fn;
-cl_int (CL_API_CALL*clGetSupportedImageFormats)(cl_context, cl_mem_flags, cl_mem_object_type, cl_uint, cl_image_format*, cl_uint*) = opencl_fn6<OPENCL_FN_clGetSupportedImageFormats, cl_int, cl_context, cl_mem_flags, cl_mem_object_type, cl_uint, cl_image_format*, cl_uint*>::switch_fn;
-cl_int (CL_API_CALL*clGetMemObjectInfo)(cl_mem, cl_mem_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetMemObjectInfo, cl_int, cl_mem, cl_mem_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clGetImageInfo)(cl_mem, cl_image_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetImageInfo, cl_int, cl_mem, cl_image_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clSetMemObjectDestructorCallback)(cl_mem, void (CL_CALLBACK*) (cl_mem, void*), void*) = opencl_fn3<OPENCL_FN_clSetMemObjectDestructorCallback, cl_int, cl_mem, void (CL_CALLBACK*) (cl_mem, void*), void*>::switch_fn;
-cl_sampler (CL_API_CALL*clCreateSampler)(cl_context, cl_bool, cl_addressing_mode, cl_filter_mode, cl_int*) = opencl_fn5<OPENCL_FN_clCreateSampler, cl_sampler, cl_context, cl_bool, cl_addressing_mode, cl_filter_mode, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clRetainSampler)(cl_sampler) = opencl_fn1<OPENCL_FN_clRetainSampler, cl_int, cl_sampler>::switch_fn;
-cl_int (CL_API_CALL*clReleaseSampler)(cl_sampler) = opencl_fn1<OPENCL_FN_clReleaseSampler, cl_int, cl_sampler>::switch_fn;
-cl_int (CL_API_CALL*clGetSamplerInfo)(cl_sampler, cl_sampler_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetSamplerInfo, cl_int, cl_sampler, cl_sampler_info, size_t, void*, size_t*>::switch_fn;
-cl_program (CL_API_CALL*clCreateProgramWithSource)(cl_context, cl_uint, const char**, const size_t*, cl_int*) = opencl_fn5<OPENCL_FN_clCreateProgramWithSource, cl_program, cl_context, cl_uint, const char**, const size_t*, cl_int*>::switch_fn;
-cl_program (CL_API_CALL*clCreateProgramWithBinary)(cl_context, cl_uint, const cl_device_id*, const size_t*, const unsigned char**, cl_int*, cl_int*) = opencl_fn7<OPENCL_FN_clCreateProgramWithBinary, cl_program, cl_context, cl_uint, const cl_device_id*, const size_t*, const unsigned char**, cl_int*, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clRetainProgram)(cl_program) = opencl_fn1<OPENCL_FN_clRetainProgram, cl_int, cl_program>::switch_fn;
-cl_int (CL_API_CALL*clReleaseProgram)(cl_program) = opencl_fn1<OPENCL_FN_clReleaseProgram, cl_int, cl_program>::switch_fn;
-cl_int (CL_API_CALL*clBuildProgram)(cl_program, cl_uint, const cl_device_id*, const char*, void (CL_CALLBACK*) (cl_program, void*), void*) = opencl_fn6<OPENCL_FN_clBuildProgram, cl_int, cl_program, cl_uint, const cl_device_id*, const char*, void (CL_CALLBACK*) (cl_program, void*), void*>::switch_fn;
-cl_int (CL_API_CALL*clUnloadCompiler)() = opencl_fn0<OPENCL_FN_clUnloadCompiler, cl_int>::switch_fn;
-cl_int (CL_API_CALL*clGetProgramInfo)(cl_program, cl_program_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetProgramInfo, cl_int, cl_program, cl_program_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clGetProgramBuildInfo)(cl_program, cl_device_id, cl_program_build_info, size_t, void*, size_t*) = opencl_fn6<OPENCL_FN_clGetProgramBuildInfo, cl_int, cl_program, cl_device_id, cl_program_build_info, size_t, void*, size_t*>::switch_fn;
-cl_kernel (CL_API_CALL*clCreateKernel)(cl_program, const char*, cl_int*) = opencl_fn3<OPENCL_FN_clCreateKernel, cl_kernel, cl_program, const char*, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clCreateKernelsInProgram)(cl_program, cl_uint, cl_kernel*, cl_uint*) = opencl_fn4<OPENCL_FN_clCreateKernelsInProgram, cl_int, cl_program, cl_uint, cl_kernel*, cl_uint*>::switch_fn;
-cl_int (CL_API_CALL*clRetainKernel)(cl_kernel) = opencl_fn1<OPENCL_FN_clRetainKernel, cl_int, cl_kernel>::switch_fn;
-cl_int (CL_API_CALL*clReleaseKernel)(cl_kernel) = opencl_fn1<OPENCL_FN_clReleaseKernel, cl_int, cl_kernel>::switch_fn;
-cl_int (CL_API_CALL*clSetKernelArg)(cl_kernel, cl_uint, size_t, const void*) = opencl_fn4<OPENCL_FN_clSetKernelArg, cl_int, cl_kernel, cl_uint, size_t, const void*>::switch_fn;
-cl_int (CL_API_CALL*clGetKernelInfo)(cl_kernel, cl_kernel_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetKernelInfo, cl_int, cl_kernel, cl_kernel_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clGetKernelWorkGroupInfo)(cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void*, size_t*) = opencl_fn6<OPENCL_FN_clGetKernelWorkGroupInfo, cl_int, cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clWaitForEvents)(cl_uint, const cl_event*) = opencl_fn2<OPENCL_FN_clWaitForEvents, cl_int, cl_uint, const cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clGetEventInfo)(cl_event, cl_event_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetEventInfo, cl_int, cl_event, cl_event_info, size_t, void*, size_t*>::switch_fn;
-cl_event (CL_API_CALL*clCreateUserEvent)(cl_context, cl_int*) = opencl_fn2<OPENCL_FN_clCreateUserEvent, cl_event, cl_context, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clRetainEvent)(cl_event) = opencl_fn1<OPENCL_FN_clRetainEvent, cl_int, cl_event>::switch_fn;
-cl_int (CL_API_CALL*clReleaseEvent)(cl_event) = opencl_fn1<OPENCL_FN_clReleaseEvent, cl_int, cl_event>::switch_fn;
-cl_int (CL_API_CALL*clSetUserEventStatus)(cl_event, cl_int) = opencl_fn2<OPENCL_FN_clSetUserEventStatus, cl_int, cl_event, cl_int>::switch_fn;
-cl_int (CL_API_CALL*clSetEventCallback)(cl_event, cl_int, void (CL_CALLBACK*) (cl_event, cl_int, void*), void*) = opencl_fn4<OPENCL_FN_clSetEventCallback, cl_int, cl_event, cl_int, void (CL_CALLBACK*) (cl_event, cl_int, void*), void*>::switch_fn;
-cl_int (CL_API_CALL*clGetEventProfilingInfo)(cl_event, cl_profiling_info, size_t, void*, size_t*) = opencl_fn5<OPENCL_FN_clGetEventProfilingInfo, cl_int, cl_event, cl_profiling_info, size_t, void*, size_t*>::switch_fn;
-cl_int (CL_API_CALL*clFlush)(cl_command_queue) = opencl_fn1<OPENCL_FN_clFlush, cl_int, cl_command_queue>::switch_fn;
-cl_int (CL_API_CALL*clFinish)(cl_command_queue) = opencl_fn1<OPENCL_FN_clFinish, cl_int, cl_command_queue>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueReadBuffer)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueReadBuffer, cl_int, cl_command_queue, cl_mem, cl_bool, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueReadBufferRect)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*) = opencl_fn14<OPENCL_FN_clEnqueueReadBufferRect, cl_int, cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueWriteBuffer)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueWriteBuffer, cl_int, cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueWriteBufferRect)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*) = opencl_fn14<OPENCL_FN_clEnqueueWriteBufferRect, cl_int, cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueCopyBuffer)(cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueCopyBuffer, cl_int, cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueCopyBufferRect)(cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, cl_uint, const cl_event*, cl_event*) = opencl_fn13<OPENCL_FN_clEnqueueCopyBufferRect, cl_int, cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, const size_t*, size_t, size_t, size_t, size_t, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueReadImage)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*) = opencl_fn11<OPENCL_FN_clEnqueueReadImage, cl_int, cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, size_t, size_t, void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueWriteImage)(cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*) = opencl_fn11<OPENCL_FN_clEnqueueWriteImage, cl_int, cl_command_queue, cl_mem, cl_bool, const size_t*, const size_t*, size_t, size_t, const void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueCopyImage)(cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueCopyImage, cl_int, cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueCopyImageToBuffer)(cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, size_t, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueCopyImageToBuffer, cl_int, cl_command_queue, cl_mem, cl_mem, const size_t*, const size_t*, size_t, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueCopyBufferToImage)(cl_command_queue, cl_mem, cl_mem, size_t, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueCopyBufferToImage, cl_int, cl_command_queue, cl_mem, cl_mem, size_t, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-void* (CL_API_CALL*clEnqueueMapBuffer)(cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event*, cl_event*, cl_int*) = opencl_fn10<OPENCL_FN_clEnqueueMapBuffer, void*, cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event*, cl_event*, cl_int*>::switch_fn;
-void* (CL_API_CALL*clEnqueueMapImage)(cl_command_queue, cl_mem, cl_bool, cl_map_flags, const size_t*, const size_t*, size_t*, size_t*, cl_uint, const cl_event*, cl_event*, cl_int*) = opencl_fn12<OPENCL_FN_clEnqueueMapImage, void*, cl_command_queue, cl_mem, cl_bool, cl_map_flags, const size_t*, const size_t*, size_t*, size_t*, cl_uint, const cl_event*, cl_event*, cl_int*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueUnmapMemObject)(cl_command_queue, cl_mem, void*, cl_uint, const cl_event*, cl_event*) = opencl_fn6<OPENCL_FN_clEnqueueUnmapMemObject, cl_int, cl_command_queue, cl_mem, void*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueNDRangeKernel)(cl_command_queue, cl_kernel, cl_uint, const size_t*, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*) = opencl_fn9<OPENCL_FN_clEnqueueNDRangeKernel, cl_int, cl_command_queue, cl_kernel, cl_uint, const size_t*, const size_t*, const size_t*, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueTask)(cl_command_queue, cl_kernel, cl_uint, const cl_event*, cl_event*) = opencl_fn5<OPENCL_FN_clEnqueueTask, cl_int, cl_command_queue, cl_kernel, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueNativeKernel)(cl_command_queue, void (CL_CALLBACK*) (void*), void*, size_t, cl_uint, const cl_mem*, const void**, cl_uint, const cl_event*, cl_event*) = opencl_fn10<OPENCL_FN_clEnqueueNativeKernel, cl_int, cl_command_queue, void (CL_CALLBACK*) (void*), void*, size_t, cl_uint, const cl_mem*, const void**, cl_uint, const cl_event*, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueMarker)(cl_command_queue, cl_event*) = opencl_fn2<OPENCL_FN_clEnqueueMarker, cl_int, cl_command_queue, cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueWaitForEvents)(cl_command_queue, cl_uint, const cl_event*) = opencl_fn3<OPENCL_FN_clEnqueueWaitForEvents, cl_int, cl_command_queue, cl_uint, const cl_event*>::switch_fn;
-cl_int (CL_API_CALL*clEnqueueBarrier)(cl_command_queue) = opencl_fn1<OPENCL_FN_clEnqueueBarrier, cl_int, cl_command_queue>::switch_fn;
-void* (CL_API_CALL*clGetExtensionFunctionAddress)(const char*) = opencl_fn1<OPENCL_FN_clGetExtensionFunctionAddress, void*, const char*>::switch_fn;
-
-// generated by parser_cl.py
-void* opencl_fn_ptrs[] = {
-    &clGetPlatformIDs,
-    &clGetPlatformInfo,
-    &clGetDeviceIDs,
-    &clGetDeviceInfo,
-    &clCreateContext,
-    &clCreateContextFromType,
-    &clRetainContext,
-    &clReleaseContext,
-    &clGetContextInfo,
-    &clCreateCommandQueue,
-    &clRetainCommandQueue,
-    &clReleaseCommandQueue,
-    &clGetCommandQueueInfo,
-    &clSetCommandQueueProperty,
-    &clCreateBuffer,
-    &clCreateSubBuffer,
-    &clCreateImage2D,
-    &clCreateImage3D,
-    &clRetainMemObject,
-    &clReleaseMemObject,
-    &clGetSupportedImageFormats,
-    &clGetMemObjectInfo,
-    &clGetImageInfo,
-    &clSetMemObjectDestructorCallback,
-    &clCreateSampler,
-    &clRetainSampler,
-    &clReleaseSampler,
-    &clGetSamplerInfo,
-    &clCreateProgramWithSource,
-    &clCreateProgramWithBinary,
-    &clRetainProgram,
-    &clReleaseProgram,
-    &clBuildProgram,
-    &clUnloadCompiler,
-    &clGetProgramInfo,
-    &clGetProgramBuildInfo,
-    &clCreateKernel,
-    &clCreateKernelsInProgram,
-    &clRetainKernel,
-    &clReleaseKernel,
-    &clSetKernelArg,
-    &clGetKernelInfo,
-    &clGetKernelWorkGroupInfo,
-    &clWaitForEvents,
-    &clGetEventInfo,
-    &clCreateUserEvent,
-    &clRetainEvent,
-    &clReleaseEvent,
-    &clSetUserEventStatus,
-    &clSetEventCallback,
-    &clGetEventProfilingInfo,
-    &clFlush,
-    &clFinish,
-    &clEnqueueReadBuffer,
-    &clEnqueueReadBufferRect,
-    &clEnqueueWriteBuffer,
-    &clEnqueueWriteBufferRect,
-    &clEnqueueCopyBuffer,
-    &clEnqueueCopyBufferRect,
-    &clEnqueueReadImage,
-    &clEnqueueWriteImage,
-    &clEnqueueCopyImage,
-    &clEnqueueCopyImageToBuffer,
-    &clEnqueueCopyBufferToImage,
-    &clEnqueueMapBuffer,
-    &clEnqueueMapImage,
-    &clEnqueueUnmapMemObject,
-    &clEnqueueNDRangeKernel,
-    &clEnqueueTask,
-    &clEnqueueNativeKernel,
-    &clEnqueueMarker,
-    &clEnqueueWaitForEvents,
-    &clEnqueueBarrier,
-    &clGetExtensionFunctionAddress,
-};

modules/ocl/src/cl_runtime/generator/generate.sh

@@ -1,7 +1,6 @@
 #!/bin/bash -e
 echo "Generate files for CL runtime..."
-cat sources/opencl11/cl.h | python parser_cl.py cl_runtime_opencl11
-cat sources/opencl12/cl.h | python parser_cl.py cl_runtime_opencl12
+cat sources/cl.h | python parser_cl.py cl_runtime_opencl
 cat sources/clAmdBlas.h | python parser_clamdblas.py
 cat sources/clAmdFft.h | python parser_clamdfft.py
 echo "Generate files for CL runtime... Done"

modules/ocl/src/cl_runtime/generator/template/cl_runtime_opencl.hpp.in

@@ -1,12 +1,6 @@
 #ifndef __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__
 #define __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__
 
-#ifdef HAVE_OPENCL
-
-#if defined __APPLE__ && !defined(IOS)
-#include <OpenCL/cl.h>
-#else
-
 @CL_REMAP_ORIGIN@
 
 #if defined __APPLE__
@@ -27,8 +21,4 @@
 
 @CL_FN_DECLARATIONS@
 
-#endif
-
-#endif
-
 #endif // __OPENCV_OCL_CL_RUNTIME_OPENCL_HPP__

modules/ocl/src/columnsum.cpp

@@ -52,25 +52,24 @@ using namespace cv::ocl;
 void cv::ocl::columnSum(const oclMat &src, oclMat &dst)
 {
     CV_Assert(src.type() == CV_32FC1);
-
     dst.create(src.size(), src.type());
 
-    Context *clCxt = src.clCxt;
-
-    const String kernelName = "columnSum";
+    int src_step = src.step / src.elemSize(), src_offset = src.offset / src.elemSize();
+    int dst_step = dst.step / dst.elemSize(), dst_offset = dst.offset / dst.elemSize();
 
     std::vector< std::pair<size_t, const void *> > args;
-
     args.push_back( std::make_pair( sizeof(cl_mem), (void *)&src.data));
     args.push_back( std::make_pair( sizeof(cl_mem), (void *)&dst.data));
     args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.cols));
     args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.rows));
-    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.step));
-    args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.step));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src_step));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst_step));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src_offset));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst_offset));
 
     size_t globalThreads[3] = {dst.cols, 1, 1};
     size_t localThreads[3]  = {256, 1, 1};
 
-    openCLExecuteKernel(clCxt, &imgproc_columnsum, kernelName, globalThreads, localThreads, args, src.channels(), src.depth());
+    openCLExecuteKernel(src.clCxt, &imgproc_columnsum, "columnSum", globalThreads, localThreads, args, src.oclchannels(), src.depth());
 
 }

modules/ocl/src/filtering.cpp

@@ -197,10 +197,10 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
               (src.rows == dst.rows));
     CV_Assert((src.oclchannels() == dst.oclchannels()));
 
-    int srcStep = src.step1() / src.oclchannels();
-    int dstStep = dst.step1() / dst.oclchannels();
-    int srcOffset = src.offset /  src.elemSize();
-    int dstOffset = dst.offset /  dst.elemSize();
+    int srcStep = src.step / src.elemSize();
+    int dstStep = dst.step / dst.elemSize();
+    int srcOffset = src.offset / src.elemSize();
+    int dstOffset = dst.offset / dst.elemSize();
 
     int srcOffset_x = srcOffset % srcStep;
     int srcOffset_y = srcOffset / srcStep;
@@ -248,6 +248,7 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
         anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
         rectKernel?"-D RECTKERNEL":"",
         s);
+
     std::vector< std::pair<size_t, const void *> > args;
     args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data));
     args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data));
@@ -261,6 +262,7 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
     args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.wholecols));
     args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.wholerows));
     args.push_back(std::make_pair(sizeof(cl_int), (void *)&dstOffset));
+
     openCLExecuteKernel(clCxt, &filtering_morph, kernelName, globalThreads, localThreads, args, -1, -1, compile_option);
 }
 
@@ -352,7 +354,7 @@ Ptr<BaseFilter_GPU> cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat
     };
 
     CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE);
-    CV_Assert(type == CV_8UC1 || type == CV_8UC3 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC1 || type == CV_32FC4);
+    CV_Assert(type == CV_8UC1 || type == CV_8UC3 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC3 || type == CV_32FC4);
 
     oclMat gpu_krnl;
     normalizeKernel(kernel, gpu_krnl);
@@ -362,9 +364,11 @@ Ptr<BaseFilter_GPU> cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat
     for(int i = 0; i < kernel.rows * kernel.cols; ++i)
         if(kernel.data[i] != 1)
             noZero = false;
-    MorphFilter_GPU* mfgpu=new MorphFilter_GPU(ksize, anchor, gpu_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]);
+
+    MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, gpu_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]);
     if(noZero)
         mfgpu->rectKernel = true;
+
     return Ptr<BaseFilter_GPU>(mfgpu);
 }
 
@@ -446,9 +450,7 @@ void morphOp(int op, const oclMat &src, oclMat &dst, const Mat &_kernel, Point a
         iterations = 1;
     }
     else
-    {
         kernel = _kernel;
-    }
 
     Ptr<FilterEngine_GPU> f = createMorphologyFilter_GPU(op, src.type(), kernel, anchor, iterations);
 
@@ -463,14 +465,10 @@ void cv::ocl::erode(const oclMat &src, oclMat &dst, const Mat &kernel, Point anc
 
     for (int i = 0; i < kernel.rows * kernel.cols; ++i)
         if (kernel.data[i] != 0)
-        {
             allZero = false;
-        }
 
     if (allZero)
-    {
         kernel.data[0] = 1;
-    }
 
     morphOp(MORPH_ERODE, src, dst, kernel, anchor, iterations, borderType, borderValue);
 }
@@ -559,7 +557,7 @@ static void GPUFilter2D(const oclMat &src, oclMat &dst, const oclMat &mat_kernel
     Context *clCxt = src.clCxt;
 
     int filterWidth = ksize.width;
-    bool ksize_3x3 = filterWidth == 3 && src.type() != CV_32FC4; // CV_32FC4 is not tuned up with filter2d_3x3 kernel
+    bool ksize_3x3 = filterWidth == 3 && src.type() != CV_32FC4 && src.type() != CV_32FC3; // CV_32FC4 is not tuned up with filter2d_3x3 kernel
 
     String kernelName = ksize_3x3 ? "filter2D_3x3" : "filter2D";
 
@@ -650,9 +648,7 @@ Ptr<BaseFilter_GPU> cv::ocl::getLinearFilter_GPU(int srcType, int dstType, const
 Ptr<FilterEngine_GPU> cv::ocl::createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Point &anchor,
         int borderType)
 {
-
     Size ksize = kernel.size();
-
     Ptr<BaseFilter_GPU> linearFilter = getLinearFilter_GPU(srcType, dstType, kernel, ksize, anchor, borderType);
 
     return createFilter2D_GPU(linearFilter);
@@ -660,11 +656,8 @@ Ptr<FilterEngine_GPU> cv::ocl::createLinearFilter_GPU(int srcType, int dstType,
 
 void cv::ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor, int borderType)
 {
-
     if (ddepth < 0)
-    {
         ddepth = src.depth();
-    }
 
     dst.create(src.size(), CV_MAKETYPE(ddepth, src.channels()));
 
@@ -1424,7 +1417,7 @@ void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, d
 
     CV_Assert(ksize == 1 || ksize == 3);
 
-    static const int K[2][9] =
+    int K[2][9] =
     {
         {0, 1, 0, 1, -4, 1, 0, 1, 0},
         {2, 0, 2, 0, -8, 0, 2, 0, 2}
@@ -1445,9 +1438,7 @@ Ptr<FilterEngine_GPU> cv::ocl::createGaussianFilter_GPU(int type, Size ksize, do
     int depth = CV_MAT_DEPTH(type);
 
     if (sigma2 <= 0)
-    {
         sigma2 = sigma1;
-    }
 
     // automatic detection of kernel size from sigma
     if (ksize.width <= 0 && sigma1 > 0)

modules/ocl/src/imgproc.cpp

@@ -183,112 +183,90 @@ namespace cv
         void remap( const oclMat &src, oclMat &dst, oclMat &map1, oclMat &map2, int interpolation, int borderType, const Scalar &borderValue )
         {
             Context *clCxt = src.clCxt;
+            bool supportsDouble = clCxt->supportsFeature(FEATURE_CL_DOUBLE);
+            if (!supportsDouble && src.depth() == CV_64F)
+            {
+                CV_Error(CV_OpenCLDoubleNotSupported, "Selected device does not support double");
+                return;
+            }
+
             CV_Assert(interpolation == INTER_LINEAR || interpolation == INTER_NEAREST
                       || interpolation == INTER_CUBIC || interpolation == INTER_LANCZOS4);
-            CV_Assert((map1.type() == CV_16SC2 && !map2.data) || (map1.type() == CV_32FC2 && !map2.data) || (map1.type() == CV_32FC1 && map2.type() == CV_32FC1));
+            CV_Assert((map1.type() == CV_16SC2 && !map2.data) || (map1.type() == CV_32FC2 && !map2.data) ||
+                      (map1.type() == CV_32FC1 && map2.type() == CV_32FC1));
             CV_Assert(!map2.data || map2.size() == map1.size());
-            CV_Assert(dst.size() == map1.size());
+            CV_Assert(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || borderType == BORDER_WRAP
+                      || borderType == BORDER_REFLECT_101 || borderType == BORDER_REFLECT);
 
             dst.create(map1.size(), src.type());
 
-            String kernelName;
+            const char * const typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" };
+            const char * const channelMap[] = { "", "", "2", "4", "4" };
+            const char * const interMap[] = { "INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_LINEAR", "INTER_LANCZOS" };
+            const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP",
+                                   "BORDER_REFLECT_101", "BORDER_TRANSPARENT" };
 
+            String kernelName = "remap";
             if ( map1.type() == CV_32FC2 && !map2.data )
-            {
-                if (interpolation == INTER_LINEAR && borderType == BORDER_CONSTANT)
-                    kernelName = "remapLNFConstant";
-                else if (interpolation == INTER_NEAREST && borderType == BORDER_CONSTANT)
-                    kernelName = "remapNNFConstant";
-            }
+                kernelName = kernelName + "_32FC2";
             else if (map1.type() == CV_16SC2 && !map2.data)
-            {
-                if (interpolation == INTER_LINEAR && borderType == BORDER_CONSTANT)
-                    kernelName = "remapLNSConstant";
-                else if (interpolation == INTER_NEAREST && borderType == BORDER_CONSTANT)
-                    kernelName = "remapNNSConstant";
-
-            }
+                kernelName = kernelName + "_16SC2";
             else if (map1.type() == CV_32FC1 && map2.type() == CV_32FC1)
-            {
-                if (interpolation == INTER_LINEAR && borderType == BORDER_CONSTANT)
-                    kernelName = "remapLNF1Constant";
-                else if (interpolation == INTER_NEAREST && borderType == BORDER_CONSTANT)
-                    kernelName = "remapNNF1Constant";
-            }
-
-            size_t blkSizeX = 16, blkSizeY = 16;
-            size_t glbSizeX;
-            int cols = dst.cols;
-            if (src.type() == CV_8UC1)
-            {
-                cols = (dst.cols + dst.offset % 4 + 3) / 4;
-                glbSizeX = cols % blkSizeX == 0 ? cols : (cols / blkSizeX + 1) * blkSizeX;
-
-            }
-            else if (src.type() == CV_32FC1 && interpolation == INTER_LINEAR)
-            {
-                cols = (dst.cols + (dst.offset >> 2) % 4 + 3) / 4;
-                glbSizeX = cols % blkSizeX == 0 ? cols : (cols / blkSizeX + 1) * blkSizeX;
-            }
+                kernelName = kernelName + "_2_32FC1";
             else
-                glbSizeX = dst.cols % blkSizeX == 0 ? dst.cols : (dst.cols / blkSizeX + 1) * blkSizeX;
+                CV_Error(Error::StsBadArg, "Unsupported map types");
 
-            size_t glbSizeY = dst.rows % blkSizeY == 0 ? dst.rows : (dst.rows / blkSizeY + 1) * blkSizeY;
-            size_t globalThreads[3] = {glbSizeX, glbSizeY, 1};
-            size_t localThreads[3] = {blkSizeX, blkSizeY, 1};
+            int ocn = dst.oclchannels();
+            size_t localThreads[3] = { 16, 16, 1};
+            size_t globalThreads[3] = { dst.cols, dst.rows, 1};
+
+            Mat scalar(1, 1, CV_MAKE_TYPE(dst.depth(), ocn), borderValue);
+            String buildOptions = format("-D %s -D %s -D T=%s%s", interMap[interpolation],
+                                         borderMap[borderType], typeMap[src.depth()], channelMap[ocn]);
+
+            if (interpolation != INTER_NEAREST)
+            {
+                int wdepth = std::max(CV_32F, dst.depth());
+                if (!supportsDouble)
+                    wdepth = std::min(CV_32F, wdepth);
+
+                buildOptions = buildOptions
+                    + format(" -D WT=%s%s -D convertToT=convert_%s%s%s -D convertToWT=convert_%s%s"
+                             " -D convertToWT2=convert_%s2 -D WT2=%s2",
+                             typeMap[wdepth], channelMap[ocn],
+                             typeMap[src.depth()], channelMap[ocn], src.depth() < CV_32F ? "_sat_rte" : "",
+                             typeMap[wdepth], channelMap[ocn],
+                             typeMap[wdepth], typeMap[wdepth]);
+            }
+
+            int src_step = src.step / src.elemSize(), src_offset = src.offset / src.elemSize();
+            int map1_step = map1.step / map1.elemSize(), map1_offset = map1.offset / map1.elemSize();
+            int map2_step = map2.step / map2.elemSize(), map2_offset = map2.offset / map2.elemSize();
+            int dst_step = dst.step / dst.elemSize(), dst_offset = dst.offset / dst.elemSize();
 
-            float borderFloat[4] = {(float)borderValue[0], (float)borderValue[1], (float)borderValue[2], (float)borderValue[3]};
             std::vector< std::pair<size_t, const void *> > args;
-            if (map1.channels() == 2)
-            {
-                args.push_back( std::make_pair(sizeof(cl_mem), (void *)&dst.data));
-                args.push_back( std::make_pair(sizeof(cl_mem), (void *)&src.data));
-                args.push_back( std::make_pair(sizeof(cl_mem), (void *)&map1.data));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.offset));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.offset));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.offset));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.step));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.step));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.step));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.cols));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.rows));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.cols));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.rows));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.cols));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.rows));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&cols));
-                float borderFloat[4] = {(float)borderValue[0], (float)borderValue[1], (float)borderValue[2], (float)borderValue[3]};
-
-                if (src.clCxt->supportsFeature(FEATURE_CL_DOUBLE))
-                    args.push_back( std::make_pair(sizeof(cl_double4), (void *)&borderValue));
-                else
-                    args.push_back( std::make_pair(sizeof(cl_float4), (void *)&borderFloat));
-            }
-            if (map1.channels() == 1)
-            {
-                args.push_back( std::make_pair(sizeof(cl_mem), (void *)&dst.data));
-                args.push_back( std::make_pair(sizeof(cl_mem), (void *)&src.data));
-                args.push_back( std::make_pair(sizeof(cl_mem), (void *)&map1.data));
+            args.push_back( std::make_pair(sizeof(cl_mem), (void *)&src.data));
+            args.push_back( std::make_pair(sizeof(cl_mem), (void *)&dst.data));
+            args.push_back( std::make_pair(sizeof(cl_mem), (void *)&map1.data));
+            if (!map2.empty())
                 args.push_back( std::make_pair(sizeof(cl_mem), (void *)&map2.data));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.offset));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.offset));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.offset));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.step));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.step));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.step));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.cols));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.rows));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.cols));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.rows));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.cols));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1.rows));
-                args.push_back( std::make_pair(sizeof(cl_int), (void *)&cols));
-                if (src.clCxt->supportsFeature(FEATURE_CL_DOUBLE))
-                    args.push_back( std::make_pair(sizeof(cl_double4), (void *)&borderValue));
-                else
-                    args.push_back( std::make_pair(sizeof(cl_float4), (void *)&borderFloat));
-            }
-            openCLExecuteKernel(clCxt, &imgproc_remap, kernelName, globalThreads, localThreads, args, src.oclchannels(), src.depth());
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&src_offset));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst_offset));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1_offset));
+            if (!map2.empty())
+                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map2_offset));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&src_step));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst_step));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&map1_step));
+            if (!map2.empty())
+                args.push_back( std::make_pair(sizeof(cl_int), (void *)&map2_step));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.cols));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&src.rows));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.cols));
+            args.push_back( std::make_pair(sizeof(cl_int), (void *)&dst.rows));
+            args.push_back( std::make_pair(scalar.elemSize(), (void *)scalar.data));
+
+            openCLExecuteKernel(clCxt, &imgproc_remap, kernelName, globalThreads, localThreads, args, -1, -1, buildOptions.c_str());
         }
 
         ////////////////////////////////////////////////////////////////////////////////////////////
@@ -409,20 +387,11 @@ namespace cv
         void medianFilter(const oclMat &src, oclMat &dst, int m)
         {
             CV_Assert( m % 2 == 1 && m > 1 );
-            CV_Assert( m <= 5 || src.depth() == CV_8U );
-            CV_Assert( src.cols <= dst.cols && src.rows <= dst.rows );
+            CV_Assert( (src.depth() == CV_8U || src.depth() == CV_32F) && (src.channels() == 1 || src.channels() == 4));
+            dst.create(src.size(), src.type());
 
-            if (src.data == dst.data)
-            {
-                oclMat src1;
-                src.copyTo(src1);
-                return medianFilter(src1, dst, m);
-            }
-
-            int srcStep = src.step1() / src.oclchannels();
-            int dstStep = dst.step1() / dst.oclchannels();
-            int srcOffset = src.offset / src.oclchannels() / src.elemSize1();
-            int dstOffset = dst.offset / dst.oclchannels() / dst.elemSize1();
+            int srcStep = src.step / src.elemSize(), dstStep = dst.step / dst.elemSize();
+            int srcOffset = src.offset /  src.elemSize(), dstOffset = dst.offset / dst.elemSize();
 
             Context *clCxt = src.clCxt;
 
@@ -458,31 +427,47 @@ namespace cv
 
         void copyMakeBorder(const oclMat &src, oclMat &dst, int top, int bottom, int left, int right, int bordertype, const Scalar &scalar)
         {
-            CV_Assert(top >= 0 && bottom >= 0 && left >= 0 && right >= 0);
-            if ((dst.cols != dst.wholecols) || (dst.rows != dst.wholerows)) //has roi
+            if (!src.clCxt->supportsFeature(FEATURE_CL_DOUBLE) && src.depth() == CV_64F)
             {
-                if (((bordertype & cv::BORDER_ISOLATED) == 0) &&
-                        (bordertype != cv::BORDER_CONSTANT) &&
-                        (bordertype != cv::BORDER_REPLICATE))
-                {
-                    CV_Error(Error::StsBadArg, "Unsupported border type");
-                }
+                CV_Error(Error::OpenCLDoubleNotSupported, "Selected device does not support double");
+                return;
             }
 
+            oclMat _src = src;
+
+            CV_Assert(top >= 0 && bottom >= 0 && left >= 0 && right >= 0);
+
+            if( _src.offset != 0 && (bordertype & BORDER_ISOLATED) == 0 )
+            {
+                Size wholeSize;
+                Point ofs;
+                _src.locateROI(wholeSize, ofs);
+                int dtop = std::min(ofs.y, top);
+                int dbottom = std::min(wholeSize.height - _src.rows - ofs.y, bottom);
+                int dleft = std::min(ofs.x, left);
+                int dright = std::min(wholeSize.width - _src.cols - ofs.x, right);
+                _src.adjustROI(dtop, dbottom, dleft, dright);
+                top -= dtop;
+                left -= dleft;
+                bottom -= dbottom;
+                right -= dright;
+            }
             bordertype &= ~cv::BORDER_ISOLATED;
+
+            // TODO need to remove this conditions and fix the code
             if (bordertype == cv::BORDER_REFLECT || bordertype == cv::BORDER_WRAP)
             {
-                CV_Assert((src.cols >= left) && (src.cols >= right) && (src.rows >= top) && (src.rows >= bottom));
+                CV_Assert((_src.cols >= left) && (_src.cols >= right) && (_src.rows >= top) && (_src.rows >= bottom));
             }
             else if (bordertype == cv::BORDER_REFLECT_101)
             {
-                CV_Assert((src.cols > left) && (src.cols > right) && (src.rows > top) && (src.rows > bottom));
+                CV_Assert((_src.cols > left) && (_src.cols > right) && (_src.rows > top) && (_src.rows > bottom));
             }
 
-            dst.create(src.rows + top + bottom, src.cols + left + right, src.type());
-            int srcStep = src.step1() / src.oclchannels(),  dstStep = dst.step1() / dst.oclchannels();
-            int srcOffset = src.offset / src.elemSize(), dstOffset = dst.offset / dst.elemSize();
-            int depth = src.depth(), ochannels = src.oclchannels();
+            dst.create(_src.rows + top + bottom, _src.cols + left + right, _src.type());
+            int srcStep = _src.step1() / _src.oclchannels(),  dstStep = dst.step1() / dst.oclchannels();
+            int srcOffset = _src.offset / _src.elemSize(), dstOffset = dst.offset / dst.elemSize();
+            int depth = _src.depth(), ochannels = _src.oclchannels();
 
             int __bordertype[] = {cv::BORDER_CONSTANT, cv::BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101};
             const char *borderstr[] = {"BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101"};
@@ -493,19 +478,19 @@ namespace cv
                     break;
 
             if (bordertype_index == sizeof(__bordertype) / sizeof(int))
-                CV_Error(Error::StsBadArg, "unsupported border type");
+                CV_Error(Error::StsBadArg, "Unsupported border type");
 
             String kernelName = "copymakeborder";
             size_t localThreads[3] = {16, 16, 1};
             size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
 
             std::vector< std::pair<size_t, const void *> > args;
-            args.push_back( std::make_pair( sizeof(cl_mem), (void *)&src.data));
+            args.push_back( std::make_pair( sizeof(cl_mem), (void *)&_src.data));
             args.push_back( std::make_pair( sizeof(cl_mem), (void *)&dst.data));
             args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.cols));
             args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.rows));
-            args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.cols));
-            args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.rows));
+            args.push_back( std::make_pair( sizeof(cl_int), (void *)&_src.cols));
+            args.push_back( std::make_pair( sizeof(cl_int), (void *)&_src.rows));
             args.push_back( std::make_pair( sizeof(cl_int), (void *)&srcStep));
             args.push_back( std::make_pair( sizeof(cl_int), (void *)&srcOffset));
             args.push_back( std::make_pair( sizeof(cl_int), (void *)&dstStep));
@@ -1324,6 +1309,8 @@ namespace cv
                 args.push_back( std::make_pair( sizeof(cl_int), (void *)&tilesX ));
                 args.push_back( std::make_pair( sizeof(cl_int), (void *)&clipLimit ));
                 args.push_back( std::make_pair( sizeof(cl_float), (void *)&lutScale ));
+                args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.offset ));
+                args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.offset ));
 
                 String kernelName = "calcLut";
                 size_t localThreads[3]  = { 32, 8, 1 };
@@ -1343,7 +1330,7 @@ namespace cv
             }
 
             static void transform(const oclMat &src, oclMat &dst, const oclMat &lut,
-                const int tilesX, const int tilesY, const cv::Size tileSize)
+                const int tilesX, const int tilesY, const Size & tileSize)
             {
                 cl_int2 tile_size;
                 tile_size.s[0] = tileSize.width;
@@ -1361,6 +1348,9 @@ namespace cv
                 args.push_back( std::make_pair( sizeof(cl_int2), (void *)&tile_size ));
                 args.push_back( std::make_pair( sizeof(cl_int), (void *)&tilesX ));
                 args.push_back( std::make_pair( sizeof(cl_int), (void *)&tilesY ));
+                args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.offset ));
+                args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.offset ));
+                args.push_back( std::make_pair( sizeof(cl_int), (void *)&lut.offset ));
 
                 size_t localThreads[3]  = { 32, 8, 1 };
                 size_t globalThreads[3] = { src.cols, src.rows, 1 };
@@ -1429,9 +1419,10 @@ namespace cv
                 }
                 else
                 {
-                    cv::ocl::copyMakeBorder(src, srcExt_, 0, tilesY_ - (src.rows % tilesY_), 0, tilesX_ - (src.cols % tilesX_), cv::BORDER_REFLECT_101, cv::Scalar());
+                    ocl::copyMakeBorder(src, srcExt_, 0, tilesY_ - (src.rows % tilesY_), 0,
+                                            tilesX_ - (src.cols % tilesX_), BORDER_REFLECT_101, Scalar::all(0));
 
-                    tileSize = cv::Size(srcExt_.cols / tilesX_, srcExt_.rows / tilesY_);
+                    tileSize = Size(srcExt_.cols / tilesX_, srcExt_.rows / tilesY_);
                     srcForLut = srcExt_;
                 }
 
@@ -1519,6 +1510,7 @@ namespace cv
             float *color_weight = &_color_weight[0];
             float *space_weight = &_space_weight[0];
             int *space_ofs = &_space_ofs[0];
+
             int dst_step_in_pixel = dst.step / dst.elemSize();
             int dst_offset_in_pixel = dst.offset / dst.elemSize();
             int temp_step_in_pixel = temp.step / temp.elemSize();
@@ -1549,7 +1541,7 @@ namespace cv
             if ((dst.type() == CV_8UC1) && ((dst.offset & 3) == 0) && ((dst.cols & 3) == 0))
             {
                 kernelName = "bilateral2";
-                globalThreads[0] = dst.cols / 4;
+                globalThreads[0] = dst.cols >> 2;
             }
 
             std::vector<std::pair<size_t , const void *> > args;
@@ -1567,15 +1559,17 @@ namespace cv
             args.push_back( std::make_pair( sizeof(cl_mem), (void *)&oclcolor_weight.data ));
             args.push_back( std::make_pair( sizeof(cl_mem), (void *)&oclspace_weight.data ));
             args.push_back( std::make_pair( sizeof(cl_mem), (void *)&oclspace_ofs.data ));
+
             openCLExecuteKernel(src.clCxt, &imgproc_bilateral, kernelName, globalThreads, localThreads, args, dst.oclchannels(), dst.depth());
         }
+
         void bilateralFilter(const oclMat &src, oclMat &dst, int radius, double sigmaclr, double sigmaspc, int borderType)
         {
             dst.create( src.size(), src.type() );
             if ( src.depth() == CV_8U )
                 oclbilateralFilter_8u( src, dst, radius, sigmaclr, sigmaspc, borderType );
             else
-                CV_Error(Error::StsUnsupportedFormat, "Bilateral filtering is only implemented for 8uimages");
+                CV_Error(Error::StsUnsupportedFormat, "Bilateral filtering is only implemented for CV_8U images");
         }
 
     }
@@ -1732,32 +1726,31 @@ static void convolve_run(const oclMat &src, const oclMat &temp1, oclMat &dst, St
     CV_Assert(src.cols == dst.cols && src.rows == dst.rows);
     CV_Assert(src.type() == dst.type());
 
-    Context  *clCxt = src.clCxt;
-    int channels = dst.oclchannels();
-    int depth = dst.depth();
-
-    size_t vector_length = 1;
-    int offset_cols = ((dst.offset % dst.step) / dst.elemSize1()) & (vector_length - 1);
-    int cols = divUp(dst.cols * channels + offset_cols, vector_length);
-    int rows = dst.rows;
-
     size_t localThreads[3]  = { 16, 16, 1 };
-    size_t globalThreads[3] = { cols, rows, 1 };
+    size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
+
+    int src_step = src.step / src.elemSize(), src_offset = src.offset / src.elemSize();
+    int dst_step = dst.step / dst.elemSize(), dst_offset = dst.offset / dst.elemSize();
+    int temp1_step = temp1.step / temp1.elemSize(), temp1_offset = temp1.offset / temp1.elemSize();
 
     std::vector<std::pair<size_t , const void *> > args;
     args.push_back( std::make_pair( sizeof(cl_mem), (void *)&src.data ));
     args.push_back( std::make_pair( sizeof(cl_mem), (void *)&temp1.data ));
     args.push_back( std::make_pair( sizeof(cl_mem), (void *)&dst.data ));
     args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.rows ));
-    args.push_back( std::make_pair( sizeof(cl_int), (void *)&cols ));
-    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.step ));
-    args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst.step ));
-    args.push_back( std::make_pair( sizeof(cl_int), (void *)&temp1.step ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src.cols ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src_step ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst_step ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&temp1_step ));
     args.push_back( std::make_pair( sizeof(cl_int), (void *)&temp1.rows ));
     args.push_back( std::make_pair( sizeof(cl_int), (void *)&temp1.cols ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&src_offset ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&dst_offset ));
+    args.push_back( std::make_pair( sizeof(cl_int), (void *)&temp1_offset ));
 
-    openCLExecuteKernel(clCxt, source, kernelName, globalThreads, localThreads, args, -1, depth);
+    openCLExecuteKernel(src.clCxt, source, kernelName, globalThreads, localThreads, args, -1, dst.depth());
 }
+
 void cv::ocl::convolve(const oclMat &x, const oclMat &t, oclMat &y, bool ccorr)
 {
     CV_Assert(x.depth() == CV_32F);

modules/ocl/src/opencl/filtering_laplacian.cl

@@ -169,6 +169,7 @@ __kernel void filter2D(
     int globalRow = groupStartRow + localRow;
     const int src_offset = mad24(src_offset_y, src_step, src_offset_x);
     const int dst_offset = mad24(dst_offset_y, dst_step, dst_offset_x);
+
 #ifdef BORDER_CONSTANT
     for(int i = localRow; i < LOCAL_HEIGHT; i += get_local_size(1))
     {
@@ -208,6 +209,7 @@ __kernel void filter2D(
         }
     }
 #endif
+
     barrier(CLK_LOCAL_MEM_FENCE);
     if(globalRow < rows && globalCol < cols)
     {
@@ -231,6 +233,7 @@ __kernel void filter2D(
 //////////////////////////////////////////////////////////////////////////////////////////////////////
 /////////////////////////////Macro for define elements number per thread/////////////////////////////
 ////////////////////////////////////////////////////////////////////////////////////////////////////
+
 #define ANX                     1
 #define ANY                     1
 
@@ -249,6 +252,7 @@ __kernel void filter2D(
 ///////////////////////////////////////////////////////////////////////////////////////////////////
 /////////////////////////////////////////8uC1////////////////////////////////////////////////////////
 ////////////////////////////////////////////////////////////////////////////////////////////////////
+
 __kernel void filter2D_3x3(
     __global T_IMG *src,
     __global T_IMG *dst,
@@ -359,6 +363,7 @@ __kernel void filter2D_3x3(
                 }
             }
         }
+
         if(dst_rows_index < dst_rows_end)
         {
             T_IMGx4 tmp_dst = CONVERT_TYPEx4(sum);

modules/ocl/src/opencl/filtering_morph.cl

@@ -45,6 +45,7 @@
 //BORDER_CONSTANT:      iiiiii|abcdefgh|iiiiiii
 #define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
 #ifndef GENTYPE
+
 __kernel void morph_C1_D0(__global const uchar * restrict src,
                           __global uchar *dst,
                           int src_offset_x, int src_offset_y,
@@ -150,7 +151,9 @@ __kernel void morph_C1_D0(__global const uchar * restrict src,
         }
     }
 }
+
 #else
+
 __kernel void morph(__global const GENTYPE * restrict src,
                     __global GENTYPE *dst,
                     int src_offset_x, int src_offset_y,
@@ -221,4 +224,5 @@ __kernel void morph(__global const GENTYPE * restrict src,
         dst[out_addr] = res;
     }
 }
+
 #endif

modules/ocl/src/opencl/imgproc_bilateral.cl

@@ -47,25 +47,27 @@ __kernel void bilateral_C1_D0(__global uchar *dst,
         __constant float *space_weight,
         __constant int *space_ofs)
 {
-    int gidx = get_global_id(0);
-    int gidy = get_global_id(1);
-    if((gidy<dst_rows) && (gidx<dst_cols))
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < dst_rows && x < dst_cols)
     {
-        int src_addr = mad24(gidy+radius,src_step,gidx+radius);
-        int dst_addr = mad24(gidy,dst_step,gidx+dst_offset);
+        int src_index = mad24(y + radius, src_step, x + radius);
+        int dst_index = mad24(y, dst_step, x + dst_offset);
         float sum = 0.f, wsum = 0.f;
 
-        int val0 = (int)src[src_addr];
+        int val0 = (int)src[src_index];
         for(int k = 0; k < maxk; k++ )
         {
-            int val = (int)src[src_addr + space_ofs[k]];
-            float w = space_weight[k]*color_weight[abs(val - val0)];
-            sum += (float)(val)*w;
+            int val = (int)src[src_index + space_ofs[k]];
+            float w = space_weight[k] * color_weight[abs(val - val0)];
+            sum += (float)(val) * w;
             wsum += w;
         }
-        dst[dst_addr] = convert_uchar_rtz(sum/wsum+0.5f);
+        dst[dst_index] = convert_uchar_rtz(sum / wsum + 0.5f);
     }
 }
+
 __kernel void bilateral2_C1_D0(__global uchar *dst,
         __global const uchar *src,
         const int dst_rows,
@@ -81,25 +83,28 @@ __kernel void bilateral2_C1_D0(__global uchar *dst,
         __constant float *space_weight,
         __constant int *space_ofs)
 {
-    int gidx = get_global_id(0)<<2;
-    int gidy = get_global_id(1);
-    if((gidy<dst_rows) && (gidx<dst_cols))
+    int x = get_global_id(0) << 2;
+    int y = get_global_id(1);
+
+    if (y < dst_rows && x < dst_cols)
     {
-        int src_addr = mad24(gidy+radius,src_step,gidx+radius);
-        int dst_addr = mad24(gidy,dst_step,gidx+dst_offset);
+        int src_index = mad24(y + radius, src_step, x + radius);
+        int dst_index = mad24(y, dst_step, x + dst_offset);
         float4 sum = (float4)(0.f), wsum = (float4)(0.f);
 
-        int4 val0 = convert_int4(vload4(0,src+src_addr));
+        int4 val0 = convert_int4(vload4(0,src + src_index));
         for(int k = 0; k < maxk; k++ )
         {
-            int4 val = convert_int4(vload4(0,src+src_addr + space_ofs[k]));
-            float4 w = (float4)(space_weight[k])*(float4)(color_weight[abs(val.x - val0.x)],color_weight[abs(val.y - val0.y)],color_weight[abs(val.z - val0.z)],color_weight[abs(val.w - val0.w)]);
-            sum += convert_float4(val)*w;
+            int4 val = convert_int4(vload4(0,src+src_index + space_ofs[k]));
+            float4 w = (float4)(space_weight[k]) * (float4)(color_weight[abs(val.x - val0.x)], color_weight[abs(val.y - val0.y)],
+                color_weight[abs(val.z - val0.z)], color_weight[abs(val.w - val0.w)]);
+            sum += convert_float4(val) * w;
             wsum += w;
         }
-        *(__global uchar4*)(dst+dst_addr) = convert_uchar4_rtz(sum/wsum+0.5f);
+        *(__global uchar4*)(dst+dst_index) = convert_uchar4_rtz(sum/wsum+0.5f);
     }
 }
+
 __kernel void bilateral_C4_D0(__global uchar4 *dst,
         __global const uchar4 *src,
         const int dst_rows,
@@ -115,24 +120,26 @@ __kernel void bilateral_C4_D0(__global uchar4 *dst,
         __constant float *space_weight,
         __constant int *space_ofs)
 {
-    int gidx = get_global_id(0);
-    int gidy = get_global_id(1);
-    if((gidy<dst_rows) && (gidx<dst_cols))
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < dst_rows && x < dst_cols)
     {
-        int src_addr = mad24(gidy+radius,src_step,gidx+radius);
-        int dst_addr = mad24(gidy,dst_step,gidx+dst_offset);
+        int src_index = mad24(y + radius, src_step, x + radius);
+        int dst_index = mad24(y, dst_step, x + dst_offset);
         float4 sum = (float4)0.f;
         float wsum = 0.f;
 
-        int4 val0 = convert_int4(src[src_addr]);
+        int4 val0 = convert_int4(src[src_index]);
         for(int k = 0; k < maxk; k++ )
         {
-            int4 val = convert_int4(src[src_addr + space_ofs[k]]);
-            float w = space_weight[k]*color_weight[abs(val.x - val0.x)+abs(val.y - val0.y)+abs(val.z - val0.z)];
-            sum += convert_float4(val)*(float4)w;
+            int4 val = convert_int4(src[src_index + space_ofs[k]]);
+            float w = space_weight[k] * color_weight[abs(val.x - val0.x) + abs(val.y - val0.y) + abs(val.z - val0.z)];
+            sum += convert_float4(val) * (float4)w;
             wsum += w;
         }
-        wsum=1.f/wsum;
-        dst[dst_addr] = convert_uchar4_rtz(sum*(float4)wsum+(float4)0.5f);
+
+        wsum = 1.f / wsum;
+        dst[dst_index] = convert_uchar4_rtz(sum * (float4)wsum + (float4)0.5f);
     }
 }

modules/ocl/src/opencl/imgproc_clahe.cl

@@ -53,12 +53,8 @@ int calc_lut(__local int* smem, int val, int tid)
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid == 0)
-    {
         for (int i = 1; i < 256; ++i)
-        {
             smem[i] += smem[i - 1];
-        }
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     return smem[tid];
@@ -71,69 +67,51 @@ void reduce(volatile __local int* smem, int val, int tid)
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 128)
-    {
         smem[tid] = val += smem[tid + 128];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 64)
-    {
         smem[tid] = val += smem[tid + 64];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 32)
-    {
         smem[tid] += smem[tid + 32];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 16)
-    {
         smem[tid] += smem[tid + 16];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 8)
-    {
         smem[tid] += smem[tid + 8];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 4)
-    {
         smem[tid] += smem[tid + 4];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 2)
-    {
         smem[tid] += smem[tid + 2];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 1)
-    {
         smem[256] = smem[tid] + smem[tid + 1];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 }
+
 #else
+
 void reduce(__local volatile int* smem, int val, int tid)
 {
     smem[tid] = val;
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 128)
-    {
         smem[tid] = val += smem[tid + 128];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 64)
-    {
         smem[tid] = val += smem[tid + 64];
-    }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (tid < 32)
@@ -141,12 +119,17 @@ void reduce(__local volatile int* smem, int val, int tid)
         smem[tid] += smem[tid + 32];
 #if WAVE_SIZE < 32
     } barrier(CLK_LOCAL_MEM_FENCE);
-    if (tid < 16) {
+
+    if (tid < 16)
+    {
 #endif
         smem[tid] += smem[tid + 16];
 #if WAVE_SIZE < 16
-    } barrier(CLK_LOCAL_MEM_FENCE);
-    if (tid < 8) {
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (tid < 8)
+    {
 #endif
         smem[tid] += smem[tid + 8];
         smem[tid] += smem[tid + 4];
@@ -159,7 +142,8 @@ void reduce(__local volatile int* smem, int val, int tid)
 __kernel void calcLut(__global __const uchar * src, __global uchar * lut,
                       const int srcStep, const int dstStep,
                       const int2 tileSize, const int tilesX,
-                      const int clipLimit, const float lutScale)
+                      const int clipLimit, const float lutScale,
+                      const int src_offset, const int dst_offset)
 {
     __local int smem[512];
 
@@ -173,25 +157,21 @@ __kernel void calcLut(__global __const uchar * src, __global uchar * lut,
 
     for (int i = get_local_id(1); i < tileSize.y; i += get_local_size(1))
     {
-        __global const uchar* srcPtr = src + mad24( ty * tileSize.y + i,
-                                                    srcStep, tx * tileSize.x );
+        __global const uchar* srcPtr = src + mad24(ty * tileSize.y + i, srcStep, tx * tileSize.x + src_offset);
         for (int j = get_local_id(0); j < tileSize.x; j += get_local_size(0))
         {
             const int data = srcPtr[j];
             atomic_inc(&smem[data]);
         }
     }
-
     barrier(CLK_LOCAL_MEM_FENCE);
 
     int tHistVal = smem[tid];
-
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (clipLimit > 0)
     {
         // clip histogram bar
-
         int clipped = 0;
         if (tHistVal > clipLimit)
         {
@@ -200,7 +180,6 @@ __kernel void calcLut(__global __const uchar * src, __global uchar * lut,
         }
 
         // find number of overall clipped samples
-
         reduce(smem, clipped, tid);
         barrier(CLK_LOCAL_MEM_FENCE);
 #ifdef CPU
@@ -229,7 +208,7 @@ __kernel void calcLut(__global __const uchar * src, __global uchar * lut,
 
     const int lutVal = calc_lut(smem, tHistVal, tid);
     uint ires = (uint)convert_int_rte(lutScale * lutVal);
-    lut[(ty * tilesX + tx) * dstStep + tid] =
+    lut[(ty * tilesX + tx) * dstStep + tid + dst_offset] =
         convert_uchar(clamp(ires, (uint)0, (uint)255));
 }
 
@@ -239,7 +218,8 @@ __kernel void transform(__global __const uchar * src,
                         const int srcStep, const int dstStep, const int lutStep,
                         const int cols, const int rows,
                         const int2 tileSize,
-                        const int tilesX, const int tilesY)
+                        const int tilesX, const int tilesY,
+                        const int src_offset, const int dst_offset, int lut_offset)
 {
     const int x = get_global_id(0);
     const int y = get_global_id(1);
@@ -261,15 +241,15 @@ __kernel void transform(__global __const uchar * src,
     tx1 = max(tx1, 0);
     tx2 = min(tx2, tilesX - 1);
 
-    const int srcVal = src[mad24(y, srcStep, x)];
+    const int srcVal = src[mad24(y, srcStep, x + src_offset)];
 
     float res = 0;
 
-    res += lut[mad24(ty1 * tilesX + tx1, lutStep, srcVal)] * ((1.0f - xa) * (1.0f - ya));
-    res += lut[mad24(ty1 * tilesX + tx2, lutStep, srcVal)] * ((xa) * (1.0f - ya));
-    res += lut[mad24(ty2 * tilesX + tx1, lutStep, srcVal)] * ((1.0f - xa) * (ya));
-    res += lut[mad24(ty2 * tilesX + tx2, lutStep, srcVal)] * ((xa) * (ya));
+    res += lut[mad24(ty1 * tilesX + tx1, lutStep, srcVal + lut_offset)] * ((1.0f - xa) * (1.0f - ya));
+    res += lut[mad24(ty1 * tilesX + tx2, lutStep, srcVal + lut_offset)] * ((xa) * (1.0f - ya));
+    res += lut[mad24(ty2 * tilesX + tx1, lutStep, srcVal + lut_offset)] * ((1.0f - xa) * (ya));
+    res += lut[mad24(ty2 * tilesX + tx2, lutStep, srcVal + lut_offset)] * ((xa) * (ya));
 
     uint ires = (uint)convert_int_rte(res);
-    dst[mad24(y, dstStep, x)] = convert_uchar(clamp(ires, (uint)0, (uint)255));
+    dst[mad24(y, dstStep, x + dst_offset)] = convert_uchar(clamp(ires, (uint)0, (uint)255));
 }

modules/ocl/src/opencl/imgproc_columnsum.cl

@@ -43,38 +43,28 @@
 //
 //M*/
 
-#pragma OPENCL EXTENSION cl_amd_printf : enable
-#if defined (__ATI__)
-#pragma OPENCL EXTENSION cl_amd_fp64:enable
-
-#elif defined (__NVIDIA__)
-#pragma OPENCL EXTENSION cl_khr_fp64:enable
-#endif
-
 ////////////////////////////////////////////////////////////////////
 ///////////////////////// columnSum ////////////////////////////////
 ////////////////////////////////////////////////////////////////////
-/// CV_32FC1
-__kernel void columnSum_C1_D5(__global float* src,__global float* dst,int srcCols,int srcRows,int srcStep,int dstStep)
+
+__kernel void columnSum_C1_D5(__global float * src, __global float * dst,
+    int cols, int rows, int src_step, int dst_step, int src_offset, int dst_offset)
 {
     const int x = get_global_id(0);
 
-    srcStep >>= 2;
-    dstStep >>= 2;
-
-    if (x < srcCols)
+    if (x < cols)
     {
-        int srcIdx = x ;
-        int dstIdx = x ;
+        int srcIdx = x + src_offset;
+        int dstIdx = x + dst_offset;
 
         float sum = 0;
 
-        for (int y = 0; y < srcRows; ++y)
+        for (int y = 0; y < rows; ++y)
         {
             sum += src[srcIdx];
             dst[dstIdx] = sum;
-            srcIdx += srcStep;
-            dstIdx += dstStep;
+            srcIdx += src_step;
+            dstIdx += dst_step;
         }
     }
 }

modules/ocl/src/opencl/imgproc_convolve.cl

@@ -48,9 +48,12 @@
 #elif defined (__NVIDIA__)
 #pragma OPENCL EXTENSION cl_khr_fp64:enable
 #endif
+
 /************************************** convolve **************************************/
-__kernel void convolve_D5 (__global float *src, __global float *temp1, __global float *dst,
-                                  int rows, int cols, int src_step, int dst_step,int k_step, int kWidth, int kHeight)
+
+__kernel void convolve_D5(__global float *src, __global float *temp1, __global float *dst,
+                          int rows, int cols, int src_step, int dst_step,int k_step, int kWidth, int kHeight,
+                          int src_offset, int dst_offset, int koffset)
 {
     __local float smem[16 + 2 * 8][16 + 2 * 8];
 
@@ -65,7 +68,7 @@ __kernel void convolve_D5 (__global float *src, __global float *temp1, __global
             // 0 | 0 0 | 0
             // -----------
             // 0 | 0 0 | 0
-    smem[y][x] = src[min(max(gy - 8, 0), rows - 1)*(src_step >> 2) + min(max(gx - 8, 0), cols - 1)];
+    smem[y][x] = src[min(max(gy - 8, 0), rows - 1) * src_step + min(max(gx - 8, 0), cols - 1) + src_offset];
 
             // 0 | 0 x | x
             // -----------
@@ -73,7 +76,7 @@ __kernel void convolve_D5 (__global float *src, __global float *temp1, __global
             // 0 | 0 0 | 0
             // -----------
             // 0 | 0 0 | 0
-    smem[y][x + 16] = src[min(max(gy - 8, 0), rows - 1)*(src_step >> 2) + min(gx + 8, cols - 1)];
+    smem[y][x + 16] = src[min(max(gy - 8, 0), rows - 1) * src_step + min(gx + 8, cols - 1) + src_offset];
 
             // 0 | 0 0 | 0
             // -----------
@@ -81,7 +84,7 @@ __kernel void convolve_D5 (__global float *src, __global float *temp1, __global
             // x | x 0 | 0
             // -----------
             // x | x 0 | 0
-    smem[y + 16][x] = src[min(gy + 8, rows - 1)*(src_step >> 2) + min(max(gx - 8, 0), cols - 1)];
+    smem[y + 16][x] = src[min(gy + 8, rows - 1) * src_step + min(max(gx - 8, 0), cols - 1) + src_offset];
 
             // 0 | 0 0 | 0
             // -----------
@@ -89,21 +92,18 @@ __kernel void convolve_D5 (__global float *src, __global float *temp1, __global
             // 0 | 0 x | x
             // -----------
             // 0 | 0 x | x
-    smem[y + 16][x + 16] = src[min(gy + 8, rows - 1)*(src_step >> 2) + min(gx + 8, cols - 1)];
+    smem[y + 16][x + 16] = src[min(gy + 8, rows - 1) * src_step + min(gx + 8, cols - 1) + src_offset];
 
     barrier(CLK_LOCAL_MEM_FENCE);
 
     if (gx < cols && gy < rows)
     {
-       float res = 0;
+        float res = 0;
 
         for (int i = 0; i < kHeight; ++i)
-        {
             for (int j = 0; j < kWidth; ++j)
-            {
-                res += smem[y + 8 - kHeight / 2 + i][x + 8 - kWidth / 2 + j] * temp1[i * (k_step>>2) + j];
-            }
-        }
-        dst[gy*(dst_step >> 2)+gx] = res;
-   }
+                res += smem[y + 8 - kHeight / 2 + i][x + 8 - kWidth / 2 + j] * temp1[i * k_step + j + koffset];
+
+        dst[gy * dst_step + gx + dst_offset] = res;
+    }
 }

modules/ocl/src/opencl/imgproc_copymakeboder.cl

@@ -34,6 +34,13 @@
 //
 //
 
+#if defined (DOUBLE_SUPPORT)
+#ifdef cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#elif defined (cl_amd_fp64)
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#endif
+#endif
 
 #ifdef BORDER_CONSTANT
 //BORDER_CONSTANT:      iiiiii|abcdefgh|iiiiiii

modules/ocl/src/svm.cpp

@@ -46,7 +46,7 @@
 #include "opencl_kernels.hpp"
 
 // TODO Remove this after HAVE_CLAMDBLAS eliminating
-#ifdef __GNUC__
+#if defined(__GNUC__) && (__GNUC__ == 4) && (__GNUC_MINOR__ == 8)
 #  pragma GCC diagnostic ignored "-Wunused-but-set-variable"
 #endif
 

modules/ocl/test/main.cpp

@@ -55,4 +55,24 @@
 
 #include "opencv2/ocl/private/opencl_dumpinfo.hpp"
 
-CV_TEST_MAIN(".", dumpOpenCLDevice())
+int LOOP_TIMES = 1;
+
+void readLoopTimes(int argc, char ** argv)
+{
+    const char * const command_line_keys =
+            "{   test_loop_times             |1        |count of iterations per each test}"
+            "{h  help                        |false    |print help info}";
+
+    cv::CommandLineParser parser(argc, argv, command_line_keys);
+    if (parser.has("help"))
+    {
+        std::cout << "\nAvailable options besides google test option: \n";
+        parser.printMessage();
+    }
+
+    LOOP_TIMES = parser.get<int>("test_loop_times");
+    CV_Assert(LOOP_TIMES > 0);
+}
+
+CV_TEST_MAIN(".", dumpOpenCLDevice(),
+                  readLoopTimes(argc, argv))

modules/ocl/test/test_arithm.cpp

@@ -63,7 +63,7 @@ using namespace std;
 
 //////////////////////////////// LUT /////////////////////////////////////////////////
 
-PARAM_TEST_CASE(Lut, int, int, bool, bool)
+PARAM_TEST_CASE(Lut, MatDepth, MatDepth, bool, bool)
 {
     int lut_depth;
     int cn;
@@ -141,7 +141,7 @@ OCL_TEST_P(Lut, Mat)
 
 ///////////////////////// ArithmTestBase ///////////////////////////
 
-PARAM_TEST_CASE(ArithmTestBase, int, int, bool)
+PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
 {
     int depth;
     int cn;
@@ -1490,38 +1490,38 @@ OCL_TEST_P(Norm, NORM_L2)
 
 //////////////////////////////////////// Instantiation /////////////////////////////////////////
 
-INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool(), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(testing::Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(testing::Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Sub, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Mul, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Div, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Min, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Max, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Abs, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Absdiff, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Transpose, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Flip, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, MinMax, Combine(testing::Range(CV_8U, CV_USRTYPE1), Values(1), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, MinMaxLoc, Combine(testing::Range(CV_8U, CV_USRTYPE1), Values(1), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Sum, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, SqrSum, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, AbsSum, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, CountNonZero, Combine(testing::Range(CV_8U, CV_USRTYPE1), Values(1), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_and, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_or, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_xor, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_not, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Compare, Combine(testing::Range(CV_8U, CV_USRTYPE1), Values(1), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Pow, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, SetIdentity, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
-INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool(), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(testing::Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(testing::Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Sub, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Mul, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Div, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Min, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Max, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Abs, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Absdiff, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine(Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine(Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Transpose, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Flip, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, MinMax, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(Channels(1)), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, MinMaxLoc, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(Channels(1)), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Sum, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, SqrSum, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, AbsSum, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, CountNonZero, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(Channels(1)), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_and, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_or, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_xor, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_not, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Compare, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(Channels(1)), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Pow, Combine(Values(CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, SetIdentity, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
 
 #endif // HAVE_OPENCL

modules/ocl/test/test_color.cpp

@@ -49,148 +49,194 @@ using namespace cv;
 
 #ifdef HAVE_OPENCL
 
-//#define MAT_DEBUG
-#ifdef MAT_DEBUG
-#define MAT_DIFF(mat, mat2)\
-{\
-    for(int i = 0; i < mat.rows; i ++)\
-    {\
-        for(int j = 0; j < mat.cols; j ++)\
-        {\
-            cv::Vec4b s = mat.at<cv::Vec4b>(i, j);\
-            cv::Vec4b s2 = mat2.at<cv::Vec4b>(i, j);\
-            if(s != s2) printf("*");\
-            else printf(".");\
-        }\
-        puts("\n");\
-    }\
-}
-#else
-#define MAT_DIFF(mat, mat2)
-#endif
-
-
 namespace
 {
+using namespace testing;
 
 ///////////////////////////////////////////////////////////////////////////////////////////////////////
 // cvtColor
-PARAM_TEST_CASE(CvtColor, cv::Size, MatDepth)
-{
-    cv::Size size;
-    int depth;
-    bool useRoi;
 
-    cv::Mat img;
+PARAM_TEST_CASE(CvtColor, MatDepth, bool)
+{
+    int depth;
+    bool use_roi;
+
+    // src mat
+    cv::Mat src1;
+    cv::Mat dst1;
+
+    // src mat with roi
+    cv::Mat src1_roi;
+    cv::Mat dst1_roi;
+
+    // ocl dst mat for testing
+    cv::ocl::oclMat gsrc1_whole;
+    cv::ocl::oclMat gdst1_whole;
+
+    // ocl mat with roi
+    cv::ocl::oclMat gsrc1_roi;
+    cv::ocl::oclMat gdst1_roi;
 
     virtual void SetUp()
     {
-        size = GET_PARAM(0);
-        depth = GET_PARAM(1);
+        depth = GET_PARAM(0);
+        use_roi = GET_PARAM(1);
+    }
 
-        img = randomMat(size, CV_MAKE_TYPE(depth, 3), 0.0, depth == CV_32F ? 1.0 : 255.0);
+    virtual void random_roi(int channelsIn, int channelsOut)
+    {
+        const int srcType = CV_MAKE_TYPE(depth, channelsIn);
+        const int dstType = CV_MAKE_TYPE(depth, channelsOut);
+
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, roiSize, srcBorder, srcType, 2, 100);
+
+        Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst1, dst1_roi, roiSize, dst1Border, dstType, 5, 16);
+
+        generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder);
+        generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border);
+    }
+
+    void Near(double threshold = 1e-3)
+    {
+        EXPECT_MAT_NEAR(dst1, gdst1_whole, threshold);
+        EXPECT_MAT_NEAR(dst1_roi, gdst1_roi, threshold);
+    }
+
+    void doTest(int channelsIn, int channelsOut, int code)
+    {
+        for (int j = 0; j < LOOP_TIMES; j++)
+        {
+            random_roi(channelsIn, channelsOut);
+
+            cv::cvtColor(src1_roi, dst1_roi, code);
+            cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code);
+
+            Near();
+        }
     }
 };
 
 #define CVTCODE(name) cv::COLOR_ ## name
-#define OCL_TEST_P_CVTCOLOR(name) OCL_TEST_P(CvtColor, name)\
-{\
-    cv::Mat src = img;\
-    cv::ocl::oclMat ocl_img, dst;\
-    ocl_img.upload(img);\
-    cv::ocl::cvtColor(ocl_img, dst, CVTCODE(name));\
-    cv::Mat dst_gold;\
-    cv::cvtColor(src, dst_gold, CVTCODE(name));\
-    cv::Mat dst_mat;\
-    dst.download(dst_mat);\
-    EXPECT_MAT_NEAR(dst_gold, dst_mat, 1e-5);\
-}
 
-//add new ones here using macro
-OCL_TEST_P_CVTCOLOR(RGB2GRAY)
-OCL_TEST_P_CVTCOLOR(BGR2GRAY)
-OCL_TEST_P_CVTCOLOR(RGBA2GRAY)
-OCL_TEST_P_CVTCOLOR(BGRA2GRAY)
-
-OCL_TEST_P_CVTCOLOR(RGB2YUV)
-OCL_TEST_P_CVTCOLOR(BGR2YUV)
-OCL_TEST_P_CVTCOLOR(YUV2RGB)
-OCL_TEST_P_CVTCOLOR(YUV2BGR)
-OCL_TEST_P_CVTCOLOR(RGB2YCrCb)
-OCL_TEST_P_CVTCOLOR(BGR2YCrCb)
-
-PARAM_TEST_CASE(CvtColor_Gray2RGB, cv::Size, MatDepth, int)
+OCL_TEST_P(CvtColor, RGB2GRAY)
 {
-    cv::Size size;
-    int code;
-    int depth;
-    cv::Mat img;
-
-    virtual void SetUp()
-    {
-        size  = GET_PARAM(0);
-        depth = GET_PARAM(1);
-        code  = GET_PARAM(2);
-        img   = randomMat(size, CV_MAKETYPE(depth, 1), 0.0, depth == CV_32F ? 1.0 : 255.0);
-    }
+    doTest(3, 1, CVTCODE(RGB2GRAY));
+}
+OCL_TEST_P(CvtColor, GRAY2RGB)
+{
+    doTest(1, 3, CVTCODE(GRAY2RGB));
 };
-OCL_TEST_P(CvtColor_Gray2RGB, Accuracy)
+
+OCL_TEST_P(CvtColor, BGR2GRAY)
 {
-    cv::Mat src = img;
-    cv::ocl::oclMat ocl_img, dst;
-    ocl_img.upload(src);
-    cv::ocl::cvtColor(ocl_img, dst, code);
-    cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, code);
-    cv::Mat dst_mat;
-    dst.download(dst_mat);
-    EXPECT_MAT_NEAR(dst_gold, dst_mat, 1e-5);
+    doTest(3, 1, CVTCODE(BGR2GRAY));
+}
+OCL_TEST_P(CvtColor, GRAY2BGR)
+{
+    doTest(1, 3, CVTCODE(GRAY2BGR));
+};
+
+OCL_TEST_P(CvtColor, RGBA2GRAY)
+{
+    doTest(3, 1, CVTCODE(RGBA2GRAY));
+}
+OCL_TEST_P(CvtColor, GRAY2RGBA)
+{
+    doTest(1, 3, CVTCODE(GRAY2RGBA));
+};
+
+OCL_TEST_P(CvtColor, BGRA2GRAY)
+{
+    doTest(3, 1, CVTCODE(BGRA2GRAY));
+}
+OCL_TEST_P(CvtColor, GRAY2BGRA)
+{
+    doTest(1, 3, CVTCODE(GRAY2BGRA));
+};
+
+OCL_TEST_P(CvtColor, RGB2YUV)
+{
+    doTest(3, 3, CVTCODE(RGB2YUV));
+}
+OCL_TEST_P(CvtColor, BGR2YUV)
+{
+    doTest(3, 3, CVTCODE(BGR2YUV));
+}
+OCL_TEST_P(CvtColor, YUV2RGB)
+{
+    doTest(3, 3, CVTCODE(YUV2RGB));
+}
+OCL_TEST_P(CvtColor, YUV2BGR)
+{
+    doTest(3, 3, CVTCODE(YUV2BGR));
+}
+OCL_TEST_P(CvtColor, RGB2YCrCb)
+{
+    doTest(3, 3, CVTCODE(RGB2YCrCb));
+}
+OCL_TEST_P(CvtColor, BGR2YCrCb)
+{
+    doTest(3, 3, CVTCODE(BGR2YCrCb));
 }
 
-
-PARAM_TEST_CASE(CvtColor_YUV420, cv::Size, int)
+struct CvtColor_YUV420 : CvtColor
 {
-    cv::Size size;
-    int code;
-
-    cv::Mat img;
-
-    virtual void SetUp()
+    void random_roi(int channelsIn, int channelsOut)
     {
-        size = GET_PARAM(0);
-        code = GET_PARAM(1);
-        img  = randomMat(size, CV_8UC1, 0.0, 255.0);
+        const int srcType = CV_MAKE_TYPE(depth, channelsIn);
+        const int dstType = CV_MAKE_TYPE(depth, channelsOut);
+
+        Size roiSize = randomSize(1, MAX_VALUE);
+        roiSize.width *= 2;
+        roiSize.height *= 3;
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, roiSize, srcBorder, srcType, 2, 100);
+
+        Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst1, dst1_roi, roiSize, dst1Border, dstType, 5, 16);
+
+        generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder);
+        generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border);
     }
 };
 
-OCL_TEST_P(CvtColor_YUV420, Accuracy)
+OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12)
 {
-    cv::Mat src = img;
-    cv::ocl::oclMat ocl_img, dst;
-    ocl_img.upload(src);
-    cv::ocl::cvtColor(ocl_img, dst, code);
-    cv::Mat dst_gold;
-    cv::cvtColor(src, dst_gold, code);
-    cv::Mat dst_mat;
-    dst.download(dst_mat);
-    MAT_DIFF(dst_mat, dst_gold);
-    EXPECT_MAT_NEAR(dst_gold, dst_mat, 1e-5);
-}
+    doTest(1, 4, COLOR_YUV2RGBA_NV12);
+};
 
-INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine(
-                            DIFFERENT_SIZES,
-                            testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F))
-                        ));
+OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12)
+{
+    doTest(1, 4, COLOR_YUV2BGRA_NV12);
+};
 
-INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420, testing::Combine(
-                            testing::Values(cv::Size(128, 45), cv::Size(46, 132), cv::Size(1024, 1023)),
-                            testing::Values((int)COLOR_YUV2RGBA_NV12, (int)COLOR_YUV2BGRA_NV12, (int)COLOR_YUV2RGB_NV12, (int)COLOR_YUV2BGR_NV12)
-                        ));
+OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12)
+{
+    doTest(1, 3, COLOR_YUV2RGB_NV12);
+};
+
+OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12)
+{
+    doTest(1, 3, COLOR_YUV2BGR_NV12);
+};
+
+
+INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor,
+                            testing::Combine(
+                                testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),
+                                Bool()
+                            )
+                        );
+
+INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420,
+                            testing::Combine(
+                                testing::Values(MatDepth(CV_8U)),
+                                Bool()
+                            )
+                        );
 
-INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_Gray2RGB, testing::Combine(
-                            DIFFERENT_SIZES,
-                            testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),
-                            testing::Values((int)COLOR_GRAY2BGR, (int)COLOR_GRAY2BGRA, (int)COLOR_GRAY2RGB, (int)COLOR_GRAY2RGBA)
-                        ));
 }
 #endif

modules/ocl/test/test_filters.cpp

@@ -52,424 +52,401 @@
 
 #ifdef HAVE_OPENCL
 
-using namespace cvtest;
 using namespace testing;
 using namespace std;
+using namespace cv;
 
-
-PARAM_TEST_CASE(FilterTestBase,
-                MatType,
-                cv::Size, // kernel size
-                cv::Size, // dx,dy
-                int       // border type, or iteration
-                )
+PARAM_TEST_CASE(FilterTestBase, MatType,
+                int, // kernel size
+                Size, // dx, dy
+                int, // border type, or iteration
+                bool) // roi or not
 {
-    //src mat
-    cv::Mat mat1;
-    cv::Mat dst;
+    int type, borderType, ksize;
+    bool useRoi;
 
-    // set up roi
-    int roicols;
-    int roirows;
-    int src1x;
-    int src1y;
-    int dstx;
-    int dsty;
+    Mat src, dst_whole, src_roi, dst_roi;
+    ocl::oclMat gsrc_whole, gsrc_roi, gdst_whole, gdst_roi;
 
-    //src mat with roi
-    cv::Mat mat1_roi;
-    cv::Mat dst_roi;
-
-    //ocl dst mat for testing
-    cv::ocl::oclMat gdst_whole;
-
-    //ocl mat with roi
-    cv::ocl::oclMat gmat1;
-    cv::ocl::oclMat gdst;
+    virtual void SetUp()
+    {
+        type = GET_PARAM(0);
+        ksize = GET_PARAM(1);
+        borderType = GET_PARAM(3);
+        useRoi = GET_PARAM(4);
+    }
 
     void random_roi()
     {
-#ifdef RANDOMROI
-        //randomize ROI
-        roicols = rng.uniform(2, mat1.cols);
-        roirows = rng.uniform(2, mat1.rows);
-        src1x   = rng.uniform(0, mat1.cols - roicols);
-        src1y   = rng.uniform(0, mat1.rows - roirows);
-        dstx    = rng.uniform(0, dst.cols  - roicols);
-        dsty    = rng.uniform(0, dst.rows  - roirows);
-#else
-        roicols = mat1.cols;
-        roirows = mat1.rows;
-        src1x = 0;
-        src1y = 0;
-        dstx = 0;
-        dsty = 0;
-#endif
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256);
 
-        mat1_roi = mat1(Rect(src1x, src1y, roicols, roirows));
-        dst_roi  = dst(Rect(dstx, dsty, roicols, roirows));
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, 5, 16);
 
-        gdst_whole = dst;
-        gdst = gdst_whole(Rect(dstx, dsty, roicols, roirows));
-
-        gmat1 = mat1_roi;
+        generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder);
+        generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
     }
 
-    void Init(int mat_type)
+    void Near(double threshold = 0.0)
     {
-        cv::Size size(MWIDTH, MHEIGHT);
-        mat1 = randomMat(size, mat_type, 5, 16);
-        dst  = randomMat(size, mat_type, 5, 16);
-    }
+        Mat roi, whole;
+        gdst_whole.download(whole);
+        gdst_roi.download(roi);
 
-    void Near(double threshold)
-    {
-        EXPECT_MAT_NEAR(dst, Mat(gdst_whole), threshold);
+        EXPECT_MAT_NEAR(dst_whole, whole, threshold);
+        EXPECT_MAT_NEAR(dst_roi, roi, threshold);
     }
 };
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
 // blur
-struct Blur : FilterTestBase
-{
-    int type;
-    cv::Size ksize;
-    int bordertype;
 
-    virtual void SetUp()
-    {
-        type = GET_PARAM(0);
-        ksize = GET_PARAM(1);
-        bordertype = GET_PARAM(3);
-        Init(type);
-    }
-};
+typedef FilterTestBase Blur;
 
 OCL_TEST_P(Blur, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    Size kernelSize(ksize, ksize);
+
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
-        cv::blur(mat1_roi, dst_roi, ksize, Point(-1, -1), bordertype);
-        cv::ocl::blur(gmat1, gdst, ksize, Point(-1, -1), bordertype);
+
+        blur(src_roi, dst_roi, kernelSize, Point(-1, -1), borderType);
+        ocl::blur(gsrc_roi, gdst_roi, kernelSize, Point(-1, -1), borderType); // TODO anchor
+
         Near(1.0);
     }
 }
 
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// Laplacian
+
+typedef FilterTestBase LaplacianTest;
+
+OCL_TEST_P(LaplacianTest, Accuracy)
+{
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        // border type is used as a scale factor for the Laplacian kernel
+        double scale = static_cast<double>(borderType);
+
+        Laplacian(src_roi, dst_roi, -1, ksize, scale);
+        ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale);
+
+        Near(1e-5);
+    }
+}
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
-//Laplacian
-struct Laplacian : FilterTestBase
+// erode & dilate
+
+struct ErodeDilate :
+        public FilterTestBase
 {
-    int type;
-    cv::Size ksize;
+    int iterations;
 
     virtual void SetUp()
     {
         type = GET_PARAM(0);
         ksize = GET_PARAM(1);
-        Init(type);
-    }
-};
-
-OCL_TEST_P(Laplacian, Accuracy)
-{
-    for(int j = 0; j < LOOP_TIMES; j++)
-    {
-        random_roi();
-        cv::Laplacian(mat1_roi, dst_roi, -1, ksize.width, 1);
-        cv::ocl::Laplacian(gmat1, gdst, -1, ksize.width, 1);
-        Near(1e-5);
-    }
-}
-
-
-
-/////////////////////////////////////////////////////////////////////////////////////////////////
-// erode & dilate
-struct ErodeDilate : FilterTestBase
-{
-    int type;
-    int iterations;
-
-    //erode or dilate kernel
-    cv::Mat kernel;
-
-    virtual void SetUp()
-    {
-        type = GET_PARAM(0);
         iterations = GET_PARAM(3);
-        Init(type);
-        kernel = randomMat(Size(3, 3), CV_8UC1, 0, 3);
+        useRoi = GET_PARAM(4);
     }
-
 };
 
-OCL_TEST_P(ErodeDilate, Mat)
+typedef ErodeDilate Erode;
+
+OCL_TEST_P(Erode, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    // erode or dilate kernel
+    Size kernelSize(ksize, ksize);
+    Mat kernel;
+
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
+        kernel = randomMat(kernelSize, CV_8UC1, 0, 3);
+
         random_roi();
-        cv::erode(mat1_roi, dst_roi, kernel, Point(-1, -1), iterations);
-        cv::ocl::erode(gmat1, gdst, kernel, Point(-1, -1), iterations);
-        Near(1e-5);
-    }
-    for(int j = 0; j < LOOP_TIMES; j++)
-    {
-        random_roi();
-        cv::dilate(mat1_roi, dst_roi, kernel, Point(-1, -1), iterations);
-        cv::ocl::dilate(gmat1, gdst, kernel, Point(-1, -1), iterations);
+
+        cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations);
+        ocl::erode(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations); // TODO iterations, borderType
+
         Near(1e-5);
     }
 }
 
+typedef ErodeDilate Dilate;
+
+OCL_TEST_P(Dilate, Mat)
+{
+    // erode or dilate kernel
+    Mat kernel;
+
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        kernel = randomMat(Size(3, 3), CV_8UC1, 0, 3);
+
+        random_roi();
+
+        cv::dilate(src_roi, dst_roi, kernel, Point(-1, -1), iterations);
+        ocl::dilate(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations); // TODO iterations, borderType
+
+        Near(1e-5);
+    }
+}
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
 // Sobel
-struct Sobel : FilterTestBase
+
+struct SobelTest :
+        public FilterTestBase
 {
-    int type;
-    int dx, dy, ksize, bordertype;
+    int dx, dy;
 
     virtual void SetUp()
     {
         type = GET_PARAM(0);
-        Size s = GET_PARAM(1);
-        ksize = s.width;
-        s = GET_PARAM(2);
-        dx = s.width;
-        dy = s.height;
-        bordertype = GET_PARAM(3);
-        Init(type);
+        ksize = GET_PARAM(1);
+        borderType = GET_PARAM(3);
+        useRoi = GET_PARAM(4);
+
+        Size d = GET_PARAM(2);
+        dx = d.width, dy = d.height;
     }
 };
 
-OCL_TEST_P(Sobel, Mat)
+OCL_TEST_P(SobelTest, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
-        cv::Sobel(mat1_roi, dst_roi, -1, dx, dy, ksize, /*scale*/0.00001,/*delta*/0, bordertype);
-        cv::ocl::Sobel(gmat1, gdst, -1, dx, dy, ksize,/*scale*/0.00001,/*delta*/0, bordertype);
+
+        Sobel(src_roi, dst_roi, -1, dx, dy, ksize, /* scale */ 0.00001, /* delta */0, borderType);
+        ocl::Sobel(gsrc_roi, gdst_roi, -1, dx, dy, ksize, /* scale */ 0.00001, /* delta */ 0, borderType);
+
         Near(1);
     }
 }
 
-
 /////////////////////////////////////////////////////////////////////////////////////////////////
 // Scharr
-struct Scharr : FilterTestBase
-{
-    int type;
-    int dx, dy, bordertype;
 
-    virtual void SetUp()
-    {
-        type = GET_PARAM(0);
-        Size s = GET_PARAM(2);
-        dx = s.width;
-        dy = s.height;
-        bordertype = GET_PARAM(3);
-        Init(type);
-    }
-};
+typedef SobelTest ScharrTest;
 
-OCL_TEST_P(Scharr, Mat)
+OCL_TEST_P(ScharrTest, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
-        cv::Scharr(mat1_roi, dst_roi, -1, dx, dy, /*scale*/1,/*delta*/0, bordertype);
-        cv::ocl::Scharr(gmat1, gdst, -1, dx, dy,/*scale*/1,/*delta*/0, bordertype);
+
+        Scharr(src_roi, dst_roi, -1, dx, dy, /* scale */ 1, /* delta */ 0, borderType);
+        ocl::Scharr(gsrc_roi, gdst_roi, -1, dx, dy, /* scale */ 1, /* delta */ 0, borderType);
+
         Near(1);
     }
-
 }
 
-
 /////////////////////////////////////////////////////////////////////////////////////////////////
 // GaussianBlur
-struct GaussianBlur : FilterTestBase
+
+struct GaussianBlurTest :
+        public FilterTestBase
 {
-    int type;
-    cv::Size ksize;
-    int bordertype;
     double sigma1, sigma2;
 
     virtual void SetUp()
     {
         type = GET_PARAM(0);
         ksize = GET_PARAM(1);
-        bordertype = GET_PARAM(3);
-        Init(type);
+        borderType = GET_PARAM(3);
+
         sigma1 = rng.uniform(0.1, 1.0);
         sigma2 = rng.uniform(0.1, 1.0);
     }
 };
 
-OCL_TEST_P(GaussianBlur, Mat)
+OCL_TEST_P(GaussianBlurTest, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
-        cv::GaussianBlur(mat1_roi, dst_roi, ksize, sigma1, sigma2, bordertype);
-        cv::ocl::GaussianBlur(gmat1, gdst, ksize, sigma1, sigma2, bordertype);
+
+        GaussianBlur(src_roi, dst_roi, Size(ksize, ksize), sigma1, sigma2, borderType);
+        ocl::GaussianBlur(gsrc_roi, gdst_roi, Size(ksize, ksize), sigma1, sigma2, borderType);
+
         Near(1);
     }
-
 }
 
-
-
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // Filter2D
-struct Filter2D : FilterTestBase
-{
-    int type;
-    cv::Size ksize;
-    int bordertype;
-    Point anchor;
-    virtual void SetUp()
-    {
-        type = GET_PARAM(0);
-        ksize = GET_PARAM(1);
-        bordertype = GET_PARAM(3);
-        Init(type);
-        anchor = Point(-1,-1);
-    }
-};
+
+typedef FilterTestBase Filter2D;
 
 OCL_TEST_P(Filter2D, Mat)
 {
-    cv::Mat kernel = randomMat(cv::Size(ksize.width, ksize.height), CV_32FC1, 0.0, 1.0);
-    for(int j = 0; j < LOOP_TIMES; j++)
+    const Size kernelSize(ksize, ksize);
+    Mat kernel;
+
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
+        kernel = randomMat(kernelSize, CV_32FC1, 0.0, 1.0);
+
         random_roi();
-        cv::filter2D(mat1_roi, dst_roi, -1, kernel, anchor, 0.0, bordertype);
-        cv::ocl::filter2D(gmat1, gdst, -1, kernel, anchor, bordertype);
+
+        cv::filter2D(src_roi, dst_roi, -1, kernel, Point(-1, -1), 0.0, borderType); // TODO anchor
+        ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, Point(-1, -1), borderType);
+
         Near(1);
     }
 }
+
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // Bilateral
-struct Bilateral : FilterTestBase
-{
-    int type;
-    cv::Size ksize;
-    int bordertype;
-    double sigmacolor, sigmaspace;
 
-    virtual void SetUp()
-    {
-        type = GET_PARAM(0);
-        ksize = GET_PARAM(1);
-        bordertype = GET_PARAM(3);
-        Init(type);
-        sigmacolor = rng.uniform(20, 100);
-        sigmaspace = rng.uniform(10, 40);
-    }
-};
+typedef FilterTestBase Bilateral;
 
 OCL_TEST_P(Bilateral, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
-        cv::bilateralFilter(mat1_roi, dst_roi, ksize.width, sigmacolor, sigmaspace, bordertype);
-        cv::ocl::bilateralFilter(gmat1, gdst, ksize.width, sigmacolor, sigmaspace, bordertype);
+
+        double sigmacolor = rng.uniform(20, 100);
+        double sigmaspace = rng.uniform(10, 40);
+
+        cv::bilateralFilter(src_roi, dst_roi, ksize, sigmacolor, sigmaspace, borderType);
+        ocl::bilateralFilter(gsrc_roi, gdst_roi, ksize, sigmacolor, sigmaspace, borderType);
+
         Near(1);
     }
-
 }
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // AdaptiveBilateral
-struct AdaptiveBilateral : FilterTestBase
-{
-    int type;
-    cv::Size ksize;
-    int bordertype;
-    Point anchor;
-    virtual void SetUp()
-    {
-        type = GET_PARAM(0);
-        ksize = GET_PARAM(1);
-        bordertype = GET_PARAM(3);
-        Init(type);
-        anchor = Point(-1,-1);
-    }
-};
+
+typedef FilterTestBase AdaptiveBilateral;
 
 OCL_TEST_P(AdaptiveBilateral, Mat)
 {
-    for(int j = 0; j < LOOP_TIMES; j++)
+    const Size kernelSize(ksize, ksize);
+
+    for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
-        cv::adaptiveBilateralFilter(mat1_roi, dst_roi, ksize, 5, anchor, bordertype);
-        cv::ocl::adaptiveBilateralFilter(gmat1, gdst, ksize, 5, anchor, bordertype);
+
+        adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, Point(-1, -1), borderType); // TODO anchor
+        ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, Point(-1, -1), borderType);
+
         Near(1);
     }
-
 }
 
+/////////////////////////////////////////////////////////////////////////////////////////////////////
+// MedianFilter
+
+typedef FilterTestBase MedianFilter;
+
+OCL_TEST_P(MedianFilter, Mat)
+{
+    for (int i = 0; i < LOOP_TIMES; ++i)
+    {
+        random_roi();
+
+        medianBlur(src_roi, dst_roi, ksize);
+        ocl::medianFilter(gsrc_roi, gdst_roi, ksize);
+
+        Near();
+    }
+}
+
+//////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
 INSTANTIATE_TEST_CASE_P(Filter, Blur, Combine(
-                        Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
-                        Values(cv::Size(3, 3), cv::Size(5, 5), cv::Size(7, 7)),
-                        Values(Size(0, 0)), //not use
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
+                            Values(3, 5, 7),
+                            Values(Size(0, 0)), // not used
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT_101),
+                            Bool()));
 
+INSTANTIATE_TEST_CASE_P(Filter, LaplacianTest, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values(1, 3),
+                            Values(Size(0, 0)), // not used
+                            Values(1, 2), // value is used as scale factor for kernel
+                            Bool()));
 
-INSTANTIATE_TEST_CASE_P(Filter, Laplacian, Combine(
-                        Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
-                        Values(Size(3, 3)),
-                        Values(Size(0, 0)), //not use
-                        Values(0)));        //not use
+INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values(3, 5, 7),
+                            Values(Size(0, 0)), // not used
+                            testing::Range(1, 2),
+                            Bool()));
 
-INSTANTIATE_TEST_CASE_P(Filter, ErodeDilate, Combine(
-                        Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4),
-                        Values(Size(0, 0)), //not use
-                        Values(Size(0, 0)), //not use
-                        Values(1)));
+INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values(3, 5, 7),
+                            Values(Size(0, 0)), // not used
+                            testing::Range(1, 2),
+                            Bool()));
 
+INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values(3, 5),
+                            Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)),
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
+                                   (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+                            Bool()));
 
-INSTANTIATE_TEST_CASE_P(Filter, Sobel, Combine(
-                        Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
-                        Values(Size(3, 3), Size(5, 5)),
-                        Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)),
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));
-
-
-INSTANTIATE_TEST_CASE_P(Filter, Scharr, Combine(
-                        Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
-                        Values(Size(0, 0)), //not use
-                        Values(Size(0, 1), Size(1, 0)),
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));
-
-INSTANTIATE_TEST_CASE_P(Filter, GaussianBlur, Combine(
-                        Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
-                        Values(Size(3, 3), Size(5, 5)),
-                        Values(Size(0, 0)), //not use
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));
-
+INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
+                            Values(0), // not used
+                            Values(Size(0, 1), Size(1, 0)),
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
+                                   (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+                            Bool()));
 
+INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4),
+                            Values(3, 5),
+                            Values(Size(0, 0)), // not used
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
+                                   (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+                            Bool()));
 
 INSTANTIATE_TEST_CASE_P(Filter, Filter2D, testing::Combine(
-                        Values(CV_8UC1, CV_32FC1, CV_32FC4),
-                        Values(Size(3, 3), Size(15, 15), Size(25, 25)),
-                        Values(Size(0, 0)), //not use
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REFLECT101, (MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT)));
+                            Values(CV_8UC1, CV_32FC1, CV_32FC4),
+                            Values(3, 15, 25),
+                            Values(Size(0, 0)), // not used
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
+                                   (int)BORDER_REPLICATE, (int)BORDER_REFLECT),
+                            Bool()));
 
 INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine(
-                        Values(CV_8UC1, CV_8UC3),
-                        Values(Size(5, 5), Size(9, 9)),
-                        Values(Size(0, 0)), //not use
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE,
-                               (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_WRAP, (MatType)cv::BORDER_REFLECT_101)));
+                            Values(CV_8UC1, CV_8UC3),
+                            Values(5, 9),
+                            Values(Size(0, 0)), // not used
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE,
+                                   (int)BORDER_REFLECT, (int)BORDER_WRAP, (int)BORDER_REFLECT_101),
+                            Values(false))); // TODO does not work with ROI
 
 INSTANTIATE_TEST_CASE_P(Filter, AdaptiveBilateral, Combine(
-                        Values(CV_8UC1, CV_8UC3),
-                        Values(Size(5, 5), Size(9, 9)),
-                        Values(Size(0, 0)), //not use
-                        Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE,
-                               (MatType)cv::BORDER_REFLECT,  (MatType)cv::BORDER_REFLECT_101)));
+                            Values(CV_8UC1, CV_8UC3),
+                            Values(5, 9),
+                            Values(Size(0, 0)), // not used
+                            Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE,
+                                   (int)BORDER_REFLECT, (int)BORDER_REFLECT_101),
+                            Bool()));
+
+INSTANTIATE_TEST_CASE_P(Filter, MedianFilter, Combine(
+                            Values((MatType)CV_8UC1, (MatType)CV_8UC4, (MatType)CV_32FC1, (MatType)CV_32FC4),
+                            Values(3, 5),
+                            Values(Size(0, 0)), // not used
+                            Values(0), // not used
+                            Bool()));
+
 #endif // HAVE_OPENCL

modules/ocl/test/test_matrix_operation.cpp

@@ -48,382 +48,176 @@
 
 #ifdef HAVE_OPENCL
 
-using namespace cvtest;
+using namespace cv;
 using namespace testing;
 using namespace std;
 
 ////////////////////////////////converto/////////////////////////////////////////////////
 
-PARAM_TEST_CASE(ConvertToTestBase, MatType, MatType, int, bool)
+PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, int, bool)
 {
-    int src_depth, dst_depth;
-    int cn, dst_type;
+    int src_depth, cn, dstType;
     bool use_roi;
 
-    // src mat
-    cv::Mat mat;
-    cv::Mat dst;
-
-    // set up roi
-    int roicols, roirows;
-    int srcx, srcy;
-    int dstx, dsty;
-
-    // src mat with roi
-    cv::Mat mat_roi;
-    cv::Mat dst_roi;
-
-    // ocl dst mat for testing
-    cv::ocl::oclMat gdst_whole;
-
-    // ocl mat with roi
-    cv::ocl::oclMat gsrc;
-    cv::ocl::oclMat gdst;
+    Mat src, dst, src_roi, dst_roi;
+    ocl::oclMat gdst, gsrc, gdst_roi, gsrc_roi;
 
     virtual void SetUp()
     {
         src_depth = GET_PARAM(0);
-        dst_depth = GET_PARAM(1);
         cn = GET_PARAM(2);
-        int src_type = CV_MAKE_TYPE(src_depth, cn);
-        dst_type = CV_MAKE_TYPE(dst_depth, cn);
+        dstType = CV_MAKE_TYPE(GET_PARAM(1), cn);
 
         use_roi = GET_PARAM(3);
-
-        mat = randomMat(randomSize(MIN_VALUE, MAX_VALUE), src_type, 5, 136, false);
-        dst = randomMat(use_roi ? randomSize(MIN_VALUE, MAX_VALUE) : mat.size(), dst_type, 5, 136, false);
     }
 
-    void random_roi()
+    virtual void random_roi()
     {
-        if (use_roi)
-        {
-            // randomize ROI
-            roicols = rng.uniform(1, MIN_VALUE);
-            roirows = rng.uniform(1, MIN_VALUE);
-            srcx = rng.uniform(0, mat.cols - roicols);
-            srcy = rng.uniform(0, mat.rows - roirows);
-            dstx = rng.uniform(0, dst.cols - roicols);
-            dsty = rng.uniform(0, dst.rows - roirows);
-        }
-        else
-        {
-            roicols = mat.cols;
-            roirows = mat.rows;
-            srcx = srcy = 0;
-            dstx = dsty = 0;
-        }
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, CV_MAKE_TYPE(src_depth, cn), -MAX_VALUE, MAX_VALUE);
 
-        mat_roi = mat(Rect(srcx, srcy, roicols, roirows));
-        dst_roi = dst(Rect(dstx, dsty, roicols, roirows));
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16);
 
-        gdst_whole = dst;
-        gdst = gdst_whole(Rect(dstx, dsty, roicols, roirows));
-
-        gsrc = mat_roi;
+        generateOclMat(gsrc, gsrc_roi, src, roiSize, srcBorder);
+        generateOclMat(gdst, gdst_roi, dst, roiSize, dstBorder);
     }
 };
 
-typedef ConvertToTestBase ConvertTo;
+typedef MatrixTestBase ConvertTo;
 
 OCL_TEST_P(ConvertTo, Accuracy)
 {
-    if((src_depth == CV_64F || dst_depth == CV_64F) &&
-            !cv::ocl::Context::getContext()->supportsFeature(cv::ocl::FEATURE_CL_DOUBLE))
-    {
-        return; // returns silently
-    }
     for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
 
-        mat_roi.convertTo(dst_roi, dst_type);
-        gsrc.convertTo(gdst, dst_type);
+        src_roi.convertTo(dst_roi, dstType);
+        gsrc_roi.convertTo(gdst_roi, dstType);
 
-        EXPECT_MAT_NEAR(dst, Mat(gdst_whole), src_depth == CV_64F ? 1.0 : 0.0);
-        EXPECT_MAT_NEAR(dst_roi, Mat(gdst), src_depth == CV_64F ? 1.0 : 0.0);
+        EXPECT_MAT_NEAR(dst, Mat(gdst), src_depth == CV_64F ? 1.0 : 0.0);
+        EXPECT_MAT_NEAR(dst_roi, Mat(gdst_roi), src_depth == CV_64F ? 1.0 : 0.0);
     }
 }
 
 ///////////////////////////////////////////copyto/////////////////////////////////////////////////////////////
 
-PARAM_TEST_CASE(CopyToTestBase, MatType, int, bool)
+struct CopyTo :
+        public MatrixTestBase
 {
-    bool use_roi;
+    Mat mask, mask_roi;
+    ocl::oclMat gmask, gmask_roi;
 
-    cv::Mat src, mask, dst;
-
-    // set up roi
-    int roicols,roirows;
-    int srcx, srcy;
-    int dstx, dsty;
-    int maskx,masky;
-
-    // src mat with roi
-    cv::Mat src_roi;
-    cv::Mat mask_roi;
-    cv::Mat dst_roi;
-
-    // ocl dst mat for testing
-    cv::ocl::oclMat gdst_whole;
-
-    // ocl mat with roi
-    cv::ocl::oclMat gsrc, gdst, gmask;
-
-    virtual void SetUp()
+    virtual void random_roi()
     {
-        int type = CV_MAKETYPE(GET_PARAM(0), GET_PARAM(1));
-        use_roi = GET_PARAM(2);
+        int type = CV_MAKE_TYPE(src_depth, cn);
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
 
-        src = randomMat(randomSize(MIN_VALUE, MAX_VALUE), type, 5, 16, false);
-        dst = randomMat(use_roi ? randomSize(MIN_VALUE, MAX_VALUE) : src.size(), type, 5, 16, false);
-        mask = randomMat(use_roi ? randomSize(MIN_VALUE, MAX_VALUE) : src.size(), CV_8UC1, 0, 2,  false);
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16);
 
-        cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
-    }
+        Border maskBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(mask, mask_roi, roiSize, maskBorder, CV_8UC1, 5, 16);
 
-    void random_roi()
-    {
-        if (use_roi)
-        {
-            // randomize ROI
-            roicols = rng.uniform(1, MIN_VALUE);
-            roirows = rng.uniform(1, MIN_VALUE);
-            srcx = rng.uniform(0, src.cols - roicols);
-            srcy = rng.uniform(0, src.rows - roirows);
-            dstx = rng.uniform(0, dst.cols - roicols);
-            dsty = rng.uniform(0, dst.rows - roirows);
-            maskx = rng.uniform(0, mask.cols - roicols);
-            masky = rng.uniform(0, mask.rows - roirows);
-        }
-        else
-        {
-            roicols = src.cols;
-            roirows = src.rows;
-            srcx = srcy = 0;
-            dstx = dsty = 0;
-            maskx = masky = 0;
-        }
-
-        src_roi = src(Rect(srcx, srcy, roicols, roirows));
-        mask_roi = mask(Rect(maskx, masky, roicols, roirows));
-        dst_roi = dst(Rect(dstx, dsty, roicols, roirows));
-
-        gdst_whole = dst;
-        gdst = gdst_whole(Rect(dstx, dsty, roicols, roirows));
-
-        gsrc = src_roi;
-        gmask = mask_roi;
+        generateOclMat(gsrc, gsrc_roi, src, roiSize, srcBorder);
+        generateOclMat(gdst, gdst_roi, dst, roiSize, dstBorder);
+        generateOclMat(gmask, gmask_roi, mask, roiSize, maskBorder);
     }
 };
 
-typedef CopyToTestBase CopyTo;
-
 OCL_TEST_P(CopyTo, Without_mask)
 {
-    if((src.depth() == CV_64F) &&
-            !cv::ocl::Context::getContext()->supportsFeature(cv::ocl::FEATURE_CL_DOUBLE))
-    {
-        return; // returns silently
-    }
     for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
 
         src_roi.copyTo(dst_roi);
-        gsrc.copyTo(gdst);
+        gsrc_roi.copyTo(gdst_roi);
 
-        EXPECT_MAT_NEAR(dst, Mat(gdst_whole), 0.0);
+        EXPECT_MAT_NEAR(dst, Mat(gdst), 0.0);
+        EXPECT_MAT_NEAR(dst_roi, Mat(gdst_roi), 0.0);
     }
 }
 
 OCL_TEST_P(CopyTo, With_mask)
 {
-    if(src.depth() == CV_64F &&
-        !cv::ocl::Context::getContext()->supportsFeature(cv::ocl::FEATURE_CL_DOUBLE))
-    {
-        return; // returns silently
-    }
     for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
 
         src_roi.copyTo(dst_roi, mask_roi);
-        gsrc.copyTo(gdst, gmask);
+        gsrc_roi.copyTo(gdst_roi, gmask_roi);
 
-        EXPECT_MAT_NEAR(dst, Mat(gdst_whole), 0.0);
+        EXPECT_MAT_NEAR(dst, Mat(gdst), 0.0);
+        EXPECT_MAT_NEAR(dst_roi, Mat(gdst_roi), 0.0);
     }
 }
 
 /////////////////////////////////////////// setTo /////////////////////////////////////////////////////////////
 
-PARAM_TEST_CASE(SetToTestBase, MatType, int, bool)
-{
-    int depth, channels;
-    bool use_roi;
-
-    cv::Scalar val;
-
-    cv::Mat src;
-    cv::Mat mask;
-
-    // set up roi
-    int roicols, roirows;
-    int srcx, srcy;
-    int maskx, masky;
-
-    // src mat with roi
-    cv::Mat src_roi;
-    cv::Mat mask_roi;
-
-    // ocl dst mat for testing
-    cv::ocl::oclMat gsrc_whole;
-
-    // ocl mat with roi
-    cv::ocl::oclMat gsrc;
-    cv::ocl::oclMat gmask;
-
-    virtual void SetUp()
-    {
-        depth = GET_PARAM(0);
-        channels = GET_PARAM(1);
-        use_roi = GET_PARAM(2);
-
-        int type = CV_MAKE_TYPE(depth, channels);
-
-        src = randomMat(randomSize(MIN_VALUE, MAX_VALUE), type, 5, 16, false);
-        mask = randomMat(use_roi ? randomSize(MIN_VALUE, MAX_VALUE) : src.size(), CV_8UC1, 0, 2,  false);
-
-        cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
-        val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0),
-                         rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
-    }
-
-    void random_roi()
-    {
-        if (use_roi)
-        {
-            // randomize ROI
-            roicols = rng.uniform(1, MIN_VALUE);
-            roirows = rng.uniform(1, MIN_VALUE);
-            srcx = rng.uniform(0, src.cols - roicols);
-            srcy = rng.uniform(0, src.rows - roirows);
-            maskx = rng.uniform(0, mask.cols - roicols);
-            masky = rng.uniform(0, mask.rows - roirows);
-        }
-        else
-        {
-            roicols = src.cols;
-            roirows = src.rows;
-            srcx = srcy = 0;
-            maskx = masky = 0;
-        }
-
-        src_roi = src(Rect(srcx, srcy, roicols, roirows));
-        mask_roi = mask(Rect(maskx, masky, roicols, roirows));
-
-        gsrc_whole = src;
-        gsrc = gsrc_whole(Rect(srcx, srcy, roicols, roirows));
-
-        gmask = mask_roi;
-    }
-};
-
-typedef SetToTestBase SetTo;
+typedef CopyTo SetTo;
 
 OCL_TEST_P(SetTo, Without_mask)
 {
-    if(depth == CV_64F &&
-            !cv::ocl::Context::getContext()->supportsFeature(cv::ocl::FEATURE_CL_DOUBLE))
-    {
-        return; // returns silently
-    }
     for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
+        Scalar scalar = randomScalar(-MAX_VALUE, MAX_VALUE);
 
-        src_roi.setTo(val);
-        gsrc.setTo(val);
+        src_roi.setTo(scalar);
+        gsrc_roi.setTo(scalar);
 
-        EXPECT_MAT_NEAR(src, Mat(gsrc_whole), 1.);
+        EXPECT_MAT_NEAR(dst, Mat(gdst), 0.0);
+        EXPECT_MAT_NEAR(dst_roi, Mat(gdst_roi), 0.0);;
     }
 }
 
 OCL_TEST_P(SetTo, With_mask)
 {
-    if(depth == CV_64F &&
-            !cv::ocl::Context::getContext()->supportsFeature(cv::ocl::FEATURE_CL_DOUBLE))
-    {
-        return; // returns silently
-    }
     for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
+        Scalar scalar = randomScalar(-MAX_VALUE, MAX_VALUE);
 
-        src_roi.setTo(val, mask_roi);
-        gsrc.setTo(val, gmask);
+        src_roi.setTo(scalar, mask_roi);
+        gsrc_roi.setTo(scalar, gmask_roi);
 
-        EXPECT_MAT_NEAR(src, Mat(gsrc_whole), 1.);
+        EXPECT_MAT_NEAR(src, Mat(gsrc), 1.);
+        EXPECT_MAT_NEAR(src_roi, Mat(gsrc_roi), 1.);
     }
 }
 
 // convertC3C4
 
-PARAM_TEST_CASE(convertC3C4, MatType, bool)
+PARAM_TEST_CASE(convertC3C4, MatDepth, bool)
 {
     int depth;
     bool use_roi;
 
-    //src mat
-    cv::Mat src;
-
-    // set up roi
-    int roicols, roirows;
-    int srcx, srcy;
-
-    //src mat with roi
-    cv::Mat src_roi;
-
-    //ocl mat with roi
-    cv::ocl::oclMat gsrc_roi;
+    Mat src, src_roi;
+    ocl::oclMat gsrc, gsrc_roi;
 
     virtual void SetUp()
     {
         depth = GET_PARAM(0);
         use_roi = GET_PARAM(1);
-        int type = CV_MAKE_TYPE(depth, 3);
-
-        src = randomMat(randomSize(1, MAX_VALUE), type, 0, 40, false);
     }
 
     void random_roi()
     {
-        if (use_roi)
-        {
-            //randomize ROI
-            roicols = rng.uniform(1, src.cols);
-            roirows = rng.uniform(1, src.rows);
-            srcx = rng.uniform(0, src.cols - roicols);
-            srcy = rng.uniform(0, src.rows - roirows);
-        }
-        else
-        {
-            roicols = src.cols;
-            roirows = src.rows;
-            srcx = srcy = 0;
-        }
-
-        src_roi = src(Rect(srcx, srcy, roicols, roirows));
+        int type = CV_MAKE_TYPE(depth, 3);
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+        generateOclMat(gsrc, gsrc_roi, src, roiSize, srcBorder);
     }
 };
 
 OCL_TEST_P(convertC3C4, Accuracy)
 {
-    if(depth == CV_64F &&
-        !cv::ocl::Context::getContext()->supportsFeature(cv::ocl::FEATURE_CL_DOUBLE))
-    {
-        return; // returns silently
-    }
     for (int j = 0; j < LOOP_TIMES; j++)
     {
         random_roi();
@@ -431,20 +225,23 @@ OCL_TEST_P(convertC3C4, Accuracy)
         gsrc_roi = src_roi;
 
         EXPECT_MAT_NEAR(src_roi, Mat(gsrc_roi), 0.0);
+        EXPECT_MAT_NEAR(src, Mat(gsrc), 0.0);
     }
 }
 
 INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine(
                             Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F),
                             Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F),
-                            Range(1, 5), Bool()));
+                            testing::Range(1, 5), Bool()));
 
 INSTANTIATE_TEST_CASE_P(MatrixOperation, CopyTo, Combine(
                             Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F),
+                            Values(MatDepth(0)), // not used
                             testing::Range(1, 5), Bool()));
 
 INSTANTIATE_TEST_CASE_P(MatrixOperation, SetTo, Combine(
                             Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F),
+                            Values((MatDepth)0), // not used
                             testing::Range(1, 5), Bool()));
 
 INSTANTIATE_TEST_CASE_P(MatrixOperation, convertC3C4, Combine(

modules/ocl/test/test_mean_shift.cpp

@@ -0,0 +1,408 @@
+/*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) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Niko Li, newlife20080214@gmail.com
+//    Jia Haipeng, jiahaipeng95@gmail.com
+//    Shengen Yan, yanshengen@gmail.com
+//    Jiang Liyuan, lyuan001.good@163.com
+//    Rock Li, Rock.Li@amd.com
+//    Wu Zailong, bullet@yeah.net
+//    Xu Pang, pangxu010@163.com
+//    Sen Liu, swjtuls1987@126.com
+//
+// 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 oclMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "test_precomp.hpp"
+
+#ifdef HAVE_OPENCL
+
+using namespace testing;
+using namespace std;
+using namespace cv;
+
+typedef struct
+{
+    short x;
+    short y;
+} COOR;
+
+COOR do_meanShift(int x0, int y0, uchar *sptr, uchar *dptr, int sstep, Size size, int sp, int sr, int maxIter, float eps, int *tab)
+{
+
+    int isr2 = sr * sr;
+    int c0, c1, c2, c3;
+    int iter;
+    uchar *ptr = NULL;
+    uchar *pstart = NULL;
+    int revx = 0, revy = 0;
+    c0 = sptr[0];
+    c1 = sptr[1];
+    c2 = sptr[2];
+    c3 = sptr[3];
+    // iterate meanshift procedure
+    for(iter = 0; iter < maxIter; iter++ )
+    {
+        int count = 0;
+        int s0 = 0, s1 = 0, s2 = 0, sx = 0, sy = 0;
+
+        //mean shift: process pixels in window (p-sigmaSp)x(p+sigmaSp)
+        int minx = x0 - sp;
+        int miny = y0 - sp;
+        int maxx = x0 + sp;
+        int maxy = y0 + sp;
+
+        //deal with the image boundary
+        if(minx < 0) minx = 0;
+        if(miny < 0) miny = 0;
+        if(maxx >= size.width) maxx = size.width - 1;
+        if(maxy >= size.height) maxy = size.height - 1;
+        if(iter == 0)
+        {
+            pstart = sptr;
+        }
+        else
+        {
+            pstart = pstart + revy * sstep + (revx << 2); //point to the new position
+        }
+        ptr = pstart;
+        ptr = ptr + (miny - y0) * sstep + ((minx - x0) << 2); //point to the start in the row
+
+        for( int y = miny; y <= maxy; y++, ptr += sstep - ((maxx - minx + 1) << 2))
+        {
+            int rowCount = 0;
+            int x = minx;
+#if CV_ENABLE_UNROLLED
+            for( ; x + 4 <= maxx; x += 4, ptr += 16)
+            {
+                int t0, t1, t2;
+                t0 = ptr[0], t1 = ptr[1], t2 = ptr[2];
+                if(tab[t0 - c0 + 255] + tab[t1 - c1 + 255] + tab[t2 - c2 + 255] <= isr2)
+                {
+                    s0 += t0;
+                    s1 += t1;
+                    s2 += t2;
+                    sx += x;
+                    rowCount++;
+                }
+                t0 = ptr[4], t1 = ptr[5], t2 = ptr[6];
+                if(tab[t0 - c0 + 255] + tab[t1 - c1 + 255] + tab[t2 - c2 + 255] <= isr2)
+                {
+                    s0 += t0;
+                    s1 += t1;
+                    s2 += t2;
+                    sx += x + 1;
+                    rowCount++;
+                }
+                t0 = ptr[8], t1 = ptr[9], t2 = ptr[10];
+                if(tab[t0 - c0 + 255] + tab[t1 - c1 + 255] + tab[t2 - c2 + 255] <= isr2)
+                {
+                    s0 += t0;
+                    s1 += t1;
+                    s2 += t2;
+                    sx += x + 2;
+                    rowCount++;
+                }
+                t0 = ptr[12], t1 = ptr[13], t2 = ptr[14];
+                if(tab[t0 - c0 + 255] + tab[t1 - c1 + 255] + tab[t2 - c2 + 255] <= isr2)
+                {
+                    s0 += t0;
+                    s1 += t1;
+                    s2 += t2;
+                    sx += x + 3;
+                    rowCount++;
+                }
+            }
+#endif
+            for(; x <= maxx; x++, ptr += 4)
+            {
+                int t0 = ptr[0], t1 = ptr[1], t2 = ptr[2];
+                if(tab[t0 - c0 + 255] + tab[t1 - c1 + 255] + tab[t2 - c2 + 255] <= isr2)
+                {
+                    s0 += t0;
+                    s1 += t1;
+                    s2 += t2;
+                    sx += x;
+                    rowCount++;
+                }
+            }
+            if(rowCount == 0)
+                continue;
+            count += rowCount;
+            sy += y * rowCount;
+        }
+
+        if( count == 0 )
+            break;
+
+        int x1 = sx / count;
+        int y1 = sy / count;
+        s0 = s0 / count;
+        s1 = s1 / count;
+        s2 = s2 / count;
+
+        bool stopFlag = (x0 == x1 && y0 == y1) || (abs(x1 - x0) + abs(y1 - y0) +
+                        tab[s0 - c0 + 255] + tab[s1 - c1 + 255] + tab[s2 - c2 + 255] <= eps);
+
+        //revise the pointer corresponding to the new (y0,x0)
+        revx = x1 - x0;
+        revy = y1 - y0;
+
+        x0 = x1;
+        y0 = y1;
+        c0 = s0;
+        c1 = s1;
+        c2 = s2;
+
+        if( stopFlag )
+            break;
+    } //for iter
+
+    dptr[0] = (uchar)c0;
+    dptr[1] = (uchar)c1;
+    dptr[2] = (uchar)c2;
+    dptr[3] = (uchar)c3;
+
+    COOR coor;
+    coor.x = (short)x0;
+    coor.y = (short)y0;
+    return coor;
+}
+
+void meanShiftFiltering_(const Mat &src_roi, Mat &dst_roi, int sp, int sr, TermCriteria crit)
+{
+    if( src_roi.empty() )
+        CV_Error( CV_StsBadArg, "The input image is empty" );
+
+    if( src_roi.depth() != CV_8U || src_roi.channels() != 4 )
+        CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );
+
+    CV_Assert( (src_roi.cols == dst_roi.cols) && (src_roi.rows == dst_roi.rows) );
+    CV_Assert( !(dst_roi.step & 0x3) );
+
+    if( !(crit.type & TermCriteria::MAX_ITER) )
+        crit.maxCount = 5;
+    int maxIter = std::min(std::max(crit.maxCount, 1), 100);
+    float eps;
+    if( !(crit.type & TermCriteria::EPS) )
+        eps = 1.f;
+    eps = (float)std::max(crit.epsilon, 0.0);
+
+    int tab[512];
+    for(int i = 0; i < 512; i++)
+        tab[i] = (i - 255) * (i - 255);
+    uchar *sptr = src_roi.data;
+    uchar *dptr = dst_roi.data;
+    int sstep = (int)src_roi.step;
+    int dstep = (int)dst_roi.step;
+    Size size = src_roi.size();
+
+    for(int i = 0; i < size.height; i++, sptr += sstep - (size.width << 2),
+            dptr += dstep - (size.width << 2))
+    {
+        for(int j = 0; j < size.width; j++, sptr += 4, dptr += 4)
+        {
+            do_meanShift(j, i, sptr, dptr, sstep, size, sp, sr, maxIter, eps, tab);
+        }
+    }
+}
+
+void meanShiftProc_(const Mat &src_roi, Mat &dst_roi, Mat &dstCoor_roi, int sp, int sr, TermCriteria crit)
+{
+    if( src_roi.empty() )
+        CV_Error( CV_StsBadArg, "The input image is empty" );
+    if( src_roi.depth() != CV_8U || src_roi.channels() != 4 )
+        CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );
+    CV_Assert( (src_roi.cols == dst_roi.cols) && (src_roi.rows == dst_roi.rows) &&
+               (src_roi.cols == dstCoor_roi.cols) && (src_roi.rows == dstCoor_roi.rows));
+    CV_Assert( !(dstCoor_roi.step & 0x3) );
+
+    if( !(crit.type & TermCriteria::MAX_ITER) )
+        crit.maxCount = 5;
+    int maxIter = std::min(std::max(crit.maxCount, 1), 100);
+    float eps;
+    if( !(crit.type & TermCriteria::EPS) )
+        eps = 1.f;
+    eps = (float)std::max(crit.epsilon, 0.0);
+
+    int tab[512];
+    for(int i = 0; i < 512; i++)
+        tab[i] = (i - 255) * (i - 255);
+    uchar *sptr = src_roi.data;
+    uchar *dptr = dst_roi.data;
+    short *dCoorptr = (short *)dstCoor_roi.data;
+    int sstep = (int)src_roi.step;
+    int dstep = (int)dst_roi.step;
+    int dCoorstep = (int)dstCoor_roi.step >> 1;
+    Size size = src_roi.size();
+
+    for(int i = 0; i < size.height; i++, sptr += sstep - (size.width << 2),
+            dptr += dstep - (size.width << 2), dCoorptr += dCoorstep - (size.width << 1))
+    {
+        for(int j = 0; j < size.width; j++, sptr += 4, dptr += 4, dCoorptr += 2)
+        {
+            *((COOR *)dCoorptr) = do_meanShift(j, i, sptr, dptr, sstep, size, sp, sr, maxIter, eps, tab);
+        }
+    }
+
+}
+
+//////////////////////////////// meanShift //////////////////////////////////////////
+
+PARAM_TEST_CASE(meanShiftTestBase, MatType, MatType, int, int, TermCriteria, bool)
+{
+    int type, typeCoor;
+    int sp, sr;
+    TermCriteria crit;
+    bool useRoi;
+
+    // src mat
+    Mat src, src_roi;
+    Mat dst, dst_roi;
+    Mat dstCoor, dstCoor_roi;
+
+    // ocl dst mat
+    ocl::oclMat gsrc, gsrc_roi;
+    ocl::oclMat gdst, gdst_roi;
+    ocl::oclMat gdstCoor, gdstCoor_roi;
+
+    virtual void SetUp()
+    {
+        type = GET_PARAM(0);
+        typeCoor = GET_PARAM(1);
+        sp = GET_PARAM(2);
+        sr = GET_PARAM(3);
+        crit = GET_PARAM(4);
+        useRoi = GET_PARAM(5);
+    }
+
+    void random_roi()
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256);
+        generateOclMat(gsrc, gsrc_roi, src, roiSize, srcBorder);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 256);
+        generateOclMat(gdst, gdst_roi, dst, roiSize, dstBorder);
+
+        randomSubMat(dstCoor, dstCoor_roi, roiSize, dstBorder, typeCoor, 5, 256);
+        generateOclMat(gdstCoor, gdstCoor_roi, dstCoor, roiSize, dstBorder);
+    }
+
+    void Near(double threshold = 0.0)
+    {
+        Mat whole, roi;
+        gdst.download(whole);
+        gdst_roi.download(roi);
+
+        EXPECT_MAT_NEAR(dst, whole, threshold);
+        EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+    }
+
+    void Near1(double threshold = 0.0)
+    {
+        Mat whole, roi;
+        gdstCoor.download(whole);
+        gdstCoor_roi.download(roi);
+
+        EXPECT_MAT_NEAR(dstCoor, whole, threshold);
+        EXPECT_MAT_NEAR(dstCoor_roi, roi, threshold);
+    }
+};
+
+/////////////////////////meanShiftFiltering/////////////////////////////
+
+typedef meanShiftTestBase meanShiftFiltering;
+
+OCL_TEST_P(meanShiftFiltering, Mat)
+{
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        meanShiftFiltering_(src_roi, dst_roi, sp, sr, crit);
+        ocl::meanShiftFiltering(gsrc_roi, gdst_roi, sp, sr, crit);
+
+        Near();
+    }
+}
+
+///////////////////////////meanShiftProc//////////////////////////////////
+
+typedef meanShiftTestBase meanShiftProc;
+
+OCL_TEST_P(meanShiftProc, Mat)
+{
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        meanShiftProc_(src_roi, dst_roi, dstCoor_roi, sp, sr, crit);
+        ocl::meanShiftProc(gsrc_roi, gdst_roi, gdstCoor_roi, sp, sr, crit);
+
+        Near();
+        Near1();
+    }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////
+
+INSTANTIATE_TEST_CASE_P(Imgproc, meanShiftFiltering, Combine(
+                            Values((MatType)CV_8UC4),
+                            Values((MatType)CV_16SC2),
+                            Values(5),
+                            Values(6),
+                            Values(TermCriteria(TermCriteria::COUNT + TermCriteria::EPS, 5, 1)),
+                            Bool()
+                        ));
+
+INSTANTIATE_TEST_CASE_P(Imgproc, meanShiftProc, Combine(
+                            Values((MatType)CV_8UC4),
+                            Values((MatType)CV_16SC2),
+                            Values(5),
+                            Values(6),
+                            Values(TermCriteria(TermCriteria::COUNT + TermCriteria::EPS, 5, 1)),
+                            Bool()
+                        ));
+
+#endif // HAVE_OPENCL

modules/ocl/test/test_pyramids.cpp

@@ -50,18 +50,16 @@
 #ifdef HAVE_OPENCL
 
 using namespace cv;
-using namespace cv::ocl;
-using namespace cvtest;
 using namespace testing;
 using namespace std;
 
-PARAM_TEST_CASE(PyrBase, MatType, int)
+PARAM_TEST_CASE(PyrBase, MatDepth, Channels)
 {
     int depth;
     int channels;
 
     Mat dst_cpu;
-    oclMat gdst;
+    ocl::oclMat gdst;
 
     virtual void SetUp()
     {
@@ -80,10 +78,10 @@ OCL_TEST_P(PyrDown, Mat)
     {
         Size size(MWIDTH, MHEIGHT);
         Mat src = randomMat(size, CV_MAKETYPE(depth, channels), 0, 255);
-        oclMat gsrc(src);
+        ocl::oclMat gsrc(src);
 
         pyrDown(src, dst_cpu);
-        pyrDown(gsrc, gdst);
+        ocl::pyrDown(gsrc, gdst);
 
         EXPECT_MAT_NEAR(dst_cpu, Mat(gdst), depth == CV_32F ? 1e-4f : 1.0f);
     }
@@ -103,10 +101,10 @@ OCL_TEST_P(PyrUp, Accuracy)
     {
         Size size(MWIDTH, MHEIGHT);
         Mat src = randomMat(size, CV_MAKETYPE(depth, channels), 0, 255);
-        oclMat gsrc(src);
+        ocl::oclMat gsrc(src);
 
         pyrUp(src, dst_cpu);
-        pyrUp(gsrc, gdst);
+        ocl::pyrUp(gsrc, gdst);
 
         EXPECT_MAT_NEAR(dst_cpu, Mat(gdst), (depth == CV_32F ? 1e-4f : 1.0));
     }

modules/ocl/test/test_split_merge.cpp

@@ -54,7 +54,7 @@ using namespace std;
 
 #define MAX_CHANNELS 4
 
-PARAM_TEST_CASE(MergeTestBase, MatType, int, bool)
+PARAM_TEST_CASE(MergeTestBase, MatDepth, Channels, bool)
 {
     int type;
     int channels;
@@ -254,11 +254,11 @@ OCL_TEST_P(Split, Accuracy)
 
 
 INSTANTIATE_TEST_CASE_P(SplitMerge, Merge, Combine(
-                            Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), Range(1, 5), Bool()));
+                            Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), Values(1, 2, 3, 4), Bool()));
 
 
 INSTANTIATE_TEST_CASE_P(SplitMerge, Split , Combine(
-                            Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), Range(1, 5), Bool()));
+                            Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), Values(1, 2, 3, 4), Bool()));
 
 
 #endif // HAVE_OPENCL

modules/ocl/test/test_warp.cpp

@@ -0,0 +1,371 @@
+/*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) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Niko Li, newlife20080214@gmail.com
+//    Jia Haipeng, jiahaipeng95@gmail.com
+//    Shengen Yan, yanshengen@gmail.com
+//    Jiang Liyuan, lyuan001.good@163.com
+//    Rock Li, Rock.Li@amd.com
+//    Wu Zailong, bullet@yeah.net
+//    Xu Pang, pangxu010@163.com
+//    Sen Liu, swjtuls1987@126.com
+//
+// 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 oclMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "test_precomp.hpp"
+
+#ifdef HAVE_OPENCL
+
+using namespace cv;
+using namespace testing;
+using namespace std;
+
+static MatType noType = -1;
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// warpAffine  & warpPerspective
+
+PARAM_TEST_CASE(WarpTestBase, MatType, Interpolation, bool, bool)
+{
+    int type, interpolation;
+    Size dsize;
+    bool useRoi, mapInverse;
+
+    Mat src, dst_whole, src_roi, dst_roi;
+    ocl::oclMat gsrc_whole, gsrc_roi, gdst_whole, gdst_roi;
+
+    virtual void SetUp()
+    {
+        type = GET_PARAM(0);
+        interpolation = GET_PARAM(1);
+        mapInverse = GET_PARAM(2);
+        useRoi = GET_PARAM(3);
+
+        if (mapInverse)
+            interpolation |= WARP_INVERSE_MAP;
+    }
+
+    void random_roi()
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder);
+        generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
+
+        dsize = randomSize(1, MAX_VALUE);
+    }
+
+    void Near(double threshold = 0.0)
+    {
+        Mat whole, roi;
+        gdst_whole.download(whole);
+        gdst_roi.download(roi);
+
+        EXPECT_MAT_NEAR(dst_whole, whole, threshold);
+        EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+    }
+};
+
+/////warpAffine
+
+typedef WarpTestBase WarpAffine;
+
+OCL_TEST_P(WarpAffine, Mat)
+{
+    static const double coeffs[2][3] =
+    {
+        { cos(CV_PI / 6), -sin(CV_PI / 6),  100.0 },
+        { sin(CV_PI / 6),  cos(CV_PI / 6), -100.0 }
+    };
+
+    static Mat M(2, 3, CV_64FC1, (void *)coeffs);
+
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        warpAffine(src_roi, dst_roi, M, dsize, interpolation);
+        ocl::warpAffine(gsrc_roi, gdst_roi, M, dsize, interpolation);
+
+        Near(1.0);
+    }
+}
+
+// warpPerspective
+
+typedef WarpTestBase WarpPerspective;
+
+OCL_TEST_P(WarpPerspective, Mat)
+{
+    static const double coeffs[3][3] =
+    {
+        { cos(CV_PI / 6), -sin(CV_PI / 6),  100.0 },
+        { sin(CV_PI / 6),  cos(CV_PI / 6), -100.0 },
+        { 0.0,             0.0,             1.0   }
+    };
+
+    static Mat M(3, 3, CV_64FC1, (void *)coeffs);
+
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        warpPerspective(src_roi, dst_roi, M, dsize, interpolation);
+        ocl::warpPerspective(gsrc_roi, gdst_roi, M, dsize, interpolation);
+
+        Near(1.0);
+    }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// remap
+
+PARAM_TEST_CASE(Remap, MatDepth, Channels, pair<MatType, MatType>, Border, bool)
+{
+    int srcType, map1Type, map2Type;
+    int borderType;
+    bool useRoi;
+
+    Scalar val;
+
+    Mat src, src_roi;
+    Mat dst, dst_roi;
+    Mat map1, map1_roi;
+    Mat map2, map2_roi;
+
+    // ocl mat with roi
+    ocl::oclMat gsrc, gsrc_roi;
+    ocl::oclMat gdst, gdst_roi;
+    ocl::oclMat gmap1, gmap1_roi;
+    ocl::oclMat gmap2, gmap2_roi;
+
+    virtual void SetUp()
+    {
+        srcType = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
+        map1Type = GET_PARAM(2).first;
+        map2Type = GET_PARAM(2).second;
+        borderType = GET_PARAM(3);
+        useRoi = GET_PARAM(4);
+    }
+
+    void random_roi()
+    {
+        val = randomScalar(-MAX_VALUE, MAX_VALUE);
+        Size srcROISize = randomSize(1, MAX_VALUE);
+        Size dstROISize = randomSize(1, MAX_VALUE);
+
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, srcROISize, srcBorder, srcType, 5, 256);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, dstROISize, dstBorder, srcType, -MAX_VALUE, MAX_VALUE);
+
+        int mapMaxValue = MAX_VALUE << 2;
+        Border map1Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(map1, map1_roi, dstROISize, map1Border, map1Type, -mapMaxValue, mapMaxValue);
+
+        Border map2Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        if (map2Type != noType)
+            randomSubMat(map2, map2_roi, dstROISize, map2Border, map2Type, -mapMaxValue, mapMaxValue);
+
+        generateOclMat(gsrc, gsrc_roi, src, srcROISize, srcBorder);
+        generateOclMat(gdst, gdst_roi, dst, dstROISize, dstBorder);
+        generateOclMat(gmap1, gmap1_roi, map1, dstROISize, map1Border);
+        if (noType != map2Type)
+            generateOclMat(gmap2, gmap2_roi, map2, dstROISize, map2Border);
+    }
+
+    void Near(double threshold = 0.0)
+    {
+        Mat whole, roi;
+        gdst.download(whole);
+        gdst_roi.download(roi);
+
+        EXPECT_MAT_NEAR(dst, whole, threshold);
+        EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+    }
+};
+
+typedef Remap Remap_INTER_NEAREST;
+
+OCL_TEST_P(Remap_INTER_NEAREST, Mat)
+{
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        remap(src_roi, dst_roi, map1_roi, map2_roi, INTER_NEAREST, borderType, val);
+        ocl::remap(gsrc_roi, gdst_roi, gmap1_roi, gmap2_roi, INTER_NEAREST, borderType, val);
+
+        Near(1.0);
+    }
+}
+
+typedef Remap Remap_INTER_LINEAR;
+
+OCL_TEST_P(Remap_INTER_LINEAR, Mat)
+{
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        cv::remap(src_roi, dst_roi, map1_roi, map2_roi, INTER_LINEAR, borderType, val);
+        ocl::remap(gsrc_roi, gdst_roi, gmap1_roi, gmap2_roi, INTER_LINEAR, borderType, val);
+
+        Near(2.0);
+    }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// resize
+
+PARAM_TEST_CASE(Resize, MatType, double, double, Interpolation, bool)
+{
+    int type, interpolation;
+    double fx, fy;
+    bool useRoi;
+
+    Mat src, dst_whole, src_roi, dst_roi;
+    ocl::oclMat gsrc_whole, gsrc_roi, gdst_whole, gdst_roi;
+
+    virtual void SetUp()
+    {
+        type = GET_PARAM(0);
+        fx = GET_PARAM(1);
+        fy = GET_PARAM(2);
+        interpolation = GET_PARAM(3);
+        useRoi = GET_PARAM(4);
+    }
+
+    void random_roi()
+    {
+        Size srcRoiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, srcRoiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        Size dstRoiSize;
+        dstRoiSize.width = cvRound(srcRoiSize.width * fx);
+        dstRoiSize.height = cvRound(srcRoiSize.height * fy);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst_whole, dst_roi, dstRoiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        generateOclMat(gsrc_whole, gsrc_roi, src, srcRoiSize, srcBorder);
+        generateOclMat(gdst_whole, gdst_roi, dst_whole, dstRoiSize, dstBorder);
+    }
+
+    void Near(double threshold = 0.0)
+    {
+        Mat whole, roi;
+        gdst_whole.download(whole);
+        gdst_roi.download(roi);
+
+        EXPECT_MAT_NEAR(dst_whole, whole, threshold);
+        EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+    }
+};
+
+OCL_TEST_P(Resize, Mat)
+{
+    for (int j = 0; j < LOOP_TIMES; j++)
+    {
+        random_roi();
+
+        resize(src_roi, dst_roi, Size(), fx, fy, interpolation);
+        ocl::resize(gsrc_roi, gdst_roi, Size(), fx, fy, interpolation);
+
+        Near(1.0);
+    }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////
+
+INSTANTIATE_TEST_CASE_P(ImgprocWarp, WarpAffine, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR, (Interpolation)INTER_CUBIC),
+                            Bool(),
+                            Bool()));
+
+INSTANTIATE_TEST_CASE_P(ImgprocWarp, WarpPerspective, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR, (Interpolation)INTER_CUBIC),
+                            Bool(),
+                            Bool()));
+
+INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_LINEAR, Combine(
+                            Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F),
+                            Values(1, 2, 3, 4),
+                            Values(pair<MatType, MatType>((MatType)CV_32FC1, (MatType)CV_32FC1),
+                                   pair<MatType, MatType>((MatType)CV_32FC2, noType)),
+                            Values((Border)BORDER_CONSTANT,
+                                   (Border)BORDER_REPLICATE,
+                                   (Border)BORDER_WRAP,
+                                   (Border)BORDER_REFLECT,
+                                   (Border)BORDER_REFLECT_101),
+                            Bool()));
+
+INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine(
+                            Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F),
+                            Values(1, 2, 3, 4),
+                            Values(pair<MatType, MatType>((MatType)CV_32FC1, (MatType)CV_32FC1),
+                                   pair<MatType, MatType>((MatType)CV_32FC2, noType),
+                                   pair<MatType, MatType>((MatType)CV_16SC2, noType)),
+                            Values((Border)BORDER_CONSTANT,
+                                   (Border)BORDER_REPLICATE,
+                                   (Border)BORDER_WRAP,
+                                   (Border)BORDER_REFLECT,
+                                   (Border)BORDER_REFLECT_101),
+                            Bool()));
+
+INSTANTIATE_TEST_CASE_P(ImgprocWarp, Resize, Combine(
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
+                            Values(0.5, 1.5, 2.0),
+                            Values(0.5, 1.5, 2.0),
+                            Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR),
+                            Bool()));
+
+#endif // HAVE_OPENCL

modules/ocl/test/utility.hpp

@@ -44,7 +44,7 @@
 #include "opencv2/core.hpp"
 
 
-#define LOOP_TIMES 1
+extern int LOOP_TIMES;
 
 #define MWIDTH 256
 #define MHEIGHT 256

modules/ts/misc/run.py

@@ -759,7 +759,10 @@ class TestSuite(object):
                 return hostlogpath
             return None
         elif path == "java":
-            cmd = [self.ant_executable, "-DjavaLibraryPath=" + self.tests_dir, "buildAndTest"]
+            cmd = [self.ant_executable,
+                   "-Dopencv.build.type="
+                     + (self.options.configuration if self.options.configuration else self.build_type),
+                   "buildAndTest"]
 
             print >> _stderr, "Run command:", " ".join(cmd)
             try:

modules/ts/src/ts_func.cpp

@@ -2968,6 +2968,16 @@ void printVersionInfo(bool useStdOut)
         if(useStdOut) std::cout << "Inner VCS version: " << ver << std::endl;
     }
 
+    const char * build_type =
+#ifdef _DEBUG
+        "debug";
+#else
+        "release";
+#endif
+
+    ::testing::Test::RecordProperty("cv_build_type", build_type);
+    if (useStdOut) std::cout << "Build type: " << build_type << std::endl;
+
     const char* parallel_framework = currentParallelFramework();
 
     if (parallel_framework) {