added cv::mixChannels to T-API

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<UMat> & 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<UMat> & src = *(const std::vector<UMat> *)_src.getObj();
+    std::vector<UMat> & dst = *(std::vector<UMat> *)_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<UMat> 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 &&
@@ -641,6 +734,14 @@ void cv::mixChannels(InputArrayOfArrays src, InputOutputArrayOfArrays dst,
 {
     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 &&

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
+    }
+}

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()));
+//////////////////////////////////////// MixChannels ///////////////////////////////////////////////
-OCL_INSTANTIATE_TEST_CASE_P(SplitMerge, Split, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+
+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<Mat> src_roi, dst_roi, dst;
+    std::vector<UMat> usrc_roi, udst_roi, udst;
+    std::vector<int> 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