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