Merge pull request #2936 from ilya-lavrenov:tapi_minmaxloc_opt

modules/core/src/opencl/minmaxloc.cl

@@ -42,9 +42,13 @@
 #if wdepth <= 4
 #define MIN_ABS(a) convertFromU(abs(a))
 #define MIN_ABS2(a, b) convertFromU(abs_diff(a, b))
+#define MIN(a, b) min(a, b)
+#define MAX(a, b) max(a, b)
 #else
 #define MIN_ABS(a) fabs(a)
 #define MIN_ABS2(a, b) fabs(a - b)
+#define MIN(a, b) fmin(a, b)
+#define MAX(a, b) fmax(a, b)
 #endif
 
 #if kercn != 3
@@ -60,44 +64,41 @@
 #define srcTSIZE (int)sizeof(srcT1)
 #endif
 
-#ifdef NEED_MINLOC
-#define CALC_MINLOC(inc) minloc = id + inc
-#else
-#define CALC_MINLOC(inc)
-#endif
-
-#ifdef NEED_MAXLOC
-#define CALC_MAXLOC(inc) maxloc = id + inc
-#else
-#define CALC_MAXLOC(inc)
-#endif
-
 #ifdef NEED_MINVAL
+#ifdef NEED_MINLOC
 #define CALC_MIN(p, inc) \
     if (minval > temp.p) \
     { \
         minval = temp.p; \
-        CALC_MINLOC(inc); \
+        minloc = id + inc; \
     }
 #else
+#define CALC_MIN(p, inc) \
+    minval = MIN(minval, temp.p);
+#endif
+#else
 #define CALC_MIN(p, inc)
 #endif
 
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
 #define CALC_MAX(p, inc) \
     if (maxval < temp.p) \
     { \
         maxval = temp.p; \
-        CALC_MAXLOC(inc); \
+        maxloc = id + inc; \
     }
 #else
+#define CALC_MAX(p, inc) \
+    maxval = MAX(maxval, temp.p);
+#endif
+#else
 #define CALC_MAX(p, inc)
 #endif
 
 #ifdef OP_CALC2
 #define CALC_MAX2(p) \
-    if (maxval2 < temp.p) \
-        maxval2 = temp.p;
+    maxval2 = MAX(maxval2, temp.p);
 #else
 #define CALC_MAX2(p)
 #endif
@@ -208,25 +209,28 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
 
 #if kercn == 1
 #ifdef NEED_MINVAL
+#if NEED_MINLOC
             if (minval > temp)
             {
                 minval = temp;
-#ifdef NEED_MINLOC
                 minloc = id;
-#endif
             }
+#else
+            minval = MIN(minval, temp);
+#endif
 #endif
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
             if (maxval < temp)
             {
                 maxval = temp;
-#ifdef NEED_MAXLOC
                 maxloc = id;
-#endif
             }
+#else
+            maxval = MAX(maxval, temp);
+#endif
 #ifdef OP_CALC2
-            if (maxval2 < temp2)
-                maxval2 = temp2;
+            maxval2 = MAX(maxval2, temp2);
 #endif
 #endif
 #elif kercn >= 2
@@ -282,32 +286,35 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
     {
         int lid3 = lid - WGS2_ALIGNED;
 #ifdef NEED_MINVAL
+#ifdef NEED_MINLOC
         if (localmem_min[lid3] >= minval)
         {
-#ifdef NEED_MINLOC
             if (localmem_min[lid3] == minval)
                 localmem_minloc[lid3] = min(localmem_minloc[lid3], minloc);
             else
                 localmem_minloc[lid3] = minloc,
-#endif
-                localmem_min[lid3] = minval;
+            localmem_min[lid3] = minval;
         }
+#else
+        localmem_min[lid3] = MIN(localmem_min[lid3], minval);
+#endif
 #endif
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
         if (localmem_max[lid3] <= maxval)
         {
-#ifdef NEED_MAXLOC
             if (localmem_max[lid3] == maxval)
                 localmem_maxloc[lid3] = min(localmem_maxloc[lid3], maxloc);
             else
                 localmem_maxloc[lid3] = maxloc,
-#endif
-                localmem_max[lid3] = maxval;
+            localmem_max[lid3] = maxval;
         }
+#else
+        localmem_max[lid3] = MAX(localmem_max[lid3], maxval);
+#endif
 #endif
 #ifdef OP_CALC2
-        if (localmem_max2[lid3] < maxval2)
-            localmem_max2[lid3] = maxval2;
+        localmem_max2[lid3] = MAX(localmem_max2[lid3], maxval2);
 #endif
     }
     barrier(CLK_LOCAL_MEM_FENCE);
@@ -319,32 +326,35 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
             int lid2 = lsize + lid;
 
 #ifdef NEED_MINVAL
+#ifdef NEED_MAXLOC
             if (localmem_min[lid] >= localmem_min[lid2])
             {
-#ifdef NEED_MINLOC
                 if (localmem_min[lid] == localmem_min[lid2])
                     localmem_minloc[lid] = min(localmem_minloc[lid2], localmem_minloc[lid]);
                 else
                     localmem_minloc[lid] = localmem_minloc[lid2],
-#endif
-                    localmem_min[lid] = localmem_min[lid2];
+                localmem_min[lid] = localmem_min[lid2];
             }
+#else
+            localmem_min[lid] = MIN(localmem_min[lid], localmem_min[lid2]);
+#endif
 #endif
 #ifdef NEED_MAXVAL
+#ifdef NEED_MAXLOC
             if (localmem_max[lid] <= localmem_max[lid2])
             {
-#ifdef NEED_MAXLOC
                 if (localmem_max[lid] == localmem_max[lid2])
                     localmem_maxloc[lid] = min(localmem_maxloc[lid2], localmem_maxloc[lid]);
                 else
                     localmem_maxloc[lid] = localmem_maxloc[lid2],
-#endif
-                    localmem_max[lid] = localmem_max[lid2];
+                localmem_max[lid] = localmem_max[lid2];
             }
+#else
+            localmem_max[lid] = MAX(localmem_max[lid], localmem_max[lid2]);
+#endif
 #endif
 #ifdef OP_CALC2
-            if (localmem_max2[lid] < localmem_max2[lid2])
-                localmem_max2[lid] = localmem_max2[lid2];
+            localmem_max2[lid] = MAX(localmem_max2[lid], localmem_max2[lid2]);
 #endif
         }
         barrier(CLK_LOCAL_MEM_FENCE);