modules/imgproc/src/opencl/hough_lines.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 #define ACCUM(ptr) *((__global int*)(ptr))
   9
  10 #ifdef MAKE_POINTS_LIST
  11
  12 __kernel void make_point_list(__global const uchar * src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
  13                               __global uchar * list_ptr, int list_step, int list_offset, __global int* global_offset)
  14 {
  15     int x = get_local_id(0);
  16     int y = get_group_id(1);
  17
  18     __local int l_index, l_offset;
  19     __local int l_points[LOCAL_SIZE];
  20     __global const uchar * src = src_ptr + mad24(y, src_step, src_offset);
  21     __global int * list = (__global int*)(list_ptr + list_offset);
  22
  23     if (x == 0)
  24         l_index = 0;
  25
  26     barrier(CLK_LOCAL_MEM_FENCE);
  27
  28     if (y < src_rows)
  29     {
  30         y <<= 16;
  31
  32         for (int i=x; i < src_cols; i+=GROUP_SIZE)
  33         {
  34             if (src[i])
  35             {
  36                 int val = y | i;
  37                 int index = atomic_inc(&l_index);
  38                 l_points[index] = val;
  39             }
  40         }
  41     }
  42
  43     barrier(CLK_LOCAL_MEM_FENCE);
  44
  45     if (x == 0)
  46         l_offset = atomic_add(global_offset, l_index);
  47
  48     barrier(CLK_LOCAL_MEM_FENCE);
  49
  50     list += l_offset;
  51     for (int i=x; i < l_index; i+=GROUP_SIZE)
  52     {
  53         list[i] = l_points[i];
  54     }
  55 }
  56
  57 #elif defined FILL_ACCUM_GLOBAL
  58
  59 __kernel void fill_accum_global(__global const uchar * list_ptr, int list_step, int list_offset,
  60                                 __global uchar * accum_ptr, int accum_step, int accum_offset,
  61                                 int total_points, float irho, float theta, int numrho, int numangle)
  62 {
  63     int theta_idx = get_global_id(1);
  64     int count_idx = get_global_id(0);
  65     int glob_size = get_global_size(0);
  66     float cosVal;
  67     float sinVal = sincos(theta * ((float)theta_idx), &cosVal);
  68     sinVal *= irho;
  69     cosVal *= irho;
  70
  71     __global const int * list = (__global const int*)(list_ptr + list_offset);
  72     __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx + 1, accum_step, accum_offset));
  73     const int shift = (numrho - 1) / 2;
  74
  75     if (theta_idx < numangle)
  76     {
  77         for (int i = count_idx; i < total_points; i += glob_size)
  78         {
  79             const int val = list[i];
  80             const int x = (val & 0xFFFF);
  81             const int y = (val >> 16) & 0xFFFF;
  82
  83             int r = convert_int_rte(mad(x, cosVal, y * sinVal)) + shift;
  84             atomic_inc(accum + r + 1);
  85         }
  86     }
  87 }
  88
  89 #elif defined FILL_ACCUM_LOCAL
  90
  91 __kernel void fill_accum_local(__global const uchar * list_ptr, int list_step, int list_offset,
  92                                __global uchar * accum_ptr, int accum_step, int accum_offset,
  93                                int total_points, float irho, float theta, int numrho, int numangle)
  94 {
  95     int theta_idx = get_group_id(1);
  96     int count_idx = get_local_id(0);
  97
  98     if (theta_idx > 0 && theta_idx < numangle + 1)
  99     {
 100         float cosVal;
 101         float sinVal = sincos(theta * (float) (theta_idx-1), &cosVal);
 102         sinVal *= irho;
 103         cosVal *= irho;
 104
 105         __local int l_accum[BUFFER_SIZE];
 106         for (int i=count_idx; i<BUFFER_SIZE; i+=LOCAL_SIZE)
 107             l_accum[i] = 0;
 108
 109         barrier(CLK_LOCAL_MEM_FENCE);
 110
 111         __global const int * list = (__global const int*)(list_ptr + list_offset);
 112         const int shift = (numrho - 1) / 2;
 113
 114         for (int i = count_idx; i < total_points; i += LOCAL_SIZE)
 115         {
 116             const int point = list[i];
 117             const int x = (point & 0xFFFF);
 118             const int y = point >> 16;
 119
 120             int r = convert_int_rte(mad(x, cosVal, y * sinVal)) + shift;
 121             atomic_inc(l_accum + r + 1);
 122         }
 123
 124         barrier(CLK_LOCAL_MEM_FENCE);
 125
 126         __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx, accum_step, accum_offset));
 127         for (int i=count_idx; i<BUFFER_SIZE; i+=LOCAL_SIZE)
 128             accum[i] = l_accum[i];
 129     }
 130     else if (theta_idx < numangle + 2)
 131     {
 132         __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx, accum_step, accum_offset));
 133         for (int i=count_idx; i<BUFFER_SIZE; i+=LOCAL_SIZE)
 134             accum[i] = 0;
 135     }
 136 }
 137
 138 #elif defined GET_LINES
 139
 140 __kernel void get_lines(__global uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols,
 141                          __global uchar * lines_ptr, int lines_step, int lines_offset, __global int* lines_index_ptr,
 142                          int linesMax, int threshold, float rho, float theta)
 143 {
 144     int x0 = get_global_id(0);
 145     int y = get_global_id(1);
 146     int glob_size = get_global_size(0);
 147
 148     if (y < accum_rows-2)
 149     {
 150         __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x0+1, (int) sizeof(int), accum_offset));
 151         __global float2* lines = (__global float2*)(lines_ptr + lines_offset);
 152         __global int* lines_index = lines_index_ptr + 1;
 153
 154         for (int x=x0; x<accum_cols-2; x+=glob_size)
 155         {
 156             int curVote = ACCUM(accum);
 157
 158             if (curVote > threshold && curVote > ACCUM(accum - sizeof(int)) && curVote >= ACCUM(accum + sizeof(int)) &&
 159                 curVote > ACCUM(accum - accum_step) && curVote >= ACCUM(accum + accum_step))
 160             {
 161                 int index = atomic_inc(lines_index);
 162
 163                 if (index < linesMax)
 164                 {
 165                     float radius = (x - (accum_cols - 3) * 0.5f) * rho;
 166                     float angle = y * theta;
 167
 168                     lines[index] = (float2)(radius, angle);
 169                 }
 170             }
 171
 172             accum += glob_size * (int) sizeof(int);
 173         }
 174     }
 175 }
 176
 177 #elif GET_LINES_PROBABOLISTIC
 178
 179 __kernel void get_lines(__global const uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols,
 180                         __global const uchar * src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
 181                         __global uchar * lines_ptr, int lines_step, int lines_offset, __global int* lines_index_ptr,
 182                         int linesMax, int threshold, int lineLength, int lineGap, float rho, float theta)
 183 {
 184     int x = get_global_id(0);
 185     int y = get_global_id(1);
 186
 187     __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x+1, (int) sizeof(int), accum_offset));
 188     __global int4* lines = (__global int4*)(lines_ptr + lines_offset);
 189     __global int* lines_index = lines_index_ptr + 1;
 190
 191     int curVote = ACCUM(accum);
 192
 193     if (curVote >= threshold &&
 194         curVote > ACCUM(accum - accum_step - sizeof(int)) &&
 195         curVote > ACCUM(accum - accum_step) &&
 196         curVote > ACCUM(accum - accum_step + sizeof(int)) &&
 197         curVote > ACCUM(accum - sizeof(int)) &&
 198         curVote > ACCUM(accum + sizeof(int)) &&
 199         curVote > ACCUM(accum + accum_step - sizeof(int)) &&
 200         curVote > ACCUM(accum + accum_step) &&
 201         curVote > ACCUM(accum + accum_step + sizeof(int)))
 202     {
 203         const float radius = (x - (accum_cols - 2 - 1) * 0.5f) * rho;
 204         const float angle = y * theta;
 205
 206         float cosa;
 207         float sina = sincos(angle, &cosa);
 208
 209         float2 p0 = (float2)(cosa * radius, sina * radius);
 210         float2 dir = (float2)(-sina, cosa);
 211
 212         float2 pb[4] = { (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1) };
 213         float a;
 214
 215         if (dir.x != 0)
 216         {
 217             a = -p0.x / dir.x;
 218             pb[0].x = 0;
 219             pb[0].y = p0.y + a * dir.y;
 220
 221             a = (src_cols - 1 - p0.x) / dir.x;
 222             pb[1].x = src_cols - 1;
 223             pb[1].y = p0.y + a * dir.y;
 224         }
 225
 226         if (dir.y != 0)
 227         {
 228             a = -p0.y / dir.y;
 229             pb[2].x = p0.x + a * dir.x;
 230             pb[2].y = 0;
 231
 232             a = (src_rows - 1 - p0.y) / dir.y;
 233             pb[3].x = p0.x + a * dir.x;
 234             pb[3].y = src_rows - 1;
 235         }
 236
 237         if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < src_rows))
 238         {
 239             p0 = pb[0];
 240             if (dir.x < 0)
 241                 dir = -dir;
 242         }
 243         else if (pb[1].x == src_cols - 1 && (pb[1].y >= 0 && pb[1].y < src_rows))
 244         {
 245             p0 = pb[1];
 246             if (dir.x > 0)
 247                 dir = -dir;
 248         }
 249         else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < src_cols))
 250         {
 251             p0 = pb[2];
 252             if (dir.y < 0)
 253                 dir = -dir;
 254         }
 255         else if (pb[3].y == src_rows - 1 && (pb[3].x >= 0 && pb[3].x < src_cols))
 256         {
 257             p0 = pb[3];
 258             if (dir.y > 0)
 259                 dir = -dir;
 260         }
 261
 262         dir /= max(fabs(dir.x), fabs(dir.y));
 263
 264         float2 line_end[2];
 265         int gap;
 266         bool inLine = false;
 267
 268         if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows)
 269             return;
 270
 271         for (;;)
 272         {
 273             if (*(src_ptr + mad24(p0.y, src_step, p0.x + src_offset)))
 274             {
 275                 gap = 0;
 276
 277                 if (!inLine)
 278                 {
 279                     line_end[0] = p0;
 280                     line_end[1] = p0;
 281                     inLine = true;
 282                 }
 283                 else
 284                 {
 285                     line_end[1] = p0;
 286                 }
 287             }
 288             else if (inLine)
 289             {
 290                 if (++gap > lineGap)
 291                 {
 292                     bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength ||
 293                                      fabs(line_end[1].y - line_end[0].y) >= lineLength;
 294
 295                     if (good_line)
 296                     {
 297                         int index = atomic_inc(lines_index);
 298                         if (index < linesMax)
 299                             lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
 300                     }
 301
 302                     gap = 0;
 303                     inLine = false;
 304                 }
 305             }
 306
 307             p0 = p0 + dir;
 308             if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows)
 309             {
 310                 if (inLine)
 311                 {
 312                     bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength ||
 313                                      fabs(line_end[1].y - line_end[0].y) >= lineLength;
 314
 315                     if (good_line)
 316                     {
 317                         int index = atomic_inc(lines_index);
 318                         if (index < linesMax)
 319                             lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
 320                     }
 321
 322                 }
 323                 break;
 324             }
 325         }
 326
 327     }
 328 }
 329
 330 #endif