diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 6be60dc008..b83147851c 100644 --- a/modules/imgproc/src/morph.cpp +++ b/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; diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl index 69257ac36d..cb6e733ed4 100644 --- a/modules/imgproc/src/opencl/morph.cl +++ b/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