From: Ilya Lavrenov Date: Thu, 12 Dec 2013 19:42:54 +0000 (+0400) Subject: added cv::equalizeHist to T-API X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~3630^2 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=5c0dabf1ea5f6d969b9301bfdb5cc3880d75bdfb;p=platform%2Fupstream%2Fopencv.git added cv::equalizeHist to T-API --- diff --git a/modules/imgproc/src/histogram.cpp b/modules/imgproc/src/histogram.cpp index f0a7818..7849d51 100644 --- a/modules/imgproc/src/histogram.cpp +++ b/modules/imgproc/src/histogram.cpp @@ -38,7 +38,9 @@ // the use of this software, even if advised of the possibility of such damage. // //M*/ + #include "precomp.hpp" +#include "opencl_kernels.hpp" namespace cv { @@ -3129,17 +3131,80 @@ CV_IMPL void cvEqualizeHist( const CvArr* srcarr, CvArr* dstarr ) cv::equalizeHist(cv::cvarrToMat(srcarr), cv::cvarrToMat(dstarr)); } +namespace cv { + +enum +{ + BINS = 256 +}; + +static bool ocl_calcHist(InputArray _src, OutputArray _hist) +{ + int compunits = ocl::Device::getDefault().maxComputeUnits(); + size_t wgs = ocl::Device::getDefault().maxWorkGroupSize(); + + ocl::Kernel k1("calculate_histogram", ocl::imgproc::histogram_oclsrc, + format("-D BINS=%d -D HISTS_COUNT=%d -D WGS=%d", BINS, compunits, wgs)); + if (k1.empty()) + return false; + + _hist.create(1, BINS, CV_32SC1); + UMat src = _src.getUMat(), hist = _hist.getUMat(), ghist(1, BINS * compunits, CV_32SC1); + + k1.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::PtrWriteOnly(ghist), + (int)src.total()); + + size_t globalsize = compunits * wgs; + if (!k1.run(1, &globalsize, &wgs, false)) + return false; + + ocl::Kernel k2("merge_histogram", ocl::imgproc::histogram_oclsrc, + format("-D BINS=%d -D HISTS_COUNT=%d -D WGS=%d", BINS, compunits, (int)wgs)); + if (k2.empty()) + return false; + + k2.args(ocl::KernelArg::PtrReadOnly(ghist), ocl::KernelArg::PtrWriteOnly(hist)); + return k2.run(1, &wgs, &wgs, false); +} + +static bool ocl_equalizeHist(InputArray _src, OutputArray _dst) +{ + size_t wgs = ocl::Device::getDefault().maxWorkGroupSize(); + + // calculation of histogram + UMat hist; + if (!ocl_calcHist(_src, hist)) + return false; + + UMat lut(1, 256, CV_8UC1); + ocl::Kernel k("calcLUT", ocl::imgproc::histogram_oclsrc, format("-D BINS=%d -D HISTS_COUNT=1 -D WGS=%d", BINS, (int)wgs)); + k.args(ocl::KernelArg::PtrWriteOnly(lut), ocl::KernelArg::PtrReadOnly(hist), (int)_src.total()); + + // calculation of LUT + if (!k.run(1, &wgs, &wgs, false)) + return false; + + // execute LUT transparently + LUT(_src, lut, _dst); + return true; +} + +} + void cv::equalizeHist( InputArray _src, OutputArray _dst ) { - Mat src = _src.getMat(); - CV_Assert( src.type() == CV_8UC1 ); + CV_Assert( _src.type() == CV_8UC1 ); - _dst.create( src.size(), src.type() ); - Mat dst = _dst.getMat(); + if (_src.empty()) + return; - if(src.empty()) + if (ocl::useOpenCL() && _dst.isUMat() && ocl_equalizeHist(_src, _dst)) return; + Mat src = _src.getMat(); + _dst.create( src.size(), src.type() ); + Mat dst = _dst.getMat(); + Mutex histogramLockInstance; const int hist_sz = EqualizeHistCalcHist_Invoker::HIST_SZ; diff --git a/modules/imgproc/src/opencl/histogram.cl b/modules/imgproc/src/opencl/histogram.cl new file mode 100644 index 0000000..481bdcf --- /dev/null +++ b/modules/imgproc/src/opencl/histogram.cl @@ -0,0 +1,120 @@ +// 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 +// Xu Pang, pangxu010@163.com +// Wenju He, wenju@multicorewareinc.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. +// +// + +__kernel void calculate_histogram(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols, + __global uchar * hist, int total) +{ + int lid = get_local_id(0); + int id = get_global_id(0); + int gid = get_group_id(0); + + __local int localhist[BINS]; + + for (int i = lid; i < BINS; i += WGS) + localhist[i] = 0; + barrier(CLK_LOCAL_MEM_FENCE); + + for (int grain = HISTS_COUNT * WGS; id < total; id += grain) + { + int src_index = mad24(id / src_cols, src_step, src_offset + id % src_cols); + atomic_inc(localhist + (int)src[src_index]); + } + barrier(CLK_LOCAL_MEM_FENCE); + + for (int i = lid; i < BINS; i += WGS) + *(__global int *)(hist + mad24(gid, BINS * (int)sizeof(int), i * (int)sizeof(int))) = localhist[i]; +} + +__kernel void merge_histogram(__global const int * ghist, __global int * hist) +{ + int lid = get_local_id(0); + + #pragma unroll + for (int i = lid; i < BINS; i += WGS) + hist[i] = 0; + barrier(CLK_LOCAL_MEM_FENCE); + + #pragma unroll + for (int i = 0; i < HISTS_COUNT; ++i) + { + #pragma unroll + for (int j = lid; j < BINS; j += WGS) + hist[j] += ghist[mad24(i, BINS, j)]; + barrier(CLK_LOCAL_MEM_FENCE); + } +} + +__kernel void calcLUT(__global uchar * dst, __constant int * hist, int total) +{ + int lid = get_local_id(0); + __local int sumhist[BINS]; + __local float scale; + + sumhist[lid] = hist[lid]; + barrier(CLK_LOCAL_MEM_FENCE); + + if (lid == 0) + { + int sum = 0, i = 0; + while (!sumhist[i]) + ++i; + + if (total == sumhist[i]) + { + scale = 1; + for (int j = 0; j < BINS; ++j) + sumhist[i] = i; + } + else + { + scale = 255.f / (total - sumhist[i]); + + for (sumhist[i++] = 0; i < BINS; i++) + { + sum += sumhist[i]; + sumhist[i] = sum; + } + } + } + barrier(CLK_LOCAL_MEM_FENCE); + + #pragma unroll + for (int i = lid; i < BINS; i += WGS) + dst[i]= convert_uchar_sat_rte(convert_float(sumhist[i]) * scale); +} diff --git a/modules/imgproc/test/ocl/test_imgproc.cpp b/modules/imgproc/test/ocl/test_imgproc.cpp index d6799ce..bf6f8e6 100644 --- a/modules/imgproc/test/ocl/test_imgproc.cpp +++ b/modules/imgproc/test/ocl/test_imgproc.cpp @@ -184,7 +184,7 @@ OCL_TEST_P(EqualizeHist, Mat) OCL_OFF(cv::equalizeHist(src_roi, dst_roi)); OCL_ON(cv::equalizeHist(usrc_roi, udst_roi)); - Near(1.1); + Near(1); } } diff --git a/modules/imgproc/test/test_imgproc_umat.cpp b/modules/imgproc/test/test_imgproc_umat.cpp index 5237038..464f0b6 100644 --- a/modules/imgproc/test/test_imgproc_umat.cpp +++ b/modules/imgproc/test/test_imgproc_umat.cpp @@ -76,6 +76,8 @@ protected: destroyWindow("equalized gray"); #endif ts->set_failed_test_info(cvtest::TS::OK); + + (void)uresult.getMat(ACCESS_READ); } };