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
