← 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/filterSep_singlePass.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/filterSep_singlePass.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) 2014, Intel Corporation, all rights reserved.
 
14
// Third party copyrights are property of their respective owners.
 
15
//
 
16
// Redistribution and use in source and binary forms, with or without modification,
 
17
// are permitted provided that the following conditions are met:
 
18
//
 
19
//   * Redistribution's of source code must retain the above copyright notice,
 
20
//     this list of conditions and the following disclaimer.
 
21
//
 
22
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
23
//     this list of conditions and the following disclaimer in the documentation
 
24
//     and/or other materials provided with the distribution.
 
25
//
 
26
//   * The name of the copyright holders may not be used to endorse or promote products
 
27
//     derived from this software without specific prior written permission.
 
28
//
 
29
// This software is provided by the copyright holders and contributors "as is" and
 
30
// any express or implied warranties, including, but not limited to, the implied
 
31
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
32
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
33
// indirect, incidental, special, exemplary, or consequential damages
 
34
// (including, but not limited to, procurement of substitute goods or services;
 
35
// loss of use, data, or profits; or business interruption) however caused
 
36
// and on any theory of liability, whether in contract, strict liability,
 
37
// or tort (including negligence or otherwise) arising in any way out of
 
38
// the use of this software, even if advised of the possibility of such damage.
 
39
//
 
40
//M*/
 
41
 
 
42
///////////////////////////////////////////////////////////////////////////////////////////////////
 
43
/////////////////////////////////Macro for border type////////////////////////////////////////////
 
44
/////////////////////////////////////////////////////////////////////////////////////////////////
 
45
 
 
46
#ifdef BORDER_CONSTANT
 
47
// CCCCCC|abcdefgh|CCCCCCC
 
48
#define EXTRAPOLATE(x, maxV)
 
49
#elif defined BORDER_REPLICATE
 
50
// aaaaaa|abcdefgh|hhhhhhh
 
51
#define EXTRAPOLATE(x, maxV) \
 
52
    { \
 
53
        (x) = clamp((x), 0, (maxV)-1); \
 
54
    }
 
55
#elif defined BORDER_WRAP
 
56
// cdefgh|abcdefgh|abcdefg
 
57
#define EXTRAPOLATE(x, maxV) \
 
58
    { \
 
59
        (x) = ( (x) + (maxV) ) % (maxV); \
 
60
    }
 
61
#elif defined BORDER_REFLECT
 
62
// fedcba|abcdefgh|hgfedcb
 
63
#define EXTRAPOLATE(x, maxV) \
 
64
    { \
 
65
        (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \
 
66
    }
 
67
#elif defined BORDER_REFLECT_101 || defined BORDER_REFLECT101
 
68
// gfedcb|abcdefgh|gfedcba
 
69
#define EXTRAPOLATE(x, maxV) \
 
70
    { \
 
71
        (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \
 
72
    }
 
73
#else
 
74
#error No extrapolation method
 
75
#endif
 
76
 
 
77
#if CN != 3
 
78
#define loadpix(addr) *(__global const srcT *)(addr)
 
79
#define storepix(val, addr)  *(__global dstT *)(addr) = val
 
80
#define SRCSIZE (int)sizeof(srcT)
 
81
#define DSTSIZE (int)sizeof(dstT)
 
82
#else
 
83
#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
 
84
#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
 
85
#define SRCSIZE (int)sizeof(srcT1)*3
 
86
#define DSTSIZE (int)sizeof(dstT1)*3
 
87
#endif
 
88
 
 
89
#define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x)))
 
90
 
 
91
#ifdef BORDER_CONSTANT
 
92
// CCCCCC|abcdefgh|CCCCCCC
 
93
#define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y))
 
94
#else
 
95
#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y))
 
96
#endif
 
97
 
 
98
#define noconvert
 
99
 
 
100
// horizontal and vertical filter kernels
 
101
// should be defined on host during compile time to avoid overhead
 
102
#define DIG(a) a,
 
103
__constant WT1 mat_kernelX[] = { KERNEL_MATRIX_X };
 
104
__constant WT1 mat_kernelY[] = { KERNEL_MATRIX_Y };
 
105
 
 
106
__kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width,
 
107
                         __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, float delta)
 
108
{
 
109
    // RADIUSX, RADIUSY are filter dimensions
 
110
    // BLK_X, BLK_Y are local wrogroup sizes
 
111
    // all these should be defined on host during compile time
 
112
    // first lsmem array for source pixels used in first pass,
 
113
    // second lsmemDy for storing first pass results
 
114
    __local WT lsmem[BLK_Y + 2 * RADIUSY][BLK_X + 2 * RADIUSX];
 
115
    __local WT lsmemDy[BLK_Y][BLK_X + 2 * RADIUSX];
 
116
 
 
117
    // get local and global ids - used as image and local memory array indexes
 
118
    int lix = get_local_id(0);
 
119
    int liy = get_local_id(1);
 
120
 
 
121
    int x = get_global_id(0);
 
122
 
 
123
    // calculate pixel position in source image taking image offset into account
 
124
    int srcX = x + srcOffsetX - RADIUSX;
 
125
 
 
126
    // extrapolate coordinates, if needed
 
127
    // and read my own source pixel into local memory
 
128
    // with account for extra border pixels, which will be read by starting workitems
 
129
    int clocY = liy;
 
130
    do
 
131
    {
 
132
        int yb = clocY + srcOffsetY - RADIUSY;
 
133
        EXTRAPOLATE(yb, (height));
 
134
 
 
135
        int clocX = lix;
 
136
        int cSrcX = srcX;
 
137
        do
 
138
        {
 
139
            int xb = cSrcX;
 
140
            EXTRAPOLATE(xb,(width));
 
141
            lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 );
 
142
 
 
143
            clocX += BLK_X;
 
144
            cSrcX += BLK_X;
 
145
        }
 
146
        while(clocX < BLK_X+(RADIUSX*2));
 
147
 
 
148
        clocY += BLK_Y;
 
149
    }
 
150
    while (clocY < BLK_Y+(RADIUSY*2));
 
151
    barrier(CLK_LOCAL_MEM_FENCE);
 
152
 
 
153
    for (int y = 0; y < dst_rows; y+=BLK_Y)
 
154
    {
 
155
        // do vertical filter pass
 
156
        // and store intermediate results to second local memory array
 
157
        int i, clocX = lix;
 
158
        WT sum = (WT) 0;
 
159
        do
 
160
        {
 
161
            sum = (WT) 0;
 
162
            for (i=0; i<=2*RADIUSY; i++)
 
163
#if (defined(INTEGER_ARITHMETIC) && !INTEL_DEVICE)
 
164
                sum = mad24(lsmem[liy + i][clocX], mat_kernelY[i], sum);
 
165
#else
 
166
                sum = mad(lsmem[liy + i][clocX], mat_kernelY[i], sum);
 
167
#endif
 
168
            lsmemDy[liy][clocX] = sum;
 
169
            clocX += BLK_X;
 
170
        }
 
171
        while(clocX < BLK_X+(RADIUSX*2));
 
172
        barrier(CLK_LOCAL_MEM_FENCE);
 
173
 
 
174
        // if this pixel happened to be out of image borders because of global size rounding,
 
175
        // then just return
 
176
        if ((x < dst_cols) && (y + liy < dst_rows))
 
177
        {
 
178
            // do second horizontal filter pass
 
179
            // and calculate final result
 
180
            sum = 0.0f;
 
181
            for (i=0; i<=2*RADIUSX; i++)
 
182
#if (defined(INTEGER_ARITHMETIC) && !INTEL_DEVICE)
 
183
                sum = mad24(lsmemDy[liy][lix+i], mat_kernelX[i], sum);
 
184
#else
 
185
                sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum);
 
186
#endif
 
187
 
 
188
#ifdef INTEGER_ARITHMETIC
 
189
#ifdef INTEL_DEVICE
 
190
            sum = (sum + (1 << (SHIFT_BITS-1))) / (1 << SHIFT_BITS);
 
191
#else
 
192
            sum = (sum + (1 << (SHIFT_BITS-1))) >> SHIFT_BITS;
 
193
#endif
 
194
#endif
 
195
            // store result into destination image
 
196
            storepix(convertToDstT(sum + (WT)(delta)), Dst + mad24(y + liy, dst_step, mad24(x, DSTSIZE, dst_offset)));
 
197
        }
 
198
 
 
199
        for (int i = liy * BLK_X + lix; i < (RADIUSY*2) * (BLK_X+(RADIUSX*2)); i += BLK_X * BLK_Y)
 
200
        {
 
201
            int clocX = i % (BLK_X+(RADIUSX*2));
 
202
            int clocY = i / (BLK_X+(RADIUSX*2));
 
203
            lsmem[clocY][clocX] = lsmem[clocY + BLK_Y][clocX];
 
204
        }
 
205
        barrier(CLK_LOCAL_MEM_FENCE);
 
206
 
 
207
        int yb = y + liy + BLK_Y + srcOffsetY + RADIUSY;
 
208
        EXTRAPOLATE(yb, (height));
 
209
 
 
210
        clocX = lix;
 
211
        int cSrcX = x + srcOffsetX - RADIUSX;
 
212
        do
 
213
        {
 
214
            int xb = cSrcX;
 
215
            EXTRAPOLATE(xb,(width));
 
216
            lsmem[liy + 2*RADIUSY][clocX] = ELEM(xb, yb, (width), (height), 0 );
 
217
 
 
218
            clocX += BLK_X;
 
219
            cSrcX += BLK_X;
 
220
        }
 
221
        while(clocX < BLK_X+(RADIUSX*2));
 
222
        barrier(CLK_LOCAL_MEM_FENCE);
 
223
    }
 
224
 
 
225
}