← 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/core/src/opencl/minmaxloc.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/core/src/opencl/minmaxloc.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
#ifdef DOUBLE_SUPPORT
 
9
#ifdef cl_amd_fp64
 
10
#pragma OPENCL EXTENSION cl_amd_fp64:enable
 
11
#elif defined (cl_khr_fp64)
 
12
#pragma OPENCL EXTENSION cl_khr_fp64:enable
 
13
#endif
 
14
#endif
 
15
 
 
16
static inline int align(int pos)
 
17
{
 
18
    return (pos + (MINMAX_STRUCT_ALIGNMENT - 1)) & (~(MINMAX_STRUCT_ALIGNMENT - 1));
 
19
}
 
20
 
 
21
#ifdef DEPTH_0
 
22
#define MIN_VAL 0
 
23
#define MAX_VAL UCHAR_MAX
 
24
#elif defined DEPTH_1
 
25
#define MIN_VAL SCHAR_MIN
 
26
#define MAX_VAL SCHAR_MAX
 
27
#elif defined DEPTH_2
 
28
#define MIN_VAL 0
 
29
#define MAX_VAL USHRT_MAX
 
30
#elif defined DEPTH_3
 
31
#define MIN_VAL SHRT_MIN
 
32
#define MAX_VAL SHRT_MAX
 
33
#elif defined DEPTH_4
 
34
#define MIN_VAL INT_MIN
 
35
#define MAX_VAL INT_MAX
 
36
#elif defined DEPTH_5
 
37
#define MIN_VAL (-FLT_MAX)
 
38
#define MAX_VAL FLT_MAX
 
39
#elif defined DEPTH_6
 
40
#define MIN_VAL (-DBL_MAX)
 
41
#define MAX_VAL DBL_MAX
 
42
#endif
 
43
 
 
44
#define noconvert
 
45
#define INDEX_MAX UINT_MAX
 
46
 
 
47
#if wdepth <= 4
 
48
#define MIN_ABS(a) convertFromU(abs(a))
 
49
#define MIN_ABS2(a, b) convertFromU(abs_diff(a, b))
 
50
#define MIN(a, b) min(a, b)
 
51
#define MAX(a, b) max(a, b)
 
52
#else
 
53
#define MIN_ABS(a) fabs(a)
 
54
#define MIN_ABS2(a, b) fabs(a - b)
 
55
#define MIN(a, b) fmin(a, b)
 
56
#define MAX(a, b) fmax(a, b)
 
57
#endif
 
58
 
 
59
#if kercn != 3
 
60
#define loadpix(addr) *(__global const srcT *)(addr)
 
61
#define srcTSIZE (int)sizeof(srcT)
 
62
#else
 
63
#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
 
64
#define srcTSIZE ((int)sizeof(srcT1) * 3)
 
65
#endif
 
66
 
 
67
#ifndef HAVE_MASK
 
68
#undef srcTSIZE
 
69
#define srcTSIZE (int)sizeof(srcT1)
 
70
#endif
 
71
 
 
72
#ifdef NEED_MINVAL
 
73
#ifdef NEED_MINLOC
 
74
#define CALC_MIN(p, inc) \
 
75
    if (minval > temp.p) \
 
76
    { \
 
77
        minval = temp.p; \
 
78
        minloc = id + inc; \
 
79
    }
 
80
#else
 
81
#define CALC_MIN(p, inc) \
 
82
    minval = MIN(minval, temp.p);
 
83
#endif
 
84
#else
 
85
#define CALC_MIN(p, inc)
 
86
#endif
 
87
 
 
88
#ifdef NEED_MAXVAL
 
89
#ifdef NEED_MAXLOC
 
90
#define CALC_MAX(p, inc) \
 
91
    if (maxval < temp.p) \
 
92
    { \
 
93
        maxval = temp.p; \
 
94
        maxloc = id + inc; \
 
95
    }
 
96
#else
 
97
#define CALC_MAX(p, inc) \
 
98
    maxval = MAX(maxval, temp.p);
 
99
#endif
 
100
#else
 
101
#define CALC_MAX(p, inc)
 
102
#endif
 
103
 
 
104
#ifdef OP_CALC2
 
105
#define CALC_MAX2(p) \
 
106
    maxval2 = MAX(maxval2, temp2.p);
 
107
#else
 
108
#define CALC_MAX2(p)
 
109
#endif
 
110
 
 
111
#define CALC_P(p, inc) \
 
112
    CALC_MIN(p, inc) \
 
113
    CALC_MAX(p, inc) \
 
114
    CALC_MAX2(p)
 
115
 
 
116
__kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_offset, int cols,
 
117
                        int total, int groupnum, __global uchar * dstptr
 
118
#ifdef HAVE_MASK
 
119
                        , __global const uchar * mask, int mask_step, int mask_offset
 
120
#endif
 
121
#ifdef HAVE_SRC2
 
122
                        , __global const uchar * src2ptr, int src2_step, int src2_offset
 
123
#endif
 
124
                        )
 
125
{
 
126
    int lid = get_local_id(0);
 
127
    int gid = get_group_id(0);
 
128
    int  id = get_global_id(0)
 
129
#ifndef HAVE_MASK
 
130
    * kercn;
 
131
#else
 
132
    ;
 
133
#endif
 
134
 
 
135
    srcptr += src_offset;
 
136
#ifdef HAVE_MASK
 
137
    mask += mask_offset;
 
138
#endif
 
139
#ifdef HAVE_SRC2
 
140
    src2ptr += src2_offset;
 
141
#endif
 
142
 
 
143
#ifdef NEED_MINVAL
 
144
    __local dstT1 localmem_min[WGS2_ALIGNED];
 
145
    dstT1 minval = MAX_VAL;
 
146
#ifdef NEED_MINLOC
 
147
    __local uint localmem_minloc[WGS2_ALIGNED];
 
148
    uint minloc = INDEX_MAX;
 
149
#endif
 
150
#endif
 
151
#ifdef NEED_MAXVAL
 
152
    dstT1 maxval = MIN_VAL;
 
153
    __local dstT1 localmem_max[WGS2_ALIGNED];
 
154
#ifdef NEED_MAXLOC
 
155
    __local uint localmem_maxloc[WGS2_ALIGNED];
 
156
    uint maxloc = INDEX_MAX;
 
157
#endif
 
158
#endif
 
159
#ifdef OP_CALC2
 
160
    __local dstT1 localmem_max2[WGS2_ALIGNED];
 
161
    dstT1 maxval2 = MIN_VAL;
 
162
#endif
 
163
 
 
164
    int src_index;
 
165
#ifdef HAVE_MASK
 
166
    int mask_index;
 
167
#endif
 
168
#ifdef HAVE_SRC2
 
169
    int src2_index;
 
170
#endif
 
171
 
 
172
    dstT temp;
 
173
#ifdef HAVE_SRC2
 
174
    dstT temp2;
 
175
#endif
 
176
 
 
177
    for (int grain = groupnum * WGS
 
178
#ifndef HAVE_MASK
 
179
        * kercn
 
180
#endif
 
181
        ; id < total; id += grain)
 
182
    {
 
183
#ifdef HAVE_MASK
 
184
#ifdef HAVE_MASK_CONT
 
185
        mask_index = id;
 
186
#else
 
187
        mask_index = mad24(id / cols, mask_step, id % cols);
 
188
#endif
 
189
        if (mask[mask_index])
 
190
#endif
 
191
        {
 
192
#ifdef HAVE_SRC_CONT
 
193
            src_index = id * srcTSIZE;//mul24(id, srcTSIZE);
 
194
#else
 
195
            src_index = mad24(id / cols, src_step, mul24(id % cols, srcTSIZE));
 
196
#endif
 
197
            temp = convertToDT(loadpix(srcptr + src_index));
 
198
#ifdef OP_ABS
 
199
            temp = MIN_ABS(temp);
 
200
#endif
 
201
 
 
202
#ifdef HAVE_SRC2
 
203
#ifdef HAVE_SRC2_CONT
 
204
            src2_index = id * srcTSIZE; //mul24(id, srcTSIZE);
 
205
#else
 
206
            src2_index = mad24(id / cols, src2_step, mul24(id % cols, srcTSIZE));
 
207
#endif
 
208
            temp2 = convertToDT(loadpix(src2ptr + src2_index));
 
209
            temp = MIN_ABS2(temp, temp2);
 
210
#ifdef OP_CALC2
 
211
            temp2 = MIN_ABS(temp2);
 
212
#endif
 
213
#endif
 
214
 
 
215
#if kercn == 1
 
216
#ifdef NEED_MINVAL
 
217
#ifdef NEED_MINLOC
 
218
            if (minval > temp)
 
219
            {
 
220
                minval = temp;
 
221
                minloc = id;
 
222
            }
 
223
#else
 
224
            minval = MIN(minval, temp);
 
225
#endif
 
226
#endif
 
227
#ifdef NEED_MAXVAL
 
228
#ifdef NEED_MAXLOC
 
229
            if (maxval < temp)
 
230
            {
 
231
                maxval = temp;
 
232
                maxloc = id;
 
233
            }
 
234
#else
 
235
            maxval = MAX(maxval, temp);
 
236
#endif
 
237
#ifdef OP_CALC2
 
238
            maxval2 = MAX(maxval2, temp2);
 
239
#endif
 
240
#endif
 
241
#elif kercn >= 2
 
242
            CALC_P(s0, 0)
 
243
            CALC_P(s1, 1)
 
244
#if kercn >= 3
 
245
            CALC_P(s2, 2)
 
246
#if kercn >= 4
 
247
            CALC_P(s3, 3)
 
248
#if kercn >= 8
 
249
            CALC_P(s4, 4)
 
250
            CALC_P(s5, 5)
 
251
            CALC_P(s6, 6)
 
252
            CALC_P(s7, 7)
 
253
#if kercn == 16
 
254
            CALC_P(s8, 8)
 
255
            CALC_P(s9, 9)
 
256
            CALC_P(sA, 10)
 
257
            CALC_P(sB, 11)
 
258
            CALC_P(sC, 12)
 
259
            CALC_P(sD, 13)
 
260
            CALC_P(sE, 14)
 
261
            CALC_P(sF, 15)
 
262
#endif
 
263
#endif
 
264
#endif
 
265
#endif
 
266
#endif
 
267
        }
 
268
    }
 
269
 
 
270
    if (lid < WGS2_ALIGNED)
 
271
    {
 
272
#ifdef NEED_MINVAL
 
273
        localmem_min[lid] = minval;
 
274
#endif
 
275
#ifdef NEED_MAXVAL
 
276
        localmem_max[lid] = maxval;
 
277
#endif
 
278
#ifdef NEED_MINLOC
 
279
        localmem_minloc[lid] = minloc;
 
280
#endif
 
281
#ifdef NEED_MAXLOC
 
282
        localmem_maxloc[lid] = maxloc;
 
283
#endif
 
284
#ifdef OP_CALC2
 
285
        localmem_max2[lid] = maxval2;
 
286
#endif
 
287
    }
 
288
    barrier(CLK_LOCAL_MEM_FENCE);
 
289
 
 
290
    if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
 
291
    {
 
292
        int lid3 = lid - WGS2_ALIGNED;
 
293
#ifdef NEED_MINVAL
 
294
#ifdef NEED_MINLOC
 
295
        if (localmem_min[lid3] >= minval)
 
296
        {
 
297
            if (localmem_min[lid3] == minval)
 
298
                localmem_minloc[lid3] = min(localmem_minloc[lid3], minloc);
 
299
            else
 
300
                localmem_minloc[lid3] = minloc,
 
301
            localmem_min[lid3] = minval;
 
302
        }
 
303
#else
 
304
        localmem_min[lid3] = MIN(localmem_min[lid3], minval);
 
305
#endif
 
306
#endif
 
307
#ifdef NEED_MAXVAL
 
308
#ifdef NEED_MAXLOC
 
309
        if (localmem_max[lid3] <= maxval)
 
310
        {
 
311
            if (localmem_max[lid3] == maxval)
 
312
                localmem_maxloc[lid3] = min(localmem_maxloc[lid3], maxloc);
 
313
            else
 
314
                localmem_maxloc[lid3] = maxloc,
 
315
            localmem_max[lid3] = maxval;
 
316
        }
 
317
#else
 
318
        localmem_max[lid3] = MAX(localmem_max[lid3], maxval);
 
319
#endif
 
320
#endif
 
321
#ifdef OP_CALC2
 
322
        localmem_max2[lid3] = MAX(localmem_max2[lid3], maxval2);
 
323
#endif
 
324
    }
 
325
    barrier(CLK_LOCAL_MEM_FENCE);
 
326
 
 
327
    for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
 
328
    {
 
329
        if (lid < lsize)
 
330
        {
 
331
            int lid2 = lsize + lid;
 
332
 
 
333
#ifdef NEED_MINVAL
 
334
#ifdef NEED_MINLOC
 
335
            if (localmem_min[lid] >= localmem_min[lid2])
 
336
            {
 
337
                if (localmem_min[lid] == localmem_min[lid2])
 
338
                    localmem_minloc[lid] = min(localmem_minloc[lid2], localmem_minloc[lid]);
 
339
                else
 
340
                    localmem_minloc[lid] = localmem_minloc[lid2],
 
341
                localmem_min[lid] = localmem_min[lid2];
 
342
            }
 
343
#else
 
344
            localmem_min[lid] = MIN(localmem_min[lid], localmem_min[lid2]);
 
345
#endif
 
346
#endif
 
347
#ifdef NEED_MAXVAL
 
348
#ifdef NEED_MAXLOC
 
349
            if (localmem_max[lid] <= localmem_max[lid2])
 
350
            {
 
351
                if (localmem_max[lid] == localmem_max[lid2])
 
352
                    localmem_maxloc[lid] = min(localmem_maxloc[lid2], localmem_maxloc[lid]);
 
353
                else
 
354
                    localmem_maxloc[lid] = localmem_maxloc[lid2],
 
355
                localmem_max[lid] = localmem_max[lid2];
 
356
            }
 
357
#else
 
358
            localmem_max[lid] = MAX(localmem_max[lid], localmem_max[lid2]);
 
359
#endif
 
360
#endif
 
361
#ifdef OP_CALC2
 
362
            localmem_max2[lid] = MAX(localmem_max2[lid], localmem_max2[lid2]);
 
363
#endif
 
364
        }
 
365
        barrier(CLK_LOCAL_MEM_FENCE);
 
366
    }
 
367
 
 
368
    if (lid == 0)
 
369
    {
 
370
        int pos = 0;
 
371
#ifdef NEED_MINVAL
 
372
        *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_min[0];
 
373
        pos = mad24(groupnum, (int)sizeof(dstT1), pos);
 
374
        pos = align(pos);
 
375
#endif
 
376
#ifdef NEED_MAXVAL
 
377
        *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max[0];
 
378
        pos = mad24(groupnum, (int)sizeof(dstT1), pos);
 
379
        pos = align(pos);
 
380
#endif
 
381
#ifdef NEED_MINLOC
 
382
        *(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_minloc[0];
 
383
        pos = mad24(groupnum, (int)sizeof(uint), pos);
 
384
        pos = align(pos);
 
385
#endif
 
386
#ifdef NEED_MAXLOC
 
387
        *(__global uint *)(dstptr + mad24(gid, (int)sizeof(uint), pos)) = localmem_maxloc[0];
 
388
#ifdef OP_CALC2
 
389
        pos = mad24(groupnum, (int)sizeof(uint), pos);
 
390
        pos = align(pos);
 
391
#endif
 
392
#endif
 
393
#ifdef OP_CALC2
 
394
        *(__global dstT1 *)(dstptr + mad24(gid, (int)sizeof(dstT1), pos)) = localmem_max2[0];
 
395
#endif
 
396
    }
 
397
}