generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Merge pull request #1956 from ilya-lavrenov:tapi_split_merge

Browse Source
This commit is contained in:
Roman Donchenko 2013-12-09 16:56:06 +04:00 committed by OpenCV Buildbot
commit 9d87f9c974
4 changed files with 417 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

92
modules/core/src/convert.cpp
View File

@ -264,8 +264,50 @@ void cv::split(const Mat& src, Mat* mv)
}
}
namespace cv {
static bool ocl_split( InputArray _m, OutputArrayOfArrays _mv )
{
int type = _m.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
String dstargs, dstdecl, processelem;
for (int i = 0; i < cn; ++i)
{
dstargs += format("DECLARE_DST_PARAM(%d)", i);
dstdecl += format("DECLARE_DATA(%d)", i);
processelem += format("PROCESS_ELEM(%d)", i);
}
ocl::Kernel k("split", ocl::core::split_merge_oclsrc,
format("-D T=%s -D OP_SPLIT -D cn=%d -D DECLARE_DST_PARAMS=%s "
"-D DECLARE_DATA_N=%s -D PROCESS_ELEMS_N=%s",
ocl::memopTypeToStr(depth), cn, dstargs.c_str(),
dstdecl.c_str(), processelem.c_str()));
if (k.empty())
return false;
Size size = _m.size();
std::vector<UMat> & dst = *(std::vector<UMat> *)_mv.getObj();
dst.resize(cn);
for (int i = 0; i < cn; ++i)
dst[i].create(size, depth);
int argidx = k.set(0, ocl::KernelArg::ReadOnly(_m.getUMat()));
for (int i = 0; i < cn; ++i)
argidx = k.set(argidx, ocl::KernelArg::WriteOnlyNoSize(dst[i]));
size_t globalsize[2] = { size.width, size.height };
return k.run(2, globalsize, NULL, false);
}
}
void cv::split(InputArray _m, OutputArrayOfArrays _mv)
{
if (ocl::useOpenCL() && _m.dims() <= 2 && _mv.isUMatVector() &&
ocl_split(_m, _mv))
return;
Mat m = _m.getMat();
if( m.empty() )
{
@ -353,8 +395,58 @@ void cv::merge(const Mat* mv, size_t n, OutputArray _dst)
}
}
namespace cv {
static bool ocl_merge( InputArrayOfArrays _mv, OutputArray _dst )
{
const std::vector<UMat> & src = *(const std::vector<UMat> *)(_mv.getObj());
CV_Assert(!src.empty());
int type = src[0].type(), depth = CV_MAT_DEPTH(type);
Size size = src[0].size();
size_t srcsize = src.size();
for (size_t i = 0; i < srcsize; ++i)
{
int itype = src[i].type(), icn = CV_MAT_CN(itype), idepth = CV_MAT_DEPTH(itype);
if (src[i].dims > 2 || icn != 1)
return false;
CV_Assert(size == src[i].size() && depth == idepth);
}
String srcargs, srcdecl, processelem;
for (size_t i = 0; i < srcsize; ++i)
{
srcargs += format("DECLARE_SRC_PARAM(%d)", i);
srcdecl += format("DECLARE_DATA(%d)", i);
processelem += format("PROCESS_ELEM(%d)", i);
}
ocl::Kernel k("merge", ocl::core::split_merge_oclsrc,
format("-D OP_MERGE -D cn=%d -D T=%s -D DECLARE_SRC_PARAMS_N=%s -D DECLARE_DATA_N=%s -D PROCESS_ELEMS_N=%s",
(int)srcsize, ocl::memopTypeToStr(depth), srcargs.c_str(), srcdecl.c_str(), processelem.c_str()));
if (k.empty())
return false;
_dst.create(size, CV_MAKE_TYPE(depth, (int)srcsize));
UMat dst = _dst.getUMat();
int argidx = 0;
for (size_t i = 0; i < srcsize; ++i)
argidx = k.set(argidx, ocl::KernelArg::ReadOnlyNoSize(src[i]));
k.set(argidx, ocl::KernelArg::WriteOnly(dst));
size_t globalsize[2] = { dst.cols, dst.rows };
return k.run(2, globalsize, NULL, false);
}
}
void cv::merge(InputArrayOfArrays _mv, OutputArray _dst)
{
if (ocl::useOpenCL() && _mv.isUMatVector() && _dst.isUMat() && ocl_merge(_mv, _dst))
return;
std::vector<Mat> mv;
_mv.getMatVector(mv);
merge(!mv.empty() ? &mv[0] : 0, mv.size(), _dst);

14
modules/core/src/matrix.cpp
View File

@ -1822,6 +1822,13 @@ size_t _InputArray::offset(int i) const
return (size_t)(vv[i].data - vv[i].datastart);
}
if( k == STD_VECTOR_UMAT )
{
const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
CV_Assert((size_t)i < vv.size());
return vv[i].offset;
}
if( k == GPU_MAT )
{
CV_Assert( i < 0 );
@ -1861,6 +1868,13 @@ size_t _InputArray::step(int i) const
return vv[i].step;
}
if( k == STD_VECTOR_UMAT )
{
const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
CV_Assert((size_t)i < vv.size());
return vv[i].step;
}
if( k == GPU_MAT )
{
CV_Assert( i < 0 );

88
modules/core/src/opencl/split_merge.cl Normal file
View File

@ -0,0 +1,88 @@
/*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, Institute Of Software Chinese Academy Of Science, all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// 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 copyright holders 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*/
#ifdef OP_MERGE
#define DECLARE_SRC_PARAM(index) __global const uchar * src##index##ptr, int src##index##_step, int src##index##_offset,
#define DECLARE_DATA(index) __global const T * src##index = \
(__global T *)(src##index##ptr + mad24(src##index##_step, y, x * (int)sizeof(T) + src##index##_offset));
#define PROCESS_ELEM(index) dst[index] = src##index[0];
__kernel void merge(DECLARE_SRC_PARAMS_N
__global uchar * dstptr, int dst_step, int dst_offset,
int rows, int cols)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
DECLARE_DATA_N
__global T * dst = (__global T *)(dstptr + mad24(dst_step, y, x * (int)sizeof(T) * cn + dst_offset));
PROCESS_ELEMS_N
}
}
#elif defined OP_SPLIT
#define DECLARE_DST_PARAM(index) , __global uchar * dst##index##ptr, int dst##index##_step, int dst##index##_offset
#define DECLARE_DATA(index) __global T * dst##index = \
(__global T *)(dst##index##ptr + mad24(y, dst##index##_step, x * (int)sizeof(T) + dst##index##_offset));
#define PROCESS_ELEM(index) dst##index[0] = src[index];
__kernel void split(__global uchar* srcptr, int src_step, int src_offset, int rows, int cols DECLARE_DST_PARAMS)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
DECLARE_DATA_N
__global const T * src = (__global const T *)(srcptr + mad24(y, src_step, x * cn * (int)sizeof(T) + src_offset));
PROCESS_ELEMS_N
}
}
#else
#error "No operation"
#endif

223
modules/core/test/ocl/test_split_merge.cpp Normal file
View File

@ -0,0 +1,223 @@
/*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, Institute Of Software Chinese Academy Of Science, all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jia Haipeng, jiahaipeng95@gmail.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 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 "test_precomp.hpp"
#include "opencv2/ts/ocl_test.hpp"
#ifdef HAVE_OPENCL
namespace cvtest {
namespace ocl {
PARAM_TEST_CASE(MergeTestBase, MatDepth, Channels, bool)
{
int depth, cn;
bool use_roi;
TEST_DECLARE_INPUT_PARAMETER(src1)
TEST_DECLARE_INPUT_PARAMETER(src2)
TEST_DECLARE_INPUT_PARAMETER(src3)
TEST_DECLARE_INPUT_PARAMETER(src4)
TEST_DECLARE_OUTPUT_PARAMETER(dst)
std::vector<Mat> src_roi;
std::vector<UMat> usrc_roi;
virtual void SetUp()
{
depth = GET_PARAM(0);
cn = GET_PARAM(1);
use_roi = GET_PARAM(2);
CV_Assert(cn >= 1 && cn <= 4);
}
void random_roi()
{
Size roiSize = randomSize(1, MAX_VALUE);
{
Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src1, src1_roi, roiSize, src1Border, depth, 2, 11);
Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src2, src2_roi, roiSize, src2Border, depth, -1540, 1740);
Border src3Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src3, src3_roi, roiSize, src3Border, depth, -1540, 1740);
Border src4Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src4, src4_roi, roiSize, src4Border, depth, -1540, 1740);
}
Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst, dst_roi, roiSize, dstBorder, CV_MAKE_TYPE(depth, cn), 5, 16);
UMAT_UPLOAD_INPUT_PARAMETER(src1)
UMAT_UPLOAD_INPUT_PARAMETER(src2)
UMAT_UPLOAD_INPUT_PARAMETER(src3)
UMAT_UPLOAD_INPUT_PARAMETER(src4)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
src_roi.push_back(src1_roi), usrc_roi.push_back(usrc1_roi);
if (cn >= 2)
src_roi.push_back(src2_roi), usrc_roi.push_back(usrc2_roi);
if (cn >= 3)
src_roi.push_back(src3_roi), usrc_roi.push_back(usrc3_roi);
if (cn >= 4)
src_roi.push_back(src4_roi), usrc_roi.push_back(usrc4_roi);
}
void Near(double threshold = 0.)
{
EXPECT_MAT_NEAR(dst, udst, threshold);
EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
}
};
typedef MergeTestBase Merge;
OCL_TEST_P(Merge, Accuracy)
{
for(int j = 0; j < test_loop_times; j++)
{
random_roi();
OCL_OFF(cv::merge(src_roi, dst_roi));
OCL_ON(cv::merge(usrc_roi, udst_roi));
Near();
}
}
PARAM_TEST_CASE(SplitTestBase, MatType, Channels, bool)
{
int depth, cn;
bool use_roi;
TEST_DECLARE_INPUT_PARAMETER(src)
TEST_DECLARE_OUTPUT_PARAMETER(dst1)
TEST_DECLARE_OUTPUT_PARAMETER(dst2)
TEST_DECLARE_OUTPUT_PARAMETER(dst3)
TEST_DECLARE_OUTPUT_PARAMETER(dst4)
std::vector<Mat> dst_roi, dst;
std::vector<UMat> udst_roi, udst;
virtual void SetUp()
{
depth = GET_PARAM(0);
cn = GET_PARAM(1);
use_roi = GET_PARAM(2);
CV_Assert(cn >= 1 && cn <= 4);
}
void random_roi()
{
Size roiSize = randomSize(1, MAX_VALUE);
Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src, src_roi, roiSize, srcBorder, CV_MAKE_TYPE(depth, cn), 5, 16);
{
Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst1, dst1_roi, roiSize, dst1Border, depth, 2, 11);
Border dst2Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst2, dst2_roi, roiSize, dst2Border, depth, -1540, 1740);
Border dst3Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst3, dst3_roi, roiSize, dst3Border, depth, -1540, 1740);
Border dst4Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst4, dst4_roi, roiSize, dst4Border, depth, -1540, 1740);
}
UMAT_UPLOAD_INPUT_PARAMETER(src)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst1)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst2)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst3)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst4)
dst_roi.push_back(dst1_roi), udst_roi.push_back(udst1_roi),
dst.push_back(dst1), udst.push_back(udst1);
if (cn >= 2)
dst_roi.push_back(dst2_roi), udst_roi.push_back(udst2_roi),
dst.push_back(dst2), udst.push_back(udst2);
if (cn >= 3)
dst_roi.push_back(dst3_roi), udst_roi.push_back(udst3_roi),
dst.push_back(dst3), udst.push_back(udst3);
if (cn >= 4)
dst_roi.push_back(dst4_roi), udst_roi.push_back(udst4_roi),
dst.push_back(dst4), udst.push_back(udst4);
}
};
typedef SplitTestBase Split;
OCL_TEST_P(Split, Accuracy)
{
for (int j = 0; j < test_loop_times; j++)
{
random_roi();
OCL_OFF(cv::split(src_roi, dst_roi));
OCL_ON(cv::split(usrc_roi, udst_roi));
for (int i = 0; i < cn; ++i)
{
EXPECT_MAT_NEAR(dst[i], udst[i], 0.0);
EXPECT_MAT_NEAR(dst_roi[i], udst_roi[i], 0.0);
}
}
}
OCL_INSTANTIATE_TEST_CASE_P(SplitMerge, Merge, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
OCL_INSTANTIATE_TEST_CASE_P(SplitMerge, Split, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
} } // namespace cvtest::ocl
#endif // HAVE_OPENCL