Merge remote-tracking branch 'origin/2.4' into merge-2.4

Conflicts:
	modules/core/include/opencv2/core/version.hpp
	modules/ocl/include/opencv2/ocl/ocl.hpp
	modules/ocl/src/initialization.cpp
	modules/ocl/test/main.cpp
	modules/superres/CMakeLists.txt
	modules/superres/src/input_array_utility.cpp
	modules/superres/src/input_array_utility.hpp
	modules/superres/src/optical_flow.cpp

modules/core/doc/xml_yaml_persistence.rst

@@ -11,7 +11,7 @@ You can store and then restore various OpenCV data structures to/from XML (http:
 
 Use the following procedure to write something to XML or YAML:
  #. Create new :ocv:class:`FileStorage` and open it for writing. It can be done with a single call to :ocv:func:`FileStorage::FileStorage` constructor that takes a filename, or you can use the default constructor and then call :ocv:func:`FileStorage::open`. Format of the file (XML or YAML) is determined from the filename extension (".xml" and ".yml"/".yaml", respectively)
- #. Write all the data you want using the streaming operator ``>>``, just like in the case of STL streams.
+ #. Write all the data you want using the streaming operator ``<<``, just like in the case of STL streams.
  #. Close the file using :ocv:func:`FileStorage::release`. ``FileStorage`` destructor also closes the file.
 
 Here is an example: ::

modules/features2d/src/brisk.cpp

@@ -525,7 +525,11 @@ BRISK::operator()( InputArray _image, InputArray _mask, std::vector<KeyPoint>& k
   bool doOrientation=true;
   if (useProvidedKeypoints)
     doOrientation = false;
-  computeDescriptorsAndOrOrientation(_image, _mask, keypoints, _descriptors, true, doOrientation,
+
+  // If the user specified cv::noArray(), this will yield false. Otherwise it will return true.
+  bool doDescriptors = _descriptors.needed();
+
+  computeDescriptorsAndOrOrientation(_image, _mask, keypoints, _descriptors, doDescriptors, doOrientation,
                                        useProvidedKeypoints);
 }
 

modules/highgui/src/cap.cpp

@@ -220,8 +220,8 @@ CV_IMPL CvCapture * cvCreateCameraCapture (int index)
                 return capture;
             break;
 #endif
-#ifdef HAVE_VFW
         case CV_CAP_VFW:
+#ifdef HAVE_VFW
             capture = cvCreateCameraCapture_VFW (index);
             if (capture)
                 return capture;

modules/imgproc/src/templmatch.cpp

@@ -248,6 +248,8 @@ void cv::matchTemplate( InputArray _img, InputArray _templ, OutputArray _result,
     CV_Assert( (img.depth() == CV_8U || img.depth() == CV_32F) &&
                img.type() == templ.type() );
 
+    CV_Assert( img.rows >= templ.rows && img.cols >= templ.cols);
+
     Size corrSize(img.cols - templ.cols + 1, img.rows - templ.rows + 1);
     _result.create(corrSize, CV_32F);
     Mat result = _result.getMat();

modules/ocl/include/opencv2/ocl.hpp

@@ -853,6 +853,19 @@ namespace cv
         CV_EXPORTS void cornerMinEigenVal_dxdy(const oclMat &src, oclMat &dst, oclMat &Dx, oclMat &Dy,
             int blockSize, int ksize, int bordertype = cv::BORDER_DEFAULT);
 
+
+        /////////////////////////////////// ML ///////////////////////////////////////////
+
+        //! Compute closest centers for each lines in source and lable it after center's index
+        // supports CV_32FC1/CV_32FC2/CV_32FC4 data type
+        CV_EXPORTS void distanceToCenters(oclMat &dists, oclMat &labels, const oclMat &src, const oclMat &centers);
+
+        //!Does k-means procedure on GPU
+        // supports CV_32FC1/CV_32FC2/CV_32FC4 data type
+        CV_EXPORTS double kmeans(const oclMat &src, int K, oclMat &bestLabels,
+                                     TermCriteria criteria, int attemps, int flags, oclMat &centers);
+
+
         ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
         ///////////////////////////////////////////CascadeClassifier//////////////////////////////////////////////////////////////////
         ///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

modules/ocl/src/gfft.cpp → modules/ocl/src/gftt.cpp

@@ -55,7 +55,7 @@ namespace cv
     namespace ocl
     {
         ///////////////////////////OpenCL kernel strings///////////////////////////
-        extern const char *imgproc_gfft;
+        extern const char *imgproc_gftt;
     }
 }
 
@@ -133,7 +133,7 @@ struct Sorter<BITONIC>
             for(int passOfStage = 0; passOfStage < stage + 1; ++passOfStage)
             {
                 args[4] = std::make_pair(sizeof(cl_int), (void *)&passOfStage);
-                openCLExecuteKernel(cxt, &imgproc_gfft, kernelname, globalThreads, localThreads, args, -1, -1);
+                openCLExecuteKernel(cxt, &imgproc_gftt, kernelname, globalThreads, localThreads, args, -1, -1);
             }
         }
     }
@@ -160,12 +160,12 @@ struct Sorter<SELECTION>
         args.push_back( std::make_pair( sizeof(cl_int), (void*)&count) );
         args.push_back( std::make_pair( lds_size,       (void*)NULL) );
 
-        openCLExecuteKernel(cxt, &imgproc_gfft, kernelname, globalThreads, localThreads, args, -1, -1);
+        openCLExecuteKernel(cxt, &imgproc_gftt, kernelname, globalThreads, localThreads, args, -1, -1);
 
         //final
         kernelname = "sortCorners_selectionSortFinal";
         args.pop_back();
-        openCLExecuteKernel(cxt, &imgproc_gfft, kernelname, globalThreads, localThreads, args, -1, -1);
+        openCLExecuteKernel(cxt, &imgproc_gftt, kernelname, globalThreads, localThreads, args, -1, -1);
     }
 };
 
@@ -201,7 +201,7 @@ int findCorners_caller(
     size_t localThreads[3]  = {16, 16, 1};
 
     const char * opt = mask.empty() ? "" : "-D WITH_MASK";
-    openCLExecuteKernel(cxt, &imgproc_gfft, kernelname, globalThreads, localThreads, args, -1, -1, opt);
+    openCLExecuteKernel(cxt, &imgproc_gftt, kernelname, globalThreads, localThreads, args, -1, -1, opt);
     return std::min(Mat(g_counter).at<int>(0), max_count);
 }
 }//unnamed namespace

modules/ocl/src/initialization.cpp

@@ -319,8 +319,7 @@ namespace cv
             char clVersion[256];
             for (unsigned i = 0; i < numPlatforms; ++i)
             {
-
-                cl_uint numsdev;
+                cl_uint numsdev = 0;
                 cl_int status = clGetDeviceIDs(platforms[i], devicetype, 0, NULL, &numsdev);
                 if(status != CL_DEVICE_NOT_FOUND)
                     openCLVerifyCall(status);

modules/ocl/src/opencl/kmeans_kernel.cl

+/*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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Xiaopeng Fu, fuxiaopeng2222@163.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 GpuMaterials 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*/
+
+__kernel void distanceToCenters(
+    int label_step, int K,
+    __global float *src,
+    __global int *labels, int dims, int rows,
+    __global float *centers,
+    __global float *dists)
+{
+    int gid = get_global_id(1);
+
+    float dist, euDist, min;
+    int minCentroid;
+
+    if(gid >= rows)
+        return;
+
+    for(int i = 0 ; i < K; i++)
+    {
+        euDist = 0;
+        for(int j = 0; j < dims; j++)
+        {
+            dist = (src[j + gid * dims]
+                    - centers[j + i * dims]);
+            euDist += dist * dist;
+        }
+
+        if(i == 0)
+        {
+            min = euDist;
+            minCentroid = 0;
+        }
+        else if(euDist < min)
+        {
+            min = euDist;
+            minCentroid = i;
+        }
+    }
+    dists[gid] = min;
+    labels[label_step * gid] = minCentroid;
+}

modules/ocl/test/main.cpp

@@ -73,14 +73,12 @@ void print_info()
 #endif
 
 }
-std::string workdir;
 int main(int argc, char **argv)
 {
-    TS::ptr()->init("ocl");
+    TS::ptr()->init(".");
     InitGoogleTest(&argc, argv);
     const char *keys =
         "{ h | false              | print help message }"
-		"{ w | ../../../samples/c/| set working directory i.e. -w=C:\\}"
         "{ t | gpu                | set device type:i.e. -t=cpu or gpu}"
         "{ p | 0                  | set platform id i.e. -p=0}"
         "{ d | 0                  | set device id i.e. -d=0}";
@@ -92,7 +90,6 @@ int main(int argc, char **argv)
         cmd.printMessage();
         return 0;
     }
-    workdir = cmd.get<string>("w");
     string type = cmd.get<string>("t");
     unsigned int pid = cmd.get<unsigned int>("p");
     int device = cmd.get<int>("d");

modules/ocl/test/test_calib3d.cpp

@@ -50,7 +50,6 @@
 
 using namespace cv;
 
-extern std::string workdir;
 PARAM_TEST_CASE(StereoMatchBM, int, int)
 {
     int n_disp;
@@ -66,9 +65,9 @@ PARAM_TEST_CASE(StereoMatchBM, int, int)
 TEST_P(StereoMatchBM, Regression)
 {
 
-    Mat left_image  = readImage("stereobm/aloe-L.png", IMREAD_GRAYSCALE);
-    Mat right_image = readImage("stereobm/aloe-R.png", IMREAD_GRAYSCALE);
-    Mat disp_gold   = readImage("stereobm/aloe-disp.png", IMREAD_GRAYSCALE);
+    Mat left_image  = readImage("gpu/stereobm/aloe-L.png", IMREAD_GRAYSCALE);
+    Mat right_image = readImage("gpu/stereobm/aloe-R.png", IMREAD_GRAYSCALE);
+    Mat disp_gold   = readImage("gpu/stereobm/aloe-disp.png", IMREAD_GRAYSCALE);
     ocl::oclMat d_left, d_right;
     ocl::oclMat d_disp(left_image.size(), CV_8U);
     Mat  disp;
@@ -113,9 +112,9 @@ PARAM_TEST_CASE(StereoMatchBP, int, int, int, float, float, float, float)
 };
 TEST_P(StereoMatchBP, Regression)
 {
-    Mat left_image  = readImage("stereobp/aloe-L.png");
-    Mat right_image = readImage("stereobp/aloe-R.png");
-    Mat disp_gold   = readImage("stereobp/aloe-disp.png", IMREAD_GRAYSCALE);
+    Mat left_image  = readImage("gpu/stereobp/aloe-L.png");
+    Mat right_image = readImage("gpu/stereobp/aloe-R.png");
+    Mat disp_gold   = readImage("gpu/stereobp/aloe-disp.png", IMREAD_GRAYSCALE);
     ocl::oclMat d_left, d_right;
     ocl::oclMat d_disp;
     Mat  disp;
@@ -166,9 +165,9 @@ PARAM_TEST_CASE(StereoMatchConstSpaceBP, int, int, int, int, float, float, float
 };
 TEST_P(StereoMatchConstSpaceBP, Regression)
 {
-    Mat left_image  = readImage("csstereobp/aloe-L.png");
-    Mat right_image = readImage("csstereobp/aloe-R.png");
-    Mat disp_gold   = readImage("csstereobp/aloe-disp.png", IMREAD_GRAYSCALE);
+    Mat left_image  = readImage("gpu/csstereobp/aloe-L.png");
+    Mat right_image = readImage("gpu/csstereobp/aloe-R.png");
+    Mat disp_gold   = readImage("gpu/csstereobp/aloe-disp.png", IMREAD_GRAYSCALE);
 
     ocl::oclMat d_left, d_right;
     ocl::oclMat d_disp;

modules/ocl/test/test_canny.cpp

@@ -48,7 +48,6 @@
 
 ////////////////////////////////////////////////////////
 // Canny
-extern std::string workdir;
 IMPLEMENT_PARAM_CLASS(AppertureSize, int);
 IMPLEMENT_PARAM_CLASS(L2gradient, bool);
 
@@ -67,7 +66,7 @@ PARAM_TEST_CASE(Canny, AppertureSize, L2gradient)
 
 TEST_P(Canny, Accuracy)
 {
-    cv::Mat img = readImage(workdir + "fruits.jpg", cv::IMREAD_GRAYSCALE);
+    cv::Mat img = readImage("cv/shared/fruits.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(img.empty());
 
     double low_thresh = 50.0;

modules/ocl/test/test_kmeans.cpp

+/*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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Erping Pang,   pang_er_ping@163.com
+//    Xiaopeng Fu,   fuxiaopeng2222@163.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 "precomp.hpp"
+
+#ifdef HAVE_OPENCL
+
+using namespace cvtest;
+using namespace testing;
+using namespace std;
+using namespace cv;
+
+#define OCL_KMEANS_USE_INITIAL_LABELS 1
+#define OCL_KMEANS_PP_CENTERS         2
+
+PARAM_TEST_CASE(Kmeans, int, int, int)
+{
+    int type;
+    int K;
+    int flags;
+    cv::Mat src ;
+    ocl::oclMat d_src, d_dists;
+
+    Mat labels, centers;
+    ocl::oclMat d_labels, d_centers;
+    cv::RNG rng ;
+    virtual void SetUp(){
+        K = GET_PARAM(0);
+        type = GET_PARAM(1);
+        flags = GET_PARAM(2);
+        rng = TS::ptr()->get_rng();
+
+        // MWIDTH=256, MHEIGHT=256. defined in utility.hpp
+        cv::Size size = cv::Size(MWIDTH, MHEIGHT);
+        src.create(size, type);
+        int row_idx = 0;
+        const int max_neighbour = MHEIGHT / K - 1;
+        CV_Assert(K <= MWIDTH);
+        for(int i = 0; i < K; i++ )
+        {
+            Mat center_row_header = src.row(row_idx);
+            center_row_header.setTo(0);
+            int nchannel = center_row_header.channels();
+            for(int j = 0; j < nchannel; j++)
+                center_row_header.at<float>(0, i*nchannel+j) = 50000.0;
+
+            for(int j = 0; (j < max_neighbour) || 
+                           (i == K-1 && j < max_neighbour + MHEIGHT%K); j ++)
+            {
+                Mat cur_row_header = src.row(row_idx + 1 + j);
+                center_row_header.copyTo(cur_row_header);
+                Mat tmpmat = randomMat(rng, cur_row_header.size(), cur_row_header.type(), -200, 200, false);
+                cur_row_header += tmpmat;
+            }
+            row_idx += 1 + max_neighbour;
+        }
+    }
+};
+TEST_P(Kmeans, Mat){
+
+    if(flags & KMEANS_USE_INITIAL_LABELS)
+    {
+        // inital a given labels
+        labels.create(src.rows, 1, CV_32S);
+        int *label = labels.ptr<int>();
+        for(int i = 0; i < src.rows; i++)
+            label[i] = rng.uniform(0, K);
+        d_labels.upload(labels);
+    }
+    d_src.upload(src);
+
+    for(int j = 0; j < LOOP_TIMES; j++)
+    {
+        kmeans(src, K, labels,
+            TermCriteria( TermCriteria::EPS + TermCriteria::MAX_ITER, 100, 0),
+            1, flags, centers);
+
+        ocl::kmeans(d_src, K, d_labels,
+            TermCriteria( TermCriteria::EPS + TermCriteria::MAX_ITER, 100, 0),
+            1, flags, d_centers);
+  
+        Mat dd_labels(d_labels);
+        Mat dd_centers(d_centers);
+        if(flags & KMEANS_USE_INITIAL_LABELS)
+        {
+            EXPECT_MAT_NEAR(labels, dd_labels, 0);
+            EXPECT_MAT_NEAR(centers, dd_centers, 1e-3);
+        } 
+        else 
+        {
+            int row_idx = 0;
+            for(int i = 0; i < K; i++)
+            {
+                // verify lables with ground truth resutls
+                int label = labels.at<int>(row_idx);
+                int header_label = dd_labels.at<int>(row_idx);
+                for(int j = 0; (j < MHEIGHT/K)||(i == K-1 && j < MHEIGHT/K+MHEIGHT%K); j++)
+                {
+                    ASSERT_NEAR(labels.at<int>(row_idx+j), label, 0);
+                    ASSERT_NEAR(dd_labels.at<int>(row_idx+j), header_label, 0);
+                }
+
+                // verify centers
+                float *center = centers.ptr<float>(label);
+                float *header_center = dd_centers.ptr<float>(header_label);
+                for(int t = 0; t < centers.cols; t++)
+                    ASSERT_NEAR(center[t], header_center[t], 1e-3);
+
+                row_idx += MHEIGHT/K;
+            }
+        }
+    }
+}
+INSTANTIATE_TEST_CASE_P(OCL_ML, Kmeans, Combine(
+    Values(3, 5, 8),
+    Values(CV_32FC1, CV_32FC2, CV_32FC4),
+    Values(OCL_KMEANS_USE_INITIAL_LABELS/*, OCL_KMEANS_PP_CENTERS*/))); 
+
+#endif

modules/ocl/test/test_moments.cpp

@@ -44,7 +44,7 @@ TEST_P(MomentsTest, Mat)
     {
         if(test_contours)
         {
-            Mat src = imread( workdir + "../cpp/pic3.png", IMREAD_GRAYSCALE );
+            Mat src = readImage( "cv/shared/pic3.png", IMREAD_GRAYSCALE );
             ASSERT_FALSE(src.empty());
             Mat canny_output;
             vector<vector<Point> > contours;

modules/ocl/test/test_objdetect.cpp

@@ -66,11 +66,8 @@ PARAM_TEST_CASE(HOG, Size, int)
     {
         winSize = GET_PARAM(0);
         type = GET_PARAM(1);
-        img_rgb = readImage(workdir + "../gpu/road.png");
-        if(img_rgb.empty())
-        {
-            std::cout << "Couldn't read road.png" << std::endl;
-        }
+        img_rgb = readImage("gpu/hog/road.png");
+        ASSERT_FALSE(img_rgb.empty());
     }
 };
 
@@ -215,18 +212,11 @@ PARAM_TEST_CASE(Haar, int, CascadeName)
     virtual void SetUp()
     {
         flags = GET_PARAM(0);
-        cascadeName = (workdir + "../../data/haarcascades/").append(GET_PARAM(1));
-        if( (!cascade.load( cascadeName )) || (!cpucascade.load(cascadeName)) )
-        {
-            std::cout << "ERROR: Could not load classifier cascade" << std::endl;
-            return;
-        }
-        img = readImage(workdir + "lena.jpg", IMREAD_GRAYSCALE);
-        if(img.empty())
-        {
-            std::cout << "Couldn't read lena.jpg" << std::endl;
-            return ;
-        }
+        cascadeName = (string(cvtest::TS::ptr()->get_data_path()) + "cv/cascadeandhog/cascades/").append(GET_PARAM(1));
+        ASSERT_TRUE(cascade.load( cascadeName ));
+        ASSERT_TRUE(cpucascade.load(cascadeName));
+        img = readImage("cv/shared/lena.png", IMREAD_GRAYSCALE);
+        ASSERT_FALSE(img.empty());
         equalizeHist(img, img);
         d_img.upload(img);
     }

modules/ocl/test/test_optflow.cpp

@@ -75,7 +75,7 @@ PARAM_TEST_CASE(GoodFeaturesToTrack, MinDistance)
 
 TEST_P(GoodFeaturesToTrack, Accuracy)
 {
-    cv::Mat frame = readImage(workdir + "../gpu/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame.empty());
 
     int maxCorners = 1000;
@@ -146,10 +146,10 @@ PARAM_TEST_CASE(TVL1, bool)
 
 TEST_P(TVL1, Accuracy)
 {
-    cv::Mat frame0 = readImage(workdir + "../gpu/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame0.empty());
 
-    cv::Mat frame1 = readImage(workdir + "../gpu/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame1.empty());
 
     cv::ocl::OpticalFlowDual_TVL1_OCL d_alg;
@@ -188,10 +188,10 @@ PARAM_TEST_CASE(Sparse, bool, bool)
 
 TEST_P(Sparse, Mat)
 {
-    cv::Mat frame0 = readImage(workdir + "../gpu/rubberwhale1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
     ASSERT_FALSE(frame0.empty());
 
-    cv::Mat frame1 = readImage(workdir + "../gpu/rubberwhale2.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
     ASSERT_FALSE(frame1.empty());
 
     cv::Mat gray_frame;
@@ -301,10 +301,10 @@ PARAM_TEST_CASE(Farneback, PyrScale, PolyN, FarnebackOptFlowFlags, UseInitFlow)
 
 TEST_P(Farneback, Accuracy)
 {
-    cv::Mat frame0 = imread(workdir + "/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame0.empty());
 
-    cv::Mat frame1 = imread(workdir + "/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame1.empty());
 
     double polySigma = polyN <= 5 ? 1.1 : 1.5;

modules/superres/CMakeLists.txt

@@ -5,5 +5,5 @@ endif()
 set(the_description "Super Resolution")
 ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 -Wundef)
 ocv_define_module(superres opencv_imgproc opencv_video
-                  OPTIONAL opencv_highgui
+                  OPTIONAL opencv_highgui opencv_ocl
                   opencv_gpuarithm opencv_gpufilters opencv_gpuwarping opencv_gpuimgproc opencv_gpuoptflow opencv_gpucodec)

modules/superres/include/opencv2/superres.hpp

@@ -92,6 +92,7 @@ namespace cv
         // Dennis Mitzel, Thomas Pock, Thomas Schoenemann, Daniel Cremers. Video Super Resolution using Duality Based TV-L1 Optical Flow.
         CV_EXPORTS Ptr<SuperResolution> createSuperResolution_BTVL1();
         CV_EXPORTS Ptr<SuperResolution> createSuperResolution_BTVL1_GPU();
+        CV_EXPORTS Ptr<SuperResolution> createSuperResolution_BTVL1_OCL();
     }
 }
 

modules/superres/include/opencv2/superres/optical_flow.hpp

@@ -63,10 +63,12 @@ namespace cv
 
         CV_EXPORTS Ptr<DenseOpticalFlowExt> createOptFlow_DualTVL1();
         CV_EXPORTS Ptr<DenseOpticalFlowExt> createOptFlow_DualTVL1_GPU();
+        CV_EXPORTS Ptr<DenseOpticalFlowExt> createOptFlow_DualTVL1_OCL();
 
         CV_EXPORTS Ptr<DenseOpticalFlowExt> createOptFlow_Brox_GPU();
 
         CV_EXPORTS Ptr<DenseOpticalFlowExt> createOptFlow_PyrLK_GPU();
+        CV_EXPORTS Ptr<DenseOpticalFlowExt> createOptFlow_PyrLK_OCL();
     }
 }
 

modules/superres/perf/perf_superres_ocl.cpp

+/*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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "perf_precomp.hpp"
+
+#ifdef HAVE_OPENCL
+
+#include "opencv2/ocl.hpp"
+using namespace std;
+using namespace std::tr1;
+using namespace testing;
+using namespace perf;
+using namespace cv;
+using namespace cv::superres;
+
+namespace
+{
+    class OneFrameSource_OCL : public FrameSource
+    {
+    public:
+        explicit OneFrameSource_OCL(const ocl::oclMat& frame) : frame_(frame) {}
+
+        void nextFrame(OutputArray frame)
+        {
+            ocl::getOclMatRef(frame) = frame_;
+        }
+        void reset()
+        {
+        }
+
+    private:
+        ocl::oclMat frame_;
+    };
+
+
+    class ZeroOpticalFlowOCL : public DenseOpticalFlowExt
+    {
+    public:
+        void calc(InputArray frame0, InputArray, OutputArray flow1, OutputArray flow2)
+        {
+            ocl::oclMat& frame0_ = ocl::getOclMatRef(frame0);
+            ocl::oclMat& flow1_ = ocl::getOclMatRef(flow1);
+            ocl::oclMat& flow2_ = ocl::getOclMatRef(flow2);
+
+            cv::Size size = frame0_.size();
+
+            if(!flow2.needed())
+            {
+                flow1_.create(size, CV_32FC2);
+                flow1_.setTo(Scalar::all(0));
+            }
+            else
+            {
+                flow1_.create(size, CV_32FC1);
+                flow2_.create(size, CV_32FC1);
+
+                flow1_.setTo(Scalar::all(0));
+                flow2_.setTo(Scalar::all(0));
+            }
+        }
+
+        void collectGarbage()
+        {
+        }
+    };
+}
+
+PERF_TEST_P(Size_MatType, SuperResolution_BTVL1_OCL,
+    Combine(Values(szSmall64, szSmall128),
+    Values(MatType(CV_8UC1), MatType(CV_8UC3))))
+{
+    std::vector<cv::ocl::Info>info;
+    cv::ocl::getDevice(info);
+
+    declare.time(5 * 60);
+
+    const Size size = get<0>(GetParam());
+    const int type = get<1>(GetParam());
+
+    Mat frame(size, type);
+    declare.in(frame, WARMUP_RNG);
+
+    ocl::oclMat frame_ocl;
+    frame_ocl.upload(frame);
+
+
+    const int scale = 2;
+    const int iterations = 50;
+    const int temporalAreaRadius = 1;
+    Ptr<DenseOpticalFlowExt> opticalFlowOcl(new ZeroOpticalFlowOCL);
+
+    Ptr<SuperResolution> superRes_ocl = createSuperResolution_BTVL1_OCL();
+
+    superRes_ocl->set("scale", scale);
+    superRes_ocl->set("iterations", iterations);
+    superRes_ocl->set("temporalAreaRadius", temporalAreaRadius);
+    superRes_ocl->set("opticalFlow", opticalFlowOcl);
+
+    superRes_ocl->setInput(new OneFrameSource_OCL(frame_ocl));
+
+    ocl::oclMat dst_ocl;
+    superRes_ocl->nextFrame(dst_ocl);
+
+    TEST_CYCLE_N(10) superRes_ocl->nextFrame(dst_ocl);
+    frame_ocl.release();
+    CPU_SANITY_CHECK(dst_ocl);
+}
+#endif

modules/superres/src/frame_source.cpp

@@ -118,12 +118,24 @@ namespace
         {
             vc_ >> _frame.getMatRef();
         }
-        else
+        else if(_frame.kind() == _InputArray::GPU_MAT)
+        {
+            vc_ >> frame_;
+            arrCopy(frame_, _frame);
+        }
+        else if(_frame.kind() == _InputArray::OCL_MAT)
         {
             vc_ >> frame_;
+            if(!frame_.empty())
+            {
                 arrCopy(frame_, _frame);
             }
         }
+        else
+        {
+            //should never get here
+        }
+    }
 
     class VideoFrameSource : public CaptureFrameSource
     {

modules/superres/src/input_array_utility.cpp

@@ -108,30 +108,59 @@ namespace
     {
         src.getGpuMat().copyTo(dst.getGpuMatRef());
     }
+#ifdef HAVE_OPENCV_OCL
+    void ocl2mat(InputArray src, OutputArray dst)
+    {
+        dst.getMatRef() = (Mat)ocl::getOclMatRef(src);
+    }
+    void mat2ocl(InputArray src, OutputArray dst)
+    {
+        Mat m = src.getMat();
+        ocl::getOclMatRef(dst) = (ocl::oclMat)m;
+    }
+    void ocl2ocl(InputArray src, OutputArray dst)
+    {
+        ocl::getOclMatRef(src).copyTo(ocl::getOclMatRef(dst));
+    }
+#else
+    void ocl2mat(InputArray, OutputArray)
+    {
+        CV_Error(Error::StsNotImplemented, "The called functionality is disabled for current build or platform");;
+    }
+    void mat2ocl(InputArray, OutputArray)
+    {
+        CV_Error(Error::StsNotImplemented, "The called functionality is disabled for current build or platform");;
+    }
+    void ocl2ocl(InputArray, OutputArray)
+    {
+        CV_Error(Error::StsNotImplemented, "The called functionality is disabled for current build or platform");
+    }
+#endif
 }
 
 void cv::superres::arrCopy(InputArray src, OutputArray dst)
 {
     typedef void (*func_t)(InputArray src, OutputArray dst);
-    static const func_t funcs[10][10] =
+    static const func_t funcs[11][11] =
     {
         {0, 0, 0, 0, 0, 0, 0, 0, 0, 0},
-        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu},
-        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu},
-        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu},
-        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu},
-        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu},
-        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu},
-        {0, buf2arr, buf2arr, buf2arr, buf2arr, buf2arr, buf2arr, buf2arr, 0 /*buf2arr*/, buf2arr},
-        {0, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/},
-        {0, gpu2mat, gpu2mat, gpu2mat, gpu2mat, gpu2mat, gpu2mat, arr2buf, 0 /*arr2tex*/, gpu2gpu}
+        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu, mat2ocl},
+        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu, mat2ocl},
+        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu, mat2ocl},
+        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu, mat2ocl},
+        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu, mat2ocl},
+        {0, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, mat2mat, arr2buf, 0 /*arr2tex*/, mat2gpu, mat2ocl},
+        {0, buf2arr, buf2arr, buf2arr, buf2arr, buf2arr, buf2arr, buf2arr, 0 /*buf2arr*/, buf2arr, 0      },
+        {0, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0 /*tex2arr*/, 0},
+        {0, gpu2mat, gpu2mat, gpu2mat, gpu2mat, gpu2mat, gpu2mat, arr2buf, 0 /*arr2tex*/, gpu2gpu, 0      },
+        {0, ocl2mat, ocl2mat, ocl2mat, ocl2mat, ocl2mat, ocl2mat, 0,       0,             0,       ocl2ocl}
     };
 
     const int src_kind = src.kind() >> _InputArray::KIND_SHIFT;
     const int dst_kind = dst.kind() >> _InputArray::KIND_SHIFT;
 
-    CV_DbgAssert( src_kind >= 0 && src_kind < 10 );
-    CV_DbgAssert( dst_kind >= 0 && dst_kind < 10 );
+    CV_DbgAssert( src_kind >= 0 && src_kind < 11 );
+    CV_DbgAssert( dst_kind >= 0 && dst_kind < 11 );
 
     const func_t func = funcs[src_kind][dst_kind];
     CV_DbgAssert( func != 0 );
@@ -173,7 +202,6 @@ namespace
             break;
         }
     }
-
     void convertToDepth(InputArray src, OutputArray dst, int depth)
     {
         CV_Assert( src.depth() <= CV_64F );
@@ -254,3 +282,70 @@ GpuMat cv::superres::convertToType(const GpuMat& src, int type, GpuMat& buf0, Gp
     convertToDepth(buf0, buf1, depth);
     return buf1;
 }
+#ifdef HAVE_OPENCV_OCL
+namespace
+{
+    // TODO(pengx17): remove these overloaded functions until IntputArray fully supports oclMat
+    void convertToCn(const ocl::oclMat& src, ocl::oclMat& dst, int cn)
+    {
+        CV_Assert( src.channels() == 1 || src.channels() == 3 || src.channels() == 4 );
+        CV_Assert( cn == 1 || cn == 3 || cn == 4 );
+
+        static const int codes[5][5] =
+        {
+            {-1, -1, -1, -1, -1},
+            {-1, -1, -1, COLOR_GRAY2BGR, COLOR_GRAY2BGRA},
+            {-1, -1, -1, -1, -1},
+            {-1, COLOR_BGR2GRAY, -1, -1, COLOR_BGR2BGRA},
+            {-1, COLOR_BGRA2GRAY, -1, COLOR_BGRA2BGR, -1},
+        };
+
+        const int code = codes[src.channels()][cn];
+        CV_DbgAssert( code >= 0 );
+
+        ocl::cvtColor(src, dst, code, cn);
+    }
+    void convertToDepth(const ocl::oclMat& src, ocl::oclMat& dst, int depth)
+    {
+        CV_Assert( src.depth() <= CV_64F );
+        CV_Assert( depth == CV_8U || depth == CV_32F );
+
+        static const double maxVals[] =
+        {
+            std::numeric_limits<uchar>::max(),
+            std::numeric_limits<schar>::max(),
+            std::numeric_limits<ushort>::max(),
+            std::numeric_limits<short>::max(),
+            std::numeric_limits<int>::max(),
+            1.0,
+            1.0,
+        };
+        const double scale = maxVals[depth] / maxVals[src.depth()];
+        src.convertTo(dst, depth, scale);
+    }
+}
+ocl::oclMat cv::superres::convertToType(const ocl::oclMat& src, int type, ocl::oclMat& buf0, ocl::oclMat& buf1)
+{
+    if (src.type() == type)
+        return src;
+
+    const int depth = CV_MAT_DEPTH(type);
+    const int cn = CV_MAT_CN(type);
+
+    if (src.depth() == depth)
+    {
+        convertToCn(src, buf0, cn);
+        return buf0;
+    }
+
+    if (src.channels() == cn)
+    {
+        convertToDepth(src, buf1, depth);
+        return buf1;
+    }
+
+    convertToCn(src, buf0, cn);
+    convertToDepth(buf0, buf1, depth);
+    return buf1;
+}
+#endif

modules/superres/src/input_array_utility.hpp

@@ -45,6 +45,9 @@
 
 #include "opencv2/core.hpp"
 #include "opencv2/core/gpu.hpp"
+#ifdef HAVE_OPENCV_OCL
+#include "opencv2/ocl.hpp"
+#endif
 
 namespace cv
 {
@@ -57,6 +60,10 @@ namespace cv
 
         CV_EXPORTS Mat convertToType(const Mat& src, int type, Mat& buf0, Mat& buf1);
         CV_EXPORTS gpu::GpuMat convertToType(const gpu::GpuMat& src, int type, gpu::GpuMat& buf0, gpu::GpuMat& buf1);
+
+#ifdef HAVE_OPENCV_OCL
+        CV_EXPORTS ocl::oclMat convertToType(const ocl::oclMat& src, int type, ocl::oclMat& buf0, ocl::oclMat& buf1);
+#endif
     }
 }
 

modules/superres/src/opencl/superres_btvl1.cl

+/*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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Jin Ma jin@multicorewareinc.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*/
+
+__kernel void buildMotionMapsKernel(__global float* forwardMotionX,
+    __global float* forwardMotionY,
+    __global float* backwardMotionX,
+    __global float* backwardMotionY,
+    __global float* forwardMapX,
+    __global float* forwardMapY,
+    __global float* backwardMapX,
+    __global float* backwardMapY,
+    int forwardMotionX_row,
+    int forwardMotionX_col,
+    int forwardMotionX_step,
+    int forwardMotionY_step,
+    int backwardMotionX_step,
+    int backwardMotionY_step,
+    int forwardMapX_step,
+    int forwardMapY_step,
+    int backwardMapX_step,
+    int backwardMapY_step
+    )
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if(x < forwardMotionX_col && y < forwardMotionX_row)
+    {
+        float fx = forwardMotionX[y * forwardMotionX_step + x];
+        float fy = forwardMotionY[y * forwardMotionY_step + x];
+
+        float bx = backwardMotionX[y * backwardMotionX_step + x];
+        float by = backwardMotionY[y * backwardMotionY_step + x];
+
+        forwardMapX[y * forwardMapX_step + x] = x + bx;
+        forwardMapY[y * forwardMapY_step + x] = y + by;
+
+        backwardMapX[y * backwardMapX_step + x] = x + fx;
+        backwardMapY[y * backwardMapY_step + x] = y + fy;
+    }
+}
+
+__kernel void upscaleKernel(__global float* src,
+    __global float* dst,
+    int src_step,
+    int dst_step,
+    int src_row,
+    int src_col,
+    int scale,
+    int channels
+    )
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if(x < src_col && y < src_row)
+    {
+        if(channels == 1)
+        {
+            dst[y * scale * dst_step + x * scale] = src[y * src_step + x];
+        }else if(channels == 3)
+        {
+            dst[y * channels * scale * dst_step + 3 * x * scale + 0] = src[y * channels * src_step + 3 * x + 0];
+            dst[y * channels * scale * dst_step + 3 * x * scale + 1] = src[y * channels * src_step + 3 * x + 1];
+            dst[y * channels * scale * dst_step + 3 * x * scale + 2] = src[y * channels * src_step + 3 * x + 2];
+        }else
+        {
+            dst[y * channels * scale * dst_step + 4 * x * scale + 0] = src[y * channels * src_step + 4 * x + 0];
+            dst[y * channels * scale * dst_step + 4 * x * scale + 1] = src[y * channels * src_step + 4 * x + 1];
+            dst[y * channels * scale * dst_step + 4 * x * scale + 2] = src[y * channels * src_step + 4 * x + 2];
+            dst[y * channels * scale * dst_step + 4 * x * scale + 3] = src[y * channels * src_step + 4 * x + 3];            
+        }
+    }
+}
+
+
+float diffSign(float a, float b)
+{
+    return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
+}
+
+float3 diffSign3(float3 a, float3 b)
+{
+    float3 pos;
+    pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
+    pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
+    pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
+    return pos;
+}
+
+float4 diffSign4(float4 a, float4 b)
+{
+    float4 pos;
+    pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
+    pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
+    pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
+    pos.w = 0.0f;
+    return pos;
+}
+
+__kernel void diffSignKernel(__global float* src1,
+    __global float* src2,
+    __global float* dst,
+    int src1_row,
+    int src1_col,
+    int dst_step,
+    int src1_step,
+    int src2_step)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if(x < src1_col && y < src1_row)
+    {
+        dst[y * dst_step + x] = diffSign(src1[y * src1_step + x], src2[y * src2_step + x]);
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+}
+
+__kernel void calcBtvRegularizationKernel(__global float* src,
+    __global float* dst,
+    int src_step,
+    int dst_step,
+    int src_row,
+    int src_col,
+    int ksize,
+    int channels,
+    __global float* c_btvRegWeights
+    )
+{
+    int x = get_global_id(0) + ksize;
+    int y = get_global_id(1) + ksize;
+
+    if ((y < src_row - ksize) && (x < src_col - ksize))
+    {
+        if(channels == 1)
+        {
+            const float srcVal = src[y * src_step + x];
+            float dstVal = 0.0f;
+
+            for (int m = 0, count = 0; m <= ksize; ++m)
+            {
+                for (int l = ksize; l + m >= 0; --l, ++count)
+                    dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal));
+            }
+            dst[y * dst_step + x] = dstVal;
+        }else if(channels == 3)
+        {
+            float3 srcVal;
+            srcVal.x = src[y * src_step + 3 * x + 0];
+            srcVal.y = src[y * src_step + 3 * x + 1];
+            srcVal.z = src[y * src_step + 3 * x + 2];
+
+            float3 dstVal;
+            dstVal.x = 0.0f;
+            dstVal.y = 0.0f;
+            dstVal.z = 0.0f;
+
+            for (int m = 0, count = 0; m <= ksize; ++m)
+            {
+                for (int l = ksize; l + m >= 0; --l, ++count)
+                {
+                    float3 src1;
+                    src1.x = src[(y + m) * src_step + 3 * (x + l) + 0];
+                    src1.y = src[(y + m) * src_step + 3 * (x + l) + 1];
+                    src1.z = src[(y + m) * src_step + 3 * (x + l) + 2];
+
+                    float3 src2;
+                    src2.x = src[(y - m) * src_step + 3 * (x - l) + 0];
+                    src2.y = src[(y - m) * src_step + 3 * (x - l) + 1];
+                    src2.z = src[(y - m) * src_step + 3 * (x - l) + 2];
+
+                    dstVal = dstVal + c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal));
+                }
+            }
+            dst[y * dst_step + 3 * x + 0] = dstVal.x;
+            dst[y * dst_step + 3 * x + 1] = dstVal.y;
+            dst[y * dst_step + 3 * x + 2] = dstVal.z;
+        }else
+        {
+            float4 srcVal;
+            srcVal.x = src[y * src_step + 4 * x + 0];//r type =float
+            srcVal.y = src[y * src_step + 4 * x + 1];//g
+            srcVal.z = src[y * src_step + 4 * x + 2];//b
+            srcVal.w = src[y * src_step + 4 * x + 3];//a
+
+            float4 dstVal;
+            dstVal.x = 0.0f;
+            dstVal.y = 0.0f;
+            dstVal.z = 0.0f;
+            dstVal.w = 0.0f;
+
+            for (int m = 0, count = 0; m <= ksize; ++m)
+            {
+                for (int l = ksize; l + m >= 0; --l, ++count)
+                {
+                    float4 src1;
+                    src1.x = src[(y + m) * src_step + 4 * (x + l) + 0];
+                    src1.y = src[(y + m) * src_step + 4 * (x + l) + 1];
+                    src1.z = src[(y + m) * src_step + 4 * (x + l) + 2];
+                    src1.w = src[(y + m) * src_step + 4 * (x + l) + 3];
+
+                    float4 src2;
+                    src2.x = src[(y - m) * src_step + 4 * (x - l) + 0];
+                    src2.y = src[(y - m) * src_step + 4 * (x - l) + 1];
+                    src2.z = src[(y - m) * src_step + 4 * (x - l) + 2];
+                    src2.w = src[(y - m) * src_step + 4 * (x - l) + 3];
+
+                    dstVal = dstVal + c_btvRegWeights[count] * (diffSign4(srcVal, src1) - diffSign4(src2, srcVal));
+
+                }
+            }
+            dst[y * dst_step + 4 * x + 0] = dstVal.x;
+            dst[y * dst_step + 4 * x + 1] = dstVal.y;
+            dst[y * dst_step + 4 * x + 2] = dstVal.z;
+            dst[y * dst_step + 4 * x + 3] = dstVal.w;
+        }
+    }
+}
\ No newline at end of file

modules/superres/src/optical_flow.cpp

@@ -718,3 +718,195 @@ Ptr<DenseOpticalFlowExt> cv::superres::createOptFlow_DualTVL1_GPU()
 }
 
 #endif // HAVE_OPENCV_GPUOPTFLOW
+#ifdef HAVE_OPENCV_OCL
+
+namespace
+{
+    class oclOpticalFlow : public DenseOpticalFlowExt
+    {
+    public:
+        explicit oclOpticalFlow(int work_type);
+
+        void calc(InputArray frame0, InputArray frame1, OutputArray flow1, OutputArray flow2);
+        void collectGarbage();
+
+    protected:
+        virtual void impl(const cv::ocl::oclMat& input0, const cv::ocl::oclMat& input1, cv::ocl::oclMat& dst1, cv::ocl::oclMat& dst2) = 0;
+
+    private:
+        int work_type_;
+        cv::ocl::oclMat buf_[6];
+        cv::ocl::oclMat u_, v_, flow_;
+    };
+
+    oclOpticalFlow::oclOpticalFlow(int work_type) : work_type_(work_type)
+    {
+    }
+
+    void oclOpticalFlow::calc(InputArray frame0, InputArray frame1, OutputArray flow1, OutputArray flow2)
+    {
+        ocl::oclMat& _frame0 = ocl::getOclMatRef(frame0);
+        ocl::oclMat& _frame1 = ocl::getOclMatRef(frame1);
+        ocl::oclMat& _flow1  = ocl::getOclMatRef(flow1);
+        ocl::oclMat& _flow2  = ocl::getOclMatRef(flow2);
+
+        CV_Assert( _frame1.type() == _frame0.type() );
+        CV_Assert( _frame1.size() == _frame0.size() );
+
+        cv::ocl::oclMat input0_ = convertToType(_frame0, work_type_, buf_[2], buf_[3]);
+        cv::ocl::oclMat input1_ = convertToType(_frame1, work_type_, buf_[4], buf_[5]);
+
+        impl(input0_, input1_, u_, v_);//go to tvl1 algorithm
+
+        u_.copyTo(_flow1);
+        v_.copyTo(_flow2);
+    }
+
+    void oclOpticalFlow::collectGarbage()
+    {
+        for (int i = 0; i < 6; ++i)
+            buf_[i].release();
+        u_.release();
+        v_.release();
+        flow_.release();
+    }
+}
+///////////////////////////////////////////////////////////////////
+// PyrLK_OCL
+
+namespace
+{
+    class PyrLK_OCL : public oclOpticalFlow
+    {
+    public:
+        AlgorithmInfo* info() const;
+
+        PyrLK_OCL();
+
+        void collectGarbage();
+
+    protected:
+        void impl(const ocl::oclMat& input0, const ocl::oclMat& input1, ocl::oclMat& dst1, ocl::oclMat& dst2);
+
+    private:
+        int winSize_;
+        int maxLevel_;
+        int iterations_;
+
+        ocl::PyrLKOpticalFlow alg_;
+    };
+
+    CV_INIT_ALGORITHM(PyrLK_OCL, "DenseOpticalFlowExt.PyrLK_OCL",
+        obj.info()->addParam(obj, "winSize", obj.winSize_);
+    obj.info()->addParam(obj, "maxLevel", obj.maxLevel_);
+    obj.info()->addParam(obj, "iterations", obj.iterations_));
+
+    PyrLK_OCL::PyrLK_OCL() : oclOpticalFlow(CV_8UC1)
+    {
+        winSize_ = alg_.winSize.width;
+        maxLevel_ = alg_.maxLevel;
+        iterations_ = alg_.iters;
+    }
+
+    void PyrLK_OCL::impl(const cv::ocl::oclMat& input0, const cv::ocl::oclMat& input1, cv::ocl::oclMat& dst1, cv::ocl::oclMat& dst2)
+    {
+        alg_.winSize.width = winSize_;
+        alg_.winSize.height = winSize_;
+        alg_.maxLevel = maxLevel_;
+        alg_.iters = iterations_;
+
+        alg_.dense(input0, input1, dst1, dst2);
+    }
+
+    void PyrLK_OCL::collectGarbage()
+    {
+        alg_.releaseMemory();
+        oclOpticalFlow::collectGarbage();
+    }
+}
+
+Ptr<DenseOpticalFlowExt> cv::superres::createOptFlow_PyrLK_OCL()
+{
+    return new PyrLK_OCL;
+}
+
+///////////////////////////////////////////////////////////////////
+// DualTVL1_OCL
+
+namespace
+{
+    class DualTVL1_OCL : public oclOpticalFlow
+    {
+    public:
+        AlgorithmInfo* info() const;
+
+        DualTVL1_OCL();
+
+        void collectGarbage();
+
+    protected:
+        void impl(const cv::ocl::oclMat& input0, const cv::ocl::oclMat& input1, cv::ocl::oclMat& dst1, cv::ocl::oclMat& dst2);
+
+    private:
+        double tau_;
+        double lambda_;
+        double theta_;
+        int nscales_;
+        int warps_;
+        double epsilon_;
+        int iterations_;
+        bool useInitialFlow_;
+
+        ocl::OpticalFlowDual_TVL1_OCL alg_;
+    };
+
+    CV_INIT_ALGORITHM(DualTVL1_OCL, "DenseOpticalFlowExt.DualTVL1_OCL",
+    obj.info()->addParam(obj, "tau", obj.tau_);
+    obj.info()->addParam(obj, "lambda", obj.lambda_);
+    obj.info()->addParam(obj, "theta", obj.theta_);
+    obj.info()->addParam(obj, "nscales", obj.nscales_);
+    obj.info()->addParam(obj, "warps", obj.warps_);
+    obj.info()->addParam(obj, "epsilon", obj.epsilon_);
+    obj.info()->addParam(obj, "iterations", obj.iterations_);
+    obj.info()->addParam(obj, "useInitialFlow", obj.useInitialFlow_));
+
+    DualTVL1_OCL::DualTVL1_OCL() : oclOpticalFlow(CV_8UC1)
+    {
+        tau_ = alg_.tau;
+        lambda_ = alg_.lambda;
+        theta_ = alg_.theta;
+        nscales_ = alg_.nscales;
+        warps_ = alg_.warps;
+        epsilon_ = alg_.epsilon;
+        iterations_ = alg_.iterations;
+        useInitialFlow_ = alg_.useInitialFlow;
+    }
+
+    void DualTVL1_OCL::impl(const cv::ocl::oclMat& input0, const cv::ocl::oclMat& input1, cv::ocl::oclMat& dst1, cv::ocl::oclMat& dst2)
+    {
+        alg_.tau = tau_;
+        alg_.lambda = lambda_;
+        alg_.theta = theta_;
+        alg_.nscales = nscales_;
+        alg_.warps = warps_;
+        alg_.epsilon = epsilon_;
+        alg_.iterations = iterations_;
+        alg_.useInitialFlow = useInitialFlow_;
+
+        alg_(input0, input1, dst1, dst2);
+
+    }
+
+    void DualTVL1_OCL::collectGarbage()
+    {
+        alg_.collectGarbage();
+        oclOpticalFlow::collectGarbage();
+    }
+}
+
+Ptr<DenseOpticalFlowExt> cv::superres::createOptFlow_DualTVL1_OCL()
+{
+    return new DualTVL1_OCL;
+}
+
+#endif

modules/superres/src/precomp.hpp

@@ -81,6 +81,10 @@
 #  include "opencv2/gpucodec.hpp"
 #endif
 
+#ifdef HAVE_OPENCV_OCL
+    #include "opencv2/ocl/private/util.hpp"
+#endif
+
 #ifdef HAVE_OPENCV_HIGHGUI
     #include "opencv2/highgui.hpp"
 #endif

modules/superres/test/test_superres.cpp

@@ -274,5 +274,12 @@ TEST_F(SuperResolution, BTVL1_GPU)
 {
     RunTest(cv::superres::createSuperResolution_BTVL1_GPU());
 }
-
+#endif
+#if defined(HAVE_OPENCV_OCL) && defined(HAVE_OPENCL)
+TEST_F(SuperResolution, BTVL1_OCL)
+{
+    std::vector<cv::ocl::Info> infos;
+    cv::ocl::getDevice(infos);
+    RunTest(cv::superres::createSuperResolution_BTVL1_OCL());
+}
 #endif 

modules/ts/include/opencv2/ts.hpp

@@ -551,6 +551,13 @@ int main(int argc, char **argv) \
     return RUN_ALL_TESTS(); \
 }
 
+// This usually only makes sense in perf tests with several implementations,
+// some of which are not available.
+#define CV_TEST_FAIL_NO_IMPL() do { \
+    ::testing::Test::RecordProperty("custom_status", "noimpl"); \
+    FAIL() << "No equivalent implementation."; \
+} while (0)
+
 #endif
 
 #include "opencv2/ts/ts_perf.hpp"

modules/ts/misc/testlog_parser.py

@@ -13,10 +13,17 @@ class TestInfo(object):
         self.name = xmlnode.getAttribute("name")
         self.value_param = xmlnode.getAttribute("value_param")
         self.type_param = xmlnode.getAttribute("type_param")
-        if xmlnode.getElementsByTagName("failure"):
+
+        custom_status = xmlnode.getAttribute("custom_status")
+        failures = xmlnode.getElementsByTagName("failure")
+
+        if len(custom_status) > 0:
+            self.status = custom_status
+        elif len(failures) > 0:
             self.status = "failed"
         else:
             self.status = xmlnode.getAttribute("status")
+
         if self.name.startswith("DISABLED_"):
             self.status = "disabled"
             self.fixture = self.fixture.replace("DISABLED_", "")

modules/ts/misc/xls-report.py

@@ -64,6 +64,10 @@
         Name for the sheet. If this parameter is missing, the name of sheet's directory
         will be used.
 
+    * 'sheet_properties': [(string, string)]
+        List of arbitrary (key, value) pairs that somehow describe the sheet. Will be
+        dumped into the first row of the sheet in string form.
+
     Note that all keys are optional, although to get useful results, you'll want to
     specify at least 'configurations' and 'configuration_matchers'.
 
@@ -100,6 +104,7 @@ bad_speedup_style = xlwt.easyxf('font: color red', num_format_str='#0.00')
 no_speedup_style = no_time_style
 error_speedup_style = xlwt.easyxf('pattern: pattern solid, fore_color orange')
 header_style = xlwt.easyxf('font: bold true; alignment: horizontal centre, vertical top, wrap True')
+subheader_style = xlwt.easyxf('alignment: horizontal centre, vertical top')
 
 class Collector(object):
     def __init__(self, config_match_func, include_unmatched):
@@ -189,6 +194,8 @@ def main():
     arg_parser.add_argument('-c', '--config', metavar='CONF', help='global configuration file')
     arg_parser.add_argument('--include-unmatched', action='store_true',
         help='include results from XML files that were not recognized by configuration matchers')
+    arg_parser.add_argument('--show-times-per-pixel', action='store_true',
+        help='for tests that have an image size parameter, show per-pixel time, as well as total time')
 
     args = arg_parser.parse_args()
 
@@ -231,24 +238,64 @@ def main():
 
         sheet = wb.add_sheet(sheet_conf.get('sheet_name', os.path.basename(os.path.abspath(sheet_path))))
 
-        sheet.row(0).height = 800
+        sheet_properties = sheet_conf.get('sheet_properties', [])
+
+        sheet.write(0, 0, 'Properties:')
+
+        sheet.write(0, 1,
+          'N/A' if len(sheet_properties) == 0 else
+          ' '.join(str(k) + '=' + repr(v) for (k, v) in sheet_properties))
+
+        sheet.row(2).height = 800
         sheet.panes_frozen = True
         sheet.remove_splits = True
-        sheet.horz_split_pos = 1
-        sheet.horz_split_first_visible = 1
 
         sheet_comparisons = sheet_conf.get('comparisons', [])
 
-        for i, w in enumerate([2000, 15000, 2500, 2000, 15000]
-                + (len(config_names) + 1 + len(sheet_comparisons)) * [4000]):
-            sheet.col(i).width = w
+        row = 2
+
+        col = 0
 
-        for i, caption in enumerate(['Module', 'Test', 'Image\nsize', 'Data\ntype', 'Parameters']
-                + config_names + [None]
-                + [comp['to'] + '\nvs\n' + comp['from'] for comp in sheet_comparisons]):
-            sheet.row(0).write(i, caption, header_style)
+        for (w, caption) in [
+                (2500, 'Module'),
+                (10000, 'Test'),
+                (2500, 'Image\nsize'),
+                (2000, 'Data\ntype'),
+                (7500, 'Other parameters')]:
+            sheet.col(col).width = w
+            if args.show_times_per_pixel:
+                sheet.write_merge(row, row + 1, col, col, caption, header_style)
+            else:
+                sheet.write(row, col, caption, header_style)
+            col += 1
+
+        for config_name in config_names:
+            if args.show_times_per_pixel:
+                sheet.col(col).width = 3000
+                sheet.col(col + 1).width = 3000
+                sheet.write_merge(row, row, col, col + 1, config_name, header_style)
+                sheet.write(row + 1, col, 'total, ms', subheader_style)
+                sheet.write(row + 1, col + 1, 'per pixel, ns', subheader_style)
+                col += 2
+            else:
+                sheet.col(col).width = 4000
+                sheet.write(row, col, config_name, header_style)
+                col += 1
+
+        col += 1 # blank column between configurations and comparisons
+
+        for comp in sheet_comparisons:
+            sheet.col(col).width = 4000
+            caption = comp['to'] + '\nvs\n' + comp['from']
+            if args.show_times_per_pixel:
+                sheet.write_merge(row, row + 1, col, col, caption, header_style)
+            else:
+                sheet.write(row, col, caption, header_style)
 
-        row = 1
+        row += 2 if args.show_times_per_pixel else 1
+
+        sheet.horz_split_pos = row
+        sheet.horz_split_first_visible = row
 
         module_colors = sheet_conf.get('module_colors', {})
         module_styles = {module: xlwt.easyxf('pattern: pattern solid, fore_color {}'.format(color))
@@ -259,21 +306,49 @@ def main():
                 sheet.write(row, 0, module, module_styles.get(module, xlwt.Style.default_style))
                 sheet.write(row, 1, test)
 
-                param_list = param[1:-1].split(", ")
-                sheet.write(row, 2, next(ifilter(re_image_size.match, param_list), None))
-                sheet.write(row, 3, next(ifilter(re_data_type.match, param_list), None))
+                param_list = param[1:-1].split(', ') if param.startswith('(') and param.endswith(')') else [param]
+
+                image_size = next(ifilter(re_image_size.match, param_list), None)
+                if image_size is not None:
+                    sheet.write(row, 2, image_size)
+                    del param_list[param_list.index(image_size)]
+
+                data_type = next(ifilter(re_data_type.match, param_list), None)
+                if data_type is not None:
+                    sheet.write(row, 3, data_type)
+                    del param_list[param_list.index(data_type)]
+
+                sheet.row(row).write(4, ' | '.join(param_list))
+
+                col = 5
 
-                sheet.row(row).write(4, param)
-                for i, c in enumerate(config_names):
+                for c in config_names:
                     if c in configs:
-                        sheet.write(row, 5 + i, configs[c], time_style)
+                        sheet.write(row, col, configs[c], time_style)
                     else:
-                        sheet.write(row, 5 + i, None, no_time_style)
+                        sheet.write(row, col, None, no_time_style)
+                    col += 1
+                    if args.show_times_per_pixel:
+                        sheet.write(row, col,
+                          xlwt.Formula(
+                            '''
+                              {0} * 1000000 / (
+                                VALUE(MID({1}; 1; SEARCH("x"; {1}) - 1))
+                                  * VALUE(MID({1}; SEARCH("x"; {1}) + 1; LEN({1})))
+                              )
+                            '''.replace('\n', '').replace(' ', '').format(
+                              xlwt.Utils.rowcol_to_cell(row, col - 1),
+                              xlwt.Utils.rowcol_to_cell(row, 2)
+                            )
+                          ),
+                          time_style)
+                        col += 1
 
-                for i, comp in enumerate(sheet_comparisons):
+                col += 1 # blank column
+
+                for comp in sheet_comparisons:
                     cmp_from = configs.get(comp["from"])
                     cmp_to = configs.get(comp["to"])
-                    col = 5 + len(config_names) + 1 + i
 
                     if isinstance(cmp_from, numbers.Number) and isinstance(cmp_to, numbers.Number):
                         try:
@@ -286,6 +361,8 @@ def main():
                     else:
                         sheet.write(row, col, None, no_speedup_style)
 
+                    col += 1
+
                 row += 1
                 if row % 1000 == 0: sheet.flush_row_data()