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

Merge pull request #2459 from ilya-lavrenov:tapi_experiments

Browse Source
This commit is contained in:
Andrey Pavlenko
2014-03-13 16:49:56 +04:00
committed by OpenCV Buildbot
parent 6b6cfa893d f138b613b7
commit 0764a23eb2
12 changed files with 239 additions and 140 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
modules/imgproc/src/imgwarp.cpp
View File

@@ -4077,7 +4077,7 @@ static bool ocl_warpTransform(InputArray _src, OutputArray _dst, InputArray _M0,
matM.convertTo(M0, doubleSupport ? CV_64F : CV_32F);
k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrReadOnly(M0),
ocl::KernelArg(0, 0, 0, borderBuf, CV_ELEM_SIZE(sctype)));
ocl::KernelArg(0, 0, 0, 0, borderBuf, CV_ELEM_SIZE(sctype)));
size_t globalThreads[2] = { dst.cols, dst.rows };
return k.run(2, globalThreads, NULL, false);

16
modules/imgproc/src/opencl/threshold.cl
View File

@@ -53,29 +53,29 @@
__kernel void threshold(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
T thresh, T max_val)
T1 thresh, T1 max_val)
{
int gx = get_global_id(0);
int gy = get_global_id(1);
if (gx < cols && gy < rows)
{
int src_index = mad24(gy, src_step, src_offset + gx * (int)sizeof(T));
int dst_index = mad24(gy, dst_step, dst_offset + gx * (int)sizeof(T));
int src_index = mad24(gy, src_step, mad24(gx, (int)sizeof(T), src_offset));
int dst_index = mad24(gy, dst_step, mad24(gx, (int)sizeof(T), dst_offset));
T sdata = *(__global const T *)(srcptr + src_index);
__global T * dst = (__global T *)(dstptr + dst_index);
#ifdef THRESH_BINARY
dst[0] = sdata > thresh ? max_val : (T)(0);
dst[0] = sdata > (T)(thresh) ? (T)(max_val) : (T)(0);
#elif defined THRESH_BINARY_INV
dst[0] = sdata > thresh ? (T)(0) : max_val;
dst[0] = sdata > (T)(thresh) ? (T)(0) : (T)(max_val);
#elif defined THRESH_TRUNC
dst[0] = sdata > thresh ? thresh : sdata;
dst[0] = sdata > (T)(thresh) ? (T)(thresh) : sdata;
#elif defined THRESH_TOZERO
dst[0] = sdata > thresh ? sdata : (T)(0);
dst[0] = sdata > (T)(thresh) ? sdata : (T)(0);
#elif defined THRESH_TOZERO_INV
dst[0] = sdata > thresh ? (T)(0) : sdata;
dst[0] = sdata > (T)(thresh) ? (T)(0) : sdata;
#endif
}
}

16
modules/imgproc/src/thresh.cpp
View File

@@ -710,7 +710,8 @@ private:
static bool ocl_threshold( InputArray _src, OutputArray _dst, double & thresh, double maxval, int thresh_type )
{
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), ktype = CV_MAKE_TYPE(depth, 1);
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
kercn = ocl::predictOptimalVectorWidth(_src, _dst), ktype = CV_MAKE_TYPE(depth, kercn);
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
if ( !(thresh_type == THRESH_BINARY || thresh_type == THRESH_BINARY_INV || thresh_type == THRESH_TRUNC ||
@@ -721,8 +722,9 @@ static bool ocl_threshold( InputArray _src, OutputArray _dst, double & thresh, d
const char * const thresholdMap[] = { "THRESH_BINARY", "THRESH_BINARY_INV", "THRESH_TRUNC",
"THRESH_TOZERO", "THRESH_TOZERO_INV" };
ocl::Kernel k("threshold", ocl::imgproc::threshold_oclsrc,
format("-D %s -D T=%s%s", thresholdMap[thresh_type],
ocl::typeToStr(ktype), doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
format("-D %s -D T=%s -D T1=%s%s", thresholdMap[thresh_type],
ocl::typeToStr(ktype), ocl::typeToStr(depth),
doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
if (k.empty())
return false;
@@ -733,11 +735,11 @@ static bool ocl_threshold( InputArray _src, OutputArray _dst, double & thresh, d
if (depth <= CV_32S)
thresh = cvFloor(thresh);
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn),
ocl::KernelArg::Constant(Mat(1, 1, ktype, thresh)),
ocl::KernelArg::Constant(Mat(1, 1, ktype, maxval)));
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn, kercn),
ocl::KernelArg::Constant(Mat(1, 1, depth, Scalar::all(thresh))),
ocl::KernelArg::Constant(Mat(1, 1, depth, Scalar::all(maxval))));
size_t globalsize[2] = { dst.cols * cn, dst.rows };
size_t globalsize[2] = { dst.cols * cn / kercn, dst.rows };
return k.run(2, globalsize, NULL, false);
}