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(srcT1)
  45 #else
  46 #define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
  47 #define srcTSIZE ((int)sizeof(srcT1))
  48 #endif
  49
  50 #ifdef NEED_MINLOC
  51 #define CALC_MINLOC(inc) minloc = id + inc
  52 #else
  53 #define CALC_MINLOC(inc)
  54 #endif
  55
  56 #ifdef NEED_MAXLOC
  57 #define CALC_MAXLOC(inc) maxloc = id + inc
  58 #else
  59 #define CALC_MAXLOC(inc)
  60 #endif
  61
  62 #ifdef NEED_MINVAL
  63 #define CALC_MIN(p, inc) \
  64     if (minval > temp.p) \
  65     { \
  66         minval = temp.p; \
  67         CALC_MINLOC(inc); \
  68     }
  69 #else
  70 #define CALC_MIN(p, inc)
  71 #endif
  72
  73 #ifdef NEED_MAXVAL
  74 #define CALC_MAX(p, inc) \
  75     if (maxval < temp.p) \
  76     { \
  77         maxval = temp.p; \
  78         CALC_MAXLOC(inc); \
  79     }
  80 #else
  81 #define CALC_MAX(p, inc)
  82 #endif
  83
  84 #ifdef OP_CALC2
  85 #define CALC_MAX2(p) \
  86     if (maxval2 < temp.p) \
  87         maxval2 = temp.p;
  88 #else
  89 #define CALC_MAX2(p)
  90 #endif
  91
  92 #define CALC_P(p, inc) \
  93     CALC_MIN(p, inc) \
  94     CALC_MAX(p, inc) \
  95     CALC_MAX2(p)
  96
  97 __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_offset, int cols,
  98                         int total, int groupnum, __global uchar * dstptr
  99 #ifdef HAVE_MASK
 100                         , __global const uchar * mask, int mask_step, int mask_offset
 101 #endif
 102 #ifdef HAVE_SRC2
 103                         , __global const uchar * src2ptr, int src2_step, int src2_offset
 104 #endif
 105                         )
 106 {
 107     int lid = get_local_id(0);
 108     int gid = get_group_id(0);
 109     int  id = get_global_id(0) * kercn;
 110
 111     srcptr += src_offset;
 112 #ifdef HAVE_MASK
 113     mask += mask_offset;
 114 #endif
 115 #ifdef HAVE_SRC2
 116     src2ptr += src2_offset;
 117 #endif
 118
 119 #ifdef NEED_MINVAL
 120     __local dstT1 localmem_min[WGS2_ALIGNED];
 121     dstT1 minval = MAX_VAL;
 122 #ifdef NEED_MINLOC
 123     __local uint localmem_minloc[WGS2_ALIGNED];
 124     uint minloc = INDEX_MAX;
 125 #endif
 126 #endif
 127 #ifdef NEED_MAXVAL
 128     dstT1 maxval = MIN_VAL;
 129     __local dstT1 localmem_max[WGS2_ALIGNED];
 130 #ifdef NEED_MAXLOC
 131     __local uint localmem_maxloc[WGS2_ALIGNED];
 132     uint maxloc = INDEX_MAX;
 133 #endif
 134 #endif
 135 #ifdef OP_CALC2
 136     __local dstT1 localmem_max2[WGS2_ALIGNED];
 137     dstT1 maxval2 = MIN_VAL;
 138 #endif
 139
 140     int src_index;
 141 #ifdef HAVE_MASK
 142     int mask_index;
 143 #endif
 144 #ifdef HAVE_SRC2
 145     int src2_index;
 146 #endif
 147
 148     dstT temp;
 149 #ifdef HAVE_SRC2
 150     dstT temp2;
 151 #endif
 152
 153     for (int grain = groupnum * WGS * kercn; id < total; id += grain)
 154     {
 155 #ifdef HAVE_MASK
 156 #ifdef HAVE_MASK_CONT
 157         mask_index = id;
 158 #else
 159         mask_index = mad24(id / cols, mask_step, id % cols);
 160 #endif
 161         if (mask[mask_index])
 162 #endif
 163         {
 164 #ifdef HAVE_SRC_CONT
 165             src_index = mul24(id, srcTSIZE);
 166 #else
 167             src_index = mad24(id / cols, src_step, mul24(id % cols, srcTSIZE));
 168 #endif
 169             temp = convertToDT(loadpix(srcptr + src_index));
 170 #ifdef OP_ABS
 171             temp = temp >= (dstT)(0) ? temp : -temp;
 172 #endif
 173
 174 #ifdef HAVE_SRC2
 175 #ifdef HAVE_SRC2_CONT
 176             src2_index = mul24(id, srcTSIZE);
 177 #else
 178             src2_index = mad24(id / cols, src2_step, mul24(id % cols, srcTSIZE));
 179 #endif
 180             temp2 = convertToDT(loadpix(src2ptr + src2_index));
 181             temp = temp > temp2 ? temp - temp2 : (temp2 - temp);
 182 #ifdef OP_CALC2
 183             temp2 = temp2 >= (dstT)(0) ? temp2 : -temp2;
 184 #endif
 185 #endif
 186
 187 #if kercn == 1
 188 #ifdef NEED_MINVAL
 189             if (minval > temp)
 190             {
 191                 minval = temp;
 192 #ifdef NEED_MINLOC
 193                 minloc = id;
 194 #endif
 195             }
 196 #endif
 197 #ifdef NEED_MAXVAL
 198             if (maxval < temp)
 199             {
 200                 maxval = temp;
 201 #ifdef NEED_MAXLOC
 202                 maxloc = id;
 203 #endif
 204             }
 205 #ifdef OP_CALC2
 206             if (maxval2 < temp2)
 207                 maxval2 = temp2;
 208 #endif
 209 #endif
 210 #elif kercn >= 2
 211             CALC_P(s0, 0)
 212             CALC_P(s1, 1)
 213 #if kercn >= 3
 214             CALC_P(s2, 2)
 215 #if kercn >= 4
 216             CALC_P(s3, 3)
 217 #if kercn >= 8
 218             CALC_P(s4, 4)
 219             CALC_P(s5, 5)
 220             CALC_P(s6, 6)
 221             CALC_P(s7, 7)
 222 #if kercn == 16
 223             CALC_P(s8, 8)
 224             CALC_P(s9, 9)
 225             CALC_P(sA, 10)
 226             CALC_P(sB, 11)
 227             CALC_P(sC, 12)
 228             CALC_P(sD, 13)
 229             CALC_P(sE, 14)
 230             CALC_P(sF, 15)
 231 #endif
 232 #endif
 233 #endif
 234 #endif
 235 #endif
 236         }
 237     }
 238
 239     if (lid < WGS2_ALIGNED)
 240     {
 241 #ifdef NEED_MINVAL
 242         localmem_min[lid] = minval;
 243 #endif
 244 #ifdef NEED_MAXVAL
 245         localmem_max[lid] = maxval;
 246 #endif
 247 #ifdef NEED_MINLOC
 248         localmem_minloc[lid] = minloc;
 249 #endif
 250 #ifdef NEED_MAXLOC
 251         localmem_maxloc[lid] = maxloc;
 252 #endif
 253 #ifdef OP_CALC2
 254         localmem_max2[lid] = maxval2;
 255 #endif
 256     }
 257     barrier(CLK_LOCAL_MEM_FENCE);
 258
 259     if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
 260     {
 261         int lid3 = lid - WGS2_ALIGNED;
 262 #ifdef NEED_MINVAL
 263         if (localmem_min[lid3] >= minval)
 264         {
 265 #ifdef NEED_MINLOC
 266             if (localmem_min[lid3] == minval)
 267                 localmem_minloc[lid3] = min(localmem_minloc[lid3], minloc);
 268             else
 269                 localmem_minloc[lid3] = minloc,
 270 #endif
 271                 localmem_min[lid3] = minval;
 272         }
 273 #endif
 274 #ifdef NEED_MAXVAL
 275         if (localmem_max[lid3] <= maxval)
 276         {
 277 #ifdef NEED_MAXLOC
 278             if (localmem_max[lid3] == maxval)
 279                 localmem_maxloc[lid3] = min(localmem_maxloc[lid3], maxloc);
 280             else
 281                 localmem_maxloc[lid3] = maxloc,
 282 #endif
 283                 localmem_max[lid3] = maxval;
 284         }
 285 #endif
 286 #ifdef OP_CALC2
 287         if (localmem_max2[lid3] < maxval2)
 288             localmem_max2[lid3] = maxval2;
 289 #endif
 290     }
 291     barrier(CLK_LOCAL_MEM_FENCE);
 292
 293     for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
 294     {
 295         if (lid < lsize)
 296         {
 297             int lid2 = lsize + lid;
 298
 299 #ifdef NEED_MINVAL
 300             if (localmem_min[lid] >= localmem_min[lid2])
 301             {
 302 #ifdef NEED_MINLOC
 303                 if (localmem_min[lid] == localmem_min[lid2])
 304                     localmem_minloc[lid] = min(localmem_minloc[lid2], localmem_minloc[lid]);
 305                 else
 306                     localmem_minloc[lid] = localmem_minloc[lid2],
 307 #endif
 308                     localmem_min[lid] = localmem_min[lid2];
 309             }
 310 #endif
 311 #ifdef NEED_MAXVAL
 312             if (localmem_max[lid] <= localmem_max[lid2])
 313             {
 314 #ifdef NEED_MAXLOC
 315                 if (localmem_max[lid] == localmem_max[lid2])
 316                     localmem_maxloc[lid] = min(localmem_maxloc[lid2], localmem_maxloc[lid]);
 317                 else
 318                     localmem_maxloc[lid] = localmem_maxloc[lid2],
 319 #endif
 320                     localmem_max[lid] = localmem_max[lid2];
 321             }
 322 #endif
 323 #ifdef OP_CALC2
 324             if (localmem_max2[lid] < localmem_max2[lid2])
 325                 localmem_max2[lid] = localmem_max2[lid2];
 326 #endif
 327         }
 328         barrier(CLK_LOCAL_MEM_FENCE);
 329     }
 330
 331     if (lid == 0)
 332     {
 333         int pos = 0;
 334 #ifdef NEED_MINVAL
 335         *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_min[0];
 336         pos = mad24(groupnum, (int)sizeof(dstT1), pos);
 337 #endif
 338 #ifdef NEED_MAXVAL
 339         *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max[0];
 340         pos = mad24(groupnum, (int)sizeof(dstT1), pos);
 341 #endif
 342 #ifdef NEED_MINLOC
 343         *(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_minloc[0];
 344         pos = mad24(groupnum, (int)sizeof(uint), pos);
 345 #endif
 346 #ifdef NEED_MAXLOC
 347         *(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_maxloc[0];
 348 #ifdef OP_CALC2
 349         pos = mad24(groupnum, (int)sizeof(uint), pos);
 350 #endif
 351 #endif
 352 #ifdef OP_CALC2
 353         *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max2[0];
 354 #endif
 355     }
 356 }