class CV_EXPORTS MatExpr;
class CV_EXPORTS UMat;
-class CV_EXPORTS UMatExpr;
class CV_EXPORTS SparseMat;
typedef Mat MatND;
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();
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;
typedef Mat_<Vec3d> Mat3d;
typedef Mat_<Vec4d> Mat4d;
-
-class CV_EXPORTS UMatExpr;
-
class CV_EXPORTS UMat
{
public:
~UMat();
//! assignment operators
UMat& operator = (const UMat& m);
- UMat& operator = (const UMatExpr& expr);
Mat getMat(int flags) const;
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;
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.
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;
}
#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 \
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]); \
#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
#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);
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;
}
return sumSqrTab[depth];
}
+#ifdef HAVE_OPENCL
+
template <typename T> Scalar ocl_part_sum(Mat m)
{
CV_Assert(m.rows == 1);
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() )
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),
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. */
}
}
+//////////////////////////////// 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;
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),
--- /dev/null
+/*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
projects / profile / ivi / opencv.git / commitdiff