modules/imgproc/src/opencl/gftt.cl

   1 /*M///////////////////////////////////////////////////////////////////////////////////////
   2 //
   3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
   4 //
   5 //  By downloading, copying, installing or using the software you agree to this license.
   6 //  If you do not agree to this license, do not download, install,
   7 //  copy or use the software.
   8 //
   9 //
  10 //                           License Agreement
  11 //                For Open Source Computer Vision Library
  12 //
  13 // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
  14 // Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
  15 // Third party copyrights are property of their respective owners.
  16 //
  17 // @Authors
  18 //    Zhang Ying, zhangying913@gmail.com
  19 //
  20 // Redistribution and use in source and binary forms, with or without modification,
  21 // are permitted provided that the following conditions are met:
  22 //
  23 //   * Redistribution's of source code must retain the above copyright notice,
  24 //     this list of conditions and the following disclaimer.
  25 //
  26 //   * Redistribution's in binary form must reproduce the above copyright notice,
  27 //     this list of conditions and the following disclaimer in the documentation
  28 //     and/or other materials provided with the distribution.
  29 //
  30 //   * The name of the copyright holders may not be used to endorse or promote products
  31 //     derived from this software without specific prior written permission.
  32 //
  33 // This software is provided by the copyright holders and contributors as is and
  34 // any express or implied warranties, including, but not limited to, the implied
  35 // warranties of merchantability and fitness for a particular purpose are disclaimed.
  36 // In no event shall the Intel Corporation or contributors be liable for any direct,
  37 // indirect, incidental, special, exemplary, or consequential damages
  38 // (including, but not limited to, procurement of substitute goods or services;
  39 // loss of use, data, or profits; or business interruption) however caused
  40 // and on any theory of liability, whether in contract, strict liability,
  41 // or tort (including negligence or otherwise) arising in any way out of
  42 // the use of this software, even if advised of the possibility of such damage.
  43 //
  44 //M*/
  45
  46 #ifdef OP_MAX_EIGEN_VAL
  47
  48 __kernel void maxEigenVal(__global const uchar * srcptr, int src_step, int src_offset, int cols,
  49                           int total, __global uchar * dstptr
  50 #ifdef HAVE_MASK
  51                           , __global const uchar * maskptr, int mask_step, int mask_offset
  52 #endif
  53                           )
  54 {
  55     int lid = get_local_id(0);
  56     int gid = get_group_id(0);
  57     int  id = get_global_id(0);
  58
  59     __local float localmem_max[WGS2_ALIGNED];
  60     float maxval = -FLT_MAX;
  61
  62     for (int grain = groupnum * WGS; id < total; id += grain)
  63     {
  64         int src_index = mad24(id / cols, src_step, mad24((id % cols), (int)sizeof(float), src_offset));
  65 #ifdef HAVE_MASK
  66         int mask_index = mad24(id / cols, mask_step, id % cols + mask_offset);
  67         if (mask[mask_index])
  68 #endif
  69             maxval = max(maxval, *(__global const float *)(srcptr + src_index));
  70     }
  71
  72     if (lid < WGS2_ALIGNED)
  73         localmem_max[lid] = maxval;
  74     barrier(CLK_LOCAL_MEM_FENCE);
  75
  76     if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
  77         localmem_max[lid - WGS2_ALIGNED] = max(maxval, localmem_max[lid - WGS2_ALIGNED]);
  78     barrier(CLK_LOCAL_MEM_FENCE);
  79
  80     for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
  81     {
  82         if (lid < lsize)
  83         {
  84            int lid2 = lsize + lid;
  85            localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]);
  86         }
  87         barrier(CLK_LOCAL_MEM_FENCE);
  88     }
  89
  90     if (lid == 0)
  91         *(__global float *)(dstptr + (int)sizeof(float) * gid) = localmem_max[0];
  92 }
  93
  94 __kernel void maxEigenValTask(__global float * dst, float qualityLevel,
  95                               __global int * cornersptr)
  96 {
  97     float maxval = -FLT_MAX;
  98
  99     #pragma unroll
 100     for (int x = 0; x < groupnum; ++x)
 101         maxval = max(maxval, dst[x]);
 102
 103     dst[0] = maxval * qualityLevel;
 104     cornersptr[0] = 0;
 105 }
 106
 107 #elif OP_FIND_CORNERS
 108
 109 #define GET_SRC_32F(_y, _x) *(__global const float *)(eigptr + (_y) * eig_step + (_x) * (int)sizeof(float) )
 110
 111 __kernel void findCorners(__global const uchar * eigptr, int eig_step, int eig_offset,
 112 #ifdef HAVE_MASK
 113                           __global const uchar * mask, int mask_step, int mask_offset,
 114 #endif
 115                           __global uchar * cornersptr, int rows, int cols,
 116                           __constant float * threshold, int max_corners)
 117 {
 118     int x = get_global_id(0);
 119     int y = get_global_id(1);
 120
 121     __global int* counter = (__global int*) cornersptr;
 122     __global float2 * corners = (__global float2 *)(cornersptr + (int)sizeof(float2));
 123
 124     if (y < rows && x < cols
 125 #ifdef HAVE_MASK
 126             && mask[mad24(y, mask_step, x + mask_offset)]
 127 #endif
 128         )
 129     {
 130         ++x, ++y;
 131         float val = GET_SRC_32F(y, x);
 132
 133         if (val > threshold[0])
 134         {
 135             float maxVal = val;
 136             maxVal = max(GET_SRC_32F(y - 1, x - 1), maxVal);
 137             maxVal = max(GET_SRC_32F(y - 1, x    ), maxVal);
 138             maxVal = max(GET_SRC_32F(y - 1, x + 1), maxVal);
 139
 140             maxVal = max(GET_SRC_32F(y    , x - 1), maxVal);
 141             maxVal = max(GET_SRC_32F(y    , x + 1), maxVal);
 142
 143             maxVal = max(GET_SRC_32F(y + 1, x - 1), maxVal);
 144             maxVal = max(GET_SRC_32F(y + 1, x    ), maxVal);
 145             maxVal = max(GET_SRC_32F(y + 1, x + 1), maxVal);
 146
 147             if (val == maxVal)
 148             {
 149                 int ind = atomic_inc(counter);
 150                 if (ind < max_corners)
 151                 {
 152                     // pack and store eigenvalue and its coordinates
 153                     corners[ind].x = val;
 154                     corners[ind].y = as_float(y | (x << 16));
 155                 }
 156             }
 157         }
 158     }
 159 }
 160
 161 #endif