Merge pull request #2289 from ilya-lavrenov:umat_expr

modules/core/include/opencv2/core/base.hpp

@@ -502,7 +502,6 @@ class CV_EXPORTS Mat;
 class CV_EXPORTS MatExpr;
 
 class CV_EXPORTS UMat;
-class CV_EXPORTS UMatExpr;
 
 class CV_EXPORTS SparseMat;
 typedef Mat MatND;

modules/core/include/opencv2/core/mat.hpp

@@ -86,8 +86,7 @@ public:
         CUDA_MEM          = 8 << KIND_SHIFT,
         GPU_MAT           = 9 << KIND_SHIFT,
         UMAT              =10 << KIND_SHIFT,
-        STD_VECTOR_UMAT   =11 << KIND_SHIFT,
+        STD_VECTOR_UMAT   =11 << KIND_SHIFT
-        UEXPR             =12 << KIND_SHIFT
     };
 
     _InputArray();
@@ -108,7 +107,6 @@ public:
     template<typename _Tp> _InputArray(const cudev::GpuMat_<_Tp>& m);
     _InputArray(const UMat& um);
     _InputArray(const std::vector<UMat>& umv);
-    _InputArray(const UMatExpr& uexpr);
 
     virtual Mat getMat(int idx=-1) const;
     virtual UMat getUMat(int idx=-1) const;
@@ -1134,9 +1132,6 @@ typedef Mat_<Vec2d> Mat2d;
 typedef Mat_<Vec3d> Mat3d;
 typedef Mat_<Vec4d> Mat4d;
 
-
-class CV_EXPORTS UMatExpr;
-
 class CV_EXPORTS UMat
 {
 public:
@@ -1178,7 +1173,6 @@ public:
     ~UMat();
     //! assignment operators
     UMat& operator = (const UMat& m);
-    UMat& operator = (const UMatExpr& expr);
 
     Mat getMat(int flags) const;
 
@@ -1222,26 +1216,24 @@ public:
     UMat reshape(int cn, int newndims, const int* newsz) const;
 
     //! matrix transposition by means of matrix expressions
-    UMatExpr t() const;
+    UMat t() const;
     //! matrix inversion by means of matrix expressions
-    UMatExpr inv(int method=DECOMP_LU) const;
+    UMat inv(int method=DECOMP_LU) const;
     //! per-element matrix multiplication by means of matrix expressions
-    UMatExpr mul(InputArray m, double scale=1) const;
+    UMat mul(InputArray m, double scale=1) const;
 
-    //! computes cross-product of 2 3D vectors
-    UMat cross(InputArray m) const;
     //! computes dot-product
     double dot(InputArray m) const;
 
     //! Matlab-style matrix initialization
-    static UMatExpr zeros(int rows, int cols, int type);
+    static UMat zeros(int rows, int cols, int type);
-    static UMatExpr zeros(Size size, int type);
+    static UMat zeros(Size size, int type);
-    static UMatExpr zeros(int ndims, const int* sz, int type);
+    static UMat zeros(int ndims, const int* sz, int type);
-    static UMatExpr ones(int rows, int cols, int type);
+    static UMat ones(int rows, int cols, int type);
-    static UMatExpr ones(Size size, int type);
+    static UMat ones(Size size, int type);
-    static UMatExpr ones(int ndims, const int* sz, int type);
+    static UMat ones(int ndims, const int* sz, int type);
-    static UMatExpr eye(int rows, int cols, int type);
+    static UMat eye(int rows, int cols, int type);
-    static UMatExpr eye(Size size, int type);
+    static UMat eye(Size size, int type);
 
     //! allocates new matrix data unless the matrix already has specified size and type.
     // previous data is unreferenced if needed.

modules/core/perf/opencl/perf_arithm.cpp

@@ -706,6 +706,28 @@ OCL_PERF_TEST_P(NormFixture, Norm,
     SANITY_CHECK(res, 1e-5, ERROR_RELATIVE);
 }
 
+///////////// UMat::dot ////////////////////////
+
+typedef Size_MatType UMatDotFixture;
+
+OCL_PERF_TEST_P(UMatDotFixture, UMatDot,
+            ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), OCL_TEST_TYPES))
+{
+    const Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int type = get<1>(params);
+    double r = 0.0;
+
+    checkDeviceMaxMemoryAllocSize(srcSize, type);
+
+    UMat src1(srcSize, type), src2(srcSize, type);
+    declare.in(src1, src2, WARMUP_RNG);
+
+    OCL_TEST_CYCLE() r = src1.dot(src2);
+
+    SANITY_CHECK(r, 1e-5, ERROR_RELATIVE);
+}
+
 ///////////// Repeat ////////////////////////
 
 typedef Size_MatType RepeatFixture;

modules/core/src/ocl.cpp

@@ -2800,7 +2800,10 @@ int Kernel::set(int i, const void* value, size_t sz)
     CV_Assert(i >= 0);
     if( i == 0 )
         p->cleanupUMats();
-    if( clSetKernelArg(p->handle, (cl_uint)i, sz, value) < 0 )
+
+    cl_int retval = clSetKernelArg(p->handle, (cl_uint)i, sz, value);
+    CV_OclDbgAssert(retval == CL_SUCCESS);
+    if (retval != CL_SUCCESS)
         return -1;
     return i+1;
 }

modules/core/src/opencl/reduce.cl

@@ -58,12 +58,14 @@
 #define EXTRA_PARAMS
 #endif
 
-#if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR
+#if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR || defined OP_DOT
-#if OP_SUM
+#ifdef OP_DOT
+#define FUNC(a, b, c) a += b * c
+#elif defined OP_SUM
 #define FUNC(a, b) a += b
-#elif OP_SUM_ABS
+#elif defined OP_SUM_ABS
 #define FUNC(a, b) a += b >= (dstT)(0) ? b : -b
-#elif OP_SUM_SQR
+#elif defined OP_SUM_SQR
 #define FUNC(a, b) a += b * b
 #endif
 #define DECLARE_LOCAL_MEM \
@@ -76,6 +78,12 @@
     int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols)); \
     if (mask[mask_index]) \
         FUNC(accumulator, temp)
+#elif defined OP_DOT
+#define REDUCE_GLOBAL \
+    int src2_index = mad24(id / cols, src2_step, src2_offset + (id % cols) * (int)sizeof(srcT)); \
+    __global const srcT * src2 = (__global const srcT *)(src2ptr + src2_index); \
+    dstT temp = convertToDT(src[0]), temp2 = convertToDT(src2[0]); \
+    FUNC(accumulator, temp, temp2)
 #else
 #define REDUCE_GLOBAL \
     dstT temp = convertToDT(src[0]); \
@@ -112,37 +120,31 @@
 
 #elif defined OP_MIN_MAX_LOC || defined OP_MIN_MAX_LOC_MASK
 
-#if defined (DEPTH_0)
+#ifdef DEPTH_0
 #define srcT uchar
 #define MIN_VAL 0
 #define MAX_VAL 255
-#endif
+#elif defined DEPTH_1
-#if defined (DEPTH_1)
 #define srcT char
 #define MIN_VAL -128
 #define MAX_VAL 127
-#endif
+#elif defined DEPTH_2
-#if defined (DEPTH_2)
 #define srcT ushort
 #define MIN_VAL 0
 #define MAX_VAL 65535
-#endif
+#elif defined DEPTH_3
-#if defined (DEPTH_3)
 #define srcT short
 #define MIN_VAL -32768
 #define MAX_VAL 32767
-#endif
+#elif defined DEPTH_4
-#if defined (DEPTH_4)
 #define srcT int
 #define MIN_VAL INT_MIN
 #define MAX_VAL INT_MAX
-#endif
+#elif defined DEPTH_5
-#if defined (DEPTH_5)
 #define srcT float
 #define MIN_VAL (-FLT_MAX)
 #define MAX_VAL FLT_MAX
-#endif
+#elif defined DEPTH_6
-#if defined (DEPTH_6)
 #define srcT double
 #define MIN_VAL (-DBL_MAX)
 #define MAX_VAL DBL_MAX
@@ -233,17 +235,19 @@
 #error "No operation"
 #endif
 
-#if defined OP_MIN_MAX_LOC
+#ifdef OP_MIN_MAX_LOC
 #undef EXTRA_PARAMS
 #define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2
-#endif
+#elif defined OP_MIN_MAX_LOC_MASK
-#if defined OP_MIN_MAX_LOC_MASK
 #undef EXTRA_PARAMS
 #define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2, \
-    __global const uchar * maskptr, int mask_step, int mask_offset, __global int * test
+    __global const uchar * maskptr, int mask_step, int mask_offset
+#elif defined OP_DOT
+#undef EXTRA_PARAMS
+#define EXTRA_PARAMS , __global uchar * src2ptr, int src2_step, int src2_offset
 #endif
 
-__kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols,
+__kernel void reduce(__global const uchar * srcptr, int src_step, int src_offset, int cols,
                      int total, int groupnum, __global uchar * dstptr EXTRA_PARAMS)
 {
     int lid = get_local_id(0);
@@ -255,7 +259,7 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
 
     for (int grain = groupnum * WGS; id < total; id += grain)
     {
-        int src_index = mad24(id / cols, step, offset + (id % cols) * (int)sizeof(srcT));
+        int src_index = mad24(id / cols, src_step, src_offset + (id % cols) * (int)sizeof(srcT));
         __global const srcT * src = (__global const srcT *)(srcptr + src_index);
         REDUCE_GLOBAL;
     }

modules/core/src/stat.cpp

@@ -449,6 +449,8 @@ static SumSqrFunc getSumSqrTab(int depth)
     return sumSqrTab[depth];
 }
 
+#ifdef HAVE_OPENCL
+
 template <typename T> Scalar ocl_part_sum(Mat m)
 {
     CV_Assert(m.rows == 1);
@@ -464,8 +466,6 @@ template <typename T> Scalar ocl_part_sum(Mat m)
     return s;
 }
 
-#ifdef HAVE_OPENCL
-
 enum { OCL_OP_SUM = 0, OCL_OP_SUM_ABS =  1, OCL_OP_SUM_SQR = 2 };
 
 static bool ocl_sum( InputArray _src, Scalar & res, int sum_op, InputArray _mask = noArray() )
@@ -1279,7 +1279,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int*
             ocl::KernelArg::PtrWriteOnly(minloc), ocl::KernelArg::PtrWriteOnly(maxloc), ocl::KernelArg::ReadOnlyNoSize(mask));
 
     size_t globalsize = groupnum * wgs;
-    if (!k.run(1, &globalsize, &wgs, true))
+    if (!k.run(1, &globalsize, &wgs, false))
         return false;
 
     Mat minv = minval.getMat(ACCESS_READ), maxv = maxval.getMat(ACCESS_READ),

modules/core/src/umatrix.cpp

@@ -551,14 +551,6 @@ int UMat::checkVector(int _elemChannels, int _depth, bool _requireContinuous) co
     ? (int)(total()*channels()/_elemChannels) : -1;
 }
 
-
-UMat UMat::cross(InputArray) const
-{
-    CV_Error(CV_StsNotImplemented, "");
-    return UMat();
-}
-
-
 UMat UMat::reshape(int _cn, int _newndims, const int* _newsz) const
 {
     if(_newndims == dims)
@@ -798,6 +790,127 @@ UMat& UMat::operator = (const Scalar& s)
     return *this;
 }
 
+UMat UMat::t() const
+{
+    UMat m;
+    transpose(*this, m);
+    return m;
+}
+
+UMat UMat::inv(int method) const
+{
+    UMat m;
+    invert(*this, m, method);
+    return m;
+}
+
+UMat UMat::mul(InputArray m, double scale) const
+{
+    UMat dst;
+    multiply(*this, m, dst, scale);
+    return dst;
+}
+
+#ifdef HAVE_OPENCL
+
+static bool ocl_dot( InputArray _src1, InputArray _src2, double & res )
+{
+    int type = _src1.type(), depth = CV_MAT_DEPTH(type);
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+    if ( !doubleSupport && depth == CV_64F )
+        return false;
+
+    int dbsize = ocl::Device::getDefault().maxComputeUnits();
+    size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+    int ddepth = std::max(CV_32F, depth);
+
+    int wgs2_aligned = 1;
+    while (wgs2_aligned < (int)wgs)
+        wgs2_aligned <<= 1;
+    wgs2_aligned >>= 1;
+
+    char cvt[40];
+    ocl::Kernel k("reduce", ocl::core::reduce_oclsrc,
+                  format("-D srcT=%s -D dstT=%s -D convertToDT=%s -D OP_DOT -D WGS=%d -D WGS2_ALIGNED=%d%s",
+                         ocl::typeToStr(depth), ocl::typeToStr(ddepth), ocl::convertTypeStr(depth, ddepth, 1, cvt),
+                         (int)wgs, wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    UMat src1 = _src1.getUMat().reshape(1), src2 = _src2.getUMat().reshape(1), db(1, dbsize, ddepth);
+
+    ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1),
+            src2arg = ocl::KernelArg::ReadOnlyNoSize(src2),
+            dbarg = ocl::KernelArg::PtrWriteOnly(db);
+
+    k.args(src1arg, src1.cols, (int)src1.total(), dbsize, dbarg, src2arg);
+
+    size_t globalsize = dbsize * wgs;
+    if (k.run(1, &globalsize, &wgs, false))
+    {
+        res = sum(db.getMat(ACCESS_READ))[0];
+        return true;
+    }
+    return false;
+}
+
+#endif
+
+double UMat::dot(InputArray m) const
+{
+    CV_Assert(m.sameSize(*this) && m.type() == type());
+
+#ifdef HAVE_OPENCL
+    double r = 0;
+    CV_OCL_RUN_(dims <= 2, ocl_dot(*this, m, r), r)
+#endif
+
+    return getMat(ACCESS_READ).dot(m);
+}
+
+UMat UMat::zeros(int rows, int cols, int type)
+{
+    return UMat(rows, cols, type, Scalar::all(0));
+}
+
+UMat UMat::zeros(Size size, int type)
+{
+    return UMat(size, type, Scalar::all(0));
+}
+
+UMat UMat::zeros(int ndims, const int* sz, int type)
+{
+    return UMat(ndims, sz, type, Scalar::all(0));
+}
+
+UMat UMat::ones(int rows, int cols, int type)
+{
+    return UMat::ones(Size(cols, rows), type);
+}
+
+UMat UMat::ones(Size size, int type)
+{
+    return UMat(size, type, Scalar(1));
+}
+
+UMat UMat::ones(int ndims, const int* sz, int type)
+{
+    return UMat(ndims, sz, type, Scalar(1));
+}
+
+UMat UMat::eye(int rows, int cols, int type)
+{
+    return UMat::eye(Size(cols, rows), type);
+}
+
+UMat UMat::eye(Size size, int type)
+{
+    UMat m(size, type);
+    setIdentity(m);
+    return m;
+}
+
 }
 
 /* End of file. */

modules/core/test/ocl/test_arithm.cpp

@@ -1337,6 +1337,23 @@ OCL_TEST_P(Norm, NORM_L2_2args_mask)
         }
 }
 
+//////////////////////////////// UMat::dot ////////////////////////////////////////////////
+
+typedef ArithmTestBase UMatDot;
+
+OCL_TEST_P(UMatDot, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(const double cpuRes = src1_roi.dot(src2_roi));
+        OCL_ON(const double gpuRes = usrc1_roi.dot(usrc2_roi));
+
+        EXPECT_PRED3(relativeError, cpuRes, gpuRes, 1e-6);
+    }
+}
+
 //////////////////////////////// Sqrt ////////////////////////////////////////////////
 
 typedef ArithmTestBase Sqrt;
@@ -1708,6 +1725,7 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, ConvertScaleAbs, Combine(OCL_ALL_DEPTHS, OCL
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, ScaleAdd, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, PatchNaNs, Combine(OCL_ALL_CHANNELS, Bool()));
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, Psnr, Combine(::testing::Values((MatDepth)CV_8U), OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, UMatDot, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Arithm, ReduceSum, Combine(testing::Values(std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32S),
                                                                        std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32F),

modules/core/test/ocl/test_matrix_expr.cpp

@@ -0,0 +1,85 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+// Copyright (C) 2014, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#include "test_precomp.hpp"
+#include "opencv2/ts/ocl_test.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+//////////////////////////////// UMat Expressions /////////////////////////////////////////////////
+
+PARAM_TEST_CASE(UMatExpr, MatDepth, Channels)
+{
+    int type;
+    Size size;
+
+    virtual void SetUp()
+    {
+        type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
+    }
+
+    void generateTestData()
+    {
+        size = randomSize(1, MAX_VALUE);
+    }
+};
+
+//////////////////////////////// UMat::eye /////////////////////////////////////////////////
+
+OCL_TEST_P(UMatExpr, Eye)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Mat m = Mat::eye(size, type);
+        UMat um = UMat::eye(size, type);
+
+        EXPECT_MAT_NEAR(m, um, 0);
+    }
+}
+
+//////////////////////////////// UMat::zeros /////////////////////////////////////////////////
+
+OCL_TEST_P(UMatExpr, Zeros)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Mat m = Mat::zeros(size, type);
+        UMat um = UMat::zeros(size, type);
+
+        EXPECT_MAT_NEAR(m, um, 0);
+    }
+}
+
+//////////////////////////////// UMat::ones /////////////////////////////////////////////////
+
+OCL_TEST_P(UMatExpr, Ones)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        Mat m = Mat::ones(size, type);
+        UMat um = UMat::ones(size, type);
+
+        EXPECT_MAT_NEAR(m, um, 0);
+    }
+}
+
+//////////////////////////////// Instantiation /////////////////////////////////////////////////
+
+OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, UMatExpr, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS));
+
+} } // namespace cvtest::ocl
+
+#endif

modules/superres/src/input_array_utility.cpp

@@ -223,7 +223,6 @@ namespace
             break;
 
         case _InputArray::UMAT:
-        case _InputArray::UEXPR:
             src.getUMat().convertTo(dst, depth, scale);
             break;