← 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/laplacian5.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/laplacian5.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
 
 
9
#define noconvert
 
10
 
 
11
#ifdef ONLY_SUM_CONVERT
 
12
 
 
13
__kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1_offset,
 
14
                         __global const uchar * src2ptr, int src2_step, int src2_offset,
 
15
                         __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
 
16
                         coeffT scale, coeffT delta)
 
17
{
 
18
    int x = get_global_id(0);
 
19
    int y = get_global_id(1);
 
20
 
 
21
    if (y < dst_rows && x < dst_cols)
 
22
    {
 
23
        int src1_index = mad24(y, src1_step, mad24(x, (int)sizeof(srcT), src1_offset));
 
24
        int src2_index = mad24(y, src2_step, mad24(x, (int)sizeof(srcT), src2_offset));
 
25
        int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset));
 
26
 
 
27
        __global const srcT * src1 = (__global const srcT *)(src1ptr + src1_index);
 
28
        __global const srcT * src2 = (__global const srcT *)(src2ptr + src2_index);
 
29
        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
 
30
 
 
31
#if wdepth <= 4
 
32
        dst[0] = convertToDT( mad24((WT)(scale), convertToWT(src1[0]) + convertToWT(src2[0]), (WT)(delta)) );
 
33
#else
 
34
        dst[0] = convertToDT( mad((WT)(scale), convertToWT(src1[0]) + convertToWT(src2[0]), (WT)(delta)) );
 
35
#endif
 
36
    }
 
37
}
 
38
 
 
39
#else
 
40
 
 
41
///////////////////////////////////////////////////////////////////////////////////////////////////
 
42
/////////////////////////////////Macro for border type////////////////////////////////////////////
 
43
/////////////////////////////////////////////////////////////////////////////////////////////////
 
44
 
 
45
#ifdef BORDER_CONSTANT
 
46
// CCCCCC|abcdefgh|CCCCCCC
 
47
#define EXTRAPOLATE(x, maxV)
 
48
#elif defined BORDER_REPLICATE
 
49
// aaaaaa|abcdefgh|hhhhhhh
 
50
#define EXTRAPOLATE(x, maxV) \
 
51
    { \
 
52
        (x) = clamp((x), 0, (maxV)-1); \
 
53
    }
 
54
#elif defined BORDER_WRAP
 
55
// cdefgh|abcdefgh|abcdefg
 
56
#define EXTRAPOLATE(x, maxV) \
 
57
    { \
 
58
        (x) = ( (x) + (maxV) ) % (maxV); \
 
59
    }
 
60
#elif defined BORDER_REFLECT
 
61
// fedcba|abcdefgh|hgfedcb
 
62
#define EXTRAPOLATE(x, maxV) \
 
63
    { \
 
64
        (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \
 
65
    }
 
66
#elif defined BORDER_REFLECT_101
 
67
// gfedcb|abcdefgh|gfedcba
 
68
#define EXTRAPOLATE(x, maxV) \
 
69
    { \
 
70
        (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \
 
71
    }
 
72
#else
 
73
#error No extrapolation method
 
74
#endif
 
75
 
 
76
#if CN != 3
 
77
#define loadpix(addr) *(__global const srcT *)(addr)
 
78
#define storepix(val, addr)  *(__global dstT *)(addr) = val
 
79
#define SRCSIZE (int)sizeof(srcT)
 
80
#define DSTSIZE (int)sizeof(dstT)
 
81
#else
 
82
#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
 
83
#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
 
84
#define SRCSIZE (int)sizeof(srcT1)*3
 
85
#define DSTSIZE (int)sizeof(dstT1)*3
 
86
#endif
 
87
 
 
88
#define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x)))
 
89
 
 
90
#ifdef BORDER_CONSTANT
 
91
// CCCCCC|abcdefgh|CCCCCCC
 
92
#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))
 
93
#else
 
94
#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y))
 
95
#endif
 
96
 
 
97
// horizontal and vertical filter kernels
 
98
// should be defined on host during compile time to avoid overhead
 
99
#define DIG(a) a,
 
100
__constant WT1 mat_kernelX[] = { KERNEL_MATRIX_X };
 
101
__constant WT1 mat_kernelY[] = { KERNEL_MATRIX_Y };
 
102
 
 
103
__kernel void laplacian(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width,
 
104
                         __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
 
105
                         WT1 scale, WT1 delta)
 
106
{
 
107
    __local WT lsmem[BLK_Y + 2 * RADIUS][BLK_X + 2 * RADIUS];
 
108
    __local WT lsmemDy1[BLK_Y][BLK_X + 2 * RADIUS];
 
109
    __local WT lsmemDy2[BLK_Y][BLK_X + 2 * RADIUS];
 
110
 
 
111
    int lix = get_local_id(0);
 
112
    int liy = get_local_id(1);
 
113
 
 
114
    int x = get_global_id(0);
 
115
 
 
116
    int srcX = x + srcOffsetX - RADIUS;
 
117
 
 
118
    int clocY = liy;
 
119
    do
 
120
    {
 
121
        int yb = clocY + srcOffsetY - RADIUS;
 
122
        EXTRAPOLATE(yb, (height));
 
123
 
 
124
        int clocX = lix;
 
125
        int cSrcX = srcX;
 
126
        do
 
127
        {
 
128
            int xb = cSrcX;
 
129
            EXTRAPOLATE(xb,(width));
 
130
            lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 );
 
131
 
 
132
            clocX += BLK_X;
 
133
            cSrcX += BLK_X;
 
134
        }
 
135
        while(clocX < BLK_X+(RADIUS*2));
 
136
 
 
137
        clocY += BLK_Y;
 
138
    }
 
139
    while (clocY < BLK_Y+(RADIUS*2));
 
140
    barrier(CLK_LOCAL_MEM_FENCE);
 
141
 
 
142
    WT scale_v = (WT)scale;
 
143
    WT delta_v = (WT)delta;
 
144
    for (int y = 0; y < dst_rows; y+=BLK_Y)
 
145
    {
 
146
        int i, clocX = lix;
 
147
        WT sum1 = (WT) 0;
 
148
        WT sum2 = (WT) 0;
 
149
        do
 
150
        {
 
151
            sum1 = (WT) 0;
 
152
            sum2 = (WT) 0;
 
153
            for (i=0; i<=2*RADIUS; i++)
 
154
            {
 
155
                sum1 = mad(lsmem[liy + i][clocX], mat_kernelY[i], sum1);
 
156
                sum2 = mad(lsmem[liy + i][clocX], mat_kernelX[i], sum2);
 
157
            }
 
158
            lsmemDy1[liy][clocX] = sum1;
 
159
            lsmemDy2[liy][clocX] = sum2;
 
160
            clocX += BLK_X;
 
161
        }
 
162
        while(clocX < BLK_X+(RADIUS*2));
 
163
        barrier(CLK_LOCAL_MEM_FENCE);
 
164
 
 
165
        if ((x < dst_cols) && (y + liy < dst_rows))
 
166
        {
 
167
            sum1 = (WT) 0;
 
168
            sum2 = (WT) 0;
 
169
            for (i=0; i<=2*RADIUS; i++)
 
170
            {
 
171
                sum1 = mad(lsmemDy1[liy][lix+i], mat_kernelX[i], sum1);
 
172
                sum2 = mad(lsmemDy2[liy][lix+i], mat_kernelY[i], sum2);
 
173
            }
 
174
 
 
175
            WT sum = mad(scale_v, (sum1 + sum2), delta_v);
 
176
            storepix(convertToDT(sum), Dst + mad24(y + liy, dst_step, mad24(x, DSTSIZE, dst_offset)));
 
177
        }
 
178
 
 
179
        for (int i = liy * BLK_X + lix; i < (RADIUS*2) * (BLK_X+(RADIUS*2)); i += BLK_X * BLK_Y)
 
180
        {
 
181
            int clocX = i % (BLK_X+(RADIUS*2));
 
182
            int clocY = i / (BLK_X+(RADIUS*2));
 
183
            lsmem[clocY][clocX] = lsmem[clocY + BLK_Y][clocX];
 
184
        }
 
185
        barrier(CLK_LOCAL_MEM_FENCE);
 
186
 
 
187
        int yb = y + liy + BLK_Y + srcOffsetY + RADIUS;
 
188
        EXTRAPOLATE(yb, (height));
 
189
 
 
190
        clocX = lix;
 
191
        int cSrcX = x + srcOffsetX - RADIUS;
 
192
        do
 
193
        {
 
194
            int xb = cSrcX;
 
195
            EXTRAPOLATE(xb,(width));
 
196
            lsmem[liy + 2*RADIUS][clocX] = ELEM(xb, yb, (width), (height), 0 );
 
197
 
 
198
            clocX += BLK_X;
 
199
            cSrcX += BLK_X;
 
200
        }
 
201
        while(clocX < BLK_X+(RADIUS*2));
 
202
        barrier(CLK_LOCAL_MEM_FENCE);
 
203
    }
 
204
}
 
205
 
 
206
#endif
 
 
b'\\ No newline at end of file'