From e2b99a32b29dd2dbfd24d2ae294527743a9fe19c Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Tue, 3 Dec 2013 00:13:05 +0400 Subject: [PATCH] added cv::threshold to T-API --- modules/core/src/ocl.cpp | 2 +- modules/imgproc/src/opencl/threshold.cl | 79 +++++------------------------- modules/imgproc/src/thresh.cpp | 38 ++++++++++++++ modules/imgproc/test/ocl/test_warp.cpp | 22 ++++----- modules/ts/include/opencv2/ts/ocl_test.hpp | 2 +- 5 files changed, 63 insertions(+), 80 deletions(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 8b54876..f733dd1 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -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) diff --git a/modules/imgproc/src/opencl/threshold.cl b/modules/imgproc/src/opencl/threshold.cl index 63e4102..f5b6fbb 100644 --- a/modules/imgproc/src/opencl/threshold.cl +++ b/modules/imgproc/src/opencl/threshold.cl @@ -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) -{ - 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); - -#ifdef SRC_ALIGNED - VT sdata = *((__global VT *)(src + src_index)); -#else - VT sdata = VLOADN(0, src + src_index); -#endif - VT vthresh = (VT)(thresh); - -#ifdef THRESH_BINARY - VT vecValue = sdata > vthresh ? max_val : (VT)(0); -#elif defined THRESH_BINARY_INV - VT vecValue = sdata > vthresh ? (VT)(0) : max_val; -#elif defined THRESH_TRUNC - VT vecValue = sdata > vthresh ? thresh : sdata; -#elif defined THRESH_TOZERO - VT vecValue = sdata > vthresh ? sdata : (VT)(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) +__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) { - 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)); - T sdata = src[src_index]; + T sdata = *(__global const T *)(srcptr + src_index); + __global T * dst = (__global T *)(dstptr + dst_index); #ifdef THRESH_BINARY - dst[dst_index] = sdata > thresh ? max_val : (T)(0); + dst[0] = sdata > thresh ? max_val : (T)(0); #elif defined THRESH_BINARY_INV - dst[dst_index] = sdata > thresh ? (T)(0) : max_val; + dst[0] = sdata > thresh ? (T)(0) : max_val; #elif defined THRESH_TRUNC - dst[dst_index] = sdata > thresh ? thresh : sdata; + dst[0] = sdata > thresh ? thresh : sdata; #elif defined THRESH_TOZERO - dst[dst_index] = sdata > thresh ? sdata : (T)(0); + dst[0] = 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 diff --git a/modules/imgproc/src/thresh.cpp b/modules/imgproc/src/thresh.cpp index 6e80b3b..ce853a7 100644 --- a/modules/imgproc/src/thresh.cpp +++ b/modules/imgproc/src/thresh.cpp @@ -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; diff --git a/modules/imgproc/test/ocl/test_warp.cpp b/modules/imgproc/test/ocl/test_warp.cpp index 9c22e17..c05c335 100644 --- a/modules/imgproc/test/ocl/test_warp.cpp +++ b/modules/imgproc/test/ocl/test_warp.cpp @@ -224,7 +224,7 @@ OCL_TEST_P(Resize, Mat) ///////////////////////////////////////////////////////////////////////////////////////////////// // remap -PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair, Border, bool) +PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair, 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)CV_32FC1, (MatType)CV_32FC1), std::pair((MatType)CV_16SC2, (MatType)CV_16UC1), std::pair((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)CV_32FC2, noType), std::pair((MatType)CV_16SC2, (MatType)CV_16UC1), std::pair((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 diff --git a/modules/ts/include/opencv2/ts/ocl_test.hpp b/modules/ts/include/opencv2/ts/ocl_test.hpp index f6c6f5b..1b66799 100644 --- a/modules/ts/include/opencv2/ts/ocl_test.hpp +++ b/modules/ts/include/opencv2/ts/ocl_test.hpp @@ -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) -- 2.7.4