opencl/edge-laplace.cl.h


1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184

static const char* edge_laplace_cl_source =
"#define LAPLACE_RADIUS 2                                                      \n"
"#define EPSILON        1e-5f                                                  \n"
"                                                                              \n"
"void minmax(float x1, float x2, float x3,                                     \n"
"            float x4, float x5,                                               \n"
"            float *min_result,                                                \n"
"            float *max_result)                                                \n"
"{                                                                             \n"
"    float min1, min2, max1, max2;                                             \n"
"                                                                              \n"
"    if (x1 > x2)                                                              \n"
"    {                                                                         \n"
"        max1 = x1;                                                            \n"
"        min1 = x2;                                                            \n"
"    }                                                                         \n"
"    else                                                                      \n"
"    {                                                                         \n"
"        max1 = x2;                                                            \n"
"        min1 = x1;                                                            \n"
"    }                                                                         \n"
"                                                                              \n"
"    if (x3 > x4)                                                              \n"
"    {                                                                         \n"
"        max2 = x3;                                                            \n"
"        min2 = x4;                                                            \n"
"    }                                                                         \n"
"    else                                                                      \n"
"    {                                                                         \n"
"        max2 = x4;                                                            \n"
"        min2 = x3;                                                            \n"
"    }                                                                         \n"
"                                                                              \n"
"    if (min1 < min2)                                                          \n"
"        *min_result = fmin(min1, x5);                                         \n"
"    else                                                                      \n"
"        *min_result = fmin(min2, x5);                                         \n"
"    if (max1 > max2)                                                          \n"
"        *max_result = fmax(max1, x5);                                         \n"
"    else                                                                      \n"
"        *max_result = fmax(max2, x5);                                         \n"
"}                                                                             \n"
"                                                                              \n"
"kernel void pre_edgelaplace (global float4 *in,                               \n"
"                             global float4 *out)                              \n"
"{                                                                             \n"
"    int gidx = get_global_id(0);                                              \n"
"    int gidy = get_global_id(1);                                              \n"
"                                                                              \n"
"    int src_width  = get_global_size(0) + LAPLACE_RADIUS * 2;                 \n"
"    int src_height = get_global_size(1);                                      \n"
"                                                                              \n"
"    int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;                 \n"
"    int gid1d = i + j * src_width;                                            \n"
"                                                                              \n"
"    float pix_fl[4] = {                                                       \n"
"        in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y,             \n"
"        in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w              \n"
"    };                                                                        \n"
"    float pix_fm[4] = {                                                       \n"
"        in[gid1d     - src_width].x, in[gid1d     - src_width].y,             \n"
"        in[gid1d     - src_width].z, in[gid1d     - src_width].w              \n"
"    };                                                                        \n"
"    float pix_fr[4] = {                                                       \n"
"        in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y,             \n"
"        in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w              \n"
"    };                                                                        \n"
"    float pix_ml[4] = {                                                       \n"
"        in[gid1d - 1            ].x, in[gid1d - 1            ].y,             \n"
"        in[gid1d - 1            ].z, in[gid1d - 1            ].w              \n"
"    };                                                                        \n"
"    float pix_mm[4] = {                                                       \n"
"        in[gid1d                ].x, in[gid1d                ].y,             \n"
"        in[gid1d                ].z, in[gid1d                ].w              \n"
"    };                                                                        \n"
"    float pix_mr[4] = {                                                       \n"
"        in[gid1d + 1            ].x, in[gid1d + 1            ].y,             \n"
"        in[gid1d + 1            ].z, in[gid1d + 1            ].w              \n"
"    };                                                                        \n"
"    float pix_bl[4] = {                                                       \n"
"        in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y,             \n"
"        in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w              \n"
"    };                                                                        \n"
"    float pix_bm[4] = {                                                       \n"
"        in[gid1d     + src_width].x, in[gid1d     + src_width].y,             \n"
"        in[gid1d     + src_width].z, in[gid1d     + src_width].w              \n"
"    };                                                                        \n"
"    float pix_br[4] = {                                                       \n"
"        in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y,             \n"
"        in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w              \n"
"    };                                                                        \n"
"                                                                              \n"
"    int c;                                                                    \n"
"    float minval, maxval;                                                     \n"
"    float gradient[4];                                                        \n"
"                                                                              \n"
"    for (c = 0;c < 3; ++c)                                                    \n"
"    {                                                                         \n"
"        minmax(pix_fm[c], pix_bm[c], pix_ml[c], pix_mr[c],                    \n"
"            pix_mm[c], &minval, &maxval);                                     \n"
"        gradient[c] = 0.5f *                                                  \n"
"            fmax((maxval - pix_mm[c]),(pix_mm[c] - minval));                  \n"
"        gradient[c] =                                                         \n"
"            (pix_fl[c] + pix_fm[c] + pix_fr[c] +                              \n"
"             pix_bm[c] - 8.0f * pix_mm[c] + pix_br[c] +                       \n"
"             pix_ml[c] + pix_mr[c] + pix_bl[c]) <                             \n"
"             EPSILON ? -1.0f * gradient[c] : gradient[c];                     \n"
"    }                                                                         \n"
"    gradient[3] = pix_mm[3];                                                  \n"
"                                                                              \n"
"    out[gid1d] = (float4)                                                     \n"
"        (gradient[0], gradient[1], gradient[2], gradient[3]);                 \n"
"}                                                                             \n"
"                                                                              \n"
"kernel void knl_edgelaplace (global float4 *in,                               \n"
"                             global float4 *out)                              \n"
"{                                                                             \n"
"    int gidx = get_global_id(0);                                              \n"
"    int gidy = get_global_id(1);                                              \n"
"                                                                              \n"
"    int src_width  = get_global_size(0) + LAPLACE_RADIUS * 2;                 \n"
"    int src_height = get_global_size(1);                                      \n"
"                                                                              \n"
"    int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;                 \n"
"    int gid1d = i + j * src_width;                                            \n"
"                                                                              \n"
"    float pix_fl[4] = {                                                       \n"
"        in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y,             \n"
"        in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w              \n"
"    };                                                                        \n"
"    float pix_fm[4] = {                                                       \n"
"        in[gid1d     - src_width].x, in[gid1d     - src_width].y,             \n"
"        in[gid1d     - src_width].z, in[gid1d     - src_width].w              \n"
"    };                                                                        \n"
"    float pix_fr[4] = {                                                       \n"
"        in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y,             \n"
"        in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w              \n"
"    };                                                                        \n"
"    float pix_ml[4] = {                                                       \n"
"        in[gid1d - 1            ].x, in[gid1d - 1            ].y,             \n"
"        in[gid1d - 1            ].z, in[gid1d - 1            ].w              \n"
"    };                                                                        \n"
"    float pix_mm[4] = {                                                       \n"
"        in[gid1d                ].x, in[gid1d                ].y,             \n"
"        in[gid1d                ].z, in[gid1d                ].w              \n"
"    };                                                                        \n"
"    float pix_mr[4] = {                                                       \n"
"        in[gid1d + 1            ].x, in[gid1d + 1            ].y,             \n"
"        in[gid1d + 1            ].z, in[gid1d + 1            ].w              \n"
"    };                                                                        \n"
"    float pix_bl[4] = {                                                       \n"
"        in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y,             \n"
"        in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w              \n"
"    };                                                                        \n"
"    float pix_bm[4] = {                                                       \n"
"        in[gid1d     + src_width].x, in[gid1d     + src_width].y,             \n"
"        in[gid1d     + src_width].z, in[gid1d     + src_width].w              \n"
"    };                                                                        \n"
"    float pix_br[4] = {                                                       \n"
"        in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y,             \n"
"        in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w              \n"
"    };                                                                        \n"
"                                                                              \n"
"    int c;                                                                    \n"
"    float value[4];                                                           \n"
"                                                                              \n"
"    for (c = 0;c < 3; ++c)                                                    \n"
"    {                                                                         \n"
"        float current = pix_mm[c];                                            \n"
"        current =                                                             \n"
"            ((current > 0.0f) &&                                              \n"
"             (pix_fl[c] < 0.0f || pix_fm[c] < 0.0f ||                         \n"
"              pix_fr[c] < 0.0f || pix_ml[c] < 0.0f ||                         \n"
"              pix_mr[c] < 0.0f || pix_bl[c] < 0.0f ||                         \n"
"              pix_bm[c] < 0.0f || pix_br[c] < 0.0f )                          \n"
"            ) ? current : 0.0f;                                               \n"
"        value[c] = current;                                                   \n"
"    }                                                                         \n"
"    value[3] = pix_mm[3];                                                     \n"
"                                                                              \n"
"    out[gidx + gidy * get_global_size(0)] = (float4)                          \n"
"        (value[0], value[1], value[2], value[3]);                             \n"
"}                                                                             \n"
;