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     if (y < accum_rows-2)
 188     {
 189         __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x+1, (int) sizeof(int), accum_offset));
 190         __global int4* lines = (__global int4*)(lines_ptr + lines_offset);
 191         __global int* lines_index = lines_index_ptr + 1;
 192
 193         int curVote = ACCUM(accum);
 194
 195         if (curVote >= threshold &&
 196             curVote > ACCUM(accum - accum_step - sizeof(int)) &&
 197             curVote > ACCUM(accum - accum_step) &&
 198             curVote > ACCUM(accum - accum_step + sizeof(int)) &&
 199             curVote > ACCUM(accum - sizeof(int)) &&
 200             curVote > ACCUM(accum + sizeof(int)) &&
 201             curVote > ACCUM(accum + accum_step - sizeof(int)) &&
 202             curVote > ACCUM(accum + accum_step) &&
 203             curVote > ACCUM(accum + accum_step + sizeof(int)))
 204         {
 205             const float radius = (x - (accum_cols - 2 - 1) * 0.5f) * rho;
 206             const float angle = y * theta;
 207
 208             float cosa;
 209             float sina = sincos(angle, &cosa);
 210
 211             float2 p0 = (float2)(cosa * radius, sina * radius);
 212             float2 dir = (float2)(-sina, cosa);
 213
 214             float2 pb[4] = { (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1) };
 215             float a;
 216
 217             if (dir.x != 0)
 218             {
 219                 a = -p0.x / dir.x;
 220                 pb[0].x = 0;
 221                 pb[0].y = p0.y + a * dir.y;
 222
 223                 a = (src_cols - 1 - p0.x) / dir.x;
 224                 pb[1].x = src_cols - 1;
 225                 pb[1].y = p0.y + a * dir.y;
 226             }
 227
 228             if (dir.y != 0)
 229             {
 230                 a = -p0.y / dir.y;
 231                 pb[2].x = p0.x + a * dir.x;
 232                 pb[2].y = 0;
 233
 234                 a = (src_rows - 1 - p0.y) / dir.y;
 235                 pb[3].x = p0.x + a * dir.x;
 236                 pb[3].y = src_rows - 1;
 237             }
 238
 239             if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < src_rows))
 240             {
 241                 p0 = pb[0];
 242                 if (dir.x < 0)
 243                     dir = -dir;
 244             }
 245             else if (pb[1].x == src_cols - 1 && (pb[1].y >= 0 && pb[1].y < src_rows))
 246             {
 247                 p0 = pb[1];
 248                 if (dir.x > 0)
 249                     dir = -dir;
 250             }
 251             else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < src_cols))
 252             {
 253                 p0 = pb[2];
 254                 if (dir.y < 0)
 255                     dir = -dir;
 256             }
 257             else if (pb[3].y == src_rows - 1 && (pb[3].x >= 0 && pb[3].x < src_cols))
 258             {
 259                 p0 = pb[3];
 260                 if (dir.y > 0)
 261                     dir = -dir;
 262             }
 263
 264             dir /= max(fabs(dir.x), fabs(dir.y));
 265
 266             float2 line_end[2];
 267             int gap;
 268             bool inLine = false;
 269
 270             if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows)
 271                 return;
 272
 273             for (;;)
 274             {
 275                 if (*(src_ptr + mad24(p0.y, src_step, p0.x + src_offset)))
 276                 {
 277                     gap = 0;
 278
 279                     if (!inLine)
 280                     {
 281                         line_end[0] = p0;
 282                         line_end[1] = p0;
 283                         inLine = true;
 284                     }
 285                     else
 286                     {
 287                         line_end[1] = p0;
 288                     }
 289                 }
 290                 else if (inLine)
 291                 {
 292                     if (++gap > lineGap)
 293                     {
 294                         bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength ||
 295                                          fabs(line_end[1].y - line_end[0].y) >= lineLength;
 296
 297                         if (good_line)
 298                         {
 299                             int index = atomic_inc(lines_index);
 300                             if (index < linesMax)
 301                                 lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
 302                         }
 303
 304                         gap = 0;
 305                         inLine = false;
 306                     }
 307                 }
 308
 309                 p0 = p0 + dir;
 310                 if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows)
 311                 {
 312                     if (inLine)
 313                     {
 314                         bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength ||
 315                                          fabs(line_end[1].y - line_end[0].y) >= lineLength;
 316
 317                         if (good_line)
 318                         {
 319                             int index = atomic_inc(lines_index);
 320                             if (index < linesMax)
 321                                 lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
 322                         }
 323
 324                     }
 325                     break;
 326                 }
 327             }
 328
 329         }
 330     }
 331 }
 332
 333 #endif