Moved new kernels under conditional compilation to disable their build for 3-channel images; added condition to use new kernels only for images that are big enough

modules/imgproc/src/opencl/medianFilter.cl

@@ -39,6 +39,10 @@
 #define TSIZE (int)sizeof(T1) * cn
 #endif
 
+#define OP(a,b) {    mid=a; a=min(a,b); b=max(mid,b);}
+
+#ifdef USE_4OPT
+
 //Utility macros for for 1,2,4 channel images:
 //  - LOAD4/STORE4 - load/store 4-pixel groups from/to global memory
 
@@ -89,8 +93,6 @@
 
 #endif
 
-#define OP(a,b) {    mid=a; a=min(a,b); b=max(mid,b);}
-
 __kernel void medianFilter3_u(__global const uchar* srcptr, int srcStep, int srcOffset,
                             __global uchar*       dstptr, int dstStep, int dstOffset,
                             int rows, int cols)
@@ -274,6 +276,8 @@ __kernel void medianFilter5_u(__global const uchar* srcptr, int srcStep, int src
     }
 }
 
+#endif
+
 __kernel void medianFilter3(__global const uchar * srcptr, int src_step, int src_offset,
                             __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols)
 {

modules/imgproc/src/smooth.cpp

@@ -2021,13 +2021,20 @@ static bool ocl_medianFilter(InputArray _src, OutputArray _dst, int m)
     if ( !((depth == CV_8U || depth == CV_16U || depth == CV_16S || depth == CV_32F) && cn <= 4 && (m == 3 || m == 5)) )
         return false;
 
-    bool useOptimized = (1 == cn) && (ocl::Device::getDefault().isIntel());
+    Size imgSize = _src.size();
+    bool useOptimized = (1 == cn) &&
+                        imgSize.width >= localsize[0] * 8  &&
+                        imgSize.height >= localsize[1] * 8 &&
+                        (ocl::Device::getDefault().isIntel());
 
     cv::String kname = format( useOptimized ? "medianFilter%d_u" : "medianFilter%d", m) ;
+    cv::String kdefs = useOptimized ?
+                         format("-D T=%s -D T1=%s -D T4=%s%d -D cn=%d -D USE_4OPT", ocl::typeToStr(type),
+                         ocl::typeToStr(depth), ocl::typeToStr(depth), cn*4, cn)
+                         :
+                         format("-D T=%s -D T1=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn) ;
 
-    ocl::Kernel k(kname.c_str(), ocl::imgproc::medianFilter_oclsrc,
+    ocl::Kernel k(kname.c_str(), ocl::imgproc::medianFilter_oclsrc, kdefs.c_str() );
-                  format("-D T=%s -D T1=%s -D T4=%s%d -D cn=%d", ocl::typeToStr(type),
-                         ocl::typeToStr(depth), ocl::typeToStr(depth), cn*4, cn));
 
     if (k.empty())
         return false;