From d1cfcfcafd41d81522f8c4d3b64e62d3e7ecb9dc Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Tue, 18 Mar 2014 20:02:04 +0400
Subject: [PATCH] added 3-channels support to morphology operations

---
 modules/imgproc/src/morph.cpp             |  54 +++++---------
 modules/imgproc/src/opencl/morph.cl       | 117 +++++++++++++++---------------
 modules/imgproc/test/ocl/test_filters.cpp |  16 ++--
 3 files changed, 86 insertions(+), 101 deletions(-)

diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp
index ac958fc..b4011ee 100644
--- a/modules/imgproc/src/morph.cpp
+++ b/modules/imgproc/src/morph.cpp
@@ -42,7 +42,6 @@
 
 #include "precomp.hpp"
 #include <limits.h>
-#include <stdio.h>
 #include "opencl_kernels.hpp"
 
 /****************************************************************************************\
@@ -1291,9 +1290,10 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
 {
     CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE);
 
+    int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
     bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
 
-    if (_src.depth() == CV_64F && !doubleSupport)
+    if (depth == CV_64F && !doubleSupport)
         return false;
 
     UMat kernel8U;
@@ -1324,13 +1324,14 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
         return false;
 
     static const char * const op2str[] = { "ERODE", "DILATE" };
-    String buildOptions = format("-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],
+    String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s"
+                                 " -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", 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::typeToStr(_src.type()), _src.depth(), cn, ocl::typeToStr(depth));
 
     std::vector<ocl::Kernel> kernels;
-    for (int i = 0; i<iterations; i++)
+    for (int i = 0; i < iterations; i++)
     {
         ocl::Kernel k("morph", ocl::imgproc::morph_oclsrc, buildOptions);
         if (k.empty())
@@ -1341,33 +1342,27 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
     _dst.create(src.size(), src.type());
     UMat dst = _dst.getUMat();
 
-    if( iterations== 1 && src.u != dst.u)
+    if (iterations == 1 && src.u != dst.u)
     {
         Size wholesize;
         Point ofs;
         src.locateROI(wholesize, ofs);
         int wholecols = wholesize.width, wholerows = wholesize.height;
 
-        int idxArg = 0;
-        idxArg = kernels[0].set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src));
-        idxArg = kernels[0].set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst));
-        idxArg = kernels[0].set(idxArg, ofs.x);
-        idxArg = kernels[0].set(idxArg, ofs.y);
-        idxArg = kernels[0].set(idxArg, src.cols);
-        idxArg = kernels[0].set(idxArg, src.rows);
-        idxArg = kernels[0].set(idxArg, ocl::KernelArg::PtrReadOnly(kernel8U));
-        idxArg = kernels[0].set(idxArg, wholecols);
-        idxArg = kernels[0].set(idxArg, wholerows);
+        kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst),
+                        ofs.x, ofs.y, src.cols, src.rows, ocl::KernelArg::PtrReadOnly(kernel8U),
+                        wholecols, wholerows);
 
         return kernels[0].run(2, globalThreads, localThreads, false);
     }
 
-    for(int i = 0; i< iterations; i++)
+    for (int i = 0; i < iterations; i++)
     {
         UMat source;
         Size wholesize;
         Point ofs;
-        if( i == 0)
+
+        if (i == 0)
         {
             int cols =  src.cols, rows = src.rows;
             src.locateROI(wholesize,ofs);
@@ -1385,20 +1380,11 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
             dst.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);
         }
-
         source.locateROI(wholesize, ofs);
-        int wholecols = wholesize.width, wholerows = wholesize.height;
 
-        int idxArg = 0;
-        idxArg = kernels[i].set(idxArg, ocl::KernelArg::ReadOnlyNoSize(source));
-        idxArg = kernels[i].set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst));
-        idxArg = kernels[i].set(idxArg, ofs.x);
-        idxArg = kernels[i].set(idxArg, ofs.y);
-        idxArg = kernels[i].set(idxArg, source.cols);
-        idxArg = kernels[i].set(idxArg, source.rows);
-        idxArg = kernels[i].set(idxArg, ocl::KernelArg::PtrReadOnly(kernel8U));
-        idxArg = kernels[i].set(idxArg, wholecols);
-        idxArg = kernels[i].set(idxArg, wholerows);
+        kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst),
+                        ofs.x, ofs.y, source.cols, source.rows, ocl::KernelArg::PtrReadOnly(kernel8U),
+                        wholesize.width, wholesize.height);
 
         if (!kernels[i].run(2, globalThreads, localThreads, false))
             return false;
@@ -1427,13 +1413,13 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
         return;
 #endif
 
-    if( iterations == 0 || kernel.rows*kernel.cols == 1 )
+    if (iterations == 0 || kernel.rows*kernel.cols == 1)
     {
         _src.copyTo(_dst);
         return;
     }
 
-    if( !kernel.data )
+    if (!kernel.data)
     {
         kernel = getStructuringElement(MORPH_RECT, Size(1+iterations*2,1+iterations*2));
         anchor = Point(iterations, iterations);
@@ -1450,7 +1436,7 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
     }
 
     CV_OCL_RUN(_dst.isUMat() && _src.size() == _dst.size() && src_type == dst_type &&
-               _src.dims() <= 2 && (src_cn == 1 || src_cn == 4) &&
+               _src.dims() <= 2 && src_cn <= 4 &&
                (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) &&
                borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() &&
                (op == MORPH_ERODE || op == MORPH_DILATE),
diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl
index cb6e733..34c5ac9 100644
--- a/modules/imgproc/src/opencl/morph.cl
+++ b/modules/imgproc/src/opencl/morph.cl
@@ -43,6 +43,16 @@
 #endif
 #endif
 
+#if cn != 3
+#define loadpix(addr) *(__global const T *)(addr)
+#define storepix(val, addr)  *(__global T *)(addr) = val
+#define TSIZE (int)sizeof(T)
+#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)
+#endif
+
 #ifdef DEPTH_0
 #ifdef ERODE
 #define VAL 255
@@ -50,16 +60,14 @@
 #ifdef DILATE
 #define VAL 0
 #endif
-#endif
-#ifdef DEPTH_5
+#elif defined DEPTH_5
 #ifdef ERODE
 #define VAL FLT_MAX
 #endif
 #ifdef DILATE
 #define VAL -FLT_MAX
 #endif
-#endif
-#ifdef DEPTH_6
+#elif defined DEPTH_6
 #ifdef ERODE
 #define VAL DBL_MAX
 #endif
@@ -74,79 +82,70 @@
 #ifdef DILATE
 #define MORPH_OP(A,B) max((A),(B))
 #endif
-//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 uchar * restrict srcptr, int src_step, int src_offset,
+// 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 uchar * 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,
-                    __constant uchar * mat_kernel,
-                    int src_whole_cols, int src_whole_rows)
+                    int src_offset_x, int src_offset_y, int cols, int rows,
+                    __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows)
 {
-    int l_x = get_local_id(0);
-    int l_y = get_local_id(1);
-    int x = get_group_id(0)*LSIZE0;
-    int y = get_group_id(1)*LSIZE1;
-    int start_x = x+src_offset_x-RADIUSX;
-    int end_x = x + src_offset_x+LSIZE0+RADIUSX;
-    int width = end_x -(x+src_offset_x-RADIUSX)+1;
-    int start_y = y+src_offset_y-RADIUSY;
-    int point1 = mad24(l_y,LSIZE0,l_x);
-    int point2 = point1 + LSIZE0*LSIZE1;
-    int tl_x = point1 % width;
-    int tl_y = point1 / width;
-    int tl_x2 = point2 % width;
-    int tl_y2 = point2 / width;
-    int cur_x = start_x + tl_x;
-    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, 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 gidx = get_global_id(0), gidy = get_global_id(1);
+    int l_x = get_local_id(0), l_y = get_local_id(1);
+    int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1;
+    int start_x = x + src_offset_x - RADIUSX;
+    int end_x = x + src_offset_x + LSIZE0 + RADIUSX;
+    int width = end_x - (x + src_offset_x - RADIUSX) + 1;
+    int start_y = y + src_offset_y - RADIUSY;
+    int point1 = mad24(l_y, LSIZE0, l_x);
+    int point2 = point1 + LSIZE0 * LSIZE1;
+    int tl_x = point1 % width, tl_y = point1 / width;
+    int tl_x2 = point2 % width, tl_y2 = point2 / width;
+    int cur_x = start_x + tl_x, cur_y = start_y + tl_y;
+    int cur_x2 = start_x + tl_x2, cur_y2 = start_y + tl_y2;
+    int start_addr = mad24(cur_y, src_step, cur_x * TSIZE);
+    int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE);
+
+    __local T LDS_DAT[2*LSIZE1*LSIZE0];
 
-    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;
-    __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);
+    // read pixels from src
+    int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE);
+    start_addr = start_addr < end_addr && start_addr > 0 ? start_addr : 0;
+    start_addr2 = start_addr2 < end_addr && start_addr2 > 0 ? start_addr2 : 0;
 
-    temp1= ELEM(cur_x2,0,src_whole_cols,(GENTYPE)VAL,temp1);
-    temp1= ELEM(cur_y2,0,src_whole_rows,(GENTYPE)VAL,temp1);
+    T temp0 = loadpix(srcptr + start_addr);
+    T temp1 = loadpix(srcptr + start_addr2);
+
+    // judge if read out of boundary
+    temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL),temp0);
+    temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL),temp0);
+
+    temp1 = ELEM(cur_x2, 0, src_whole_cols, (T)(VAL), temp1);
+    temp1 = ELEM(cur_y2, 0, src_whole_rows, (T)(VAL), temp1);
 
     LDS_DAT[point1] = temp0;
     LDS_DAT[point2] = temp1;
     barrier(CLK_LOCAL_MEM_FENCE);
-    GENTYPE res = (GENTYPE)VAL;
-    for(int i=0; i<2*RADIUSY+1; i++)
-        for(int j=0; j<2*RADIUSX+1; j++)
+
+    T res = (T)(VAL);
+    for (int i = 0, sizey = 2 * RADIUSY + 1; i < sizey; i++)
+        for (int j = 0, sizex = 2 * RADIUSX + 1; j < sizex; j++)
         {
             res =
 #ifndef RECTKERNEL
                 mat_kernel[i*(2*RADIUSX+1)+j] ?
 #endif
-                MORPH_OP(res,LDS_DAT[mad24(l_y+i,width,l_x+j)])
+                MORPH_OP(res, LDS_DAT[mad24(l_y + i, width, l_x + j)])
 #ifndef RECTKERNEL
-                :res
+                : res
 #endif
                 ;
         }
-    int gidx = get_global_id(0);
-    int gidy = get_global_id(1);
-    if(gidx<cols && gidy<rows)
+
+    if (gidx < cols && gidy < rows)
     {
-        int dst_index = mad24(gidy, dst_step, dst_offset + gidx * (int)sizeof(GENTYPE));
-        __global GENTYPE * dst = (__global GENTYPE *)(dstptr + dst_index);
-        dst[0] = res;
+        int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset));
+        storepix(res, dstptr + dst_index);
     }
-
 }
diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp
index fe16fe8..4c68bb5 100644
--- a/modules/imgproc/test/ocl/test_filters.cpp
+++ b/modules/imgproc/test/ocl/test_filters.cpp
@@ -355,28 +355,28 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine(
                             Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
-                            Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
                             Values(3, 5, 7),
                             Values(Size(0,0)),//not used
                             Values((BorderType)BORDER_CONSTANT),//not used
                             Values(1.0, 2.0, 3.0),
-                            Bool() ) );
+                            Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(
-                            Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
                             Values(3, 5, 7),
                             Values(Size(0,0)),//not used
                             Values((BorderType)BORDER_CONSTANT),//not used
                             Values(1.0, 2.0, 3.0),
-                            Bool() ) );
+                            Bool()));
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine(
-                            Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
                             Values(3, 5, 7),
-                            Values(Size(0,0), Size(0,1), Size(0,2), Size(0,3), Size(0,4), Size(0,5),Size(0,6)),//uses as generator of operations
-                            Values((BorderType)BORDER_CONSTANT),//not used
+                            Values(Size(0, 0), Size(0, 1), Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations
+                            Values((BorderType)BORDER_CONSTANT),// not used
                             Values(1.0, 2.0, 3.0),
-                            Bool() ) );
+                            Bool()));
 
 
 } } // namespace cvtest::ocl
-- 
2.7.4