// 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
//
// */
#include "precomp.hpp"
+#include "opencl_kernels.hpp"
namespace cv
{
}
}
+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();
--- /dev/null
+/*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];
+ }
+}
}
}
+//////////////////////////////// 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(1e-5);
+ }
+}
+
+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(1e-5);
+ }
+}
+
+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(1e-5);
+ }
+}
+
//////////////////////////////////////// Instantiation /////////////////////////////////////////
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
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