← 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/filter2DSmall.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/filter2DSmall.cl
 
1
/*M///////////////////////////////////////////////////////////////////////////////////////
 
2
//
 
3
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
 
4
//
 
5
//  By downloading, copying, installing or using the software you agree to this license.
 
6
//  If you do not agree to this license, do not download, install,
 
7
//  copy or use the software.
 
8
//
 
9
//
 
10
//                           License Agreement
 
11
//                For Open Source Computer Vision Library
 
12
//
 
13
// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
 
14
// Copyright (C) 2014, Intel Corporation, all rights reserved.
 
15
// Third party copyrights are property of their respective owners.
 
16
//
 
17
// Redistribution and use in source and binary forms, with or without modification,
 
18
// are permitted provided that the following conditions are met:
 
19
//
 
20
//   * Redistribution's of source code must retain the above copyright notice,
 
21
//     this list of conditions and the following disclaimer.
 
22
//
 
23
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
24
//     this list of conditions and the following disclaimer in the documentation
 
25
//     and/or other materials provided with the distribution.
 
26
//
 
27
//   * The name of the copyright holders may not be used to endorse or promote products
 
28
//     derived from this software without specific prior written permission.
 
29
//
 
30
// This software is provided by the copyright holders and contributors as is and
 
31
// any express or implied warranties, including, but not limited to, the implied
 
32
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
33
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
34
// indirect, incidental, special, exemplary, or consequential damages
 
35
// (including, but not limited to, procurement of substitute goods or services;
 
36
// loss of use, data, or profits; or business interruption) however caused
 
37
// and on any theory of liability, whether in contract, strict liability,
 
38
// or tort (including negligence or otherwise) arising in any way out of
 
39
// the use of this software, even if advised of the possibility of such damage.
 
40
//
 
41
//M*/
 
42
 
 
43
#ifdef BORDER_REPLICATE
 
44
//BORDER_REPLICATE:     aaaaaa|abcdefgh|hhhhhhh
 
45
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? (l_edge)   : (i))
 
46
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? (r_edge)-1 : (addr))
 
47
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? (t_edge)   :(i))
 
48
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? (b_edge)-1 :(addr))
 
49
#endif
 
50
 
 
51
#ifdef BORDER_REFLECT
 
52
//BORDER_REFLECT:       fedcba|abcdefgh|hgfedcb
 
53
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? -(i)-1               : (i))
 
54
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
 
55
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? -(i)-1 : (i))
 
56
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
 
57
#endif
 
58
 
 
59
#ifdef BORDER_REFLECT_101
 
60
//BORDER_REFLECT_101:   gfedcb|abcdefgh|gfedcba
 
61
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? -(i)                 : (i))
 
62
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
 
63
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? -(i)                 : (i))
 
64
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
 
65
#endif
 
66
 
 
67
//blur function does not support BORDER_WRAP
 
68
#ifdef BORDER_WRAP
 
69
//BORDER_WRAP:          cdefgh|abcdefgh|abcdefg
 
70
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? (i)+(r_edge) : (i))
 
71
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
 
72
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? (i)+(b_edge) : (i))
 
73
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
 
74
#endif
 
75
 
 
76
#ifdef BORDER_ISOLATED
 
77
#define ISOLATED_MIN(VAL) (VAL)
 
78
#else
 
79
#define ISOLATED_MIN(VAL) 0
 
80
#endif
 
81
 
 
82
#ifdef EXTRA_EXTRAPOLATION // border > src image size
 
83
#ifdef BORDER_CONSTANT
 
84
// None
 
85
#elif defined BORDER_REPLICATE
 
86
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
 
87
    { \
 
88
        x = max(min(x, maxX - 1), minX); \
 
89
        y = max(min(y, maxY - 1), minY); \
 
90
    }
 
91
#elif defined BORDER_WRAP
 
92
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
 
93
    { \
 
94
        if (x < minX) \
 
95
            x -= ((x - maxX + 1) / maxX) * maxX; \
 
96
        if (x >= maxX) \
 
97
            x %= maxX; \
 
98
        if (y < minY) \
 
99
            y -= ((y - maxY + 1) / maxY) * maxY; \
 
100
        if (y >= maxY) \
 
101
            y %= maxY; \
 
102
    }
 
103
#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
 
104
#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \
 
105
    { \
 
106
        if (maxX - minX == 1) \
 
107
            x = minX; \
 
108
        else \
 
109
            do \
 
110
            { \
 
111
                if (x < minX) \
 
112
                    x = minX - (x - minX) - 1 + delta; \
 
113
                else \
 
114
                    x = maxX - 1 - (x - maxX) - delta; \
 
115
            } \
 
116
            while (x >= maxX || x < minX); \
 
117
        \
 
118
        if (maxY - minY == 1) \
 
119
            y = minY; \
 
120
        else \
 
121
            do \
 
122
            { \
 
123
                if (y < minY) \
 
124
                    y = minY - (y - minY) - 1 + delta; \
 
125
                else \
 
126
                    y = maxY - 1 - (y - maxY) - delta; \
 
127
            } \
 
128
            while (y >= maxY || y < minY); \
 
129
    }
 
130
#ifdef BORDER_REFLECT
 
131
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0)
 
132
#elif defined(BORDER_REFLECT_101) || defined(BORDER_REFLECT101)
 
133
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1)
 
134
#endif
 
135
#else
 
136
#error No extrapolation method
 
137
#endif
 
138
#else
 
139
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
 
140
    { \
 
141
        int _row = y - ISOLATED_MIN(minY), _col = x - ISOLATED_MIN(minX); \
 
142
        _row = ADDR_H(_row, 0, maxY - ISOLATED_MIN(minY)); \
 
143
        _row = ADDR_B(_row, maxY - ISOLATED_MIN(minY), _row); \
 
144
        y = _row + ISOLATED_MIN(minY); \
 
145
        \
 
146
        _col = ADDR_L(_col, 0, maxX - ISOLATED_MIN(minX)); \
 
147
        _col = ADDR_R(_col, maxX - ISOLATED_MIN(minX), _col); \
 
148
        x = _col + ISOLATED_MIN(minX); \
 
149
    }
 
150
#endif
 
151
 
 
152
#ifdef DOUBLE_SUPPORT
 
153
#ifdef cl_amd_fp64
 
154
#pragma OPENCL EXTENSION cl_amd_fp64:enable
 
155
#elif defined (cl_khr_fp64)
 
156
#pragma OPENCL EXTENSION cl_khr_fp64:enable
 
157
#endif
 
158
#endif
 
159
 
 
160
#if cn != 3
 
161
#define loadpix(addr) *(__global const srcT *)(addr)
 
162
#define storepix(val, addr)  *(__global dstT *)(addr) = val
 
163
#define SRCSIZE (int)sizeof(srcT)
 
164
#define DSTSIZE (int)sizeof(dstT)
 
165
#else
 
166
#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
 
167
#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
 
168
#define SRCSIZE (int)sizeof(srcT1) * cn
 
169
#define DSTSIZE (int)sizeof(dstT1) * cn
 
170
#endif
 
171
 
 
172
#define noconvert
 
173
 
 
174
struct RectCoords
 
175
{
 
176
    int x1, y1, x2, y2;
 
177
};
 
178
 
 
179
#ifdef BORDER_ISOLATED
 
180
inline bool isBorder(const struct RectCoords bounds, int2 coord, int numPixels)
 
181
{
 
182
    return (coord.x < bounds.x1 || coord.y < bounds.y1 || coord.x + numPixels > bounds.x2 || coord.y >= bounds.y2);
 
183
}
 
184
#else
 
185
inline bool isBorder(const struct RectCoords bounds, int2 coord, int numPixels)
 
186
{
 
187
    return (coord.x < 0 || coord.y < 0 || coord.x + numPixels > bounds.x2 || coord.y >= bounds.y2);
 
188
}
 
189
#endif
 
190
 
 
191
inline WT getBorderPixel(const struct RectCoords bounds, int2 coord,
 
192
                  __global const uchar* srcptr, int srcstep)
 
193
{
 
194
#ifdef BORDER_CONSTANT
 
195
    return (WT)(0);
 
196
#else
 
197
    int selected_col = coord.x;
 
198
    int selected_row = coord.y;
 
199
 
 
200
    EXTRAPOLATE(selected_col, selected_row,
 
201
            bounds.x1, bounds.y1,
 
202
            bounds.x2, bounds.y2
 
203
        );
 
204
 
 
205
    coord = (int2)(selected_col, selected_row);
 
206
    __global const uchar* ptr = srcptr + mul24(coord.y, srcstep) +
 
207
                                coord.x * SRCSIZE;
 
208
    return convertToWT(loadpix(ptr));
 
209
#endif
 
210
}
 
211
 
 
212
inline WT readSrcPixelSingle(int2 pos, __global const uchar* srcptr,
 
213
                             int srcstep, const struct RectCoords srcCoords)
 
214
{
 
215
    if (!isBorder(srcCoords, pos, 1))
 
216
    {
 
217
        __global const uchar* ptr = srcptr + mul24(pos.y, srcstep) +
 
218
                                    pos.x * SRCSIZE;
 
219
 
 
220
        return convertToWT(loadpix(ptr));
 
221
    }
 
222
    else
 
223
    {
 
224
        return getBorderPixel(srcCoords, pos, srcptr, srcstep);
 
225
    }
 
226
}
 
227
 
 
228
#define __CAT(x, y) x##y
 
229
#define CAT(x, y) __CAT(x, y)
 
230
 
 
231
#define vload1(OFFSET, PTR) (*(PTR + OFFSET))
 
232
#define PX_LOAD_VEC_TYPE CAT(srcT1, PX_LOAD_VEC_SIZE)
 
233
#define PX_LOAD_FLOAT_VEC_TYPE CAT(WT1, PX_LOAD_VEC_SIZE)
 
234
//#define PX_LOAD_FLOAT_VEC_CONV CAT(convert_, PX_LOAD_FLOAT_VEC_TYPE)
 
235
 
 
236
#if PX_LOAD_VEC_SIZE == 1
 
237
#define PX_LOAD_FLOAT_VEC_CONV (float)
 
238
#elif PX_LOAD_VEC_SIZE == 2
 
239
#define PX_LOAD_FLOAT_VEC_CONV convert_float2
 
240
#elif PX_LOAD_VEC_SIZE == 3
 
241
#define PX_LOAD_FLOAT_VEC_CONV convert_float3
 
242
#elif PX_LOAD_VEC_SIZE == 4
 
243
#define PX_LOAD_FLOAT_VEC_CONV convert_float4
 
244
#endif
 
245
 
 
246
#define PX_LOAD CAT(vload, PX_LOAD_VEC_SIZE)
 
247
#define float1 float
 
248
 
 
249
inline PX_LOAD_FLOAT_VEC_TYPE readSrcPixelGroup(int2 pos, __global const uchar* srcptr,
 
250
                                                int srcstep, const struct RectCoords srcCoords)
 
251
{
 
252
    __global const srcT1* ptr = (__global const srcT1*)
 
253
                                (srcptr + mul24(pos.y, srcstep) +
 
254
                                 pos.x * SRCSIZE);
 
255
    return PX_LOAD_FLOAT_VEC_CONV(PX_LOAD(0, ptr));
 
256
}
 
257
 
 
258
// Macros to ensure unrolled loops
 
259
#define LOOP1(VAR, STMT) (STMT); (VAR)++;
 
260
#define LOOP2(VAR, STMT) LOOP1(VAR, STMT); (STMT); (VAR)++;
 
261
#define LOOP3(VAR, STMT) LOOP2(VAR, STMT); (STMT); (VAR)++;
 
262
#define LOOP4(VAR, STMT) LOOP3(VAR, STMT); (STMT); (VAR)++;
 
263
#define LOOP5(VAR, STMT) LOOP4(VAR, STMT); (STMT); (VAR)++;
 
264
#define LOOP6(VAR, STMT) LOOP5(VAR, STMT); (STMT); (VAR)++;
 
265
#define LOOP7(VAR, STMT) LOOP6(VAR, STMT); (STMT); (VAR)++;
 
266
#define LOOP8(VAR, STMT) LOOP7(VAR, STMT); (STMT); (VAR)++;
 
267
#define LOOP9(VAR, STMT) LOOP8(VAR, STMT); (STMT); (VAR)++;
 
268
#define LOOP10(VAR, STMT) LOOP9(VAR, STMT); (STMT); (VAR)++;
 
269
#define LOOP11(VAR, STMT) LOOP10(VAR, STMT); (STMT); (VAR)++;
 
270
#define LOOP12(VAR, STMT) LOOP11(VAR, STMT); (STMT); (VAR)++;
 
271
#define LOOP13(VAR, STMT) LOOP12(VAR, STMT); (STMT); (VAR)++;
 
272
 
 
273
#define LOOP(N, VAR, STMT) CAT(LOOP, N)((VAR), (STMT))
 
274
 
 
275
#define DIG(a) a,
 
276
__constant WT1 kernelData[] = { COEFF };
 
277
 
 
278
__kernel void filter2DSmall(__global const uchar * srcptr, int src_step, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY,
 
279
                       __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols, float delta)
 
280
{
 
281
    const struct RectCoords srcCoords = { srcOffsetX, srcOffsetY, srcEndX, srcEndY }; // for non-isolated border: offsetX, offsetY, wholeX, wholeY
 
282
 
 
283
    const int startX = get_global_id(0) * PX_PER_WI_X;
 
284
    const int startY = get_global_id(1) * PX_PER_WI_Y;
 
285
 
 
286
    if ((startX >= cols) || (startY >= rows))
 
287
    {
 
288
        return;
 
289
    }
 
290
 
 
291
    WT privateData[PX_PER_WI_Y + KERNEL_SIZE_Y - 1][PRIV_DATA_WIDTH];
 
292
 
 
293
    // Load all of the pixels needed for the calculation
 
294
    int py = 0;
 
295
    LOOP(PX_LOAD_Y_ITERATIONS, py,
 
296
    {
 
297
        int y = startY + py;
 
298
        int px = 0;
 
299
        LOOP(PX_LOAD_X_ITERATIONS, px,
 
300
        {
 
301
            int x = startX + (px * PX_LOAD_NUM_PX);
 
302
            int2 srcPos = (int2)(srcCoords.x1 + x - ANCHOR_X, srcCoords.y1 + y - ANCHOR_Y);
 
303
 
 
304
            if (!isBorder(srcCoords, srcPos, PX_LOAD_NUM_PX))
 
305
            {
 
306
                PX_LOAD_FLOAT_VEC_TYPE p = readSrcPixelGroup(srcPos, srcptr, src_step, srcCoords);
 
307
                *((PX_LOAD_FLOAT_VEC_TYPE*)&privateData[py][px * PX_LOAD_NUM_PX]) = p;
 
308
            }
 
309
            else
 
310
            {
 
311
                int lx = 0;
 
312
                LOOP(PX_LOAD_NUM_PX, lx,
 
313
                {
 
314
                    WT p = readSrcPixelSingle(srcPos, srcptr, src_step, srcCoords);
 
315
                    *((WT*)&privateData[py][px * PX_LOAD_NUM_PX + lx]) = p;
 
316
                    srcPos.x++;
 
317
                });
 
318
            }
 
319
        });
 
320
    });
 
321
    // Use the stored pixels to compute the results
 
322
    py = 0;
 
323
    LOOP(PX_PER_WI_Y, py,
 
324
    {
 
325
        int y = startY + py;
 
326
        int px = 0;
 
327
        LOOP(PX_PER_WI_X, px,
 
328
        {
 
329
            int x = startX + px;
 
330
            WT total_sum = 0;
 
331
            int sy = 0;
 
332
            int kernelIndex = 0;
 
333
            LOOP(KERNEL_SIZE_Y, sy,
 
334
            {
 
335
                int sx = 0;
 
336
                LOOP(KERNEL_SIZE_X, sx,
 
337
                {
 
338
                    total_sum = mad(kernelData[kernelIndex++], privateData[py + sy][px + sx], total_sum);
 
339
                });
 
340
            });
 
341
 
 
342
            __global dstT* dstPtr = (__global dstT*)(dstptr + y * dst_step + dst_offset + x * DSTSIZE); // Pointer can be out of bounds!
 
343
            storepix(convertToDstT(total_sum + (WT)(delta)), dstPtr);
 
344
        });
 
345
    });
 
346
}