Merge pull request #1867 from ilya-lavrenov:ocl_resize_master

modules/imgproc/src/imgwarp.cpp

@@ -1908,7 +1908,7 @@ static bool ocl_resize( InputArray _src, OutputArray _dst, Size dsize,
     int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
     if( !(cn <= 4 &&
            (interpolation == INTER_NEAREST ||
-           (interpolation == INTER_LINEAR && (depth == CV_8U || depth == CV_32F)))) )
+           (interpolation == INTER_LINEAR))) )
         return false;
     UMat src = _src.getUMat();
     _dst.create(dsize, type);
@@ -1917,11 +1917,11 @@ static bool ocl_resize( InputArray _src, OutputArray _dst, Size dsize,
 
     if (interpolation == INTER_LINEAR)
     {
-        int wdepth = depth == CV_8U ? CV_32S : CV_32F;
+        int wdepth = std::max(depth, CV_32S);
         int wtype = CV_MAKETYPE(wdepth, cn);
         char buf[2][32];
         k.create("resizeLN", ocl::imgproc::resize_oclsrc,
-                 format("-D INTER_LINEAR -D depth=%s -D PIXTYPE=%s -D WORKTYPE=%s -D convertToWT=%s -D convertToDT=%s",
+                 format("-D INTER_LINEAR -D depth=%d -D PIXTYPE=%s -D WORKTYPE=%s -D convertToWT=%s -D convertToDT=%s",
                         depth, ocl::typeToStr(type), ocl::typeToStr(wtype),
                         ocl::convertTypeStr(depth, wdepth, cn, buf[0]),
                         ocl::convertTypeStr(wdepth, depth, cn, buf[1])));
@@ -1937,6 +1937,7 @@ static bool ocl_resize( InputArray _src, OutputArray _dst, Size dsize,
     k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst),
            (float)(1./fx), (float)(1./fy));
     size_t globalsize[] = { dst.cols, dst.rows };
+
     return k.run(2, globalsize, 0, false);
 }
 

modules/imgproc/src/opencl/resize.cl

@@ -89,7 +89,8 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset,
     int x_ = INC(x,srccols);
     __global const PIXTYPE* src = (__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE));
 
-#if depth == 0
+#if depth <= 4
+
     u = u * INTER_RESIZE_COEF_SCALE;
     v = v * INTER_RESIZE_COEF_SCALE;
 
@@ -102,18 +103,22 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset,
     WORKTYPE data1 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE)));
     WORKTYPE data2 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE)));
     WORKTYPE data3 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE)));
+
     WORKTYPE val = mul24((WORKTYPE)mul24(U1, V1), data0) + mul24((WORKTYPE)mul24(U, V1), data1) +
                mul24((WORKTYPE)mul24(U1, V), data2) + mul24((WORKTYPE)mul24(U, V), data3);
 
     PIXTYPE uval = convertToDT((val + (1<<(CAST_BITS-1)))>>CAST_BITS);
+
 #else
-    float u1 = 1.f-u;
+    float u1 = 1.f - u;
-    float v1 = 1.f-v;
+    float v1 = 1.f - v;
     WORKTYPE data0 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE)));
     WORKTYPE data1 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE)));
     WORKTYPE data2 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE)));
     WORKTYPE data3 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE)));
-    PIXTYPE uval = u1 * v1 * s_data1 + u * v1 * s_data2 + u1 * v *s_data3 + u * v *s_data4;
+
+    PIXTYPE uval = u1 * v1 * data0 + u * v1 * data1 + u1 * v *data2 + u * v *data3;
+
 #endif
 
     if(dx < dstcols && dy < dstrows)

modules/imgproc/test/ocl/test_warp.cpp

@@ -0,0 +1,139 @@
+/*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.
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Niko Li, newlife20080214@gmail.com
+//    Jia Haipeng, jiahaipeng95@gmail.com
+//    Shengen Yan, yanshengen@gmail.com
+//    Jiang Liyuan, lyuan001.good@163.com
+//    Rock Li, Rock.Li@amd.com
+//    Wu Zailong, bullet@yeah.net
+//    Xu Pang, pangxu010@163.com
+//    Sen Liu, swjtuls1987@126.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*/
+
+#include "test_precomp.hpp"
+#include "cvconfig.h"
+#include "opencv2/ts/ocl_test.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// resize
+
+PARAM_TEST_CASE(Resize, MatType, double, double, Interpolation, bool)
+{
+    int type, interpolation;
+    double fx, fy;
+    bool useRoi;
+
+    TEST_DECLARE_INPUT_PARATEMER(src)
+    TEST_DECLARE_OUTPUT_PARATEMER(dst)
+
+    virtual void SetUp()
+    {
+        type = GET_PARAM(0);
+        fx = GET_PARAM(1);
+        fy = GET_PARAM(2);
+        interpolation = GET_PARAM(3);
+        useRoi = GET_PARAM(4);
+    }
+
+    void random_roi()
+    {
+        CV_Assert(fx > 0 && fy > 0);
+
+        Size srcRoiSize = randomSize(1, MAX_VALUE), dstRoiSize;
+        dstRoiSize.width = cvRound(srcRoiSize.width * fx);
+        dstRoiSize.height = cvRound(srcRoiSize.height * fy);
+
+        if (dstRoiSize.area() == 0)
+        {
+            random_roi();
+            return;
+        }
+
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, srcRoiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, dstRoiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src);
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst);
+    }
+
+    void Near(double threshold = 0.0)
+    {
+        EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
+        EXPECT_MAT_NEAR(dst, udst, threshold);
+    }
+};
+
+OCL_TEST_P(Resize, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        random_roi();
+
+        OCL_OFF(cv::resize(src_roi, dst_roi, Size(), fx, fy, interpolation));
+        OCL_ON(cv::resize(usrc_roi, udst_roi, Size(), fx, fy, interpolation));
+
+        Near(1.0);
+    }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////
+
+OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Resize, Combine(
+                            Values(CV_8UC1, CV_8UC4, CV_16UC2, CV_32FC1, CV_32FC4),
+                            Values(0.5, 1.5, 2.0),
+                            Values(0.5, 1.5, 2.0),
+                            Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR),
+                            Bool()));
+
+} } // namespace cvtest::ocl
+
+#endif // HAVE_OPENCL

modules/ts/include/opencv2/ts/ocl_test.hpp

@@ -305,6 +305,8 @@ IMPLEMENT_PARAM_CLASS(Channels, int)
 #define OCL_ALL_DEPTHS Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F)
 #define OCL_ALL_CHANNELS Values(1, 2, 3, 4)
 
+CV_ENUM(Interpolation, INTER_NEAREST, INTER_LINEAR, INTER_CUBIC)
+
 #define OCL_INSTANTIATE_TEST_CASE_P(prefix, test_case_name, generator) \
     INSTANTIATE_TEST_CASE_P(OCL_ ## prefix, test_case_name, generator)