some fixes

modules/imgproc/src/morph.cpp

@@ -1314,35 +1314,10 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, InputArray _ker
     if(localThreads[0]*localThreads[1] * 2 < (localThreads[0] + ksize.width - 1) * (localThreads[1] + ksize.height - 1))
         return false;
 
-    char s[64];
-
-    switch (src.type())
-    {
-    case CV_8UC1:
-        sprintf(s, "-D VAL=%s -D GENTYPE=uchar", (op==MORPH_ERODE) ? "255" : "0");
-        break;
-    case CV_8UC4:
-        sprintf(s, "-D VAL=%s -D GENTYPE=uchar4", (op==MORPH_ERODE) ? "255" : "0");
-        break;
-    case CV_32FC1:
-        sprintf(s, "-D VAL=%s -D GENTYPE=float", (op==MORPH_ERODE) ? "FLT_MAX" : "-FLT_MAX");
-        break;
-    case CV_32FC4:
-        sprintf(s, "-D VAL=%s -D GENTYPE=float4", (op==MORPH_ERODE) ? "FLT_MAX" : "-FLT_MAX");
-        break;
-    case CV_64FC1:
-        sprintf(s, "-D VAL=%s -D GENTYPE=double", (op==MORPH_ERODE) ? "DBL_MAX" : "-DBL_MAX");
-        break;
-    case CV_64FC4:
-        sprintf(s, "-D VAL=%s -D GENTYPE=double4", (op==MORPH_ERODE) ? "DBL_MAX" : "-DBL_MAX");
-        break;
-    default:
-        CV_Error(Error::StsUnsupportedFormat, "unsupported type");
-    }
-
     char compile_option[128];
-    sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s %s %s %s",
-        anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op], doubleSupport?"-D DOUBLE_SUPPORT" :"", rectKernel?"-D RECTKERNEL":"", s);
+    sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s %s %s -D GENTYPE=%s -D DEPTH_%d",
+        anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op], doubleSupport?"-D DOUBLE_SUPPORT" :"", rectKernel?"-D RECTKERNEL":"",
+        ocl::typeToStr(_src.type()), _src.depth() );
 
     ocl::Kernel k( "morph", ocl::imgproc::morph_oclsrc, compile_option);
     if (k.empty())
@@ -1357,7 +1332,14 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, InputArray _ker
         Size wholesize;
         Point ofs;
         if( i == 0)
-            source = src;
+        {
+            int cols =  src.cols, rows = src.rows;
+            src.locateROI(wholesize,ofs);
+            src.adjustROI(ofs.y, wholesize.height - rows - ofs.y, ofs.x, wholesize.width - cols - ofs.x);
+            src.copyTo(source);
+            src.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
+            source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
+        }
         else
         {
             int cols =  dst.cols, rows = dst.rows;
@@ -1372,18 +1354,15 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, InputArray _ker
         int wholecols = wholesize.width, wholerows = wholesize.height;
 
         int idxArg = 0;
-        idxArg = k.set(idxArg, ocl::KernelArg::PtrReadOnly(source));
-        idxArg = k.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
-        idxArg = k.set(idxArg, (int)( (source.offset / source.elemSize())%(source.step / source.elemSize()) ) );
-        idxArg = k.set(idxArg, (int)( (source.offset / source.elemSize())/(source.step / source.elemSize()) ) );
+        idxArg = k.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(source));
+        idxArg = k.set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst));
+        idxArg = k.set(idxArg, ofs.x);
+        idxArg = k.set(idxArg, ofs.y);
         idxArg = k.set(idxArg, source.cols);
         idxArg = k.set(idxArg, source.rows);
-        idxArg = k.set(idxArg, (int)(source.step / source.elemSize()));
-        idxArg = k.set(idxArg, (int)(dst.step / dst.elemSize()));
         idxArg = k.set(idxArg, ocl::KernelArg::PtrReadOnly(kernel));
         idxArg = k.set(idxArg, wholecols);
         idxArg = k.set(idxArg, wholerows);
-        idxArg = k.set(idxArg, (int)( dst.offset / dst.elemSize() ) );
 
         if (!k.run(2, globalThreads, localThreads, true))
             return false;

modules/imgproc/src/opencl/morph.cl

@@ -43,6 +43,31 @@
 #endif
 #endif
 
+#ifdef DEPTH_0
+#ifdef ERODE
+#define VAL 255
+#endif
+#ifdef DILATE
+#define VAL 0
+#endif
+#endif
+#ifdef DEPTH_5
+#ifdef ERODE
+#define VAL FLT_MAX
+#endif
+#ifdef DILATE
+#define VAL -FLT_MAX
+#endif
+#endif
+#ifdef DEPTH_6
+#ifdef ERODE
+#define VAL DBL_MAX
+#endif
+#ifdef DILATE
+#define VAL -DBL_MAX
+#endif
+#endif
+
 #ifdef ERODE
 #define MORPH_OP(A,B) min((A),(B))
 #endif
@@ -52,14 +77,12 @@
 //BORDER_CONSTANT:      iiiiii|abcdefgh|iiiiiii
 #define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
 
-__kernel void morph(__global const GENTYPE * restrict src,
-                    __global GENTYPE *dst,
+__kernel void morph(__global const uchar * restrict srcptr, int src_step, int src_offset,
+                    __global uchar * dstptr, int dst_step, int dst_offset,
                     int src_offset_x, int src_offset_y,
                     int cols, int rows,
-                    int src_step_in_pixel, int dst_step_in_pixel,
                     __constant uchar * mat_kernel,
-                    int src_whole_cols, int src_whole_rows,
-                    int dst_offset_in_pixel)
+                    int src_whole_cols, int src_whole_rows)
 {
     int l_x = get_local_id(0);
     int l_y = get_local_id(1);
@@ -79,17 +102,20 @@ __kernel void morph(__global const GENTYPE * restrict src,
     int cur_y = start_y + tl_y;
     int cur_x2 = start_x + tl_x2;
     int cur_y2 = start_y + tl_y2;
-    int start_addr = mad24(cur_y,src_step_in_pixel,cur_x);
-    int start_addr2 = mad24(cur_y2,src_step_in_pixel,cur_x2);
+    int start_addr = mad24(cur_y,src_step, cur_x*(int)sizeof(GENTYPE));
+    int start_addr2 = mad24(cur_y2,src_step, cur_x2*(int)sizeof(GENTYPE));
     GENTYPE temp0,temp1;
     __local GENTYPE LDS_DAT[2*LSIZE1*LSIZE0];
 
-    int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
+    int end_addr = mad24(src_whole_rows - 1,src_step,src_whole_cols*(int)sizeof(GENTYPE));
     //read pixels from src
     start_addr = ((start_addr < end_addr) && (start_addr > 0)) ? start_addr : 0;
     start_addr2 = ((start_addr2 < end_addr) && (start_addr2 > 0)) ? start_addr2 : 0;
-    temp0 = src[start_addr];
-    temp1 = src[start_addr2];
+    __global const GENTYPE * src;
+    src = (__global const GENTYPE *)(srcptr+start_addr);
+    temp0 = src[0];
+    src = (__global const GENTYPE *)(srcptr+start_addr2);
+    temp1 = src[0];
     //judge if read out of boundary
     temp0= ELEM(cur_x,0,src_whole_cols,(GENTYPE)VAL,temp0);
     temp0= ELEM(cur_y,0,src_whole_rows,(GENTYPE)VAL,temp0);
@@ -116,10 +142,11 @@ __kernel void morph(__global const GENTYPE * restrict src,
         }
     int gidx = get_global_id(0);
     int gidy = get_global_id(1);
-    int out_addr = mad24(gidy,dst_step_in_pixel,gidx+dst_offset_in_pixel);
     if(gidx<cols && gidy<rows)
     {
-        dst[out_addr] = res;
+        int dst_index = mad24(gidy, dst_step, dst_offset + gidx * (int)sizeof(GENTYPE));
+        __global GENTYPE * dst = (__global GENTYPE *)(dstptr + dst_index);
+        dst[0] = res;
     }
 
 }