← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/imgproc/src/opencl/hough_lines.cl

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/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
 
 
b'\\ No newline at end of file'