added gpu::equalizeHist function

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -1092,6 +1092,11 @@ namespace cv
         //! Output hist will have one row, 256 cols and CV32SC1 type.
         CV_EXPORTS void calcHist(const GpuMat& src, GpuMat& hist, Stream& stream = Stream::Null());
         CV_EXPORTS void calcHist(const GpuMat& src, GpuMat& hist, GpuMat& buf, Stream& stream = Stream::Null());
+        
+        //! normalizes the grayscale image brightness and contrast by normalizing its histogram
+        CV_EXPORTS void equalizeHist(const GpuMat& src, GpuMat& dst, Stream& stream = Stream::Null());
+        CV_EXPORTS void equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, Stream& stream = Stream::Null());
+        CV_EXPORTS void equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, GpuMat& buf, Stream& stream = Stream::Null());
 
         //////////////////////////////// StereoBM_GPU ////////////////////////////////
 

modules/gpu/src/cuda/hist.cu

@@ -190,4 +190,34 @@ namespace cv { namespace gpu { namespace histograms
         if (stream == 0)
             cudaSafeCall( cudaDeviceSynchronize() );
     }
+
+    __global__ void equalizeHist(DevMem2D src, PtrStep dst, const int* lut)
+    {
+        __shared__ int s_lut[256];
+
+        const int tid = threadIdx.y * blockDim.x + threadIdx.x;
+
+        s_lut[tid] = lut[tid];
+        __syncthreads();
+
+        const int x = blockIdx.x * blockDim.x + threadIdx.x;
+        const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+        if (x < src.cols && y < src.rows)
+        {
+            dst.ptr(y)[x] = __float2int_rn(255.0f * s_lut[src.ptr(y)[x]] / (src.cols * src.rows));
+        }
+    }
+
+    void equalizeHist_gpu(DevMem2D src, DevMem2D dst, const int* lut, cudaStream_t stream)
+    {
+        dim3 block(16, 16);
+        dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
+
+        equalizeHist<<<grid, block, 0, stream>>>(src, dst, lut);
+        cudaSafeCall( cudaGetLastError() );
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
+    }
 }}}

modules/gpu/src/imgproc_gpu.cpp

@@ -73,6 +73,9 @@ void cv::gpu::histRange(const GpuMat&, GpuMat&, const GpuMat&, Stream&) { throw_
 void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*, Stream&) { throw_nogpu(); }
 void cv::gpu::calcHist(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
 void cv::gpu::calcHist(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
+void cv::gpu::equalizeHist(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
+void cv::gpu::equalizeHist(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
+void cv::gpu::equalizeHist(const GpuMat&, GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
 void cv::gpu::cornerHarris(const GpuMat&, GpuMat&, int, int, double, int) { throw_nogpu(); }
 void cv::gpu::cornerMinEigenVal(const GpuMat&, GpuMat&, int, int, int) { throw_nogpu(); }
 void cv::gpu::mulSpectrums(const GpuMat&, const GpuMat&, GpuMat&, int, bool) { throw_nogpu(); }
@@ -1066,6 +1069,57 @@ void cv::gpu::calcHist(const GpuMat& src, GpuMat& hist, GpuMat& buf, Stream& str
     histogram256_gpu(src, hist.ptr<int>(), buf.ptr<unsigned int>(), StreamAccessor::getStream(stream));
 }
 
+void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, Stream& stream)
+{
+    GpuMat hist;
+    GpuMat buf;
+    equalizeHist(src, dst, hist, buf, stream);
+}
+
+void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, Stream& stream)
+{
+    GpuMat buf;
+    equalizeHist(src, dst, hist, buf, stream);
+}
+
+namespace cv { namespace gpu { namespace histograms
+{
+    void equalizeHist_gpu(DevMem2D src, DevMem2D dst, const int* lut, cudaStream_t stream);
+}}}
+
+void cv::gpu::equalizeHist(const GpuMat& src, GpuMat& dst, GpuMat& hist, GpuMat& buf, Stream& s)
+{
+    using namespace cv::gpu::histograms;
+
+    CV_Assert(src.type() == CV_8UC1);
+
+    dst.create(src.size(), src.type());
+
+    int intBufSize;
+    nppSafeCall( nppsIntegralGetBufferSize_32s(256, &intBufSize) );
+
+    int bufSize = std::max(256 * 240 * sizeof(int), intBufSize + 256 * sizeof(int));
+
+    ensureSizeIsEnough(1, bufSize, CV_8UC1, buf);
+
+    GpuMat histBuf(1, 256 * 240, CV_32SC1, buf.ptr());
+    GpuMat intBuf(1, intBufSize, CV_8UC1, buf.ptr());
+    GpuMat lut(1, 256, CV_32S, buf.ptr() + intBufSize);
+
+    calcHist(src, hist, histBuf, s);
+
+    cudaStream_t stream = StreamAccessor::getStream(s);
+
+    NppStreamHandler h(stream);
+
+    nppSafeCall( nppsIntegral_32s(hist.ptr<Npp32s>(), lut.ptr<Npp32s>(), 256, intBuf.ptr<Npp8u>()) );
+    
+    if (stream == 0)
+        cudaSafeCall( cudaDeviceSynchronize() );
+
+    equalizeHist_gpu(src, dst, lut.ptr<int>(), stream);
+}
+
 ////////////////////////////////////////////////////////////////////////
 // cornerHarris & minEgenVal
 

modules/gpu/test/test_imgproc.cpp

@@ -1094,6 +1094,50 @@ TEST_P(CalcHist, Accuracy)
 
 INSTANTIATE_TEST_CASE_P(ImgProc, CalcHist, testing::ValuesIn(devices()));
 
+struct EqualizeHist : testing::TestWithParam<cv::gpu::DeviceInfo>
+{
+    cv::gpu::DeviceInfo devInfo;
+
+    cv::Size size;
+    cv::Mat src;
+    cv::Mat dst_gold;
+    
+    virtual void SetUp()
+    {
+        devInfo = GetParam();
+
+        cv::gpu::setDevice(devInfo.deviceID());
+
+        cv::RNG& rng = cvtest::TS::ptr()->get_rng();
+
+        size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));
+        
+        src = cvtest::randomMat(rng, size, CV_8UC1, 0, 255, false);
+
+        cv::equalizeHist(src, dst_gold);
+    }
+};
+
+TEST_P(EqualizeHist, Accuracy)
+{
+    PRINT_PARAM(devInfo);
+    PRINT_PARAM(size);
+
+    cv::Mat dst;
+    
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat gpuDst;
+
+        cv::gpu::equalizeHist(cv::gpu::GpuMat(src), gpuDst);
+
+        gpuDst.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 3.0);
+}
+
+INSTANTIATE_TEST_CASE_P(ImgProc, EqualizeHist, testing::ValuesIn(devices()));
+
 ///////////////////////////////////////////////////////////////////////////////////////////////////////
 // cornerHarris