Merge pull request #1936 from ElenaGvozdeva:ocl_flip

modules/core/src/copy.cpp

@@ -6,7 +6,6 @@
 //  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
 //
@@ -478,11 +477,59 @@ flipVert( const uchar* src0, size_t sstep, uchar* dst0, size_t dstep, Size size,
     }
 }
 
+enum { FLIP_COLS = 1 << 0, FLIP_ROWS = 1 << 1, FLIP_BOTH = FLIP_ROWS | FLIP_COLS };
+
+static bool ocl_flip(InputArray _src, OutputArray _dst, int flipCode )
+{
+    int type = _src.type(), cn = CV_MAT_CN(type);
+
+    if (cn > 4 || cn == 3)
+        return false;
+
+    const char * kernelName;
+    int flipType;
+
+    if (flipCode == 0)
+        kernelName = "arithm_flip_rows", flipType = FLIP_ROWS;
+    else if (flipCode > 0)
+        kernelName = "arithm_flip_cols", flipType = FLIP_COLS;
+    else
+        kernelName = "arithm_flip_rows_cols", flipType = FLIP_BOTH;
+
+    Size size = _src.size();
+    int cols = size.width, rows = size.height;
+    if ((cols == 1 && flipType == FLIP_COLS) ||
+            (rows == 1 && flipType == FLIP_ROWS) ||
+            (rows == 1 && cols == 1 && flipType == FLIP_BOTH))
+    {
+        _src.copyTo(_dst);
+        return true;
+    }
+
+    ocl::Kernel k(kernelName, ocl::core::flip_oclsrc,
+        format( "-D type=%s", ocl::memopTypeToStr(type)));
+    if (k.empty())
+        return false;
+
+    _dst.create(size, type);
+    UMat src = _src.getUMat(), dst = _dst.getUMat();
+
+    cols = flipType == FLIP_COLS ? ((cols+1)/2) : cols;
+    rows = flipType & FLIP_ROWS ? ((rows+1)/2) : rows;
+
+    size_t globalsize[2] = { cols, rows };
+    return k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), rows, cols).run(2, globalsize, NULL, false);
+}
+
 void flip( InputArray _src, OutputArray _dst, int flip_mode )
 {
-    Mat src = _src.getMat();
+    CV_Assert( _src.dims() <= 2 );
 
-    CV_Assert( src.dims <= 2 );
+    bool use_opencl = ocl::useOpenCL() && _dst.isUMat();
+    if ( use_opencl && ocl_flip(_src,_dst, flip_mode))
+        return;
+
+    Mat src = _src.getMat();
     _dst.create( src.size(), src.type() );
     Mat dst = _dst.getMat();
     size_t esz = src.elemSize();

modules/core/src/opencl/flip.cl

@@ -0,0 +1,102 @@
+/*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) 2013, OpenCV Foundation, 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 OpenCV Foundation 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*/
+
+#define sizeoftype ((int)sizeof(type))
+
+__kernel void arithm_flip_rows(__global const uchar* srcptr, int srcstep, int srcoffset,
+                               __global uchar* dstptr, int dststep, int dstoffset,
+                               int rows, int cols, int thread_rows, int thread_cols)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < thread_rows)
+    {
+        __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, srcoffset + x * sizeoftype));
+        __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, srcoffset + x * sizeoftype));
+
+        __global type* dst0 = (__global type*)(dstptr + mad24(y, dststep, dstoffset + x  * sizeoftype));
+        __global type* dst1 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, dstoffset + x  * sizeoftype));
+
+        dst0[0] = src1[0];
+        dst1[0] = src0[0];
+    }
+}
+
+__kernel void arithm_flip_rows_cols(__global const uchar* srcptr, int srcstep, int srcoffset,
+                                    __global uchar* dstptr, int dststep, int dstoffset,
+                                    int rows, int cols, int thread_rows, int thread_cols)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < thread_rows)
+    {
+        __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, x*sizeoftype + srcoffset));
+        __global const type* src1 = (__global const type*)(srcptr + mad24(rows - y - 1, srcstep, (cols - x - 1)*sizeoftype + srcoffset));
+
+        __global type* dst0 = (__global type*)(dstptr + mad24(rows - y - 1, dststep, (cols - x - 1)*sizeoftype + dstoffset));
+        __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, x * sizeoftype + dstoffset));
+
+        dst0[0] = src0[0];
+        dst1[0] = src1[0];
+    }
+}
+
+__kernel void arithm_flip_cols(__global const uchar* srcptr, int srcstep, int srcoffset,
+                               __global uchar* dstptr, int dststep, int dstoffset,
+                               int rows, int cols, int thread_rows, int thread_cols)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < thread_cols && y < rows)
+    {
+        __global const type* src0 = (__global const type*)(srcptr + mad24(y, srcstep, x * sizeoftype + srcoffset));
+        __global const type* src1 = (__global const type*)(srcptr + mad24(y, srcstep, (cols - x - 1)*sizeoftype + srcoffset));
+
+        __global type* dst0 = (__global type*)(dstptr + mad24(y, dststep, (cols - x - 1)*sizeoftype + dstoffset));
+        __global type* dst1 = (__global type*)(dstptr + mad24(y, dststep, x * sizeoftype + dstoffset));
+
+        dst1[0] = src1[0];
+        dst0[0] = src0[0];
+    }
+}

modules/core/test/ocl/test_arithm.cpp

@@ -964,6 +964,46 @@ OCL_TEST_P(Magnitude, Mat)
     }
 }
 
+//////////////////////////////// Flip /////////////////////////////////////////////////
+
+typedef ArithmTestBase Flip;
+
+OCL_TEST_P(Flip, X)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::flip(src1_roi, dst1_roi, 0));
+        OCL_ON(cv::flip(usrc1_roi, udst1_roi, 0));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Flip, Y)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::flip(src1_roi, dst1_roi, 1));
+        OCL_ON(cv::flip(usrc1_roi, udst1_roi, 1));
+        Near(0);
+    }
+}
+
+OCL_TEST_P(Flip, BOTH)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::flip(src1_roi, dst1_roi, -1));
+        OCL_ON(cv::flip(usrc1_roi, udst1_roi, -1));
+        Near(0);
+    }
+}
+
 //////////////////////////////////////// Instantiation /////////////////////////////////////////
 
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
@@ -993,6 +1033,7 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
+OCL_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()));
 
 } } // namespace cvtest::ocl