added cv::threshold to T-API

diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index 8b54876..f733dd1 100644 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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<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

diff --git a/modules/ts/include/opencv2/ts/ocl_test.hpp b/modules/ts/include/opencv2/ts/ocl_test.hpp
index f6c6f5b..1b66799 100644 (file)
--- 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)