Merge pull request #1988 from ilya-lavrenov:tapi_umat_convertTo

modules/core/src/opencl/convert.cl

+/*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 DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+#define noconvert
+
+__kernel void convertTo(__global const uchar * srcptr, int src_step, int src_offset,
+                        __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+                        float alpha, float beta )
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < dst_cols && y < dst_rows)
+    {
+        int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(srcT) );
+        int dst_index = mad24(y, dst_step, dst_offset + x * (int)sizeof(dstT) );
+
+        __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+
+        dst[0] = convertToDT( src[0] * alpha + beta );
+    }
+}

modules/core/src/umatrix.cpp

@@ -658,16 +658,56 @@ void UMat::copyTo(OutputArray _dst) const
     }
 }
 
-void UMat::convertTo(OutputArray, int, double, double) const
+void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const
 {
-    CV_Error(Error::StsNotImplemented, "");
+    bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON;
+    int stype = type(), cn = CV_MAT_CN(stype);
+
+    if( _type < 0 )
+        _type = _dst.fixedType() ? _dst.type() : stype;
+    else
+        _type = CV_MAKETYPE(CV_MAT_DEPTH(_type), cn);
+
+    int sdepth = CV_MAT_DEPTH(stype), ddepth = CV_MAT_DEPTH(_type);
+    if( sdepth == ddepth && noScale )
+    {
+        copyTo(_dst);
+        return;
+    }
+
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+    bool needDouble = sdepth == CV_64F || ddepth == CV_64F;
+    if( dims <= 2 && cn && _dst.isUMat() && ocl::useOpenCL() &&
+            ((needDouble && doubleSupport) || !needDouble) )
+    {
+        char cvt[40];
+        ocl::Kernel k("convertTo", ocl::core::convert_oclsrc,
+                      format("-D srcT=%s -D dstT=%s -D convertToDT=%s%s", ocl::typeToStr(sdepth),
+                             ocl::typeToStr(ddepth), ocl::convertTypeStr(CV_32F, ddepth, 1, cvt),
+                             doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+        if (!k.empty())
+        {
+            _dst.create( size(), _type );
+            UMat dst = _dst.getUMat();
+
+            float alphaf = (float)alpha, betaf = (float)beta;
+            k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::WriteOnly(dst, cn), alphaf, betaf);
+
+            size_t globalsize[2] = { dst.cols * cn, dst.rows };
+            if (k.run(2, globalsize, NULL, false))
+                return;
+        }
+    }
+
+    Mat m = getMat(ACCESS_READ);
+    m.convertTo(_dst, _type, alpha, beta);
 }
 
 UMat& UMat::setTo(InputArray _value, InputArray _mask)
 {
     bool haveMask = !_mask.empty();
     int tp = type(), cn = CV_MAT_CN(tp);
-    if( dims <= 2 && cn <= 4 && ocl::useOpenCL() )
+    if( dims <= 2 && cn <= 4 && cn != 3 && ocl::useOpenCL() )
     {
         Mat value = _value.getMat();
         CV_Assert( checkScalar(value, type(), _value.kind(), _InputArray::UMAT) );
@@ -707,9 +747,9 @@ UMat& UMat::setTo(InputArray _value, InputArray _mask)
     return *this;
 }
 
-UMat& UMat::operator = (const Scalar&)
+UMat& UMat::operator = (const Scalar& s)
 {
-    CV_Error(Error::StsNotImplemented, "");
+    setTo(s);
     return *this;
 }
 

modules/core/test/ocl/test_matrix_operation.cpp

+/*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 {
+
+////////////////////////////////converto/////////////////////////////////////////////////
+
+PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool)
+{
+    int src_depth, cn, dstType;
+    bool use_roi;
+
+    TEST_DECLARE_INPUT_PARAMETER(src)
+    TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+    virtual void SetUp()
+    {
+        src_depth = GET_PARAM(0);
+        cn = GET_PARAM(2);
+        dstType = CV_MAKE_TYPE(GET_PARAM(1), cn);
+
+        use_roi = GET_PARAM(3);
+    }
+
+    virtual void generateTestData()
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, CV_MAKE_TYPE(src_depth, cn), -MAX_VALUE, MAX_VALUE);
+
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+    }
+};
+
+typedef MatrixTestBase ConvertTo;
+
+OCL_TEST_P(ConvertTo, Accuracy)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        double alpha = randomDouble(-4, 4), beta = randomDouble(-4, 4);
+
+        OCL_OFF(src_roi.convertTo(dst_roi, dstType, alpha, beta));
+        OCL_ON(usrc_roi.convertTo(udst_roi, dstType, alpha, beta));
+
+        double eps = src_depth >= CV_32F || CV_MAT_DEPTH(dstType) >= CV_32F ? 1e-4 : 1;
+        OCL_EXPECT_MATS_NEAR(dst, eps);
+    }
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine(
+                            OCL_ALL_DEPTHS, OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+
+} } // namespace cvtest::ocl
+
+#endif

modules/core/test/test_umat.cpp

@@ -43,16 +43,13 @@
 
 #include <string>
 #include <iostream>
-#include <fstream>
-#include <iterator>
-#include <limits>
-#include <numeric>
 #include "opencv2/core/ocl.hpp"
 
 using namespace cv;
 using namespace std;
 
-class CV_UMatTest : public cvtest::BaseTest
+class CV_UMatTest :
+        public cvtest::BaseTest
 {
 public:
     CV_UMatTest() {}
@@ -204,37 +201,37 @@ TEST(Core_UMat, base) { CV_UMatTest test; test.safe_run(); }
 TEST(Core_UMat, getUMat)
 {
     {
-    int a[3] = { 1, 2, 3 };
-    Mat m = Mat(1, 1, CV_32SC3, a);
-    UMat u = m.getUMat(ACCESS_READ);
-    EXPECT_NE((void*)NULL, u.u);
+        int a[3] = { 1, 2, 3 };
+        Mat m = Mat(1, 1, CV_32SC3, a);
+        UMat u = m.getUMat(ACCESS_READ);
+        EXPECT_NE((void*)NULL, u.u);
     }
 
     {
-    Mat m(10, 10, CV_8UC1), ref;
-    for (int y = 0; y < m.rows; ++y)
-    {
-        uchar * const ptr = m.ptr<uchar>(y);
-        for (int x = 0; x < m.cols; ++x)
-            ptr[x] = (uchar)(x + y * 2);
-    }
+        Mat m(10, 10, CV_8UC1), ref;
+        for (int y = 0; y < m.rows; ++y)
+        {
+            uchar * const ptr = m.ptr<uchar>(y);
+            for (int x = 0; x < m.cols; ++x)
+                ptr[x] = (uchar)(x + y * 2);
+        }
 
-    ref = m.clone();
-    Rect r(1, 1, 8, 8);
-    ref(r).setTo(17);
+        ref = m.clone();
+        Rect r(1, 1, 8, 8);
+        ref(r).setTo(17);
 
-    {
-        UMat um = m(r).getUMat(ACCESS_WRITE);
-        um.setTo(17);
-    }
+        {
+            UMat um = m(r).getUMat(ACCESS_WRITE);
+            um.setTo(17);
+        }
 
-    double err = norm(m, ref, NORM_INF);
-    if(err > 0)
-    {
-        std::cout << "m: " << m << std::endl;
-        std::cout << "ref: " << ref << std::endl;
-    }
-    EXPECT_EQ(err, 0.);
+        double err = norm(m, ref, NORM_INF);
+        if (err > 0)
+        {
+            std::cout << "m: " << std::endl << m << std::endl;
+            std::cout << "ref: " << std::endl << ref << std::endl;
+        }
+        EXPECT_EQ(0., err);
     }
 }