Merge pull request #1992 from ElenaGvozdeva:ocl_medianFilter

modules/imgproc/src/opencl/medianFilter.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.
+//
+// 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.
+
+#define DATA_TYPE type
+
+#define scnbytes ((int)sizeof(type))
+
+#define op(a,b) {    mid=a; a=min(a,b); b=max(mid,b);}
+
+__kernel void medianFilter3(__global const uchar* srcptr, int srcStep, int srcOffset,
+                            __global uchar* dstptr, int dstStep, int dstOffset,
+                            int rows, int cols)
+{
+    __local DATA_TYPE data[18][18];
+
+    int x = get_local_id(0);
+    int y = get_local_id(1);
+
+    int gx= get_global_id(0);
+    int gy= get_global_id(1);
+
+    int dx = gx - x - 1;
+    int dy = gy - y - 1;
+
+    const int id = min((int)(x*16+y), 9*18-1);
+
+    int dr = id / 18;
+    int dc = id % 18;
+
+    int c = clamp(dx+dc, 0, cols-1);
+
+    int r = clamp(dy+dr, 0, rows-1);
+    int index1 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+    r = clamp(dy+dr+9, 0, rows-1);
+    int index9 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+    __global DATA_TYPE * src = (__global DATA_TYPE *)(srcptr + index1);
+    data[dr][dc] = src[0];
+
+    src = (__global DATA_TYPE *)(srcptr + index9);
+    data[dr+9][dc] = src[0];
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    DATA_TYPE p0=data[y][x], p1=data[y][(x+1)], p2=data[y][(x+2)];
+    DATA_TYPE p3=data[y+1][x], p4=data[y+1][(x+1)], p5=data[y+1][(x+2)];
+    DATA_TYPE p6=data[y+2][x], p7=data[y+2][(x+1)], p8=data[y+2][(x+2)];
+    DATA_TYPE mid;
+
+    op(p1, p2); op(p4, p5); op(p7, p8); op(p0, p1);
+    op(p3, p4); op(p6, p7); op(p1, p2); op(p4, p5);
+    op(p7, p8); op(p0, p3); op(p5, p8); op(p4, p7);
+    op(p3, p6); op(p1, p4); op(p2, p5); op(p4, p7);
+    op(p4, p2); op(p6, p4); op(p4, p2);
+
+    int dst_index = mad24( gy, dstStep, dstOffset + gx * scnbytes);
+
+    if( gy < rows && gx < cols)
+    {
+        __global DATA_TYPE* dst = (__global DATA_TYPE *)(dstptr + dst_index);
+        dst[0] = p4;
+    }
+}
+
+__kernel void medianFilter5(__global const uchar* srcptr, int srcStep, int srcOffset,
+                            __global uchar* dstptr, int dstStep, int dstOffset,
+                            int rows, int cols)
+{
+    __local DATA_TYPE data[20][20];
+
+    int x =get_local_id(0);
+    int y =get_local_id(1);
+
+    int gx=get_global_id(0);
+    int gy=get_global_id(1);
+
+    int dx = gx - x - 2;
+    int dy = gy - y - 2;
+
+    const int id = min((int)(x*16+y), 10*20-1);
+
+    int dr=id/20;
+    int dc=id%20;
+
+    int c=clamp(dx+dc, 0, cols-1);
+
+    int r = clamp(dy+dr, 0, rows-1);
+    int index1 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+    r = clamp(dy+dr+10, 0, rows-1);
+    int index10 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+    __global DATA_TYPE * src = (__global DATA_TYPE *)(srcptr + index1);
+    data[dr][dc] = src[0];
+    src = (__global DATA_TYPE *)(srcptr + index10);
+    data[dr+10][dc] = src[0];
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    DATA_TYPE p0=data[y][x], p1=data[y][x+1], p2=data[y][x+2], p3=data[y][x+3], p4=data[y][x+4];
+    DATA_TYPE p5=data[y+1][x], p6=data[y+1][x+1], p7=data[y+1][x+2], p8=data[y+1][x+3], p9=data[y+1][x+4];
+    DATA_TYPE p10=data[y+2][x], p11=data[y+2][x+1], p12=data[y+2][x+2], p13=data[y+2][x+3], p14=data[y+2][x+4];
+    DATA_TYPE p15=data[y+3][x], p16=data[y+3][x+1], p17=data[y+3][x+2], p18=data[y+3][x+3], p19=data[y+3][x+4];
+    DATA_TYPE p20=data[y+4][x], p21=data[y+4][x+1], p22=data[y+4][x+2], p23=data[y+4][x+3], p24=data[y+4][x+4];
+    DATA_TYPE mid;
+
+    op(p1, p2); op(p0, p1); op(p1, p2); op(p4, p5); op(p3, p4);
+    op(p4, p5); op(p0, p3); op(p2, p5); op(p2, p3); op(p1, p4);
+    op(p1, p2); op(p3, p4); op(p7, p8); op(p6, p7); op(p7, p8);
+    op(p10, p11); op(p9, p10); op(p10, p11); op(p6, p9); op(p8, p11);
+    op(p8, p9); op(p7, p10); op(p7, p8); op(p9, p10); op(p0, p6);
+    op(p4, p10); op(p4, p6); op(p2, p8); op(p2, p4); op(p6, p8);
+    op(p1, p7); op(p5, p11); op(p5, p7); op(p3, p9); op(p3, p5);
+    op(p7, p9); op(p1, p2); op(p3, p4); op(p5, p6); op(p7, p8);
+    op(p9, p10); op(p13, p14); op(p12, p13); op(p13, p14); op(p16, p17);
+    op(p15, p16); op(p16, p17); op(p12, p15); op(p14, p17); op(p14, p15);
+    op(p13, p16); op(p13, p14); op(p15, p16); op(p19, p20); op(p18, p19);
+    op(p19, p20); op(p21, p22); op(p23, p24); op(p21, p23); op(p22, p24);
+    op(p22, p23); op(p18, p21); op(p20, p23); op(p20, p21); op(p19, p22);
+    op(p22, p24); op(p19, p20); op(p21, p22); op(p23, p24); op(p12, p18);
+    op(p16, p22); op(p16, p18); op(p14, p20); op(p20, p24); op(p14, p16);
+    op(p18, p20); op(p22, p24); op(p13, p19); op(p17, p23); op(p17, p19);
+    op(p15, p21); op(p15, p17); op(p19, p21); op(p13, p14); op(p15, p16);
+    op(p17, p18); op(p19, p20); op(p21, p22); op(p23, p24); op(p0, p12);
+    op(p8, p20); op(p8, p12); op(p4, p16); op(p16, p24); op(p12, p16);
+    op(p2, p14); op(p10, p22); op(p10, p14); op(p6, p18); op(p6, p10);
+    op(p10, p12); op(p1, p13); op(p9, p21); op(p9, p13); op(p5, p17);
+    op(p13, p17); op(p3, p15); op(p11, p23); op(p11, p15); op(p7, p19);
+    op(p7, p11); op(p11, p13); op(p11, p12);
+
+    int dst_index = mad24( gy, dstStep, dstOffset + gx * scnbytes);
+
+    if( gy < rows && gx < cols)
+    {
+        __global DATA_TYPE* dst = (__global DATA_TYPE *)(dstptr + dst_index);
+        dst[0] = p12;
+    }
+}
\ No newline at end of file

modules/imgproc/src/smooth.cpp

@@ -1814,19 +1814,59 @@ medianBlur_SortNet( const Mat& _src, Mat& _dst, int m )
 
 }
 
+namespace cv
+{
+    static bool ocl_medianFilter ( InputArray _src, OutputArray _dst, int m)
+    {
+        int type = _src.type();
+        int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+
+        if (!((depth == CV_8U || depth == CV_16U || depth == CV_16S || depth == CV_32F) && (cn != 3 && cn <= 4)))
+            return false;
+
+        const char * kernelName;
+
+        if (m==3)
+            kernelName = "medianFilter3";
+        else if (m==5)
+            kernelName = "medianFilter5";
+        else
+            return false;
+
+        ocl::Kernel k(kernelName,ocl::imgproc::medianFilter_oclsrc,format("-D type=%s",ocl::typeToStr(type)));
+        if (k.empty())
+            return false;
+
+        _dst.create(_src.size(),type);
+        UMat src = _src.getUMat(), dst = _dst.getUMat();
+
+        size_t globalsize[2] = {(src.cols + 18) / 16 * 16, (src.rows + 15) / 16 * 16};
+        size_t localsize[2] = {16, 16};
+
+        return k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst)).run(2,globalsize,localsize,false);
+    }
+}
+
 void cv::medianBlur( InputArray _src0, OutputArray _dst, int ksize )
 {
-    Mat src0 = _src0.getMat();
-    _dst.create( src0.size(), src0.type() );
-    Mat dst = _dst.getMat();
+    CV_Assert( (ksize % 2 == 1) && (_src0.dims() <= 2 ));
 
     if( ksize <= 1 )
     {
+        Mat src0 = _src0.getMat();
+        _dst.create( src0.size(), src0.type() );
+        Mat dst = _dst.getMat();
         src0.copyTo(dst);
         return;
     }
 
-    CV_Assert( ksize % 2 == 1 );
+    bool use_opencl = ocl::useOpenCL() && _dst.isUMat();
+    if ( use_opencl && ocl_medianFilter(_src0,_dst, ksize))
+        return;
+
+    Mat src0 = _src0.getMat();
+    _dst.create( src0.size(), src0.type() );
+    Mat dst = _dst.getMat();
 
 #ifdef HAVE_TEGRA_OPTIMIZATION
     if (tegra::medianBlur(src0, dst, ksize))

modules/imgproc/test/ocl/test_medianfilter.cpp

+/*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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+//
+// 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 "opencv2/ts/ocl_test.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+/////////////////////////////////////////////medianFilter//////////////////////////////////////////////////////////
+
+PARAM_TEST_CASE(MedianFilter, MatDepth, Channels, int, bool)
+{
+    int type;
+    int ksize;
+    bool use_roi;
+
+    TEST_DECLARE_INPUT_PARAMETER(src)
+    TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+    virtual void SetUp()
+    {
+        type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
+        ksize = GET_PARAM(2);
+        use_roi = GET_PARAM(3);
+    }
+
+    virtual void generateTestData()
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+        Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, 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, udst, threshold);
+        EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
+    }
+};
+
+OCL_TEST_P(MedianFilter, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::medianBlur(src_roi, dst_roi, ksize));
+        OCL_ON(cv::medianBlur(usrc_roi, udst_roi, ksize));
+
+        Near(0);
+    }
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(ImageProc, MedianFilter, Combine(
+                                Values(CV_8U, CV_16U, CV_16S, CV_32F),
+                                Values(1, 2, 4),
+                                Values(3, 5),
+                                Bool())
+                           );
+} } // namespace cvtest::ocl
+
+#endif