Merge branch 'add_prob_result_bayes_classifier', remote-tracking branch…

Merge branch 'add_prob_result_bayes_classifier', remote-tracking branch 'upstream/master' into add_prob_result_bayes_classifier

modules/core/include/opencv2/core/opencl/ocl_defs.hpp

@@ -22,6 +22,15 @@
             fflush(stdout);                                                 \
         }                                                                   \
     }
+#elif defined CV_OPENCL_RUN_ASSERT
+#define CV_OCL_RUN_(condition, func, ...)                                   \
+    {                                                                       \
+        if (cv::ocl::useOpenCL() && (condition))                            \
+        {                                                                   \
+            CV_Assert(func);                                                \
+            return;                                                         \
+        }                                                                   \
+    }
 #else
 #define CV_OCL_RUN_(condition, func, ...)                                   \
     if (cv::ocl::useOpenCL() && (condition) && func)                        \

modules/core/src/matrix.cpp

@@ -1650,6 +1650,16 @@ int _InputArray::dims(int i) const
         return vv[i].dims;
     }
 
+    if( k == STD_VECTOR_UMAT )
+    {
+        const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
+        if( i < 0 )
+            return 1;
+        CV_Assert( i < (int)vv.size() );
+
+        return vv[i].dims;
+    }
+
     if( k == OPENGL_BUFFER )
     {
         CV_Assert( i < 0 );
@@ -1701,6 +1711,16 @@ size_t _InputArray::total(int i) const
         return vv[i].total();
     }
 
+    if( k == STD_VECTOR_UMAT )
+    {
+        const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
+        if( i < 0 )
+            return vv.size();
+
+        CV_Assert( i < (int)vv.size() );
+        return vv[i].total();
+    }
+
     return size(i).area();
 }
 
@@ -1723,6 +1743,18 @@ int _InputArray::type(int i) const
     if( k == NONE )
         return -1;
 
+    if( k == STD_VECTOR_UMAT )
+    {
+        const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
+        if( vv.empty() )
+        {
+            CV_Assert((flags & FIXED_TYPE) != 0);
+            return CV_MAT_TYPE(flags);
+        }
+        CV_Assert( i < (int)vv.size() );
+        return vv[i >= 0 ? i : 0].type();
+    }
+
     if( k == STD_VECTOR_MAT )
     {
         const std::vector<Mat>& vv = *(const std::vector<Mat>*)obj;
@@ -1793,6 +1825,12 @@ bool _InputArray::empty() const
         return vv.empty();
     }
 
+    if( k == STD_VECTOR_UMAT )
+    {
+        const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
+        return vv.empty();
+    }
+
     if( k == OPENGL_BUFFER )
         return ((const ogl::Buffer*)obj)->empty();
 

modules/imgproc/perf/opencl/perf_accumulate.cpp

+/*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.
+//
+// @Authors
+//    Nathan, liujun@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.
+//
+//M*/
+
+#include "perf_precomp.hpp"
+#include "opencv2/ts/ocl_perf.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+/////////////////////////////////// Accumulate ///////////////////////////////////
+
+typedef Size_MatType AccumulateFixture;
+
+OCL_PERF_TEST_P(AccumulateFixture, Accumulate,
+                ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES))
+{
+    Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int srcType = get<1>(params), cn = CV_MAT_CN(srcType), dstType = CV_32FC(cn);
+
+    checkDeviceMaxMemoryAllocSize(srcSize, dstType);
+
+    UMat src(srcSize, srcType), dst(srcSize, dstType);
+    declare.in(src, dst, WARMUP_RNG).out(dst);
+
+    OCL_TEST_CYCLE() cv::accumulate(src, dst);
+
+    SANITY_CHECK_NOTHING();
+}
+
+/////////////////////////////////// AccumulateSquare ///////////////////////////////////
+
+typedef Size_MatType AccumulateSquareFixture;
+
+OCL_PERF_TEST_P(AccumulateSquareFixture, AccumulateSquare,
+                ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES))
+{
+    Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int srcType = get<1>(params), cn = CV_MAT_CN(srcType), dstType = CV_32FC(cn);
+
+    checkDeviceMaxMemoryAllocSize(srcSize, dstType);
+
+    UMat src(srcSize, srcType), dst(srcSize, dstType);
+    declare.in(src, dst, WARMUP_RNG);
+
+    OCL_TEST_CYCLE() cv::accumulateSquare(src, dst);
+
+    SANITY_CHECK_NOTHING();
+}
+
+/////////////////////////////////// AccumulateProduct ///////////////////////////////////
+
+typedef Size_MatType AccumulateProductFixture;
+
+OCL_PERF_TEST_P(AccumulateProductFixture, AccumulateProduct,
+                ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES))
+{
+    Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int srcType = get<1>(params), cn = CV_MAT_CN(srcType), dstType = CV_32FC(cn);
+
+    checkDeviceMaxMemoryAllocSize(srcSize, dstType);
+
+    UMat src1(srcSize, srcType), src2(srcSize, srcType), dst(srcSize, dstType);
+    declare.in(src1, src2, dst, WARMUP_RNG);
+
+    OCL_TEST_CYCLE() cv::accumulateProduct(src1, src2, dst);
+
+    SANITY_CHECK_NOTHING();
+}
+
+/////////////////////////////////// AccumulateWeighted ///////////////////////////////////
+
+typedef Size_MatType AccumulateWeightedFixture;
+
+OCL_PERF_TEST_P(AccumulateWeightedFixture, AccumulateWeighted,
+                ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES))
+{
+    Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int srcType = get<1>(params), cn = CV_MAT_CN(srcType), dstType = CV_32FC(cn);
+
+    checkDeviceMaxMemoryAllocSize(srcSize, dstType);
+
+    UMat src(srcSize, srcType), dst(srcSize, dstType);
+    declare.in(src, dst, WARMUP_RNG);
+
+    OCL_TEST_CYCLE() cv::accumulateWeighted(src, dst, 2.0);
+
+    SANITY_CHECK_NOTHING();
+}
+
+} } // namespace cvtest::ocl
+
+#endif

modules/imgproc/perf/opencl/perf_imgproc.cpp

@@ -61,6 +61,8 @@ OCL_PERF_TEST_P(EqualizeHistFixture, EqualizeHist, OCL_TEST_SIZES)
     const Size srcSize = GetParam();
     const double eps = 1;
 
+    checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1);
+
     UMat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1);
     declare.in(src, WARMUP_RNG).out(dst);
 
@@ -69,6 +71,30 @@ OCL_PERF_TEST_P(EqualizeHistFixture, EqualizeHist, OCL_TEST_SIZES)
     SANITY_CHECK(dst, eps);
 }
 
+///////////// calcHist ////////////////////////
+
+typedef TestBaseWithParam<Size> CalcHistFixture;
+
+OCL_PERF_TEST_P(CalcHistFixture, CalcHist, OCL_TEST_SIZES)
+{
+    const Size srcSize = GetParam();
+
+    const std::vector<int> channels(1, 0);
+    std::vector<float> ranges(2);
+    std::vector<int> histSize(1, 256);
+    ranges[0] = 0;
+    ranges[1] = 256;
+
+    checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1);
+
+    UMat src(srcSize, CV_8UC1), hist(256, 1, CV_32FC1);
+    declare.in(src, WARMUP_RNG).out(hist);
+
+    OCL_TEST_CYCLE() cv::calcHist(std::vector<UMat>(1, src), channels, noArray(), hist, histSize, ranges, false);
+
+    SANITY_CHECK(hist);
+}
+
 /////////// CopyMakeBorder //////////////////////
 
 CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101)
@@ -83,6 +109,8 @@ OCL_PERF_TEST_P(CopyMakeBorderFixture, CopyMakeBorder,
     const Size srcSize = get<0>(params);
     const int type = get<1>(params), borderType = get<2>(params);
 
+    checkDeviceMaxMemoryAllocSize(srcSize, type);
+
     UMat src(srcSize, type), dst;
     const Size dstSize = srcSize + Size(12, 12);
     dst.create(dstSize, type);
@@ -105,6 +133,8 @@ OCL_PERF_TEST_P(CornerMinEigenValFixture, CornerMinEigenVal,
     const int type = get<1>(params), borderType = BORDER_REFLECT;
     const int blockSize = 7, apertureSize = 1 + 2 * 3;
 
+    checkDeviceMaxMemoryAllocSize(srcSize, type);
+
     UMat src(srcSize, type), dst(srcSize, CV_32FC1);
     declare.in(src, WARMUP_RNG).out(dst);
 
@@ -124,6 +154,8 @@ OCL_PERF_TEST_P(CornerHarrisFixture, CornerHarris,
     const Size srcSize = get<0>(params);
     const int type = get<1>(params), borderType = BORDER_REFLECT;
 
+    checkDeviceMaxMemoryAllocSize(srcSize, type);
+
     UMat src(srcSize, type), dst(srcSize, CV_32FC1);
     declare.in(src, WARMUP_RNG).out(dst);
 
@@ -143,6 +175,8 @@ OCL_PERF_TEST_P(PreCornerDetectFixture, PreCornerDetect,
     const Size srcSize = get<0>(params);
     const int type = get<1>(params), borderType = BORDER_REFLECT;
 
+    checkDeviceMaxMemoryAllocSize(srcSize, type);
+
     UMat src(srcSize, type), dst(srcSize, CV_32FC1);
     declare.in(src, WARMUP_RNG).out(dst);
 
@@ -162,6 +196,8 @@ OCL_PERF_TEST_P(IntegralFixture, Integral1, ::testing::Combine(OCL_TEST_SIZES, O
     const Size srcSize = get<0>(params);
     const int ddepth = get<1>(params);
 
+    checkDeviceMaxMemoryAllocSize(srcSize, ddepth);
+
     UMat src(srcSize, CV_8UC1), dst(srcSize + Size(1, 1), ddepth);
     declare.in(src, WARMUP_RNG).out(dst);
 
@@ -186,6 +222,8 @@ OCL_PERF_TEST_P(ThreshFixture, Threshold,
     const int threshType = get<2>(params);
     const double maxValue = 220.0, threshold = 50;
 
+    checkDeviceMaxMemoryAllocSize(srcSize, srcType);
+
     UMat src(srcSize, srcType), dst(srcSize, srcType);
     declare.in(src, WARMUP_RNG).out(dst);
 
@@ -202,6 +240,8 @@ OCL_PERF_TEST_P(CLAHEFixture, CLAHE, OCL_TEST_SIZES)
 {
     const Size srcSize = GetParam();
 
+    checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1);
+
     UMat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1);
     const double clipLimit = 40.0;
     declare.in(src, WARMUP_RNG).out(dst);

modules/imgproc/src/accum.cpp

@@ -41,6 +41,7 @@
 //M*/
 
 #include "precomp.hpp"
+#include "opencl_kernels.hpp"
 
 namespace cv
 {
@@ -352,15 +353,83 @@ inline int getAccTabIdx(int sdepth, int ddepth)
            sdepth == CV_64F && ddepth == CV_64F ? 6 : -1;
 }
 
+#ifdef HAVE_OPENCL
+
+enum
+{
+    ACCUMULATE = 0,
+    ACCUMULATE_SQUARE = 1,
+    ACCUMULATE_PRODUCT = 2,
+    ACCUMULATE_WEIGHTED = 3
+};
+
+static bool ocl_accumulate( InputArray _src, InputArray _src2, InputOutputArray _dst, double alpha,
+                            InputArray _mask, int op_type )
+{
+    CV_Assert(op_type == ACCUMULATE || op_type == ACCUMULATE_SQUARE ||
+              op_type == ACCUMULATE_PRODUCT || op_type == ACCUMULATE_WEIGHTED);
+
+    int stype = _src.type(), cn = CV_MAT_CN(stype);
+    int sdepth = CV_MAT_DEPTH(stype), ddepth = _dst.depth();
+
+    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0,
+            haveMask = !_mask.empty();
+
+    if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F))
+        return false;
+
+    const char * const opMap[4] = { "ACCUMULATE", "ACCUMULATE_SQUARE", "ACCUMULATE_PRODUCT",
+                                   "ACCUMULATE_WEIGHTED" };
+
+    ocl::Kernel k("accumulate", ocl::imgproc::accumulate_oclsrc,
+                  format("-D %s%s -D srcT=%s -D cn=%d -D dstT=%s%s",
+                         opMap[op_type], haveMask ? " -D HAVE_MASK" : "",
+                         ocl::typeToStr(sdepth), cn, ocl::typeToStr(ddepth),
+                         doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
+    UMat src = _src.getUMat(), src2 = _src2.getUMat(), dst = _dst.getUMat(), mask = _mask.getUMat();
+
+    ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
+            src2arg = ocl::KernelArg::ReadOnlyNoSize(src2),
+            dstarg = ocl::KernelArg::ReadWrite(dst),
+            maskarg = ocl::KernelArg::ReadOnlyNoSize(mask);
+
+    int argidx = k.set(0, srcarg);
+    if (op_type == ACCUMULATE_PRODUCT)
+        argidx = k.set(argidx, src2arg);
+    argidx = k.set(argidx, dstarg);
+    if (op_type == ACCUMULATE_WEIGHTED)
+    {
+        if (ddepth == CV_32F)
+            argidx = k.set(argidx, (float)alpha);
+        else
+            argidx = k.set(argidx, alpha);
+    }
+    if (haveMask)
+        argidx = k.set(argidx, maskarg);
+
+    size_t globalsize[2] = { src.cols, src.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
+#endif
+
 }
 
 void cv::accumulate( InputArray _src, InputOutputArray _dst, InputArray _mask )
 {
-    Mat src = _src.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
-    int sdepth = src.depth(), ddepth = dst.depth(), cn = src.channels();
+    int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), scn = CV_MAT_CN(stype);
+    int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), dcn = CV_MAT_CN(dtype);
+
+    CV_Assert( _src.sameSize(_dst) && dcn == scn );
+    CV_Assert( _mask.empty() || (_src.sameSize(_mask) && _mask.type() == CV_8U) );
 
-    CV_Assert( dst.size == src.size && dst.channels() == cn );
-    CV_Assert( mask.empty() || (mask.size == src.size && mask.type() == CV_8U) );
+    CV_OCL_RUN(_src.dims() <= 2 && _dst.isUMat(),
+               ocl_accumulate(_src, noArray(), _dst, 0.0, _mask, ACCUMULATE))
+
+    Mat src = _src.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
 
     int fidx = getAccTabIdx(sdepth, ddepth);
     AccFunc func = fidx >= 0 ? accTab[fidx] : 0;
@@ -372,17 +441,21 @@ void cv::accumulate( InputArray _src, InputOutputArray _dst, InputArray _mask )
     int len = (int)it.size;
 
     for( size_t i = 0; i < it.nplanes; i++, ++it )
-        func(ptrs[0], ptrs[1], ptrs[2], len, cn);
+        func(ptrs[0], ptrs[1], ptrs[2], len, scn);
 }
 
-
 void cv::accumulateSquare( InputArray _src, InputOutputArray _dst, InputArray _mask )
 {
-    Mat src = _src.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
-    int sdepth = src.depth(), ddepth = dst.depth(), cn = src.channels();
+    int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), scn = CV_MAT_CN(stype);
+    int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), dcn = CV_MAT_CN(dtype);
+
+    CV_Assert( _src.sameSize(_dst) && dcn == scn );
+    CV_Assert( _mask.empty() || (_src.sameSize(_mask) && _mask.type() == CV_8U) );
 
-    CV_Assert( dst.size == src.size && dst.channels() == cn );
-    CV_Assert( mask.empty() || (mask.size == src.size && mask.type() == CV_8U) );
+    CV_OCL_RUN(_src.dims() <= 2 && _dst.isUMat(),
+               ocl_accumulate(_src, noArray(), _dst, 0.0, _mask, ACCUMULATE_SQUARE))
+
+    Mat src = _src.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
 
     int fidx = getAccTabIdx(sdepth, ddepth);
     AccFunc func = fidx >= 0 ? accSqrTab[fidx] : 0;
@@ -394,18 +467,23 @@ void cv::accumulateSquare( InputArray _src, InputOutputArray _dst, InputArray _m
     int len = (int)it.size;
 
     for( size_t i = 0; i < it.nplanes; i++, ++it )
-        func(ptrs[0], ptrs[1], ptrs[2], len, cn);
+        func(ptrs[0], ptrs[1], ptrs[2], len, scn);
 }
 
 void cv::accumulateProduct( InputArray _src1, InputArray _src2,
                             InputOutputArray _dst, InputArray _mask )
 {
-    Mat src1 = _src1.getMat(), src2 = _src2.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
-    int sdepth = src1.depth(), ddepth = dst.depth(), cn = src1.channels();
+    int stype = _src1.type(), sdepth = CV_MAT_DEPTH(stype), scn = CV_MAT_CN(stype);
+    int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), dcn = CV_MAT_CN(dtype);
+
+    CV_Assert( _src1.sameSize(_src2) && stype == _src2.type() );
+    CV_Assert( _src1.sameSize(_dst) && dcn == scn );
+    CV_Assert( _mask.empty() || (_src1.sameSize(_mask) && _mask.type() == CV_8U) );
 
-    CV_Assert( src2.size && src1.size && src2.type() == src1.type() );
-    CV_Assert( dst.size == src1.size && dst.channels() == cn );
-    CV_Assert( mask.empty() || (mask.size == src1.size && mask.type() == CV_8U) );
+    CV_OCL_RUN(_src1.dims() <= 2 && _dst.isUMat(),
+               ocl_accumulate(_src1, _src2, _dst, 0.0, _mask, ACCUMULATE_PRODUCT))
+
+    Mat src1 = _src1.getMat(), src2 = _src2.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
 
     int fidx = getAccTabIdx(sdepth, ddepth);
     AccProdFunc func = fidx >= 0 ? accProdTab[fidx] : 0;
@@ -417,18 +495,22 @@ void cv::accumulateProduct( InputArray _src1, InputArray _src2,
     int len = (int)it.size;
 
     for( size_t i = 0; i < it.nplanes; i++, ++it )
-        func(ptrs[0], ptrs[1], ptrs[2], ptrs[3], len, cn);
+        func(ptrs[0], ptrs[1], ptrs[2], ptrs[3], len, scn);
 }
 
-
 void cv::accumulateWeighted( InputArray _src, InputOutputArray _dst,
                              double alpha, InputArray _mask )
 {
-    Mat src = _src.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
-    int sdepth = src.depth(), ddepth = dst.depth(), cn = src.channels();
+    int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), scn = CV_MAT_CN(stype);
+    int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), dcn = CV_MAT_CN(dtype);
+
+    CV_Assert( _src.sameSize(_dst) && dcn == scn );
+    CV_Assert( _mask.empty() || (_src.sameSize(_mask) && _mask.type() == CV_8U) );
 
-    CV_Assert( dst.size == src.size && dst.channels() == cn );
-    CV_Assert( mask.empty() || (mask.size == src.size && mask.type() == CV_8U) );
+    CV_OCL_RUN(_src.dims() <= 2 && _dst.isUMat(),
+               ocl_accumulate(_src, noArray(), _dst, alpha, _mask, ACCUMULATE_WEIGHTED))
+
+    Mat src = _src.getMat(), dst = _dst.getMat(), mask = _mask.getMat();
 
     int fidx = getAccTabIdx(sdepth, ddepth);
     AccWFunc func = fidx >= 0 ? accWTab[fidx] : 0;
@@ -440,7 +522,7 @@ void cv::accumulateWeighted( InputArray _src, InputOutputArray _dst,
     int len = (int)it.size;
 
     for( size_t i = 0; i < it.nplanes; i++, ++it )
-        func(ptrs[0], ptrs[1], ptrs[2], len, cn, alpha);
+        func(ptrs[0], ptrs[1], ptrs[2], len, scn, alpha);
 }
 
 

modules/imgproc/src/histogram.cpp

@@ -1399,6 +1399,61 @@ static void calcHist( const Mat* images, int nimages, const int* channels,
     }
 }
 
+#ifdef HAVE_OPENCL
+
+enum
+{
+    BINS = 256
+};
+
+static bool ocl_calcHist1(InputArray _src, OutputArray _hist, int ddepth = CV_32S)
+{
+    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(BINS, 1, ddepth);
+    UMat src = _src.getUMat(), ghist(1, BINS * compunits, CV_32SC1),
+            hist = ddepth == CV_32S ? _hist.getUMat() : UMat(BINS, 1, 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));
+    if (!k2.run(1, &wgs, &wgs, false))
+        return false;
+
+    if (hist.depth() != ddepth)
+        hist.convertTo(_hist, ddepth);
+    else
+        _hist.getUMatRef() = hist;
+
+    return true;
+}
+
+static bool ocl_calcHist(InputArrayOfArrays images, OutputArray hist)
+{
+    std::vector<UMat> v;
+    images.getUMatVector(v);
+
+    return ocl_calcHist1(v[0], hist, CV_32F);
+}
+
+#endif
+
 }
 
 void cv::calcHist( const Mat* images, int nimages, const int* channels,
@@ -1417,6 +1472,12 @@ void cv::calcHist( InputArrayOfArrays images, const std::vector<int>& channels,
                    const std::vector<float>& ranges,
                    bool accumulate )
 {
+    CV_OCL_RUN(images.total() == 1 && channels.size() == 1 && images.channels(0) == 1 &&
+               channels[0] == 0 && images.isUMatVector() && mask.empty() && !accumulate &&
+               histSize.size() == 1 && histSize[0] == BINS && ranges.size() == 2 &&
+               ranges[0] == 0 && ranges[1] == 256,
+               ocl_calcHist(images, hist))
+
     int i, dims = (int)histSize.size(), rsz = (int)ranges.size(), csz = (int)channels.size();
     int nimages = (int)images.total();
 
@@ -3290,47 +3351,13 @@ CV_IMPL void cvEqualizeHist( const CvArr* srcarr, CvArr* 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 = std::min<size_t>(ocl::Device::getDefault().maxWorkGroupSize(), BINS);
 
     // calculation of histogram
     UMat hist;
-    if (!ocl_calcHist(_src, hist))
+    if (!ocl_calcHist1(_src, hist))
         return false;
 
     UMat lut(1, 256, CV_8UC1);

modules/imgproc/src/opencl/accumulate.cl

+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+// Copyright (C) 2014, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+__kernel void accumulate(__global const uchar * srcptr, int src_step, int src_offset,
+#ifdef ACCUMULATE_PRODUCT
+                         __global const uchar * src2ptr, int src2_step, int src2_offset,
+#endif
+                         __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols
+#ifdef ACCUMULATE_WEIGHTED
+                         , dstT alpha
+#endif
+#ifdef HAVE_MASK
+                         , __global const uchar * mask, int mask_step, int mask_offset
+#endif
+                         )
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < dst_cols && y < dst_rows)
+    {
+        int src_index = mad24(y, src_step, src_offset + x * cn * (int)sizeof(srcT));
+#ifdef HAVE_MASK
+        int mask_index = mad24(y, mask_step, mask_offset + x);
+        mask += mask_index;
+#endif
+        int dst_index = mad24(y, dst_step, dst_offset + x * cn * (int)sizeof(dstT));
+
+        __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+#ifdef ACCUMULATE_PRODUCT
+        int src2_index = mad24(y, src2_step, src2_offset + x * cn * (int)sizeof(srcT));
+        __global const srcT * src2 = (__global const srcT *)(src2ptr + src2_index);
+#endif
+        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+
+        #pragma unroll
+        for (int c = 0; c < cn; ++c)
+#ifdef HAVE_MASK
+            if (mask[0])
+#endif
+#ifdef ACCUMULATE
+                dst[c] += src[c];
+#elif defined ACCUMULATE_SQUARE
+                dst[c] += src[c] * src[c];
+#elif defined ACCUMULATE_PRODUCT
+                dst[c] += src[c] * src2[c];
+#elif defined ACCUMULATE_WEIGHTED
+                dst[c] = (1 - alpha) * dst[c] + src[c] * alpha;
+#else
+#error "Unknown accumulation type"
+#endif
+    }
+}

modules/imgproc/test/ocl/test_accumulate.cpp

+/*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.
+//
+// @Authors
+//    Nathan, liujun@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.
+//
+//M*/
+
+#include "test_precomp.hpp"
+#include "cvconfig.h"
+#include "opencv2/ts/ocl_test.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+PARAM_TEST_CASE(AccumulateBase, std::pair<MatDepth, MatDepth>, Channels, bool)
+{
+    int sdepth, ddepth, channels;
+    bool useRoi;
+    double alpha;
+
+    TEST_DECLARE_INPUT_PARAMETER(src)
+    TEST_DECLARE_INPUT_PARAMETER(mask)
+    TEST_DECLARE_INPUT_PARAMETER(src2)
+    TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+    virtual void SetUp()
+    {
+        const std::pair<MatDepth, MatDepth> depths = GET_PARAM(0);
+        sdepth = depths.first, ddepth = depths.second;
+        channels = GET_PARAM(1);
+        useRoi = GET_PARAM(2);
+    }
+
+    void random_roi()
+    {
+        const int stype = CV_MAKE_TYPE(sdepth, channels),
+                dtype = CV_MAKE_TYPE(ddepth, channels);
+
+        Size roiSize = randomSize(1, 10);
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, stype, -MAX_VALUE, MAX_VALUE);
+
+        Border maskBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(mask, mask_roi, roiSize, maskBorder, CV_8UC1, -MAX_VALUE, MAX_VALUE);
+        threshold(mask, mask, 80, 255, THRESH_BINARY);
+
+        Border src2Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src2, src2_roi, roiSize, src2Border, stype, -MAX_VALUE, MAX_VALUE);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, roiSize, dstBorder, dtype, -MAX_VALUE, MAX_VALUE);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src)
+        UMAT_UPLOAD_INPUT_PARAMETER(mask)
+        UMAT_UPLOAD_INPUT_PARAMETER(src2)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+
+        alpha = randomDouble(-5, 5);
+    }
+};
+
+/////////////////////////////////// Accumulate ///////////////////////////////////
+
+typedef AccumulateBase Accumulate;
+
+OCL_TEST_P(Accumulate, Mat)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulate(src_roi, dst_roi));
+        OCL_ON(cv::accumulate(usrc_roi, udst_roi));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-6);
+    }
+}
+
+OCL_TEST_P(Accumulate, Mask)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulate(src_roi, dst_roi, mask_roi));
+        OCL_ON(cv::accumulate(usrc_roi, udst_roi, umask_roi));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-6);
+    }
+}
+
+/////////////////////////////////// AccumulateSquare ///////////////////////////////////
+
+typedef AccumulateBase AccumulateSquare;
+
+OCL_TEST_P(AccumulateSquare, Mat)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulateSquare(src_roi, dst_roi));
+        OCL_ON(cv::accumulateSquare(usrc_roi, udst_roi));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-2);
+    }
+}
+
+OCL_TEST_P(AccumulateSquare, Mask)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulateSquare(src_roi, dst_roi, mask_roi));
+        OCL_ON(cv::accumulateSquare(usrc_roi, udst_roi, umask_roi));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-2);
+    }
+}
+
+/////////////////////////////////// AccumulateProduct ///////////////////////////////////
+
+typedef AccumulateBase AccumulateProduct;
+
+OCL_TEST_P(AccumulateProduct, Mat)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulateProduct(src_roi, src2_roi, dst_roi));
+        OCL_ON(cv::accumulateProduct(usrc_roi, usrc2_roi, udst_roi));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-2);
+    }
+}
+
+OCL_TEST_P(AccumulateProduct, Mask)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulateProduct(src_roi, src2_roi, dst_roi, mask_roi));
+        OCL_ON(cv::accumulateProduct(usrc_roi, usrc2_roi, udst_roi, umask_roi));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-2);
+    }
+}
+
+/////////////////////////////////// AccumulateWeighted ///////////////////////////////////
+
+typedef AccumulateBase AccumulateWeighted;
+
+OCL_TEST_P(AccumulateWeighted, Mat)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulateWeighted(src_roi, dst_roi, alpha));
+        OCL_ON(cv::accumulateWeighted(usrc_roi, udst_roi, alpha));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-2);
+    }
+}
+
+OCL_TEST_P(AccumulateWeighted, Mask)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        random_roi();
+
+        OCL_OFF(cv::accumulateWeighted(src_roi, dst_roi, alpha));
+        OCL_ON(cv::accumulateWeighted(usrc_roi, udst_roi, alpha));
+
+        OCL_EXPECT_MATS_NEAR(dst, 1e-2);
+    }
+}
+
+/////////////////////////////////// Instantiation ///////////////////////////////////
+
+#define OCL_DEPTH_ALL_COMBINATIONS \
+    testing::Values(std::make_pair<MatDepth, MatDepth>(CV_8U, CV_32F), \
+    std::make_pair<MatDepth, MatDepth>(CV_16U, CV_32F), \
+    std::make_pair<MatDepth, MatDepth>(CV_32F, CV_32F), \
+    std::make_pair<MatDepth, MatDepth>(CV_8U, CV_64F), \
+    std::make_pair<MatDepth, MatDepth>(CV_16U, CV_64F), \
+    std::make_pair<MatDepth, MatDepth>(CV_32F, CV_64F), \
+    std::make_pair<MatDepth, MatDepth>(CV_64F, CV_64F))
+
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, Accumulate, Combine(OCL_DEPTH_ALL_COMBINATIONS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, AccumulateSquare, Combine(OCL_DEPTH_ALL_COMBINATIONS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, AccumulateProduct, Combine(OCL_DEPTH_ALL_COMBINATIONS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(ImgProc, AccumulateWeighted, Combine(OCL_DEPTH_ALL_COMBINATIONS, OCL_ALL_CHANNELS, Bool()));
+
+} } // namespace cvtest::ocl
+
+#endif

modules/imgproc/test/ocl/test_blend.cpp

@@ -75,7 +75,7 @@ PARAM_TEST_CASE(BlendLinear, MatDepth, Channels, bool)
         const int type = CV_MAKE_TYPE(depth, channels);
         const double upValue = 256;
 
-        Size roiSize = randomSize(1, 20);
+        Size roiSize = randomSize(1, MAX_VALUE);
         Border src1Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
         randomSubMat(src1, src1_roi, roiSize, src1Border, type, -upValue, upValue);
 
@@ -104,8 +104,7 @@ PARAM_TEST_CASE(BlendLinear, MatDepth, Channels, bool)
 
     void Near(double eps = 0.0)
     {
-        EXPECT_MAT_NEAR(dst, udst, eps);
-        EXPECT_MAT_NEAR(dst_roi, udst_roi, eps);
+        OCL_EXPECT_MATS_NEAR(dst, eps)
     }
 };
 

modules/imgproc/test/ocl/test_histogram.cpp

@@ -144,11 +144,6 @@ PARAM_TEST_CASE(CalcBackProject, MatDepth, int, bool)
 
         scale = randomDouble(0.1, 1);
     }
-
-    void Near()
-    {
-        OCL_EXPECT_MATS_NEAR(dst, 0.0)
-    }
 };
 
 //////////////////////////////// CalcBackProject //////////////////////////////////////////////
@@ -162,13 +157,62 @@ OCL_TEST_P(CalcBackProject, Mat)
         OCL_OFF(cv::calcBackProject(images_roi, channels, hist_roi, dst_roi, ranges, scale));
         OCL_ON(cv::calcBackProject(uimages_roi, channels, uhist_roi, udst_roi, ranges, scale));
 
-        Near();
+        OCL_EXPECT_MATS_NEAR(dst, 0.0)
+    }
+}
+
+//////////////////////////////// CalcHist //////////////////////////////////////////////
+
+PARAM_TEST_CASE(CalcHist, bool)
+{
+    bool useRoi;
+
+    TEST_DECLARE_INPUT_PARAMETER(src)
+    TEST_DECLARE_OUTPUT_PARAMETER(hist)
+
+    virtual void SetUp()
+    {
+        useRoi = GET_PARAM(0);
+    }
+
+    virtual void random_roi()
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+
+        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src, src_roi, roiSize, srcBorder, CV_8UC1, 0, 256);
+
+        Border histBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(hist, hist_roi, Size(1, 256), histBorder, CV_32SC1, 0, MAX_VALUE);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(hist)
+    }
+};
+
+OCL_TEST_P(CalcHist, Mat)
+{
+    const std::vector<int> channels(1, 0);
+    std::vector<float> ranges(2);
+    std::vector<int> histSize(1, 256);
+    ranges[0] = 0;
+    ranges[1] = 256;
+
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        random_roi();
+
+        OCL_OFF(cv::calcHist(std::vector<Mat>(1, src_roi), channels, noArray(), hist_roi, histSize, ranges, false));
+        OCL_ON(cv::calcHist(std::vector<UMat>(1, usrc_roi), channels, noArray(), uhist_roi, histSize, ranges, false));
+
+        OCL_EXPECT_MATS_NEAR(hist, 0.0)
     }
 }
 
 /////////////////////////////////////////////////////////////////////////////////////
 
 OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CalcBackProject, Combine(Values((MatDepth)CV_8U), Values(1, 2), Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CalcHist, Values(true, false));
 
 } } // namespace cvtest::ocl