diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 6259a7ada2..acc0e90046 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -612,12 +612,105 @@ void cv::mixChannels( const Mat* src, size_t nsrcs, Mat* dst, size_t ndsts, cons } } +namespace cv { + +static void getUMatIndex(const std::vector & um, int cn, int & idx, int & cnidx) +{ + int totalChannels = 0; + for (size_t i = 0, size = um.size(); i < size; ++i) + { + int ccn = um[i].channels(); + totalChannels += ccn; + + if (totalChannels == cn) + { + idx = (int)(i + 1); + cnidx = 0; + return; + } + else if (totalChannels > cn) + { + idx = (int)i; + cnidx = i == 0 ? cn : (cn - totalChannels + ccn); + return; + } + } + + idx = cnidx = -1; +} + +static bool ocl_mixChannels(InputArrayOfArrays _src, InputOutputArrayOfArrays _dst, + const int* fromTo, size_t npairs) +{ + const std::vector & src = *(const std::vector *)_src.getObj(); + std::vector & dst = *(std::vector *)_dst.getObj(); + + size_t nsrc = src.size(), ndst = dst.size(); + CV_Assert(nsrc > 0 && ndst > 0); + + Size size = src[0].size(); + int depth = src[0].depth(), esz = CV_ELEM_SIZE(depth); + + for (size_t i = 1, ssize = src.size(); i < ssize; ++i) + CV_Assert(src[i].size() == size && src[i].depth() == depth); + for (size_t i = 0, dsize = dst.size(); i < dsize; ++i) + CV_Assert(dst[i].size() == size && dst[i].depth() == depth); + + String declsrc, decldst, declproc, declcn; + std::vector srcargs(npairs), dstargs(npairs); + + for (size_t i = 0; i < npairs; ++i) + { + int scn = fromTo[i<<1], dcn = fromTo[(i<<1) + 1]; + int src_idx, src_cnidx, dst_idx, dst_cnidx; + + getUMatIndex(src, scn, src_idx, src_cnidx); + getUMatIndex(dst, dcn, dst_idx, dst_cnidx); + + CV_Assert(dst_idx >= 0 && src_idx >= 0); + + srcargs[i] = src[src_idx]; + srcargs[i].offset += src_cnidx * esz; + + dstargs[i] = dst[dst_idx]; + dstargs[i].offset += dst_cnidx * esz; + + declsrc += format("DECLARE_INPUT_MAT(%d)", i); + decldst += format("DECLARE_OUTPUT_MAT(%d)", i); + declproc += format("PROCESS_ELEM(%d)", i); + declcn += format(" -D scn%d=%d -D dcn%d=%d", i, src[src_idx].channels(), i, dst[dst_idx].channels()); + } + + ocl::Kernel k("mixChannels", ocl::core::mixchannels_oclsrc, + format("-D T=%s -D DECLARE_INPUT_MATS=%s -D DECLARE_OUTPUT_MATS=%s" + " -D PROCESS_ELEMS=%s%s", ocl::memopTypeToStr(depth), + declsrc.c_str(), decldst.c_str(), declproc.c_str(), declcn.c_str())); + if (k.empty()) + return false; + + size_t argindex = 0; + for (size_t i = 0; i < npairs; ++i) + argindex = k.set(argindex, ocl::KernelArg::ReadOnlyNoSize(srcargs[i])); + for (size_t i = 0; i < npairs; ++i) + argindex = k.set(argindex, ocl::KernelArg::ReadOnlyNoSize(dstargs[i])); + k.set(k.set(argindex, size.height), size.width); + + size_t globalsize[2] = { size.width, size.height }; + return k.run(2, globalsize, NULL, false); +} + +} void cv::mixChannels(InputArrayOfArrays src, InputOutputArrayOfArrays dst, const int* fromTo, size_t npairs) { - if(npairs == 0) + if (npairs == 0 || fromTo == NULL) return; + + if (ocl::useOpenCL() && src.isUMatVector() && dst.isUMatVector() && + ocl_mixChannels(src, dst, fromTo, npairs)) + return; + bool src_is_mat = src.kind() != _InputArray::STD_VECTOR_MAT && src.kind() != _InputArray::STD_VECTOR_VECTOR; bool dst_is_mat = dst.kind() != _InputArray::STD_VECTOR_MAT && @@ -639,8 +732,16 @@ void cv::mixChannels(InputArrayOfArrays src, InputOutputArrayOfArrays dst, void cv::mixChannels(InputArrayOfArrays src, InputOutputArrayOfArrays dst, const std::vector& fromTo) { - if(fromTo.empty()) + if (fromTo.empty()) return; + + if (ocl::useOpenCL() && src.isUMatVector() && dst.isUMatVector() /*&& + ocl_mixChannels(src, dst, &fromTo[0], fromTo.size()>>1)*/) + { + CV_Assert(ocl_mixChannels(src, dst, &fromTo[0], fromTo.size()>>1)); + return; + } + bool src_is_mat = src.kind() != _InputArray::STD_VECTOR_MAT && src.kind() != _InputArray::STD_VECTOR_VECTOR; bool dst_is_mat = dst.kind() != _InputArray::STD_VECTOR_MAT && diff --git a/modules/core/src/opencl/mixchannels.cl b/modules/core/src/opencl/mixchannels.cl new file mode 100644 index 0000000000..173421e6ce --- /dev/null +++ b/modules/core/src/opencl/mixchannels.cl @@ -0,0 +1,64 @@ +/*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*/ + +#define DECLARE_INPUT_MAT(i) \ + __global const uchar * src##i##ptr, int src##i##_step, int src##i##_offset, +#define DECLARE_OUTPUT_MAT(i) \ + __global const uchar * dst##i##ptr, int dst##i##_step, int dst##i##_offset, +#define PROCESS_ELEM(i) \ + int src##i##_index = mad24(src##i##_step, y, x * (int)sizeof(T) * scn##i + src##i##_offset); \ + __global const T * src##i = (__global const T *)(src##i##ptr + src##i##_index); \ + int dst##i##_index = mad24(dst##i##_step, y, x * (int)sizeof(T) * dcn##i + dst##i##_offset); \ + __global T * dst##i = (__global T *)(dst##i##ptr + dst##i##_index); \ + dst##i[0] = src##i[0]; + +__kernel void mixChannels(DECLARE_INPUT_MATS DECLARE_OUTPUT_MATS int rows, int cols) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < cols && y < rows) + { + PROCESS_ELEMS + } +} diff --git a/modules/core/test/ocl/test_split_merge.cpp b/modules/core/test/ocl/test_split_merge.cpp index c1c0f0e306..d7fdcea7c7 100644 --- a/modules/core/test/ocl/test_split_merge.cpp +++ b/modules/core/test/ocl/test_split_merge.cpp @@ -52,7 +52,9 @@ namespace cvtest { namespace ocl { -PARAM_TEST_CASE(MergeTestBase, MatDepth, Channels, bool) +//////////////////////////////////////// Merge /////////////////////////////////////////////// + +PARAM_TEST_CASE(Merge, MatDepth, Channels, bool) { int depth, cn; bool use_roi; @@ -75,7 +77,7 @@ PARAM_TEST_CASE(MergeTestBase, MatDepth, Channels, bool) CV_Assert(cn >= 1 && cn <= 4); } - void random_roi() + void generateTestData() { Size roiSize = randomSize(1, MAX_VALUE); @@ -117,13 +119,11 @@ PARAM_TEST_CASE(MergeTestBase, MatDepth, Channels, bool) } }; -typedef MergeTestBase Merge; - OCL_TEST_P(Merge, Accuracy) { for(int j = 0; j < test_loop_times; j++) { - random_roi(); + generateTestData(); OCL_OFF(cv::merge(src_roi, dst_roi)); OCL_ON(cv::merge(usrc_roi, udst_roi)); @@ -132,7 +132,9 @@ OCL_TEST_P(Merge, Accuracy) } } -PARAM_TEST_CASE(SplitTestBase, MatType, Channels, bool) +//////////////////////////////////////// Split /////////////////////////////////////////////// + +PARAM_TEST_CASE(Split, MatType, Channels, bool) { int depth, cn; bool use_roi; @@ -155,7 +157,7 @@ PARAM_TEST_CASE(SplitTestBase, MatType, Channels, bool) CV_Assert(cn >= 1 && cn <= 4); } - void random_roi() + void generateTestData() { Size roiSize = randomSize(1, MAX_VALUE); Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); @@ -195,13 +197,11 @@ PARAM_TEST_CASE(SplitTestBase, MatType, Channels, bool) } }; -typedef SplitTestBase Split; - OCL_TEST_P(Split, DISABLED_Accuracy) { for (int j = 0; j < test_loop_times; j++) { - random_roi(); + generateTestData(); OCL_OFF(cv::split(src_roi, dst_roi)); OCL_ON(cv::split(usrc_roi, udst_roi)); @@ -214,8 +214,150 @@ OCL_TEST_P(Split, DISABLED_Accuracy) } } -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())); +//////////////////////////////////////// MixChannels /////////////////////////////////////////////// + +PARAM_TEST_CASE(MixChannels, MatType, bool) +{ + int depth; + 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(dst1) + TEST_DECLARE_OUTPUT_PARAMETER(dst2) + TEST_DECLARE_OUTPUT_PARAMETER(dst3) + TEST_DECLARE_OUTPUT_PARAMETER(dst4) + + std::vector src_roi, dst_roi, dst; + std::vector usrc_roi, udst_roi, udst; + std::vector fromTo; + + virtual void SetUp() + { + depth = GET_PARAM(0); + use_roi = GET_PARAM(1); + } + + // generate number of channels and create type + int type() + { + int cn = randomInt(1, 5); + return CV_MAKE_TYPE(depth, cn); + } + + void generateTestData() + { + src_roi.clear(); + dst_roi.clear(); + dst.clear(); + usrc_roi.clear(); + udst_roi.clear(); + udst.clear(); + fromTo.clear(); + + Size roiSize = randomSize(1, MAX_VALUE); + + { + Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src1, src1_roi, roiSize, src1Border, type(), 2, 11); + + Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src2, src2_roi, roiSize, src2Border, type(), -1540, 1740); + + Border src3Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src3, src3_roi, roiSize, src3Border, type(), -1540, 1740); + + Border src4Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src4, src4_roi, roiSize, src4Border, type(), -1540, 1740); + } + + { + Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst1, dst1_roi, roiSize, dst1Border, type(), 2, 11); + + Border dst2Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst2, dst2_roi, roiSize, dst2Border, type(), -1540, 1740); + + Border dst3Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst3, dst3_roi, roiSize, dst3Border, type(), -1540, 1740); + + Border dst4Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst4, dst4_roi, roiSize, dst4Border, type(), -1540, 1740); + } + + UMAT_UPLOAD_INPUT_PARAMETER(src1) + UMAT_UPLOAD_INPUT_PARAMETER(src2) + UMAT_UPLOAD_INPUT_PARAMETER(src3) + UMAT_UPLOAD_INPUT_PARAMETER(src4) + + UMAT_UPLOAD_OUTPUT_PARAMETER(dst1) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst2) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst3) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst4) + + int nsrc = randomInt(1, 5), ndst = randomInt(1, 5); + + src_roi.push_back(src1_roi), usrc_roi.push_back(usrc1_roi); + if (nsrc >= 2) + src_roi.push_back(src2_roi), usrc_roi.push_back(usrc2_roi); + if (nsrc >= 3) + src_roi.push_back(src3_roi), usrc_roi.push_back(usrc3_roi); + if (nsrc >= 4) + src_roi.push_back(src4_roi), usrc_roi.push_back(usrc4_roi); + + dst_roi.push_back(dst1_roi), udst_roi.push_back(udst1_roi), + dst.push_back(dst1), udst.push_back(udst1); + if (ndst >= 2) + dst_roi.push_back(dst2_roi), udst_roi.push_back(udst2_roi), + dst.push_back(dst2), udst.push_back(udst2); + if (ndst >= 3) + dst_roi.push_back(dst3_roi), udst_roi.push_back(udst3_roi), + dst.push_back(dst3), udst.push_back(udst3); + if (ndst >= 4) + dst_roi.push_back(dst4_roi), udst_roi.push_back(udst4_roi), + dst.push_back(dst4), udst.push_back(udst4); + + int scntotal = 0, dcntotal = 0; + for (int i = 0; i < nsrc; ++i) + scntotal += src_roi[i].channels(); + for (int i = 0; i < ndst; ++i) + dcntotal += dst_roi[i].channels(); + + int npairs = randomInt(1, std::min(scntotal, dcntotal) + 1); + fromTo.resize(npairs << 1); + + for (int i = 0; i < npairs; ++i) + { + fromTo[i<<1] = randomInt(0, scntotal); + fromTo[(i<<1)+1] = randomInt(0, dcntotal); + } + } +}; + +OCL_TEST_P(MixChannels, Accuracy) +{ + for (int j = 0; j < test_loop_times + 10; j++) + { + generateTestData(); + + OCL_OFF(cv::mixChannels(src_roi, dst_roi, fromTo)); + OCL_ON(cv::mixChannels(usrc_roi, udst_roi, fromTo)); + + for (size_t i = 0, size = dst_roi.size(); i < size; ++i) + { + EXPECT_MAT_NEAR(dst[i], udst[i], 0.0); + EXPECT_MAT_NEAR(dst_roi[i], udst_roi[i], 0.0); + } + } +} + +//////////////////////////////////////// Instantiation /////////////////////////////////////////////// + +OCL_INSTANTIATE_TEST_CASE_P(Channels, Merge, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Channels, Split, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Channels, MixChannels, Combine(OCL_ALL_DEPTHS, Bool())); } } // namespace cvtest::ocl