Merge pull request #1908 from ilya-lavrenov:tapi_imgproc

modules/core/src/copy.cpp

@@ -47,6 +47,7 @@
 // */
 
 #include "precomp.hpp"
+#include "opencl_kernels.hpp"
 
 namespace cv
 {
@@ -701,12 +702,73 @@ void copyMakeConstBorder_8u( const uchar* src, size_t srcstep, cv::Size srcroi,
 
 }
 
+namespace cv {
+
+static bool ocl_copyMakeBorder( InputArray _src, OutputArray _dst, int top, int bottom,
+                                int left, int right, int borderType, const Scalar& value )
+{
+    int type = _src.type(), cn = CV_MAT_CN(type);
+    bool isolated = (borderType & BORDER_ISOLATED) != 0;
+    borderType &= ~cv::BORDER_ISOLATED;
+
+    if ( !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || borderType == BORDER_REFLECT ||
+           borderType == BORDER_WRAP || borderType == BORDER_REFLECT_101) ||
+         cn == 3 || cn > 4)
+        return false;
+
+    const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" };
+    ocl::Kernel k("copyMakeBorder", ocl::core::copymakeborder_oclsrc,
+                  format("-D T=%s -D %s", ocl::memopTypeToStr(type), borderMap[borderType]));
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat();
+    if( src.isSubmatrix() && !isolated )
+    {
+        Size wholeSize;
+        Point ofs;
+        src.locateROI(wholeSize, ofs);
+        int dtop = std::min(ofs.y, top);
+        int dbottom = std::min(wholeSize.height - src.rows - ofs.y, bottom);
+        int dleft = std::min(ofs.x, left);
+        int dright = std::min(wholeSize.width - src.cols - ofs.x, right);
+        src.adjustROI(dtop, dbottom, dleft, dright);
+        top -= dtop;
+        left -= dleft;
+        bottom -= dbottom;
+        right -= dright;
+    }
+
+    _dst.create(src.rows + top + bottom, src.cols + left + right, type);
+    UMat dst = _dst.getUMat();
+
+    if (top == 0 && left == 0 && bottom == 0 && right == 0)
+    {
+        if(src.u != dst.u || src.step != dst.step)
+            src.copyTo(dst);
+        return true;
+    }
+
+    k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst),
+           top, left, ocl::KernelArg::Constant(Mat(1, 1, type, value)));
+
+    size_t globalsize[2] = { dst.cols, dst.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
+}
+
 void cv::copyMakeBorder( InputArray _src, OutputArray _dst, int top, int bottom,
                          int left, int right, int borderType, const Scalar& value )
 {
-    Mat src = _src.getMat();
     CV_Assert( top >= 0 && bottom >= 0 && left >= 0 && right >= 0 );
 
+    if (ocl::useOpenCL() && _dst.isUMat() && _src.dims() <= 2 &&
+            ocl_copyMakeBorder(_src, _dst, top, bottom, left, right, borderType, value))
+        return;
+
+    Mat src = _src.getMat();
+
     if( src.isSubmatrix() && (borderType & BORDER_ISOLATED) == 0 )
     {
         Size wholeSize;

modules/core/src/opencl/copymakeborder.cl

+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Niko Li, newlife20080214@gmail.com
+//    Zero Lin zero.lin@amd.com
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//
+
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+#ifdef BORDER_CONSTANT
+#define EXTRAPOLATE(x, y, v) v = scalar;
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, y, v) \
+    { \
+        x = max(min(x, src_cols - 1), 0); \
+        y = max(min(y, src_rows - 1), 0); \
+        v = *(__global const T *)(srcptr + mad24(y, src_step, x * (int)sizeof(T) + src_offset)); \
+    }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, y, v) \
+    { \
+        if (x < 0) \
+            x -= ((x - src_cols + 1) / src_cols) * src_cols; \
+        if (x >= src_cols) \
+            x %= src_cols; \
+        \
+        if (y < 0) \
+            y -= ((y - src_rows + 1) / src_rows) * src_rows; \
+        if( y >= src_rows ) \
+            y %= src_rows; \
+        v = *(__global const T *)(srcptr + mad24(y, src_step, x * (int)sizeof(T) + src_offset)); \
+    }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
+#ifdef BORDER_REFLECT
+#define DELTA int delta = 0
+#else
+#define DELTA int delta = 1
+#endif
+#define EXTRAPOLATE(x, y, v) \
+    { \
+        DELTA; \
+        if (src_cols == 1) \
+            x = 0; \
+        else \
+            do \
+            { \
+                if( x < 0 ) \
+                    x = -x - 1 + delta; \
+                else \
+                    x = src_cols - 1 - (x - src_cols) - delta; \
+            } \
+            while (x >= src_cols || x < 0); \
+        \
+        if (src_rows == 1) \
+            y = 0; \
+        else \
+            do \
+            { \
+                if( y < 0 ) \
+                    y = -y - 1 + delta; \
+                else \
+                    y = src_rows - 1 - (y - src_rows) - delta; \
+            } \
+            while (y >= src_rows || y < 0); \
+        v = *(__global const T *)(srcptr + mad24(y, src_step, x * (int)sizeof(T) + src_offset)); \
+    }
+#else
+#error No extrapolation method
+#endif
+
+#define NEED_EXTRAPOLATION(gx, gy) (gx >= src_cols || gy >= src_rows || gx < 0 || gy < 0)
+
+__kernel void copyMakeBorder(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
+                             __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+                             int top, int left, T scalar)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < dst_cols && y < dst_rows)
+    {
+        int src_x = x - left;
+        int src_y = y - top;
+
+        int dst_index = mad24(y, dst_step, x * (int)sizeof(T) + dst_offset);
+        __global T * dst = (__global T *)(dstptr + dst_index);
+
+        if (NEED_EXTRAPOLATION(src_x, src_y))
+            EXTRAPOLATE(src_x, src_y, dst[0])
+        else
+        {
+            int src_index = mad24(src_y, src_step, src_x * (int)sizeof(T) + src_offset);
+            __global const T * src = (__global const T *)(srcptr + src_index);
+            dst[0] = src[0];
+        }
+    }
+}

modules/core/src/umatrix.cpp

@@ -580,6 +580,7 @@ Mat UMat::getMat(int accessFlags) const
     Mat hdr(dims, size.p, type(), u->data + offset, step.p);
     hdr.flags = flags;
     hdr.u = u;
+    hdr.flags = flags;
     hdr.datastart = u->data;
     hdr.data = hdr.datastart + offset;
     hdr.datalimit = hdr.dataend = u->data + u->size;

modules/imgproc/src/opencl/threshold.cl

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Zhang Ying, zhangying913@gmail.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+__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)sizeof(T));
+        int dst_index = mad24(gy, dst_step, dst_offset + gx * (int)sizeof(T));
+
+        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);
+#elif defined THRESH_BINARY_INV
+        dst[0] = sdata > thresh ? (T)(0) : max_val;
+#elif defined THRESH_TRUNC
+        dst[0] = sdata > thresh ? thresh : sdata;
+#elif defined THRESH_TOZERO
+        dst[0] = sdata > thresh ? sdata : (T)(0);
+#elif defined THRESH_TOZERO_INV
+        dst[0] = sdata > thresh ? (T)(0) : sdata;
+#endif
+    }
+}

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;

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

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)
@@ -306,7 +306,8 @@ IMPLEMENT_PARAM_CLASS(Channels, int)
 #define OCL_ALL_CHANNELS Values(1, 2, 3, 4)
 
 CV_ENUM(Interpolation, INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_AREA)
-CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_WRAP, BORDER_REFLECT, BORDER_REFLECT_101)
+CV_ENUM(ThreshOp, THRESH_BINARY, THRESH_BINARY_INV, THRESH_TRUNC, THRESH_TOZERO, THRESH_TOZERO_INV)
+CV_ENUM(BorderType, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101)
 
 #define OCL_INSTANTIATE_TEST_CASE_P(prefix, test_case_name, generator) \
     INSTANTIATE_TEST_CASE_P(OCL_ ## prefix, test_case_name, generator)