modules/core/src/opencl/minmaxloc.cl

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4
   5 // Copyright (C) 2014, Itseez, Inc., all rights reserved.
   6 // Third party copyrights are property of their respective owners.
   7
   8 #ifdef DOUBLE_SUPPORT
   9 #ifdef cl_amd_fp64
  10 #pragma OPENCL EXTENSION cl_amd_fp64:enable
  11 #elif defined (cl_khr_fp64)
  12 #pragma OPENCL EXTENSION cl_khr_fp64:enable
  13 #endif
  14 #endif
  15
  16 #ifdef DEPTH_0
  17 #define MIN_VAL 0
  18 #define MAX_VAL UCHAR_MAX
  19 #elif defined DEPTH_1
  20 #define MIN_VAL SCHAR_MIN
  21 #define MAX_VAL SCHAR_MAX
  22 #elif defined DEPTH_2
  23 #define MIN_VAL 0
  24 #define MAX_VAL USHRT_MAX
  25 #elif defined DEPTH_3
  26 #define MIN_VAL SHRT_MIN
  27 #define MAX_VAL SHRT_MAX
  28 #elif defined DEPTH_4
  29 #define MIN_VAL INT_MIN
  30 #define MAX_VAL INT_MAX
  31 #elif defined DEPTH_5
  32 #define MIN_VAL (-FLT_MAX)
  33 #define MAX_VAL FLT_MAX
  34 #elif defined DEPTH_6
  35 #define MIN_VAL (-DBL_MAX)
  36 #define MAX_VAL DBL_MAX
  37 #endif
  38
  39 #define noconvert
  40 #define INDEX_MAX UINT_MAX
  41
  42 #if kercn != 3
  43 #define loadpix(addr) *(__global const srcT *)(addr)
  44 #define srcTSIZE (int)sizeof(srcT)
  45 #else
  46 #define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
  47 #define srcTSIZE ((int)sizeof(srcT1) * 3)
  48 #endif
  49
  50 #ifndef HAVE_MASK
  51 #undef srcTSIZE
  52 #define srcTSIZE (int)sizeof(srcT1)
  53 #endif
  54
  55 #ifdef NEED_MINLOC
  56 #define CALC_MINLOC(inc) minloc = id + inc
  57 #else
  58 #define CALC_MINLOC(inc)
  59 #endif
  60
  61 #ifdef NEED_MAXLOC
  62 #define CALC_MAXLOC(inc) maxloc = id + inc
  63 #else
  64 #define CALC_MAXLOC(inc)
  65 #endif
  66
  67 #ifdef NEED_MINVAL
  68 #define CALC_MIN(p, inc) \
  69     if (minval > temp.p) \
  70     { \
  71         minval = temp.p; \
  72         CALC_MINLOC(inc); \
  73     }
  74 #else
  75 #define CALC_MIN(p, inc)
  76 #endif
  77
  78 #ifdef NEED_MAXVAL
  79 #define CALC_MAX(p, inc) \
  80     if (maxval < temp.p) \
  81     { \
  82         maxval = temp.p; \
  83         CALC_MAXLOC(inc); \
  84     }
  85 #else
  86 #define CALC_MAX(p, inc)
  87 #endif
  88
  89 #ifdef OP_CALC2
  90 #define CALC_MAX2(p) \
  91     if (maxval2 < temp.p) \
  92         maxval2 = temp.p;
  93 #else
  94 #define CALC_MAX2(p)
  95 #endif
  96
  97 #define CALC_P(p, inc) \
  98     CALC_MIN(p, inc) \
  99     CALC_MAX(p, inc) \
 100     CALC_MAX2(p)
 101
 102 __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_offset, int cols,
 103                         int total, int groupnum, __global uchar * dstptr
 104 #ifdef HAVE_MASK
 105                         , __global const uchar * mask, int mask_step, int mask_offset
 106 #endif
 107 #ifdef HAVE_SRC2
 108                         , __global const uchar * src2ptr, int src2_step, int src2_offset
 109 #endif
 110                         )
 111 {
 112     int lid = get_local_id(0);
 113     int gid = get_group_id(0);
 114     int  id = get_global_id(0)
 115 #ifndef HAVE_MASK
 116     * kercn;
 117 #else
 118     ;
 119 #endif
 120
 121     srcptr += src_offset;
 122 #ifdef HAVE_MASK
 123     mask += mask_offset;
 124 #endif
 125 #ifdef HAVE_SRC2
 126     src2ptr += src2_offset;
 127 #endif
 128
 129 #ifdef NEED_MINVAL
 130     __local dstT1 localmem_min[WGS2_ALIGNED];
 131     dstT1 minval = MAX_VAL;
 132 #ifdef NEED_MINLOC
 133     __local uint localmem_minloc[WGS2_ALIGNED];
 134     uint minloc = INDEX_MAX;
 135 #endif
 136 #endif
 137 #ifdef NEED_MAXVAL
 138     dstT1 maxval = MIN_VAL;
 139     __local dstT1 localmem_max[WGS2_ALIGNED];
 140 #ifdef NEED_MAXLOC
 141     __local uint localmem_maxloc[WGS2_ALIGNED];
 142     uint maxloc = INDEX_MAX;
 143 #endif
 144 #endif
 145 #ifdef OP_CALC2
 146     __local dstT1 localmem_max2[WGS2_ALIGNED];
 147     dstT1 maxval2 = MIN_VAL;
 148 #endif
 149
 150     int src_index;
 151 #ifdef HAVE_MASK
 152     int mask_index;
 153 #endif
 154 #ifdef HAVE_SRC2
 155     int src2_index;
 156 #endif
 157
 158     dstT temp;
 159 #ifdef HAVE_SRC2
 160     dstT temp2;
 161 #endif
 162
 163     for (int grain = groupnum * WGS
 164 #ifndef HAVE_MASK
 165         * kercn
 166 #endif
 167         ; id < total; id += grain)
 168     {
 169 #ifdef HAVE_MASK
 170 #ifdef HAVE_MASK_CONT
 171         mask_index = id;
 172 #else
 173         mask_index = mad24(id / cols, mask_step, id % cols);
 174 #endif
 175         if (mask[mask_index])
 176 #endif
 177         {
 178 #ifdef HAVE_SRC_CONT
 179             src_index = mul24(id, srcTSIZE);
 180 #else
 181             src_index = mad24(id / cols, src_step, mul24(id % cols, srcTSIZE));
 182 #endif
 183             temp = convertToDT(loadpix(srcptr + src_index));
 184 #ifdef OP_ABS
 185             temp = temp >= (dstT)(0) ? temp : -temp;
 186 #endif
 187
 188 #ifdef HAVE_SRC2
 189 #ifdef HAVE_SRC2_CONT
 190             src2_index = mul24(id, srcTSIZE);
 191 #else
 192             src2_index = mad24(id / cols, src2_step, mul24(id % cols, srcTSIZE));
 193 #endif
 194             temp2 = convertToDT(loadpix(src2ptr + src2_index));
 195             temp = temp > temp2 ? temp - temp2 : (temp2 - temp);
 196 #ifdef OP_CALC2
 197             temp2 = temp2 >= (dstT)(0) ? temp2 : -temp2;
 198 #endif
 199 #endif
 200
 201 #if kercn == 1
 202 #ifdef NEED_MINVAL
 203             if (minval > temp)
 204             {
 205                 minval = temp;
 206 #ifdef NEED_MINLOC
 207                 minloc = id;
 208 #endif
 209             }
 210 #endif
 211 #ifdef NEED_MAXVAL
 212             if (maxval < temp)
 213             {
 214                 maxval = temp;
 215 #ifdef NEED_MAXLOC
 216                 maxloc = id;
 217 #endif
 218             }
 219 #ifdef OP_CALC2
 220             if (maxval2 < temp2)
 221                 maxval2 = temp2;
 222 #endif
 223 #endif
 224 #elif kercn >= 2
 225             CALC_P(s0, 0)
 226             CALC_P(s1, 1)
 227 #if kercn >= 3
 228             CALC_P(s2, 2)
 229 #if kercn >= 4
 230             CALC_P(s3, 3)
 231 #if kercn >= 8
 232             CALC_P(s4, 4)
 233             CALC_P(s5, 5)
 234             CALC_P(s6, 6)
 235             CALC_P(s7, 7)
 236 #if kercn == 16
 237             CALC_P(s8, 8)
 238             CALC_P(s9, 9)
 239             CALC_P(sA, 10)
 240             CALC_P(sB, 11)
 241             CALC_P(sC, 12)
 242             CALC_P(sD, 13)
 243             CALC_P(sE, 14)
 244             CALC_P(sF, 15)
 245 #endif
 246 #endif
 247 #endif
 248 #endif
 249 #endif
 250         }
 251     }
 252
 253     if (lid < WGS2_ALIGNED)
 254     {
 255 #ifdef NEED_MINVAL
 256         localmem_min[lid] = minval;
 257 #endif
 258 #ifdef NEED_MAXVAL
 259         localmem_max[lid] = maxval;
 260 #endif
 261 #ifdef NEED_MINLOC
 262         localmem_minloc[lid] = minloc;
 263 #endif
 264 #ifdef NEED_MAXLOC
 265         localmem_maxloc[lid] = maxloc;
 266 #endif
 267 #ifdef OP_CALC2
 268         localmem_max2[lid] = maxval2;
 269 #endif
 270     }
 271     barrier(CLK_LOCAL_MEM_FENCE);
 272
 273     if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
 274     {
 275         int lid3 = lid - WGS2_ALIGNED;
 276 #ifdef NEED_MINVAL
 277         if (localmem_min[lid3] >= minval)
 278         {
 279 #ifdef NEED_MINLOC
 280             if (localmem_min[lid3] == minval)
 281                 localmem_minloc[lid3] = min(localmem_minloc[lid3], minloc);
 282             else
 283                 localmem_minloc[lid3] = minloc,
 284 #endif
 285                 localmem_min[lid3] = minval;
 286         }
 287 #endif
 288 #ifdef NEED_MAXVAL
 289         if (localmem_max[lid3] <= maxval)
 290         {
 291 #ifdef NEED_MAXLOC
 292             if (localmem_max[lid3] == maxval)
 293                 localmem_maxloc[lid3] = min(localmem_maxloc[lid3], maxloc);
 294             else
 295                 localmem_maxloc[lid3] = maxloc,
 296 #endif
 297                 localmem_max[lid3] = maxval;
 298         }
 299 #endif
 300 #ifdef OP_CALC2
 301         if (localmem_max2[lid3] < maxval2)
 302             localmem_max2[lid3] = maxval2;
 303 #endif
 304     }
 305     barrier(CLK_LOCAL_MEM_FENCE);
 306
 307     for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
 308     {
 309         if (lid < lsize)
 310         {
 311             int lid2 = lsize + lid;
 312
 313 #ifdef NEED_MINVAL
 314             if (localmem_min[lid] >= localmem_min[lid2])
 315             {
 316 #ifdef NEED_MINLOC
 317                 if (localmem_min[lid] == localmem_min[lid2])
 318                     localmem_minloc[lid] = min(localmem_minloc[lid2], localmem_minloc[lid]);
 319                 else
 320                     localmem_minloc[lid] = localmem_minloc[lid2],
 321 #endif
 322                     localmem_min[lid] = localmem_min[lid2];
 323             }
 324 #endif
 325 #ifdef NEED_MAXVAL
 326             if (localmem_max[lid] <= localmem_max[lid2])
 327             {
 328 #ifdef NEED_MAXLOC
 329                 if (localmem_max[lid] == localmem_max[lid2])
 330                     localmem_maxloc[lid] = min(localmem_maxloc[lid2], localmem_maxloc[lid]);
 331                 else
 332                     localmem_maxloc[lid] = localmem_maxloc[lid2],
 333 #endif
 334                     localmem_max[lid] = localmem_max[lid2];
 335             }
 336 #endif
 337 #ifdef OP_CALC2
 338             if (localmem_max2[lid] < localmem_max2[lid2])
 339                 localmem_max2[lid] = localmem_max2[lid2];
 340 #endif
 341         }
 342         barrier(CLK_LOCAL_MEM_FENCE);
 343     }
 344
 345     if (lid == 0)
 346     {
 347         int pos = 0;
 348 #ifdef NEED_MINVAL
 349         *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_min[0];
 350         pos = mad24(groupnum, (int)sizeof(dstT1), pos);
 351 #endif
 352 #ifdef NEED_MAXVAL
 353         *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max[0];
 354         pos = mad24(groupnum, (int)sizeof(dstT1), pos);
 355 #endif
 356 #ifdef NEED_MINLOC
 357         *(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_minloc[0];
 358         pos = mad24(groupnum, (int)sizeof(uint), pos);
 359 #endif
 360 #ifdef NEED_MAXLOC
 361         *(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_maxloc[0];
 362 #ifdef OP_CALC2
 363         pos = mad24(groupnum, (int)sizeof(uint), pos);
 364 #endif
 365 #endif
 366 #ifdef OP_CALC2
 367         *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max2[0];
 368 #endif
 369     }
 370 }