← 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/moments.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/moments.cl
 
1
/* See LICENSE file in the root OpenCV directory */
 
2
 
 
3
#if TILE_SIZE != 32
 
4
#error "TILE SIZE should be 32"
 
5
#endif
 
6
 
 
7
 
 
8
__kernel void moments(__global const uchar* src, int src_step, int src_offset,
 
9
    int src_rows, int src_cols, __global int* mom0, int xtiles)
 
10
{
 
11
    int x0 = get_global_id(0);
 
12
    int y0 = get_group_id(1);
 
13
    int x, y = get_local_id(1);
 
14
    int x_min = x0*TILE_SIZE;
 
15
    int ypix = y0*TILE_SIZE + y;
 
16
    __local int mom[TILE_SIZE][10];
 
17
 
 
18
    if (x_min < src_cols && y0*TILE_SIZE < src_rows)
 
19
    {
 
20
        if (ypix < src_rows)
 
21
        {
 
22
            int x_max = min(src_cols - x_min, TILE_SIZE);
 
23
            __global const uchar* ptr = src + src_offset + ypix*src_step + x_min;
 
24
            int4 S = (int4)(0, 0, 0, 0), p;
 
25
 
 
26
#define SUM_ELEM(elem, ofs) \
 
27
    (int4)(1, (ofs), (ofs)*(ofs), (ofs)*(ofs)*(ofs))*elem
 
28
 
 
29
            x = x_max & -4;
 
30
            if (x_max >= 4)
 
31
            {
 
32
                p = convert_int4(vload4(0, ptr));
 
33
#ifdef OP_MOMENTS_BINARY
 
34
                p = min(p, 1);
 
35
#endif
 
36
                S += (int4)(p.s0, 0, 0, 0) + (int4)(p.s1, p.s1, p.s1, p.s1) +
 
37
                    (int4)(p.s2, p.s2 * 2, p.s2 * 4, p.s2 * 8) + (int4)(p.s3, p.s3 * 3, p.s3 * 9, p.s3 * 27);
 
38
                //SUM_ELEM(p.s0, 0) + SUM_ELEM(p.s1, 1) + SUM_ELEM(p.s2, 2) + SUM_ELEM(p.s3, 3);
 
39
 
 
40
                if (x_max >= 8)
 
41
                {
 
42
                    p = convert_int4(vload4(0, ptr + 4));
 
43
#ifdef OP_MOMENTS_BINARY
 
44
                    p = min(p, 1);
 
45
#endif
 
46
                    S += (int4)(p.s0, p.s0 * 4, p.s0 * 16, p.s0 * 64) + (int4)(p.s1, p.s1 * 5, p.s1 * 25, p.s1 * 125) +
 
47
                        (int4)(p.s2, p.s2 * 6, p.s2 * 36, p.s2 * 216) + (int4)(p.s3, p.s3 * 7, p.s3 * 49, p.s3 * 343);
 
48
                    //SUM_ELEM(p.s0, 4) + SUM_ELEM(p.s1, 5) + SUM_ELEM(p.s2, 6) + SUM_ELEM(p.s3, 7);
 
49
 
 
50
                    if (x_max >= 12)
 
51
                    {
 
52
                        p = convert_int4(vload4(0, ptr + 8));
 
53
#ifdef OP_MOMENTS_BINARY
 
54
                        p = min(p, 1);
 
55
#endif
 
56
                        S += (int4)(p.s0, p.s0 * 8, p.s0 * 64, p.s0 * 512) + (int4)(p.s1, p.s1 * 9, p.s1 * 81, p.s1 * 729) +
 
57
                            (int4)(p.s2, p.s2 * 10, p.s2 * 100, p.s2 * 1000) + (int4)(p.s3, p.s3 * 11, p.s3 * 121, p.s3 * 1331);
 
58
                        //SUM_ELEM(p.s0, 8) + SUM_ELEM(p.s1, 9) + SUM_ELEM(p.s2, 10) + SUM_ELEM(p.s3, 11);
 
59
 
 
60
                        if (x_max >= 16)
 
61
                        {
 
62
                            p = convert_int4(vload4(0, ptr + 12));
 
63
#ifdef OP_MOMENTS_BINARY
 
64
                            p = min(p, 1);
 
65
#endif
 
66
                            S += (int4)(p.s0, p.s0 * 12, p.s0 * 144, p.s0 * 1728) + (int4)(p.s1, p.s1 * 13, p.s1 * 169, p.s1 * 2197) +
 
67
                                (int4)(p.s2, p.s2 * 14, p.s2 * 196, p.s2 * 2744) + (int4)(p.s3, p.s3 * 15, p.s3 * 225, p.s3 * 3375);
 
68
                            //SUM_ELEM(p.s0, 12) + SUM_ELEM(p.s1, 13) + SUM_ELEM(p.s2, 14) + SUM_ELEM(p.s3, 15);
 
69
                        }
 
70
                    }
 
71
                }
 
72
            }
 
73
 
 
74
            if (x_max >= 20)
 
75
            {
 
76
                p = convert_int4(vload4(0, ptr + 16));
 
77
#ifdef OP_MOMENTS_BINARY
 
78
                p = min(p, 1);
 
79
#endif
 
80
                S += (int4)(p.s0, p.s0 * 16, p.s0 * 256, p.s0 * 4096) + (int4)(p.s1, p.s1 * 17, p.s1 * 289, p.s1 * 4913) +
 
81
                    (int4)(p.s2, p.s2 * 18, p.s2 * 324, p.s2 * 5832) + (int4)(p.s3, p.s3 * 19, p.s3 * 361, p.s3 * 6859);
 
82
                //SUM_ELEM(p.s0, 16) + SUM_ELEM(p.s1, 17) + SUM_ELEM(p.s2, 18) + SUM_ELEM(p.s3, 19);
 
83
 
 
84
                if (x_max >= 24)
 
85
                {
 
86
                    p = convert_int4(vload4(0, ptr + 20));
 
87
#ifdef OP_MOMENTS_BINARY
 
88
                    p = min(p, 1);
 
89
#endif
 
90
                    S += (int4)(p.s0, p.s0 * 20, p.s0 * 400, p.s0 * 8000) + (int4)(p.s1, p.s1 * 21, p.s1 * 441, p.s1 * 9261) +
 
91
                        (int4)(p.s2, p.s2 * 22, p.s2 * 484, p.s2 * 10648) + (int4)(p.s3, p.s3 * 23, p.s3 * 529, p.s3 * 12167);
 
92
                    //SUM_ELEM(p.s0, 20) + SUM_ELEM(p.s1, 21) + SUM_ELEM(p.s2, 22) + SUM_ELEM(p.s3, 23);
 
93
 
 
94
                    if (x_max >= 28)
 
95
                    {
 
96
                        p = convert_int4(vload4(0, ptr + 24));
 
97
#ifdef OP_MOMENTS_BINARY
 
98
                        p = min(p, 1);
 
99
#endif
 
100
                        S += (int4)(p.s0, p.s0 * 24, p.s0 * 576, p.s0 * 13824) + (int4)(p.s1, p.s1 * 25, p.s1 * 625, p.s1 * 15625) +
 
101
                            (int4)(p.s2, p.s2 * 26, p.s2 * 676, p.s2 * 17576) + (int4)(p.s3, p.s3 * 27, p.s3 * 729, p.s3 * 19683);
 
102
                        //SUM_ELEM(p.s0, 24) + SUM_ELEM(p.s1, 25) + SUM_ELEM(p.s2, 26) + SUM_ELEM(p.s3, 27);
 
103
 
 
104
                        if (x_max >= 32)
 
105
                        {
 
106
                            p = convert_int4(vload4(0, ptr + 28));
 
107
#ifdef OP_MOMENTS_BINARY
 
108
                            p = min(p, 1);
 
109
#endif
 
110
                            S += (int4)(p.s0, p.s0 * 28, p.s0 * 784, p.s0 * 21952) + (int4)(p.s1, p.s1 * 29, p.s1 * 841, p.s1 * 24389) +
 
111
                                (int4)(p.s2, p.s2 * 30, p.s2 * 900, p.s2 * 27000) + (int4)(p.s3, p.s3 * 31, p.s3 * 961, p.s3 * 29791);
 
112
                            //SUM_ELEM(p.s0, 28) + SUM_ELEM(p.s1, 29) + SUM_ELEM(p.s2, 30) + SUM_ELEM(p.s3, 31);
 
113
                        }
 
114
                    }
 
115
                }
 
116
            }
 
117
 
 
118
            if (x < x_max)
 
119
            {
 
120
                int ps = ptr[x];
 
121
#ifdef OP_MOMENTS_BINARY
 
122
                ps = min(ps, 1);
 
123
#endif
 
124
                S += SUM_ELEM(ps, x);
 
125
                if (x + 1 < x_max)
 
126
                {
 
127
                    ps = ptr[x + 1];
 
128
#ifdef OP_MOMENTS_BINARY
 
129
                    ps = min(ps, 1);
 
130
#endif
 
131
                    S += SUM_ELEM(ps, x + 1);
 
132
                    if (x + 2 < x_max)
 
133
                    {
 
134
                        ps = ptr[x + 2];
 
135
#ifdef OP_MOMENTS_BINARY
 
136
                        ps = min(ps, 1);
 
137
#endif
 
138
                        S += SUM_ELEM(ps, x + 2);
 
139
                    }
 
140
                }
 
141
            }
 
142
 
 
143
            int sy = y*y;
 
144
 
 
145
            mom[y][0] = S.s0;
 
146
            mom[y][1] = S.s1;
 
147
            mom[y][2] = y*S.s0;
 
148
            mom[y][3] = S.s2;
 
149
            mom[y][4] = y*S.s1;
 
150
            mom[y][5] = sy*S.s0;
 
151
            mom[y][6] = S.s3;
 
152
            mom[y][7] = y*S.s2;
 
153
            mom[y][8] = sy*S.s1;
 
154
            mom[y][9] = y*sy*S.s0;
 
155
        }
 
156
        else
 
157
            mom[y][0] = mom[y][1] = mom[y][2] = mom[y][3] = mom[y][4] =
 
158
            mom[y][5] = mom[y][6] = mom[y][7] = mom[y][8] = mom[y][9] = 0;
 
159
        barrier(CLK_LOCAL_MEM_FENCE);
 
160
 
 
161
#define REDUCE(d) \
 
162
        if (y < d) \
 
163
        { \
 
164
        mom[y][0] += mom[y + d][0]; \
 
165
        mom[y][1] += mom[y + d][1]; \
 
166
        mom[y][2] += mom[y + d][2]; \
 
167
        mom[y][3] += mom[y + d][3]; \
 
168
        mom[y][4] += mom[y + d][4]; \
 
169
        mom[y][5] += mom[y + d][5]; \
 
170
        mom[y][6] += mom[y + d][6]; \
 
171
        mom[y][7] += mom[y + d][7]; \
 
172
        mom[y][8] += mom[y + d][8]; \
 
173
        mom[y][9] += mom[y + d][9]; \
 
174
        } \
 
175
        barrier(CLK_LOCAL_MEM_FENCE)
 
176
 
 
177
        REDUCE(16);
 
178
        REDUCE(8);
 
179
        REDUCE(4);
 
180
        REDUCE(2);
 
181
 
 
182
        if (y < 10)
 
183
        {
 
184
            __global int* momout = mom0 + (y0*xtiles + x0) * 10;
 
185
            momout[y] = mom[0][y] + mom[1][y];
 
186
        }
 
187
    }
 
188
}