From 563dc5aa38b8bccfe91982a007f0f3fa4f5174a4 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 5 Feb 2014 19:10:02 +0400 Subject: [PATCH] implemented umat expressions --- modules/core/include/opencv2/core/base.hpp | 1 - modules/core/include/opencv2/core/mat.hpp | 30 +++---- modules/core/src/ocl.cpp | 5 +- modules/core/src/opencl/reduce.cl | 50 ++++++------ modules/core/src/stat.cpp | 6 +- modules/core/src/umatrix.cpp | 117 +++++++++++++++++++++++++++ modules/core/test/ocl/test_arithm.cpp | 18 +++++ modules/core/test/ocl/test_matrix_expr.cpp | 124 +++++++++++++++++++++++++++++ 8 files changed, 305 insertions(+), 46 deletions(-) create mode 100644 modules/core/test/ocl/test_matrix_expr.cpp diff --git a/modules/core/include/opencv2/core/base.hpp b/modules/core/include/opencv2/core/base.hpp index 6e78300..31cae39 100644 --- a/modules/core/include/opencv2/core/base.hpp +++ b/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; diff --git a/modules/core/include/opencv2/core/mat.hpp b/modules/core/include/opencv2/core/mat.hpp index 6409222..99e7694 100644 --- a/modules/core/include/opencv2/core/mat.hpp +++ b/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, - UEXPR =12 << KIND_SHIFT + STD_VECTOR_UMAT =11 << KIND_SHIFT }; _InputArray(); @@ -108,7 +107,6 @@ public: template _InputArray(const cudev::GpuMat_<_Tp>& m); _InputArray(const UMat& um); _InputArray(const std::vector& 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_ Mat2d; typedef Mat_ Mat3d; typedef Mat_ 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,11 +1216,11 @@ 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; @@ -1234,14 +1228,14 @@ public: double dot(InputArray m) const; //! Matlab-style matrix initialization - static UMatExpr zeros(int rows, int cols, int type); - static UMatExpr zeros(Size size, int type); - static UMatExpr zeros(int ndims, const int* sz, int type); - static UMatExpr ones(int rows, int cols, int type); - static UMatExpr ones(Size size, int type); - static UMatExpr ones(int ndims, const int* sz, int type); - static UMatExpr eye(int rows, int cols, int type); - static UMatExpr eye(Size size, int type); + static UMat zeros(int rows, int cols, int type); + static UMat zeros(Size size, int type); + static UMat zeros(int ndims, const int* sz, int type); + static UMat ones(int rows, int cols, int type); + static UMat ones(Size size, int type); + static UMat ones(int ndims, const int* sz, int type); + static UMat eye(int rows, int cols, 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. diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index ce4dc17..0bbfb62 100644 --- a/modules/core/src/ocl.cpp +++ b/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; } diff --git a/modules/core/src/opencl/reduce.cl b/modules/core/src/opencl/reduce.cl index 7a35314..febc1cb 100644 --- a/modules/core/src/opencl/reduce.cl +++ b/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 OP_SUM +#if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR || defined OP_DOT +#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 -#if defined (DEPTH_1) +#elif defined DEPTH_1 #define srcT char #define MIN_VAL -128 #define MAX_VAL 127 -#endif -#if defined (DEPTH_2) +#elif defined DEPTH_2 #define srcT ushort #define MIN_VAL 0 #define MAX_VAL 65535 -#endif -#if defined (DEPTH_3) +#elif defined DEPTH_3 #define srcT short #define MIN_VAL -32768 #define MAX_VAL 32767 -#endif -#if defined (DEPTH_4) +#elif defined DEPTH_4 #define srcT int #define MIN_VAL INT_MIN #define MAX_VAL INT_MAX -#endif -#if defined (DEPTH_5) +#elif defined DEPTH_5 #define srcT float #define MIN_VAL (-FLT_MAX) #define MAX_VAL FLT_MAX -#endif -#if defined (DEPTH_6) +#elif 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 -#if defined OP_MIN_MAX_LOC_MASK +#elif 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; } diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp index 63970c5..3042ab4 100644 --- a/modules/core/src/stat.cpp +++ b/modules/core/src/stat.cpp @@ -449,6 +449,8 @@ static SumSqrFunc getSumSqrTab(int depth) return sumSqrTab[depth]; } +#ifdef HAVE_OPENCL + template Scalar ocl_part_sum(Mat m) { CV_Assert(m.rows == 1); @@ -464,8 +466,6 @@ template 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), diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp index 3e4cfa2..bb9bbb4 100644 --- a/modules/core/src/umatrix.cpp +++ b/modules/core/src/umatrix.cpp @@ -798,6 +798,123 @@ 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; +} + +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; +} + +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. */ diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index bf29c4c..e6bcf4e 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/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(CV_8U, CV_32S), std::make_pair(CV_8U, CV_32F), diff --git a/modules/core/test/ocl/test_matrix_expr.cpp b/modules/core/test/ocl/test_matrix_expr.cpp new file mode 100644 index 0000000..33be862 --- /dev/null +++ b/modules/core/test/ocl/test_matrix_expr.cpp @@ -0,0 +1,124 @@ +/*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 { + +//////////////////////////////// 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 -- 2.7.4