generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

implemented umat expressions

Browse Source
This commit is contained in:
Ilya Lavrenov
2014-02-05 19:10:02 +04:00
parent 9493a4ec4e
commit 563dc5aa38
8 changed files with 305 additions and 46 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
modules/core/src/ocl.cpp
View File

@@ -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;
}

50
modules/core/src/opencl/reduce.cl
View File

@@ -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;
}

6
modules/core/src/stat.cpp
View File

@@ -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),

117
modules/core/src/umatrix.cpp
View File

@@ -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. */