Merge branch 'master' of https://github.com/Itseez/opencv

modules/core/doc/basic_structures.rst

@@ -2981,7 +2981,7 @@ The class provides the following features for all derived classes:
 
     * so called "virtual constructor". That is, each Algorithm derivative is registered at program start and you can get the list of registered algorithms and create instance of a particular algorithm by its name (see ``Algorithm::create``). If you plan to add your own algorithms, it is good practice to add a unique prefix to your algorithms to distinguish them from other algorithms.
 
-    * setting/retrieving algorithm parameters by name. If you used video capturing functionality from OpenCV highgui module, you are probably familar with ``cvSetCaptureProperty()``, ``cvGetCaptureProperty()``, ``VideoCapture::set()`` and ``VideoCapture::get()``. ``Algorithm`` provides similar method where instead of integer id's you specify the parameter names as text strings. See ``Algorithm::set`` and ``Algorithm::get`` for details.
+    * setting/retrieving algorithm parameters by name. If you used video capturing functionality from OpenCV videoio module, you are probably familar with ``cvSetCaptureProperty()``, ``cvGetCaptureProperty()``, ``VideoCapture::set()`` and ``VideoCapture::get()``. ``Algorithm`` provides similar method where instead of integer id's you specify the parameter names as text strings. See ``Algorithm::set`` and ``Algorithm::get`` for details.
 
     * reading and writing parameters from/to XML or YAML files. Every Algorithm derivative can store all its parameters and then read them back. There is no need to re-implement it each time.
 

modules/core/doc/drawing_functions.rst

@@ -361,6 +361,37 @@ The function ``line`` draws the line segment between ``pt1`` and ``pt2`` points
 Antialiased lines are drawn using Gaussian filtering.
 
 
+arrowedLine
+----------------
+Draws a arrow segment pointing from the first point to the second one.
+
+.. ocv:function:: void arrowedLine(InputOutputArray img, Point pt1, Point pt2, const Scalar& color,          int thickness=1, int lineType=8, int shift=0, double tipLength=0.1)
+
+    :param img: Image.
+
+    :param pt1: The point the arrow starts from.
+
+    :param pt2: The point the arrow points to.
+
+    :param color: Line color.
+
+    :param thickness: Line thickness.
+
+    :param lineType: Type of the line:
+
+            * **8** (or omitted) - 8-connected line.
+
+            * **4** - 4-connected line.
+
+            * **CV_AA** - antialiased line.
+
+    :param shift: Number of fractional bits in the point coordinates.
+
+    :param tipLength: The length of the arrow tip in relation to the arrow length
+
+The function ``arrowedLine`` draws an arrow between ``pt1`` and ``pt2`` points in the image. See also :ocv:func:`line`.
+
+
 LineIterator
 ------------
 .. ocv:class:: LineIterator

modules/core/doc/intro.rst

@@ -14,7 +14,8 @@ OpenCV has a modular structure, which means that the package includes several sh
  * **calib3d** - basic multiple-view geometry algorithms, single and stereo camera calibration, object pose estimation, stereo correspondence algorithms, and elements of 3D reconstruction.
  * **features2d** - salient feature detectors, descriptors, and descriptor matchers.
  * **objdetect** - detection of objects and instances of the predefined classes (for example, faces, eyes, mugs, people, cars, and so on).
- * **highgui** - an easy-to-use interface to video capturing, image and video codecs, as well as simple UI capabilities.
+ * **highgui** - an easy-to-use interface to simple UI capabilities.
+ * **videoio** - an easy-to-use interface to video capturing and video codecs.
  * **gpu** - GPU-accelerated algorithms from different OpenCV modules.
  * ... some other helper modules, such as FLANN and Google test wrappers, Python bindings, and others.
 

modules/core/include/opencv2/core.hpp

@@ -510,6 +510,10 @@ CV_EXPORTS_W void randShuffle(InputOutputArray dst, double iterFactor = 1., RNG*
 CV_EXPORTS_W void line(InputOutputArray img, Point pt1, Point pt2, const Scalar& color,
                      int thickness = 1, int lineType = LINE_8, int shift = 0);
 
+//! draws an arrow from pt1 to pt2 in the image
+CV_EXPORTS_W void arrowedLine(InputOutputArray img, Point pt1, Point pt2, const Scalar& color,
+                     int thickness=1, int line_type=8, int shift=0, double tipLength=0.1);
+
 //! draws the rectangle outline or a solid rectangle with the opposite corners pt1 and pt2 in the image
 CV_EXPORTS_W void rectangle(InputOutputArray img, Point pt1, Point pt2,
                           const Scalar& color, int thickness = 1,

modules/core/include/opencv2/core/mat.hpp

@@ -360,7 +360,7 @@ struct CV_EXPORTS UMatData
 {
     enum { COPY_ON_MAP=1, HOST_COPY_OBSOLETE=2,
         DEVICE_COPY_OBSOLETE=4, TEMP_UMAT=8, TEMP_COPIED_UMAT=24,
-        USER_ALLOCATED=32 };
+        USER_ALLOCATED=32, DEVICE_MEM_MAPPED=64};
     UMatData(const MatAllocator* allocator);
     ~UMatData();
 
@@ -370,11 +370,13 @@ struct CV_EXPORTS UMatData
 
     bool hostCopyObsolete() const;
     bool deviceCopyObsolete() const;
+    bool deviceMemMapped() const;
     bool copyOnMap() const;
     bool tempUMat() const;
     bool tempCopiedUMat() const;
     void markHostCopyObsolete(bool flag);
     void markDeviceCopyObsolete(bool flag);
+    void markDeviceMemMapped(bool flag);
 
     const MatAllocator* prevAllocator;
     const MatAllocator* currAllocator;

modules/core/include/opencv2/core/mat.inl.hpp

@@ -3350,10 +3350,19 @@ size_t UMat::total() const
 
 inline bool UMatData::hostCopyObsolete() const { return (flags & HOST_COPY_OBSOLETE) != 0; }
 inline bool UMatData::deviceCopyObsolete() const { return (flags & DEVICE_COPY_OBSOLETE) != 0; }
+inline bool UMatData::deviceMemMapped() const { return (flags & DEVICE_MEM_MAPPED) != 0; }
 inline bool UMatData::copyOnMap() const { return (flags & COPY_ON_MAP) != 0; }
 inline bool UMatData::tempUMat() const { return (flags & TEMP_UMAT) != 0; }
 inline bool UMatData::tempCopiedUMat() const { return (flags & TEMP_COPIED_UMAT) == TEMP_COPIED_UMAT; }
 
+inline void UMatData::markDeviceMemMapped(bool flag)
+{
+  if(flag)
+    flags |= DEVICE_MEM_MAPPED;
+  else
+    flags &= ~DEVICE_MEM_MAPPED;
+}
+
 inline void UMatData::markHostCopyObsolete(bool flag)
 {
     if(flag)

modules/core/perf/opencl/perf_matop.cpp

@@ -139,6 +139,7 @@ OCL_PERF_TEST_P(CopyToFixture, CopyToWithMaskUninit,
         dst.release();
         startTimer();
         src.copyTo(dst, mask);
+        cv::ocl::finish();
         stopTimer();
     }
 

modules/core/src/arithm.cpp

@@ -2980,8 +2980,187 @@ void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op)
 namespace cv
 {
 
-template<typename T> static void
-inRange_(const T* src1, size_t step1, const T* src2, size_t step2,
+template <typename T>
+struct InRange_SSE
+{
+    int operator () (const T *, const T *, const T *, uchar *, int) const
+    {
+        return 0;
+    }
+};
+
+#if CV_SSE2
+
+template <>
+struct InRange_SSE<uchar>
+{
+    int operator () (const uchar * src1, const uchar * src2, const uchar * src3,
+                     uchar * dst, int len) const
+    {
+        int x = 0;
+
+        if (USE_SSE2)
+        {
+            __m128i v_full = _mm_set1_epi8(-1), v_128 = _mm_set1_epi8(-128);
+
+            for ( ; x <= len - 16; x += 16 )
+            {
+                __m128i v_src = _mm_add_epi8(_mm_loadu_si128((const __m128i *)(src1 + x)), v_128);
+                __m128i v_mask1 = _mm_cmpgt_epi8(_mm_add_epi8(_mm_loadu_si128((const __m128i *)(src2 + x)), v_128), v_src);
+                __m128i v_mask2 = _mm_cmpgt_epi8(v_src, _mm_add_epi8(_mm_loadu_si128((const __m128i *)(src3 + x)), v_128));
+                _mm_storeu_si128((__m128i *)(dst + x), _mm_andnot_si128(_mm_or_si128(v_mask1, v_mask2), v_full));
+            }
+        }
+
+        return x;
+    }
+};
+
+template <>
+struct InRange_SSE<schar>
+{
+    int operator () (const schar * src1, const schar * src2, const schar * src3,
+                     uchar * dst, int len) const
+    {
+        int x = 0;
+
+        if (USE_SSE2)
+        {
+            __m128i v_full = _mm_set1_epi8(-1);
+
+            for ( ; x <= len - 16; x += 16 )
+            {
+                __m128i v_src = _mm_loadu_si128((const __m128i *)(src1 + x));
+                __m128i v_mask1 = _mm_cmpgt_epi8(_mm_loadu_si128((const __m128i *)(src2 + x)), v_src);
+                __m128i v_mask2 = _mm_cmpgt_epi8(v_src, _mm_loadu_si128((const __m128i *)(src3 + x)));
+                _mm_storeu_si128((__m128i *)(dst + x), _mm_andnot_si128(_mm_or_si128(v_mask1, v_mask2), v_full));
+            }
+        }
+
+        return x;
+    }
+};
+
+template <>
+struct InRange_SSE<ushort>
+{
+    int operator () (const ushort * src1, const ushort * src2, const ushort * src3,
+                     uchar * dst, int len) const
+    {
+        int x = 0;
+
+        if (USE_SSE2)
+        {
+            __m128i v_zero = _mm_setzero_si128(), v_full = _mm_set1_epi16(-1), v_32768 = _mm_set1_epi16(-32768);
+
+            for ( ; x <= len - 8; x += 8 )
+            {
+                __m128i v_src = _mm_add_epi16(_mm_loadu_si128((const __m128i *)(src1 + x)), v_32768);
+                __m128i v_mask1 = _mm_cmpgt_epi16(_mm_add_epi16(_mm_loadu_si128((const __m128i *)(src2 + x)), v_32768), v_src);
+                __m128i v_mask2 = _mm_cmpgt_epi16(v_src, _mm_add_epi16(_mm_loadu_si128((const __m128i *)(src3 + x)), v_32768));
+                __m128i v_res = _mm_andnot_si128(_mm_or_si128(v_mask1, v_mask2), v_full);
+                _mm_storel_epi64((__m128i *)(dst + x), _mm_packus_epi16(_mm_srli_epi16(v_res, 8), v_zero));
+            }
+        }
+
+        return x;
+    }
+};
+
+template <>
+struct InRange_SSE<short>
+{
+    int operator () (const short * src1, const short * src2, const short * src3,
+                     uchar * dst, int len) const
+    {
+        int x = 0;
+
+        if (USE_SSE2)
+        {
+            __m128i v_zero = _mm_setzero_si128(), v_full = _mm_set1_epi16(-1);
+
+            for ( ; x <= len - 8; x += 8 )
+            {
+                __m128i v_src = _mm_loadu_si128((const __m128i *)(src1 + x));
+                __m128i v_mask1 = _mm_cmpgt_epi16(_mm_loadu_si128((const __m128i *)(src2 + x)), v_src);
+                __m128i v_mask2 = _mm_cmpgt_epi16(v_src, _mm_loadu_si128((const __m128i *)(src3 + x)));
+                __m128i v_res = _mm_andnot_si128(_mm_or_si128(v_mask1, v_mask2), v_full);
+                _mm_storel_epi64((__m128i *)(dst + x), _mm_packus_epi16(_mm_srli_epi16(v_res, 8), v_zero));
+            }
+        }
+
+        return x;
+    }
+};
+
+template <>
+struct InRange_SSE<int>
+{
+    int operator () (const int * src1, const int * src2, const int * src3,
+                     uchar * dst, int len) const
+    {
+        int x = 0;
+
+        if (USE_SSE2)
+        {
+            __m128i v_zero = _mm_setzero_si128(), v_full = _mm_set1_epi32(-1);
+
+            for ( ; x <= len - 8; x += 8 )
+            {
+                __m128i v_src = _mm_loadu_si128((const __m128i *)(src1 + x));
+                __m128i v_res1 = _mm_or_si128(_mm_cmpgt_epi32(_mm_loadu_si128((const __m128i *)(src2 + x)), v_src),
+                    _mm_cmpgt_epi32(v_src, _mm_loadu_si128((const __m128i *)(src3 + x))));
+
+                v_src = _mm_loadu_si128((const __m128i *)(src1 + x + 4));
+                __m128i v_res2 = _mm_or_si128(_mm_cmpgt_epi32(_mm_loadu_si128((const __m128i *)(src2 + x + 4)), v_src),
+                    _mm_cmpgt_epi32(v_src, _mm_loadu_si128((const __m128i *)(src3 + x + 4))));
+
+                __m128i v_res = _mm_packs_epi32(_mm_srli_epi32(_mm_andnot_si128(v_res1, v_full), 16),
+                                                _mm_srli_epi32(_mm_andnot_si128(v_res2, v_full), 16));
+                _mm_storel_epi64((__m128i *)(dst + x), _mm_packus_epi16(v_res, v_zero));
+            }
+        }
+
+        return x;
+    }
+};
+
+template <>
+struct InRange_SSE<float>
+{
+    int operator () (const float * src1, const float * src2, const float * src3,
+                     uchar * dst, int len) const
+    {
+        int x = 0;
+
+        if (USE_SSE2)
+        {
+            __m128i v_zero = _mm_setzero_si128();
+
+            for ( ; x <= len - 8; x += 8 )
+            {
+                __m128 v_src = _mm_loadu_ps(src1 + x);
+                __m128 v_res1 = _mm_and_ps(_mm_cmple_ps(_mm_loadu_ps(src2 + x), v_src),
+                    _mm_cmple_ps(v_src, _mm_loadu_ps(src3 + x)));
+
+                v_src = _mm_loadu_ps(src1 + x + 4);
+                __m128 v_res2 = _mm_and_ps(_mm_cmple_ps(_mm_loadu_ps(src2 + x + 4), v_src),
+                    _mm_cmple_ps(v_src, _mm_loadu_ps(src3 + x + 4)));
+
+                __m128i v_res1i = _mm_cvtps_epi32(v_res1), v_res2i = _mm_cvtps_epi32(v_res2);
+                __m128i v_res = _mm_packs_epi32(_mm_srli_epi32(v_res1i, 16), _mm_srli_epi32(v_res2i, 16));
+                _mm_storel_epi64((__m128i *)(dst + x), _mm_packus_epi16(v_res, v_zero));
+            }
+        }
+
+        return x;
+    }
+};
+
+#endif
+
+template <typename T>
+static void inRange_(const T* src1, size_t step1, const T* src2, size_t step2,
          const T* src3, size_t step3, uchar* dst, size_t step,
          Size size)
 {
@@ -2989,9 +3168,11 @@ inRange_(const T* src1, size_t step1, const T* src2, size_t step2,
     step2 /= sizeof(src2[0]);
     step3 /= sizeof(src3[0]);
 
+    InRange_SSE<T> vop;
+
     for( ; size.height--; src1 += step1, src2 += step2, src3 += step3, dst += step )
     {
-        int x = 0;
+        int x = vop(src1, src2, src3, dst, size.width);
         #if CV_ENABLE_UNROLLED
         for( ; x <= size.width - 4; x += 4 )
         {

modules/core/src/convert.cpp

@@ -1729,22 +1729,18 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
     UMat src = _src.getUMat(), lut = _lut.getUMat();
     _dst.create(src.size(), CV_MAKETYPE(ddepth, dcn));
     UMat dst = _dst.getUMat();
-    bool bAligned = (1 == lcn) && (0 == (src.offset % 4)) && (0 == ((dcn * src.cols) % 4));
-    // dst.cols == src.cols by params of dst.create
+    int kercn = lcn == 1 ? std::min(4, ocl::predictOptimalVectorWidth(_dst)) : dcn;
 
     ocl::Kernel k("LUT", ocl::core::lut_oclsrc,
-                  format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", bAligned ? 4 : dcn, lcn,
-                         ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth)
-                         ));
+                  format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", kercn, lcn,
+                         ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth)));
     if (k.empty())
         return false;
 
-    int cols = bAligned ? dcn * dst.cols / 4 : dst.cols;
-
     k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
-        ocl::KernelArg::WriteOnlyNoSize(dst), dst.rows, cols);
+        ocl::KernelArg::WriteOnly(dst, dcn, kercn));
 
-    size_t globalSize[2] = { cols, (dst.rows + 3) / 4 };
+    size_t globalSize[2] = { dst.cols * dcn / kercn, (dst.rows + 3) / 4 };
     return k.run(2, globalSize, NULL, false);
 }
 

modules/core/src/drawing.cpp

@@ -1584,6 +1584,24 @@ void line( InputOutputArray _img, Point pt1, Point pt2, const Scalar& color,
     ThickLine( img, pt1, pt2, buf, thickness, line_type, 3, shift );
 }
 
+void arrowedLine(InputOutputArray img, Point pt1, Point pt2, const Scalar& color,
+           int thickness, int line_type, int shift, double tipLength)
+{
+    const double tipSize = norm(pt1-pt2)*tipLength; // Factor to normalize the size of the tip depending on the length of the arrow
+
+    line(img, pt1, pt2, color, thickness, line_type, shift);
+
+    const double angle = atan2( (double) pt1.y - pt2.y, (double) pt1.x - pt2.x );
+
+    Point p(cvRound(pt2.x + tipSize * cos(angle + CV_PI / 4)),
+        cvRound(pt2.y + tipSize * sin(angle + CV_PI / 4)));
+    line(img, p, pt2, color, thickness, line_type, shift);
+
+    p.x = cvRound(pt2.x + tipSize * cos(angle - CV_PI / 4));
+    p.y = cvRound(pt2.y + tipSize * sin(angle - CV_PI / 4));
+    line(img, p, pt2, color, thickness, line_type, shift);
+}
+
 void rectangle( InputOutputArray _img, Point pt1, Point pt2,
                 const Scalar& color, int thickness,
                 int lineType, int shift )

modules/core/src/dxt.cpp

@@ -2080,32 +2080,32 @@ void cv::dft( InputArray _src0, OutputArray _dst, int flags, int nonzero_rows )
     {
         if ((flags & DFT_ROWS) == 0)
         {
-            if (!real_transform)
+            if (src.channels() == 2 && !(inv && (flags & DFT_REAL_OUTPUT)))
             {
-                if (ippi_DFT_C_32F(src,dst, inv, ipp_norm_flag))
+                if (ippi_DFT_C_32F(src, dst, inv, ipp_norm_flag))
                     return;
                 setIppErrorStatus();
             }
-            else if (inv || !(flags & DFT_COMPLEX_OUTPUT))
+            if (src.channels() == 1 && (inv || !(flags & DFT_COMPLEX_OUTPUT)))
             {
-                if (ippi_DFT_R_32F(src,dst, inv, ipp_norm_flag))
+                if (ippi_DFT_R_32F(src, dst, inv, ipp_norm_flag))
                     return;
                 setIppErrorStatus();
             }
         }
         else
         {
-            if (!real_transform)
+            if (src.channels() == 2 && !(inv && (flags & DFT_REAL_OUTPUT)))
             {
                 ippiDFT_C_Func ippiFunc = inv ? (ippiDFT_C_Func)ippiDFTInv_CToC_32fc_C1R : (ippiDFT_C_Func)ippiDFTFwd_CToC_32fc_C1R;
-                if (Dft_C_IPPLoop(src,dst, IPPDFT_C_Functor(ippiFunc),ipp_norm_flag))
+                if (Dft_C_IPPLoop(src, dst, IPPDFT_C_Functor(ippiFunc),ipp_norm_flag))
                     return;
                 setIppErrorStatus();
             }
-            else if (inv || !(flags & DFT_COMPLEX_OUTPUT))
+            if (src.channels() == 1 && (inv || !(flags & DFT_COMPLEX_OUTPUT)))
             {
                 ippiDFT_R_Func ippiFunc = inv ? (ippiDFT_R_Func)ippiDFTInv_PackToR_32f_C1R : (ippiDFT_R_Func)ippiDFTFwd_RToPack_32f_C1R;
-                if (Dft_R_IPPLoop(src,dst, IPPDFT_R_Functor(ippiFunc),ipp_norm_flag))
+                if (Dft_R_IPPLoop(src, dst, IPPDFT_R_Functor(ippiFunc),ipp_norm_flag))
                     return;
                 setIppErrorStatus();
             }

modules/core/src/mathfuncs.cpp

@@ -348,7 +348,18 @@ static void InvSqrt_32f(const float* src, float* dst, int len)
 
 static void InvSqrt_64f(const double* src, double* dst, int len)
 {
-    for( int i = 0; i < len; i++ )
+    int i = 0;
+
+#if CV_SSE2
+    if (USE_SSE2)
+    {
+        __m128d v_1 = _mm_set1_pd(1.0);
+        for ( ; i <= len - 2; i += 2)
+            _mm_storeu_pd(dst + i, _mm_div_pd(v_1, _mm_sqrt_pd(_mm_loadu_pd(src + i))));
+    }
+#endif
+
+    for( ; i < len; i++ )
         dst[i] = 1/std::sqrt(src[i]);
 }
 
@@ -2543,12 +2554,33 @@ void patchNaNs( InputOutputArray _a, double _val )
     NAryMatIterator it(arrays, (uchar**)ptrs);
     size_t len = it.size*a.channels();
     Cv32suf val;
-    val.f = (float)_val;
+    float fval = (float)_val;
+    val.f = fval;
+
+#if CV_SSE2
+    __m128i v_mask1 = _mm_set1_epi32(0x7fffffff), v_mask2 = _mm_set1_epi32(0x7f800000);
+    __m128i v_val = _mm_set1_epi32(val.i);
+#endif
 
     for( size_t i = 0; i < it.nplanes; i++, ++it )
     {
         int* tptr = ptrs[0];
-        for( size_t j = 0; j < len; j++ )
+        size_t j = 0;
+
+#if CV_SSE2
+        if (USE_SSE2)
+        {
+            for ( ; j < len; j += 4)
+            {
+                __m128i v_src = _mm_loadu_si128((__m128i const *)(tptr + j));
+                __m128i v_cmp_mask = _mm_cmplt_epi32(v_mask2, _mm_and_si128(v_src, v_mask1));
+                __m128i v_res = _mm_or_si128(_mm_andnot_si128(v_cmp_mask, v_src), _mm_and_si128(v_cmp_mask, v_val));
+                _mm_storeu_si128((__m128i *)(tptr + j), v_res);
+            }
+        }
+#endif
+
+        for( ; j < len; j++ )
             if( (tptr[j] & 0x7fffffff) > 0x7f800000 )
                 tptr[j] = val.i;
     }

modules/core/src/matrix.cpp

@@ -2758,21 +2758,30 @@ namespace cv {
 
 static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s )
 {
-    int type = _m.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
-            sctype = CV_MAKE_TYPE(depth, cn == 3 ? 4 : cn),
+    int type = _m.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), kercn = cn;
+    if (cn == 1)
+    {
+        kercn = std::min(ocl::predictOptimalVectorWidth(_m), 4);
+        if (kercn != 4)
+            kercn = 1;
+    }
+    int sctype = CV_MAKE_TYPE(depth, cn == 3 ? 4 : cn),
             rowsPerWI = ocl::Device::getDefault().isIntel() ? 4 : 1;
 
     ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc,
-                  format("-D T=%s -D T1=%s -D cn=%d -D ST=%s", ocl::memopTypeToStr(type),
-                         ocl::memopTypeToStr(depth), cn, ocl::memopTypeToStr(sctype)));
+                  format("-D T=%s -D T1=%s -D cn=%d -D ST=%s -D kercn=%d -D rowsPerWI=%d",
+                         ocl::memopTypeToStr(CV_MAKE_TYPE(depth, kercn)),
+                         ocl::memopTypeToStr(depth), cn,
+                         ocl::memopTypeToStr(sctype),
+                         kercn, rowsPerWI));
     if (k.empty())
         return false;
 
     UMat m = _m.getUMat();
-    k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, sctype, s)),
-           rowsPerWI);
+    k.args(ocl::KernelArg::WriteOnly(m, cn, kercn),
+           ocl::KernelArg::Constant(Mat(1, 1, sctype, s)));
 
-    size_t globalsize[2] = { m.cols, (m.rows + rowsPerWI - 1) / rowsPerWI };
+    size_t globalsize[2] = { m.cols * cn / kercn, (m.rows + rowsPerWI - 1) / rowsPerWI };
     return k.run(2, globalsize, NULL, false);
 }
 
@@ -3441,8 +3450,11 @@ static bool ocl_reduce(InputArray _src, OutputArray _dst,
     const int min_opt_cols = 128, buf_cols = 32;
     int sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype),
             ddepth = CV_MAT_DEPTH(dtype), ddepth0 = ddepth;
-    bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0,
-            useOptimized = 1 == dim && _src.cols() > min_opt_cols;
+    const ocl::Device &defDev = ocl::Device::getDefault();
+    bool doubleSupport = defDev.doubleFPConfig() > 0;
+
+    size_t wgs = defDev.maxWorkGroupSize();
+    bool useOptimized = 1 == dim && _src.cols() > min_opt_cols && (wgs >= buf_cols);
 
     if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F))
         return false;
@@ -3455,78 +3467,80 @@ static bool ocl_reduce(InputArray _src, OutputArray _dst,
 
     const char * const ops[4] = { "OCL_CV_REDUCE_SUM", "OCL_CV_REDUCE_AVG",
                                   "OCL_CV_REDUCE_MAX", "OCL_CV_REDUCE_MIN" };
-    char cvt[2][40];
-
     int wdepth = std::max(ddepth, CV_32F);
-    cv::String build_opt = format("-D %s -D dim=%d -D cn=%d -D ddepth=%d"
-                                  " -D srcT=%s -D dstT=%s -D dstT0=%s -D convertToWT=%s"
-                                  " -D convertToDT=%s -D convertToDT0=%s%s",
-                                  ops[op], dim, cn, ddepth, ocl::typeToStr(useOptimized ? ddepth : sdepth),
-                                  ocl::typeToStr(ddepth), ocl::typeToStr(ddepth0),
-                                  ocl::convertTypeStr(ddepth, wdepth, 1, cvt[0]),
-                                  ocl::convertTypeStr(sdepth, ddepth, 1, cvt[0]),
-                                  ocl::convertTypeStr(wdepth, ddepth0, 1, cvt[1]),
-                                  doubleSupport ? " -D DOUBLE_SUPPORT" : "");
-
     if (useOptimized)
     {
-        cv::String build_opt_pre = format("-D OP_REDUCE_PRE -D BUF_COLS=%d -D %s -D dim=1"
-                                          "  -D cn=%d -D ddepth=%d -D srcT=%s -D dstT=%s -D convertToDT=%s%s",
-                                          buf_cols, ops[op], cn, ddepth, ocl::typeToStr(sdepth), ocl::typeToStr(ddepth),
-                                          ocl::convertTypeStr(sdepth, ddepth, 1, cvt[0]),
-                                          doubleSupport ? " -D DOUBLE_SUPPORT" : "");
-        ocl::Kernel kpre("reduce_horz_pre", ocl::core::reduce2_oclsrc, build_opt_pre);
-        if (kpre.empty())
+        size_t tileHeight = (size_t)(wgs / buf_cols);
+        if (defDev.isIntel())
+        {
+            static const size_t maxItemInGroupCount = 16;
+            tileHeight = min(tileHeight, defDev.localMemSize() / buf_cols / CV_ELEM_SIZE(CV_MAKETYPE(wdepth, cn)) / maxItemInGroupCount);
+        }
+        char cvt[3][40];
+        cv::String build_opt = format("-D OP_REDUCE_PRE -D BUF_COLS=%d -D TILE_HEIGHT=%d -D %s -D dim=1"
+                                            " -D cn=%d -D ddepth=%d"
+                                            " -D srcT=%s -D bufT=%s -D dstT=%s"
+                                            " -D convertToWT=%s -D convertToBufT=%s -D convertToDT=%s%s",
+                                            buf_cols, tileHeight, ops[op], cn, ddepth,
+                                            ocl::typeToStr(sdepth),
+                                            ocl::typeToStr(ddepth),
+                                            ocl::typeToStr(ddepth0),
+                                            ocl::convertTypeStr(ddepth, wdepth, 1, cvt[0]),
+                                            ocl::convertTypeStr(sdepth, ddepth, 1, cvt[1]),
+                                            ocl::convertTypeStr(wdepth, ddepth0, 1, cvt[2]),
+                                            doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+        ocl::Kernel k("reduce_horz_opt", ocl::core::reduce2_oclsrc, build_opt);
+        if (k.empty())
             return false;
-
-        ocl::Kernel kmain("reduce", ocl::core::reduce2_oclsrc, build_opt);
-        if (kmain.empty())
-            return false;
-
         UMat src = _src.getUMat();
         Size dsize(1, src.rows);
         _dst.create(dsize, dtype);
         UMat dst = _dst.getUMat();
 
-        UMat buf(src.rows, buf_cols, dst.type());
-
-        kpre.args(ocl::KernelArg::ReadOnly(src),
-                  ocl::KernelArg::WriteOnlyNoSize(buf));
+        if (op0 == CV_REDUCE_AVG)
+            k.args(ocl::KernelArg::ReadOnly(src),
+                      ocl::KernelArg::WriteOnlyNoSize(dst), 1.0f / src.cols);
+        else
+            k.args(ocl::KernelArg::ReadOnly(src),
+                      ocl::KernelArg::WriteOnlyNoSize(dst));
 
+        size_t localSize[2] = { buf_cols, tileHeight};
         size_t globalSize[2] = { buf_cols, src.rows };
-        if (!kpre.run(2, globalSize, NULL, false))
+        return k.run(2, globalSize, localSize, false);
+    }
+    else
+    {
+        char cvt[2][40];
+        cv::String build_opt = format("-D %s -D dim=%d -D cn=%d -D ddepth=%d"
+                                      " -D srcT=%s -D dstT=%s -D dstT0=%s -D convertToWT=%s"
+                                      " -D convertToDT=%s -D convertToDT0=%s%s",
+                                      ops[op], dim, cn, ddepth, ocl::typeToStr(useOptimized ? ddepth : sdepth),
+                                      ocl::typeToStr(ddepth), ocl::typeToStr(ddepth0),
+                                      ocl::convertTypeStr(ddepth, wdepth, 1, cvt[0]),
+                                      ocl::convertTypeStr(sdepth, ddepth, 1, cvt[0]),
+                                      ocl::convertTypeStr(wdepth, ddepth0, 1, cvt[1]),
+                                      doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+
+        ocl::Kernel k("reduce", ocl::core::reduce2_oclsrc, build_opt);
+        if (k.empty())
             return false;
 
+        UMat src = _src.getUMat();
+        Size dsize(dim == 0 ? src.cols : 1, dim == 0 ? 1 : src.rows);
+        _dst.create(dsize, dtype);
+        UMat dst = _dst.getUMat();
+
+        ocl::KernelArg srcarg = ocl::KernelArg::ReadOnly(src),
+                temparg = ocl::KernelArg::WriteOnlyNoSize(dst);
+
         if (op0 == CV_REDUCE_AVG)
-            kmain.args(ocl::KernelArg::ReadOnly(buf),
-                       ocl::KernelArg::WriteOnlyNoSize(dst), 1.0f / src.cols);
+            k.args(srcarg, temparg, 1.0f / (dim == 0 ? src.rows : src.cols));
         else
-            kmain.args(ocl::KernelArg::ReadOnly(buf),
-                       ocl::KernelArg::WriteOnlyNoSize(dst));
+            k.args(srcarg, temparg);
 
-        globalSize[0] = src.rows;
-        return kmain.run(1, globalSize, NULL, false);
+        size_t globalsize = std::max(dsize.width, dsize.height);
+        return k.run(1, &globalsize, NULL, false);
     }
-
-    ocl::Kernel k("reduce", ocl::core::reduce2_oclsrc, build_opt);
-    if (k.empty())
-        return false;
-
-    UMat src = _src.getUMat();
-    Size dsize(dim == 0 ? src.cols : 1, dim == 0 ? 1 : src.rows);
-    _dst.create(dsize, dtype);
-    UMat dst = _dst.getUMat();
-
-    ocl::KernelArg srcarg = ocl::KernelArg::ReadOnly(src),
-            temparg = ocl::KernelArg::WriteOnlyNoSize(dst);
-
-    if (op0 == CV_REDUCE_AVG)
-        k.args(srcarg, temparg, 1.0f / (dim == 0 ? src.rows : src.cols));
-    else
-        k.args(srcarg, temparg);
-
-    size_t globalsize = std::max(dsize.width, dsize.height);
-    return k.run(1, &globalsize, NULL, false);
 }
 
 }

modules/core/src/ocl.cpp

@@ -3494,9 +3494,8 @@ public:
     OpenCLBufferPoolImpl()
         : currentReservedSize(0), maxReservedSize(0)
     {
-        // Note: Buffer pool is disabled by default,
-        //       because we didn't receive significant performance improvement
-        maxReservedSize = getConfigurationParameterForSize("OPENCV_OPENCL_BUFFERPOOL_LIMIT", 0);
+        int poolSize = ocl::Device::getDefault().isIntel() ? 1 << 27 : 0;
+        maxReservedSize = getConfigurationParameterForSize("OPENCV_OPENCL_BUFFERPOOL_LIMIT", poolSize);
     }
     virtual ~OpenCLBufferPoolImpl()
     {
@@ -3739,6 +3738,7 @@ public:
                 u->handle = clCreateBuffer(ctx_handle, CL_MEM_COPY_HOST_PTR|CL_MEM_READ_WRITE|createFlags,
                                            u->size, u->origdata, &retval);
                 tempUMatFlags = UMatData::TEMP_COPIED_UMAT;
+
             }
             if(!u->handle || retval != CL_SUCCESS)
                 return false;
@@ -3880,6 +3880,7 @@ public:
                 if(u->data && retval == CL_SUCCESS)
                 {
                     u->markHostCopyObsolete(false);
+                    u->markDeviceMemMapped(true);
                     return;
                 }
 
@@ -3908,6 +3909,7 @@ public:
         if(!u)
             return;
 
+
         CV_Assert(u->handle != 0);
 
         UMatDataAutoLock autolock(u);
@@ -3918,8 +3920,10 @@ public:
 
         cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
         cl_int retval = 0;
-        if( !u->copyOnMap() && u->data )
+        if( !u->copyOnMap() && u->deviceMemMapped() )
         {
+            CV_Assert(u->data != NULL);
+            u->markDeviceMemMapped(false);
             CV_Assert( (retval = clEnqueueUnmapMemObject(q,
                                 (cl_mem)u->handle, u->data, 0, 0, 0)) == CL_SUCCESS );
             CV_OclDbgAssert(clFinish(q) == CL_SUCCESS);

modules/core/src/opencl/lut.cl

@@ -36,114 +36,118 @@
 
 #if lcn == 1
     #if dcn == 4
-        #define LUT_OP(num)\
-            int idx = *(__global const int *)(srcptr + mad24(num, src_step, src_index));\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            dst[0] = lut_l[idx & 0xff];\
-            dst[1] = lut_l[(idx >> 8) & 0xff];\
-            dst[2] = lut_l[(idx >> 16) & 0xff];\
+        #define LUT_OP  \
+            int idx = *(__global const int *)(srcptr + src_index); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            dst[0] = lut_l[idx & 0xff]; \
+            dst[1] = lut_l[(idx >> 8) & 0xff]; \
+            dst[2] = lut_l[(idx >> 16) & 0xff]; \
             dst[3] = lut_l[(idx >> 24) & 0xff];
     #elif dcn == 3
-        #define LUT_OP(num)\
-            uchar3 idx = vload3(0, srcptr + mad24(num, src_step, src_index));\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            dst[0] = lut_l[idx.x];\
-            dst[1] = lut_l[idx.y];\
+        #define LUT_OP  \
+            uchar3 idx = vload3(0, srcptr + src_index); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            dst[0] = lut_l[idx.x]; \
+            dst[1] = lut_l[idx.y]; \
             dst[2] = lut_l[idx.z];
     #elif dcn == 2
-        #define LUT_OP(num)\
-            short idx = *(__global const short *)(srcptr + mad24(num, src_step, src_index));\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            dst[0] = lut_l[idx & 0xff];\
+        #define LUT_OP \
+            short idx = *(__global const short *)(srcptr + src_index); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            dst[0] = lut_l[idx & 0xff]; \
             dst[1] = lut_l[(idx >> 8) & 0xff];
     #elif dcn == 1
-        #define LUT_OP(num)\
-            uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+        #define LUT_OP \
+            uchar idx = (srcptr + src_index)[0]; \
+            dst = (__global dstT *)(dstptr + dst_index); \
             dst[0] = lut_l[idx];
     #else
-        #define LUT_OP(num)\
-            __global const srcT * src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            for (int cn = 0; cn < dcn; ++cn)\
+        #define LUT_OP \
+            __global const srcT * src = (__global const srcT *)(srcptr + src_index); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            for (int cn = 0; cn < dcn; ++cn) \
                 dst[cn] = lut_l[src[cn]];
     #endif
 #else
     #if dcn == 4
-        #define LUT_OP(num)\
-            __global const uchar4 *src_pixel = (__global const uchar4 *)(srcptr + mad24(num, src_step, src_index));\
-            int4 idx = convert_int4(src_pixel[0]) * lcn + (int4)(0, 1, 2, 3);\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            dst[0] = lut_l[idx.x];\
-            dst[1] = lut_l[idx.y];\
-            dst[2] = lut_l[idx.z];\
+        #define LUT_OP \
+            __global const uchar4 * src_pixel = (__global const uchar4 *)(srcptr + src_index); \
+            int4 idx = mad24(convert_int4(src_pixel[0]), (int4)(lcn), (int4)(0, 1, 2, 3)); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            dst[0] = lut_l[idx.x]; \
+            dst[1] = lut_l[idx.y]; \
+            dst[2] = lut_l[idx.z]; \
             dst[3] = lut_l[idx.w];
     #elif dcn == 3
-        #define LUT_OP(num)\
-            uchar3 src_pixel = vload3(0, srcptr + mad24(num, src_step, src_index));\
-            int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            dst[0] = lut_l[idx.x];\
-            dst[1] = lut_l[idx.y];\
+        #define LUT_OP \
+            uchar3 src_pixel = vload3(0, srcptr + src_index); \
+            int3 idx = mad24(convert_int3(src_pixel), (int3)(lcn), (int3)(0, 1, 2)); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            dst[0] = lut_l[idx.x]; \
+            dst[1] = lut_l[idx.y]; \
             dst[2] = lut_l[idx.z];
     #elif dcn == 2
-        #define LUT_OP(num)\
-            __global const uchar2 *src_pixel = (__global const uchar2 *)(srcptr + mad24(num, src_step, src_index));\
-            int2 idx = convert_int2(src_pixel[0]) * lcn + (int2)(0, 1);\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            dst[0] = lut_l[idx.x];\
+        #define LUT_OP \
+            __global const uchar2 * src_pixel = (__global const uchar2 *)(srcptr + src_index); \
+            int2 idx = mad24(convert_int2(src_pixel[0]), lcn, (int2)(0, 1)); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            dst[0] = lut_l[idx.x]; \
             dst[1] = lut_l[idx.y];
     #elif dcn == 1 //error case (1 < lcn) ==> lcn == scn == dcn
-        #define LUT_OP(num)\
-            uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+        #define LUT_OP \
+            uchar idx = (srcptr + src_index)[0]; \
+            dst = (__global dstT *)(dstptr + dst_index); \
             dst[0] = lut_l[idx];
     #else
-        #define LUT_OP(num)\
-            __global const srcT *src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\
-            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
-            for (int cn = 0; cn < dcn; ++cn)\
+        #define LUT_OP \
+            __global const srcT * src = (__global const srcT *)(srcptr + src_index); \
+            dst = (__global dstT *)(dstptr + dst_index); \
+            for (int cn = 0; cn < dcn; ++cn) \
                 dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
     #endif
 #endif
 
-#define LOCAL_LUT_INIT\
-    {\
-        __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);\
-        int init = mad24((int)get_local_id(1), (int)get_local_size(0), (int)get_local_id(0));\
-        int step = get_local_size(0) * get_local_size(1);\
-        for (int i = init; i < 256 * lcn; i += step)\
-        {\
-            lut_l[i] = lut[i];\
-        }\
-        barrier(CLK_LOCAL_MEM_FENCE);\
-    }
-
 __kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset,
                   __global const uchar * lutptr, int lut_step, int lut_offset,
                   __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols)
 {
-    __local dstT lut_l[256 * lcn];
-    LOCAL_LUT_INIT;
-
     int x = get_global_id(0);
-    int y = 4 * get_global_id(1);
+    int y = get_global_id(1) << 2;
+
+    __local dstT lut_l[256 * lcn];
+    __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
+
+    for (int i = mad24((int)get_local_id(1), (int)get_local_size(0), (int)get_local_id(0)),
+             step = get_local_size(0) * get_local_size(1); i < 256 * lcn; i += step)
+        lut_l[i] = lut[i];
+    barrier(CLK_LOCAL_MEM_FENCE);
 
     if (x < cols && y < rows)
     {
         int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
         int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
+
         __global dstT * dst;
-        LUT_OP(0);
+
+        LUT_OP;
+
         if (y < rows - 1)
         {
-            LUT_OP(1);
+            src_index += src_step;
+            dst_index += dst_step;
+            LUT_OP;
+
             if (y < rows - 2)
             {
-                LUT_OP(2);
+                src_index += src_step;
+                dst_index += dst_step;
+                LUT_OP;
+
                 if (y < rows - 3)
                 {
-                    LUT_OP(3);
+                    src_index += src_step;
+                    dst_index += dst_step;
+                    LUT_OP;
                 }
             }
         }

modules/core/src/opencl/minmaxloc.cl

@@ -42,9 +42,13 @@
 #if wdepth <= 4
 #define MIN_ABS(a) convertFromU(abs(a))
 #define MIN_ABS2(a, b) convertFromU(abs_diff(a, b))
+#define MIN(a, b) min(a, b)
+#define MAX(a, b) max(a, b)
 #else
 #define MIN_ABS(a) fabs(a)
 #define MIN_ABS2(a, b) fabs(a - b)
+#define MIN(a, b) fmin(a, b)
+#define MAX(a, b) fmax(a, b)
 #endif
 
 #if kercn != 3
@@ -60,44 +64,41 @@
 #define srcTSIZE (int)sizeof(srcT1)
 #endif
 
-#ifdef NEED_MINLOC
-#define CALC_MINLOC(inc) minloc = id + inc
-#else
-#define CALC_MINLOC(inc)
-#endif
-
-#ifdef NEED_MAXLOC
-#define CALC_MAXLOC(inc) maxloc = id + inc
-#else
-#define CALC_MAXLOC(inc)
-#endif
-
 #ifdef NEED_MINVAL
+#ifdef NEED_MINLOC
 #define CALC_MIN(p, inc) \
     if (minval > temp.p) \
     { \
         minval = temp.p; \
-        CALC_MINLOC(inc); \
+        minloc = id + inc; \
     }
 #else
+#define CALC_MIN(p, inc) \
+    minval = MIN(minval, temp.p);
+#endif
+#else
 #define CALC_MIN(p, inc)
 #endif
 
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
 #define CALC_MAX(p, inc) \
     if (maxval < temp.p) \
     { \
         maxval = temp.p; \
-        CALC_MAXLOC(inc); \
+        maxloc = id + inc; \
     }
 #else
+#define CALC_MAX(p, inc) \
+    maxval = MAX(maxval, temp.p);
+#endif
+#else
 #define CALC_MAX(p, inc)
 #endif
 
 #ifdef OP_CALC2
 #define CALC_MAX2(p) \
-    if (maxval2 < temp.p) \
-        maxval2 = temp.p;
+    maxval2 = MAX(maxval2, temp.p);
 #else
 #define CALC_MAX2(p)
 #endif
@@ -208,25 +209,28 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
 
 #if kercn == 1
 #ifdef NEED_MINVAL
+#if NEED_MINLOC
             if (minval > temp)
             {
                 minval = temp;
-#ifdef NEED_MINLOC
                 minloc = id;
-#endif
             }
+#else
+            minval = MIN(minval, temp);
+#endif
 #endif
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
             if (maxval < temp)
             {
                 maxval = temp;
-#ifdef NEED_MAXLOC
                 maxloc = id;
-#endif
             }
+#else
+            maxval = MAX(maxval, temp);
+#endif
 #ifdef OP_CALC2
-            if (maxval2 < temp2)
-                maxval2 = temp2;
+            maxval2 = MAX(maxval2, temp2);
 #endif
 #endif
 #elif kercn >= 2
@@ -282,32 +286,35 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
     {
         int lid3 = lid - WGS2_ALIGNED;
 #ifdef NEED_MINVAL
+#ifdef NEED_MINLOC
         if (localmem_min[lid3] >= minval)
         {
-#ifdef NEED_MINLOC
             if (localmem_min[lid3] == minval)
                 localmem_minloc[lid3] = min(localmem_minloc[lid3], minloc);
             else
                 localmem_minloc[lid3] = minloc,
-#endif
-                localmem_min[lid3] = minval;
+            localmem_min[lid3] = minval;
         }
+#else
+        localmem_min[lid3] = MIN(localmem_min[lid3], minval);
+#endif
 #endif
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
         if (localmem_max[lid3] <= maxval)
         {
-#ifdef NEED_MAXLOC
             if (localmem_max[lid3] == maxval)
                 localmem_maxloc[lid3] = min(localmem_maxloc[lid3], maxloc);
             else
                 localmem_maxloc[lid3] = maxloc,
-#endif
-                localmem_max[lid3] = maxval;
+            localmem_max[lid3] = maxval;
         }
+#else
+        localmem_max[lid3] = MAX(localmem_max[lid3], maxval);
+#endif
 #endif
 #ifdef OP_CALC2
-        if (localmem_max2[lid3] < maxval2)
-            localmem_max2[lid3] = maxval2;
+        localmem_max2[lid3] = MAX(localmem_max2[lid3], maxval2);
 #endif
     }
     barrier(CLK_LOCAL_MEM_FENCE);
@@ -319,32 +326,35 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
             int lid2 = lsize + lid;
 
 #ifdef NEED_MINVAL
+#ifdef NEED_MAXLOC
             if (localmem_min[lid] >= localmem_min[lid2])
             {
-#ifdef NEED_MINLOC
                 if (localmem_min[lid] == localmem_min[lid2])
                     localmem_minloc[lid] = min(localmem_minloc[lid2], localmem_minloc[lid]);
                 else
                     localmem_minloc[lid] = localmem_minloc[lid2],
-#endif
-                    localmem_min[lid] = localmem_min[lid2];
+                localmem_min[lid] = localmem_min[lid2];
             }
+#else
+            localmem_min[lid] = MIN(localmem_min[lid], localmem_min[lid2]);
+#endif
 #endif
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
             if (localmem_max[lid] <= localmem_max[lid2])
             {
-#ifdef NEED_MAXLOC
                 if (localmem_max[lid] == localmem_max[lid2])
                     localmem_maxloc[lid] = min(localmem_maxloc[lid2], localmem_maxloc[lid]);
                 else
                     localmem_maxloc[lid] = localmem_maxloc[lid2],
-#endif
-                    localmem_max[lid] = localmem_max[lid2];
+                localmem_max[lid] = localmem_max[lid2];
             }
+#else
+            localmem_max[lid] = MAX(localmem_max[lid], localmem_max[lid2]);
+#endif
 #endif
 #ifdef OP_CALC2
-            if (localmem_max2[lid] < localmem_max2[lid2])
-                localmem_max2[lid] = localmem_max2[lid2];
+            localmem_max2[lid] = MAX(localmem_max2[lid], localmem_max2[lid2]);
 #endif
         }
         barrier(CLK_LOCAL_MEM_FENCE);

modules/core/src/opencl/reduce2.cl

@@ -81,29 +81,34 @@
 #define PROCESS_ELEM(acc, value) acc += value
 #elif defined OCL_CV_REDUCE_MAX
 #define INIT_VALUE MIN_VAL
-#define PROCESS_ELEM(acc, value) acc = value > acc ? value : acc
+#define PROCESS_ELEM(acc, value) acc = max(value, acc)
 #elif defined OCL_CV_REDUCE_MIN
 #define INIT_VALUE MAX_VAL
-#define PROCESS_ELEM(acc, value) acc = value < acc ? value : acc
+#define PROCESS_ELEM(acc, value) acc = min(value, acc)
 #else
 #error "No operation is specified"
 #endif
 
 #ifdef OP_REDUCE_PRE
 
-__kernel void reduce_horz_pre(__global const uchar * srcptr, int src_step, int src_offset, int rows, int cols,
-                     __global uchar * bufptr, int buf_step, int buf_offset)
+__kernel void reduce_horz_opt(__global const uchar * srcptr, int src_step, int src_offset, int rows, int cols,
+                     __global uchar * dstptr, int dst_step, int dst_offset
+#ifdef OCL_CV_REDUCE_AVG
+                     , float fscale
+#endif
+                     )
 {
+    __local bufT lsmem[TILE_HEIGHT][BUF_COLS][cn];
+
     int x = get_global_id(0);
     int y = get_global_id(1);
-    if (x < BUF_COLS)
+    int liy = get_local_id(1);
+    if ((x < BUF_COLS) && (y < rows))
     {
         int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * cn, src_offset));
-        int buf_index = mad24(y, buf_step, mad24(x, (int)sizeof(dstT) * cn, buf_offset));
 
         __global const srcT * src = (__global const srcT *)(srcptr + src_index);
-        __global dstT * buf = (__global dstT *)(bufptr + buf_index);
-        dstT tmp[cn] = { INIT_VALUE };
+        bufT tmp[cn] = { INIT_VALUE };
 
         int src_step_mul = BUF_COLS * cn;
         for (int idx = x; idx < cols; idx += BUF_COLS, src += src_step_mul)
@@ -111,14 +116,49 @@ __kernel void reduce_horz_pre(__global const uchar * srcptr, int src_step, int s
             #pragma unroll
             for (int c = 0; c < cn; ++c)
             {
-                dstT value = convertToDT(src[c]);
+                bufT value = convertToBufT(src[c]);
                 PROCESS_ELEM(tmp[c], value);
             }
         }
 
         #pragma unroll
         for (int c = 0; c < cn; ++c)
-            buf[c] = tmp[c];
+            lsmem[liy][x][c] = tmp[c];
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+    if ((x < BUF_COLS / 2) && (y < rows))
+    {
+        #pragma unroll
+        for (int c = 0; c < cn; ++c)
+        {
+            PROCESS_ELEM(lsmem[liy][x][c], lsmem[liy][x +  BUF_COLS / 2][c]);
+        }
+    }
+    barrier(CLK_LOCAL_MEM_FENCE);
+    if ((x == 0) && (y < rows))
+    {
+        int dst_index = mad24(y, dst_step, dst_offset);
+
+        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+        bufT tmp[cn] = { INIT_VALUE };
+
+        #pragma unroll
+        for (int xin = 0; xin < BUF_COLS / 2; xin ++)
+        {
+            #pragma unroll
+            for (int c = 0; c < cn; ++c)
+            {
+                PROCESS_ELEM(tmp[c], lsmem[liy][xin][c]);
+            }
+        }
+
+        #pragma unroll
+        for (int c = 0; c < cn; ++c)
+#ifdef OCL_CV_REDUCE_AVG
+            dst[c] = convertToDT(convertToWT(tmp[c]) * fscale);
+#else
+            dst[c] = convertToDT(tmp[c]);
+#endif
     }
 }
 

modules/core/src/opencl/set_identity.cl

@@ -43,20 +43,18 @@
 //
 //M*/
 
-#if cn != 3
-#define loadpix(addr) *(__global const T *)(addr)
+#if kercn != 3
 #define storepix(val, addr)  *(__global T *)(addr) = val
 #define TSIZE (int)sizeof(T)
 #define scalar scalar_
 #else
-#define loadpix(addr) vload3(0, (__global const T1 *)(addr))
 #define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
 #define TSIZE ((int)sizeof(T1)*3)
 #define scalar (T)(scalar_.x, scalar_.y, scalar_.z)
 #endif
 
 __kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols,
-                          ST scalar_, int rowsPerWI)
+                          ST scalar_)
 {
     int x = get_global_id(0);
     int y0 = get_global_id(1) * rowsPerWI;
@@ -65,7 +63,35 @@ __kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset,
     {
         int src_index = mad24(y0, src_step, mad24(x, TSIZE, src_offset));
 
-        for (int y = y0, y1 = min(rows, y0 + rowsPerWI); y < y1; ++y, src_index += src_step)
-            storepix(x == y ? scalar : (T)(0), srcptr + src_index);
+#if kercn == cn
+        #pragma unroll
+        for (int y = y0, i = 0, y1 = min(rows, y0 + rowsPerWI); i < rowsPerWI; ++y, ++i, src_index += src_step)
+            if (y < y1)
+                storepix(x == y ? scalar : (T)(0), srcptr + src_index);
+#elif kercn == 4 && cn == 1
+        if (y0 < rows)
+        {
+            storepix(x == y0 >> 2 ? (T)(scalar, 0, 0, 0) : (T)(0), srcptr + src_index);
+            if (++y0 < rows)
+            {
+                src_index += src_step;
+                storepix(x == y0 >> 2 ? (T)(0, scalar, 0, 0) : (T)(0), srcptr + src_index);
+
+                if (++y0 < rows)
+                {
+                    src_index += src_step;
+                    storepix(x == y0 >> 2 ? (T)(0, 0, scalar, 0) : (T)(0), srcptr + src_index);
+
+                    if (++y0 < rows)
+                    {
+                        src_index += src_step;
+                        storepix(x == y0 >> 2 ? (T)(0, 0, 0, scalar) : (T)(0), srcptr + src_index);
+                    }
+                }
+            }
+        }
+#else
+#error "Incorrect combination of cn && kercn"
+#endif
     }
 }

modules/core/src/stat.cpp

@@ -918,8 +918,14 @@ static bool ocl_meanStdDev( InputArray _src, OutputArray _mean, OutputArray _sdv
         int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
         bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0,
                 isContinuous = _src.isContinuous();
-        int groups = ocl::Device::getDefault().maxComputeUnits();
-        size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+        const ocl::Device &defDev = ocl::Device::getDefault();
+        int groups = defDev.maxComputeUnits();
+        if (defDev.isIntel())
+        {
+            static const int subSliceEUCount = 10;
+            groups = (groups / subSliceEUCount) * 2;
+        }
+        size_t wgs = defDev.maxWorkGroupSize();
 
         int ddepth = std::max(CV_32S, depth), sqddepth = std::max(CV_32F, depth),
                 dtype = CV_MAKE_TYPE(ddepth, cn),