From: Elena Gvozdeva Date: Mon, 9 Dec 2013 06:34:32 +0000 (+0400) Subject: Added ocl_medianFilter using Transparent API X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~3623^2~3 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=b2d8b491c06887fe715c1266f4483ae7a2b4c09a;p=platform%2Fupstream%2Fopencv.git Added ocl_medianFilter using Transparent API --- diff --git a/modules/imgproc/src/opencl/medianFilter.cl b/modules/imgproc/src/opencl/medianFilter.cl new file mode 100644 index 0000000..3d41f7c --- /dev/null +++ b/modules/imgproc/src/opencl/medianFilter.cl @@ -0,0 +1,167 @@ +// 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 diff --git a/modules/imgproc/src/smooth.cpp b/modules/imgproc/src/smooth.cpp index 9d280c9..f7f7eed 100644 --- a/modules/imgproc/src/smooth.cpp +++ b/modules/imgproc/src/smooth.cpp @@ -1661,19 +1661,61 @@ 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 == 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)) diff --git a/modules/imgproc/test/ocl/test_medianfilter.cpp b/modules/imgproc/test/ocl/test_medianfilter.cpp new file mode 100644 index 0000000..51d8d6a --- /dev/null +++ b/modules/imgproc/test/ocl/test_medianfilter.cpp @@ -0,0 +1,111 @@ +/*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