From b93002872750b8f8844fee69180aaad57d328e85 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Thu, 5 Dec 2013 14:36:45 +0400 Subject: [PATCH 1/2] Added cv::flip using Transparent API --- modules/core/src/copy.cpp | 54 +++++++++++++- modules/core/src/opencl/flip.cl | 102 ++++++++++++++++++++++++++ modules/core/test/ocl/test_arithm.cpp | 41 +++++++++++ 3 files changed, 194 insertions(+), 3 deletions(-) create mode 100644 modules/core/src/opencl/flip.cl diff --git a/modules/core/src/copy.cpp b/modules/core/src/copy.cpp index ee34ef451..8e60227a5 100644 --- a/modules/core/src/copy.cpp +++ b/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 // @@ -47,6 +46,7 @@ // */ #include "precomp.hpp" +#include "opencl_kernels.hpp" namespace cv { @@ -477,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(); diff --git a/modules/core/src/opencl/flip.cl b/modules/core/src/opencl/flip.cl new file mode 100644 index 000000000..4e5304153 --- /dev/null +++ b/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]; + } +} diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index 10cec7bc0..4755e69c2 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/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(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())); @@ -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 From f5a01f15a59c06989801bd2e5a5aeab6f25bbdc2 Mon Sep 17 00:00:00 2001 From: Elena Gvozdeva Date: Fri, 6 Dec 2013 15:53:00 +0400 Subject: [PATCH 2/2] Fixed test for ocl_flip --- modules/core/test/ocl/test_arithm.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index 4755e69c2..844be7bdf 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/modules/core/test/ocl/test_arithm.cpp @@ -976,7 +976,7 @@ OCL_TEST_P(Flip, X) OCL_OFF(cv::flip(src1_roi, dst1_roi, 0)); OCL_ON(cv::flip(usrc1_roi, udst1_roi, 0)); - Near(1e-5); + Near(0); } } @@ -988,7 +988,7 @@ OCL_TEST_P(Flip, Y) OCL_OFF(cv::flip(src1_roi, dst1_roi, 1)); OCL_ON(cv::flip(usrc1_roi, udst1_roi, 1)); - Near(1e-5); + Near(0); } } @@ -1000,7 +1000,7 @@ OCL_TEST_P(Flip, BOTH) OCL_OFF(cv::flip(src1_roi, dst1_roi, -1)); OCL_ON(cv::flip(usrc1_roi, udst1_roi, -1)); - Near(1e-5); + Near(0); } }