From: Ilya Lavrenov Date: Mon, 9 Jun 2014 15:45:37 +0000 (+0400) Subject: eliminated convertTo X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~3171^2~2 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=eeaa4b36657838a8f79192341fc50d484d23d6e1;p=platform%2Fupstream%2Fopencv.git eliminated convertTo --- diff --git a/modules/imgproc/src/histogram.cpp b/modules/imgproc/src/histogram.cpp index 92db679..9be3f56 100644 --- a/modules/imgproc/src/histogram.cpp +++ b/modules/imgproc/src/histogram.cpp @@ -1494,7 +1494,7 @@ static bool ocl_calcHist1(InputArray _src, OutputArray _hist, int ddepth = CV_32 _hist.create(BINS, 1, ddepth); UMat src = _src.getUMat(), ghist(1, BINS * compunits, CV_32SC1), - hist = ddepth == CV_32S ? _hist.getUMat() : UMat(BINS, 1, CV_32SC1); + hist = _hist.getUMat(); k1.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::PtrWriteOnly(ghist), (int)src.total()); @@ -1503,23 +1503,18 @@ static bool ocl_calcHist1(InputArray _src, OutputArray _hist, int ddepth = CV_32 if (!k1.run(1, &globalsize, &wgs, false)) return false; + char cvt[40]; ocl::Kernel k2("merge_histogram", ocl::imgproc::histogram_oclsrc, - format("-D BINS=%d -D HISTS_COUNT=%d -D WGS=%d", - BINS, compunits, (int)wgs)); + format("-D BINS=%d -D HISTS_COUNT=%d -D WGS=%d -D convertToHT=%s -D HT=%s", + BINS, compunits, (int)wgs, ocl::convertTypeStr(CV_32S, ddepth, 1, cvt), + ocl::typeToStr(ddepth))); if (k2.empty()) return false; k2.args(ocl::KernelArg::PtrReadOnly(ghist), - ocl::KernelArg::PtrWriteOnly(hist)); - if (!k2.run(1, &wgs, &wgs, false)) - return false; - - if (hist.depth() != ddepth) - hist.convertTo(_hist, ddepth); - else - _hist.getUMatRef() = hist; + ocl::KernelArg::WriteOnlyNoSize(hist)); - return true; + return k2.run(1, &wgs, &wgs, false); } static bool ocl_calcHist(InputArrayOfArrays images, OutputArray hist) diff --git a/modules/imgproc/src/opencl/histogram.cl b/modules/imgproc/src/opencl/histogram.cl index 05cd427..2161d3b 100644 --- a/modules/imgproc/src/opencl/histogram.cl +++ b/modules/imgproc/src/opencl/histogram.cl @@ -45,6 +45,8 @@ #define T uchar #endif +#define noconvert + __kernel void calculate_histogram(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols, __global uchar * histptr, int total) { @@ -111,10 +113,20 @@ __kernel void calculate_histogram(__global const uchar * src, int src_step, int hist[i] = localhist[i]; } -__kernel void merge_histogram(__global const int * ghist, __global int * hist) +#ifndef HT +#define HT int +#endif + +#ifndef convertToHT +#define convertToHT noconvert +#endif + +__kernel void merge_histogram(__global const int * ghist, __global uchar * histptr, int hist_step, int hist_offset) { int lid = get_local_id(0); + __global HT * hist = (__global HT *)(histptr + hist_offset); + #pragma unroll for (int i = lid; i < BINS; i += WGS) hist[i] = ghist[i]; @@ -126,7 +138,7 @@ __kernel void merge_histogram(__global const int * ghist, __global int * hist) ghist += BINS; #pragma unroll for (int j = lid; j < BINS; j += WGS) - hist[j] += ghist[j]; + hist[j] += convertToHT(ghist[j]); barrier(CLK_LOCAL_MEM_FENCE); } }