--- /dev/null
+// 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
}
+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 == 1 || 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 );
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))
+ if (use_opencl)
+ { CV_Assert (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))
--- /dev/null
+/*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, 4),
+ Values(3, 5),
+ Bool())
+ );
+} } // namespace cvtest::ocl
+
+#endif
projects / profile / ivi / opencv.git / commitdiff