Merge pull request #1974 from ilya-lavrenov:tapi_equalizehist

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;

modules/imgproc/src/opencl/histogram.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.
+// 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);
+}

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);
     }
 }
 

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);
     }
 };