Merge pull request #2899 from ilya-lavrenov:tapi_ex

modules/imgproc/src/opencl/morph.cl

@@ -43,6 +43,8 @@
 #endif
 #endif
 
+#define noconvert
+
 #if cn != 3
 #define loadpix(addr) *(__global const T *)(addr)
 #define storepix(val, addr)  *(__global T *)(addr) = val
@@ -54,59 +56,75 @@
 #endif
 
 #ifdef DEPTH_0
-#ifdef ERODE
-#define VAL 255
-#endif
-#ifdef DILATE
-#define VAL 0
-#endif
+#define MIN_VAL 0
+#define MAX_VAL UCHAR_MAX
+#elif defined DEPTH_1
+#define MIN_VAL SCHAR_MIN
+#define MAX_VAL SCHAR_MAX
+#elif defined DEPTH_2
+#define MIN_VAL 0
+#define MAX_VAL USHRT_MAX
+#elif defined DEPTH_3
+#define MIN_VAL SHRT_MIN
+#define MAX_VAL SHRT_MAX
+#elif defined DEPTH_4
+#define MIN_VAL INT_MIN
+#define MAX_VAL INT_MAX
 #elif defined DEPTH_5
-#ifdef ERODE
-#define VAL FLT_MAX
-#endif
-#ifdef DILATE
-#define VAL -FLT_MAX
-#endif
+#define MIN_VAL (-FLT_MAX)
+#define MAX_VAL FLT_MAX
 #elif defined DEPTH_6
-#ifdef ERODE
-#define VAL DBL_MAX
-#endif
-#ifdef DILATE
-#define VAL -DBL_MAX
+#define MIN_VAL (-DBL_MAX)
+#define MAX_VAL DBL_MAX
 #endif
+
+#ifdef OP_ERODE
+#define VAL MAX_VAL
+#elif defined OP_DILATE
+#define VAL MIN_VAL
+#else
+#error "Unknown operation"
 #endif
 
-#ifdef ERODE
-#if defined(INTEL_DEVICE) && (DEPTH_0)
+#ifdef OP_ERODE
+#if defined INTEL_DEVICE && defined DEPTH_0
 // workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older)
 #define __CAT(x, y) x##y
 #define CAT(x, y) __CAT(x, y)
 #define WA_CONVERT_1 CAT(convert_uint, cn)
 #define WA_CONVERT_2 CAT(convert_, T)
 #define convert_uint1 convert_uint
-#define MORPH_OP(A,B) WA_CONVERT_2(min(WA_CONVERT_1(A),WA_CONVERT_1(B)))
+#define MORPH_OP(A, B) WA_CONVERT_2(min(WA_CONVERT_1(A), WA_CONVERT_1(B)))
 #else
-#define MORPH_OP(A,B) min((A),(B))
+#define MORPH_OP(A, B) min((A), (B))
 #endif
 #endif
-#ifdef DILATE
-#define MORPH_OP(A,B) max((A),(B))
+#ifdef OP_DILATE
+#define MORPH_OP(A, B) max((A), (B))
 #endif
 
+#define PROCESS(y, x) \
+    res = MORPH_OP(res, LDS_DAT[mad24(l_y + y, width, l_x + x)]);
+
 // BORDER_CONSTANT:      iiiiii|abcdefgh|iiiiiii
 #define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
 
+#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT
+#define EXTRA_PARAMS , __global const uchar * matptr, int mat_step, int mat_offset
+#else
+#define EXTRA_PARAMS
+#endif
+
 __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_whole_cols, int src_whole_rows EXTRA_PARAMS)
 {
     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 width = mad24(RADIUSX, 2, LSIZE0 + 1);
     int start_y = y + src_offset_y - RADIUSY;
     int point1 = mad24(l_y, LSIZE0, l_x);
     int point2 = point1 + LSIZE0 * LSIZE1;
@@ -117,7 +135,7 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset,
     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];
+    __local T LDS_DAT[2 * LSIZE1 * LSIZE0];
 
     // read pixels from src
     int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE);
@@ -128,8 +146,8 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset,
     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);
+    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);
@@ -138,24 +156,26 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset,
     LDS_DAT[point2] = temp1;
     barrier(CLK_LOCAL_MEM_FENCE);
 
-    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)])
-#ifndef RECTKERNEL
-                : res
-#endif
-                ;
-        }
-
     if (gidx < cols && gidy < rows)
     {
+        T res = (T)(VAL);
+        PROCESS_ELEMS;
+
         int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset));
+
+#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT
+        int mat_index =  mad24(gidy, mat_step, mad24(gidx, TSIZE, mat_offset));
+        T value = loadpix(matptr + mat_index);
+
+#ifdef OP_GRADIENT
+        storepix(convertToT(convertToWT(res) - convertToWT(value)), dstptr + dst_index);
+#elif defined OP_TOPHAT
+        storepix(convertToT(convertToWT(value) - convertToWT(res)), dstptr + dst_index);
+#elif defined OP_BLACKHAT
+        storepix(convertToT(convertToWT(res) - convertToWT(value)), dstptr + dst_index);
+#endif
+#else // erode or dilate
         storepix(res, dstptr + dst_index);
+#endif
     }
 }

modules/imgproc/test/ocl/test_filters.cpp

@@ -63,7 +63,7 @@ PARAM_TEST_CASE(FilterTestBase, MatType,
                 BorderType, // border type
                 double, // optional parameter
                 bool, // roi or not
-                int)  //width multiplier
+                int)  // width multiplier
 {
     int type, borderType, ksize;
     Size size;
@@ -244,8 +244,8 @@ OCL_TEST_P(Erode, Mat)
         random_roi();
         Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3);
 
-        OCL_OFF(cv::erode(src_roi, dst_roi, kernel, Point(-1,-1), iterations) );
-        OCL_ON(cv::erode(usrc_roi, udst_roi, kernel, Point(-1,-1), iterations) );
+        OCL_OFF(cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations) );
+        OCL_ON(cv::erode(usrc_roi, udst_roi, kernel, Point(-1, -1), iterations) );
 
         Near();
     }
@@ -266,8 +266,8 @@ OCL_TEST_P(Dilate, Mat)
         random_roi();
         Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3);
 
-        OCL_OFF(cv::dilate(src_roi, dst_roi, kernel, Point(-1,-1), iterations) );
-        OCL_ON(cv::dilate(usrc_roi, udst_roi, kernel, Point(-1,-1), iterations) );
+        OCL_OFF(cv::dilate(src_roi, dst_roi, kernel, Point(-1, -1), iterations) );
+        OCL_ON(cv::dilate(usrc_roi, udst_roi, kernel, Point(-1, -1), iterations) );
 
         Near();
     }
@@ -289,8 +289,8 @@ OCL_TEST_P(MorphologyEx, Mat)
         random_roi();
         Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3);
 
-        OCL_OFF(cv::morphologyEx(src_roi, dst_roi, op, kernel, Point(-1,-1), iterations) );
-        OCL_ON(cv::morphologyEx(usrc_roi, udst_roi, op, kernel, Point(-1,-1), iterations) );
+        OCL_OFF(cv::morphologyEx(src_roi, dst_roi, op, kernel, Point(-1, -1), iterations) );
+        OCL_ON(cv::morphologyEx(usrc_roi, udst_roi, op, kernel, Point(-1, -1), iterations) );
 
         Near();
     }
@@ -360,8 +360,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine(
 OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
                             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(Size(0, 0)), //not used
+                            Values((BorderType)BORDER_CONSTANT),
                             Values(1.0, 2.0, 3.0),
                             Bool(),
                             Values(1))); // not used
@@ -369,20 +369,20 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
 OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(
                             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(Size(0, 0)), // not used
+                            Values((BorderType)BORDER_CONSTANT),
                             Values(1.0, 2.0, 3.0),
                             Bool(),
-                            Values(1))); //not used
+                            Values(1))); // not used
 
 OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine(
-                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
+                            Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
                             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)), // used as generator of operations
-                            Values((BorderType)BORDER_CONSTANT),// not used
+                            Values(Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations
+                            Values((BorderType)BORDER_CONSTANT),
                             Values(1.0, 2.0, 3.0),
                             Bool(),
-                            Values(1))); //not used
+                            Values(1))); // not used
 
 
 } } // namespace cvtest::ocl