mirror of
https://github.com/opencv/opencv.git
synced 2024-11-29 05:29:54 +08:00
Added cv::flip using Transparent API
This commit is contained in:
parent
0736ede7e5
commit
b930028727
@ -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();
|
||||
|
102
modules/core/src/opencl/flip.cl
Normal file
102
modules/core/src/opencl/flip.cl
Normal file
@ -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];
|
||||
}
|
||||
}
|
@ -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
|
||||
|
||||
|
Loading…
Reference in New Issue
Block a user