mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-18 22:44:02 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

added cv::threshold to T-API

Browse Source
This commit is contained in:
Ilya Lavrenov 2013-12-03 00:13:05 +04:00
parent 474fc887a6
commit e2b99a32b2
5 changed files with 63 additions and 80 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1893,7 +1893,7 @@ Context2& Context2::getDefault()
// First, try to retrieve existing context of the same type.
// In its turn, Platform::getContext() may call Context2::create()
// if there is no such context.
ctx.create(Device::TYPE_CPU);
ctx.create(Device::TYPE_ACCELERATOR);
if(!ctx.p)
ctx.create(Device::TYPE_DGPU);
if(!ctx.p)

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

@ -51,86 +51,31 @@
#endif
#endif
#ifdef VECTORIZED
__kernel void threshold(__global const T * restrict src, int src_offset, int src_step,
__global T * dst, int dst_offset, int dst_step,
T thresh, T max_val, int max_index, int rows, int cols)
__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)
{
int gx = get_global_id(0);
int gy = get_global_id(1);
if (gx < cols && gy < rows)
{
gx *= VECSIZE;
int src_index = mad24(gy, src_step, src_offset + gx);
int dst_index = mad24(gy, dst_step, dst_offset + gx);
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));
#ifdef SRC_ALIGNED
VT sdata = *((__global VT *)(src + src_index));
#else
VT sdata = VLOADN(0, src + src_index);
#endif
VT vthresh = (VT)(thresh);
T sdata = *(__global const T *)(srcptr + src_index);
__global T * dst = (__global T *)(dstptr + dst_index);
#ifdef THRESH_BINARY
VT vecValue = sdata > vthresh ? max_val : (VT)(0);
dst[0] = sdata > thresh ? max_val : (T)(0);
#elif defined THRESH_BINARY_INV
VT vecValue = sdata > vthresh ? (VT)(0) : max_val;
dst[0] = sdata > thresh ? (T)(0) : max_val;
#elif defined THRESH_TRUNC
VT vecValue = sdata > vthresh ? thresh : sdata;
dst[0] = sdata > thresh ? thresh : sdata;
#elif defined THRESH_TOZERO
VT vecValue = sdata > vthresh ? sdata : (VT)(0);
dst[0] = sdata > thresh ? sdata : (T)(0);
#elif defined THRESH_TOZERO_INV
VT vecValue = sdata > vthresh ? (VT)(0) : sdata;
#endif
if (gx + VECSIZE <= max_index)
#ifdef DST_ALIGNED
*(__global VT*)(dst + dst_index) = vecValue;
#else
VSTOREN(vecValue, 0, dst + dst_index);
#endif
else
{
__attribute__(( aligned(sizeof(VT)) )) T array[VECSIZE];
*((VT*)array) = vecValue;
#pragma unroll
for (int i = 0; i < VECSIZE; ++i)
if (gx + i < max_index)
dst[dst_index + i] = array[i];
}
}
}
#else
__kernel void threshold(__global const T * restrict src, int src_offset, int src_step,
__global T * dst, int dst_offset, int dst_step,
T thresh, T max_val, int rows, int cols)
{
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 dst_index = mad24(gy, dst_step, dst_offset + gx);
T sdata = src[src_index];
#ifdef THRESH_BINARY
dst[dst_index] = sdata > thresh ? max_val : (T)(0);
#elif defined THRESH_BINARY_INV
dst[dst_index] = sdata > thresh ? (T)(0) : max_val;
#elif defined THRESH_TRUNC
dst[dst_index] = sdata > thresh ? thresh : sdata;
#elif defined THRESH_TOZERO
dst[dst_index] = sdata > thresh ? sdata : (T)(0);
#elif defined THRESH_TOZERO_INV
dst[dst_index] = sdata > thresh ? (T)(0) : sdata;
dst[0] = sdata > thresh ? (T)(0) : sdata;
#endif
}
}
#endif

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

@ -41,6 +41,7 @@
//M*/
#include "precomp.hpp"
#include "opencl_kernels.hpp"
namespace cv
{
@ -705,10 +706,47 @@ private:
int thresholdType;
};
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);
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
if ( !(thresh_type == THRESH_BINARY || thresh_type == THRESH_BINARY_INV || thresh_type == THRESH_TRUNC ||
thresh_type == THRESH_TOZERO || thresh_type == THRESH_TOZERO_INV) ||
(!doubleSupport && depth == CV_64F))
return false;
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" : ""));
if (k.empty())
return false;
UMat src = _src.getUMat();
_dst.create(src.size(), type);
UMat dst = _dst.getUMat();
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)));
size_t globalsize[2] = { dst.cols * cn, dst.rows };
return k.run(2, globalsize, NULL, false);
}
}
double cv::threshold( InputArray _src, OutputArray _dst, double thresh, double maxval, int type )
{
if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() &&
ocl_threshold(_src, _dst, thresh, maxval, type))
return thresh;
Mat src = _src.getMat();
bool use_otsu = (type & THRESH_OTSU) != 0;
type &= THRESH_MASK;

22
modules/imgproc/test/ocl/test_warp.cpp
View File

@ -224,7 +224,7 @@ OCL_TEST_P(Resize, Mat)
/////////////////////////////////////////////////////////////////////////////////////////////////
// remap
PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair<MatType, MatType>, Border, bool)
PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair<MatType, MatType>, BorderType, bool)
{
int srcType, map1Type, map2Type;
int borderType;
@ -349,11 +349,11 @@ OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_LINEAR, Combine(
Values(std::pair<MatType, MatType>((MatType)CV_32FC1, (MatType)CV_32FC1),
std::pair<MatType, MatType>((MatType)CV_16SC2, (MatType)CV_16UC1),
std::pair<MatType, MatType>((MatType)CV_32FC2, noType)),
Values((Border)BORDER_CONSTANT,
(Border)BORDER_REPLICATE,
(Border)BORDER_WRAP,
(Border)BORDER_REFLECT,
(Border)BORDER_REFLECT_101),
Values((BorderType)BORDER_CONSTANT,
(BorderType)BORDER_REPLICATE,
(BorderType)BORDER_WRAP,
(BorderType)BORDER_REFLECT,
(BorderType)BORDER_REFLECT_101),
Bool()));
OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine(
@ -363,11 +363,11 @@ OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine(
std::pair<MatType, MatType>((MatType)CV_32FC2, noType),
std::pair<MatType, MatType>((MatType)CV_16SC2, (MatType)CV_16UC1),
std::pair<MatType, MatType>((MatType)CV_16SC2, noType)),
Values((Border)BORDER_CONSTANT,
(Border)BORDER_REPLICATE,
(Border)BORDER_WRAP,
(Border)BORDER_REFLECT,
(Border)BORDER_REFLECT_101),
Values((BorderType)BORDER_CONSTANT,
(BorderType)BORDER_REPLICATE,
(BorderType)BORDER_WRAP,
(BorderType)BORDER_REFLECT,
(BorderType)BORDER_REFLECT_101),
Bool()));
} } // namespace cvtest::ocl

2
modules/ts/include/opencv2/ts/ocl_test.hpp
View File

@ -264,7 +264,7 @@ struct CV_EXPORTS TestUtils
#define UMAT_UPLOAD_INPUT_PARAMETER(name) \
{ \
name.copyTo(u ## name); \
Size wholeSize; Point ofs; name ## _roi.locateROI(wholeSize, ofs); \
Size _wholeSize; Point ofs; name ## _roi.locateROI(_wholeSize, ofs); \
u ## name ## _roi = u ## name(Rect(ofs.x, ofs.y, name ## _roi.size().width, name ## _roi.size().height)); \
}
#define UMAT_UPLOAD_OUTPUT_PARAMETER(name) UMAT_UPLOAD_INPUT_PARAMETER(name)