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
| /*
| * function: kernel_bi_filter
| * bilateral filter
| * input_y: Y channel image2d_t as read only
| * input_dark: dark channel image2d_t as read only
| * output_dark: dark channel image2d_t as write only
| *
| * data_type CL_UNSIGNED_INT16
| * channel_order CL_RGBA
| */
|
| #define PATCH_RADIUS 7
| #define PATCH_DIAMETER (2 * PATCH_RADIUS + 1)
|
| #define CALC_SUM(y1,y2,y3,dark1,dark2,dark3) \
| cur_y = (float8)(y1, y2, y3); \
| cur_dark = (float8)(dark1, dark2, dark3); \
| calc_sum (cur_y, cur_dark, center_y, &weight_sum, &data_sum);
|
| __inline void calc_sum (float8 cur_y, float8 cur_dark, float8 center_y, float8 *weight_sum, float8 *data_sum)
| {
| float8 delta = (cur_y - center_y) / 28.0f;
| delta = -0.5f * delta * delta;
|
| float8 weight = native_exp(delta);
| float8 data = cur_dark * weight;
| (*weight_sum) += weight;
| (*data_sum) += data;
| }
|
| __kernel void kernel_bi_filter (
| __read_only image2d_t input_y,
| __read_only image2d_t input_dark,
| __write_only image2d_t output_dark)
| {
| int pos_x = get_global_id (0);
| int pos_y = get_global_id (1);
| sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
|
| float8 y1, y2, dark1, dark2;
| float8 cur_y, cur_dark;
|
| float8 weight_sum = 0.0f;
| float8 data_sum = 0.0f;
| float8 center_y = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x, pos_y)))));
| for (int i = 0; i < PATCH_DIAMETER; i++) {
| y1 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x - 1, pos_y - PATCH_RADIUS + i)))));
| y2 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x, pos_y - PATCH_RADIUS + i)))));
| dark1 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x - 1, pos_y - PATCH_RADIUS + i)))));
| dark2 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x, pos_y - PATCH_RADIUS + i)))));
| CALC_SUM (y1.s1234, y1.s567, y2.s0, dark1.s1234, dark1.s567, dark2.s0);
| CALC_SUM (y1.s2345, y1.s67, y2.s01, dark1.s2345, dark1.s67, dark2.s01);
| CALC_SUM (y1.s3456, y1.s7, y2.s012, dark1.s3456, dark1.s7, dark2.s012);
| CALC_SUM (y1.s4567, y2.s01, y2.s23, dark1.s4567, dark2.s01, dark2.s23);
| CALC_SUM (y1.s567, y2.s0123, y2.s4, dark1.s567, dark2.s0123, dark2.s4);
| CALC_SUM (y1.s67, y2.s0123, y2.s45, dark1.s67, dark2.s0123, dark2.s45);
| CALC_SUM (y1.s7, y2.s0123, y2.s456, dark1.s7, dark2.s0123, dark2.s456);
| CALC_SUM (y2.s0123, y2.s45, y2.s67, dark2.s0123, dark2.s45, dark2.s67);
|
| y1 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_y, sampler, (int2)(pos_x + 1, pos_y - PATCH_RADIUS + i)))));
| dark1 = convert_float8(as_uchar8(convert_ushort4(read_imageui(input_dark, sampler, (int2)(pos_x + 1, pos_y - PATCH_RADIUS + i)))));
| CALC_SUM (y2.s1234, y2.s567, y1.s0, dark2.s1234, dark2.s567, dark1.s0);
| CALC_SUM (y2.s2345, y2.s67, y1.s01, dark2.s2345, dark2.s67, dark1.s01);
| CALC_SUM (y2.s3456, y2.s7, y1.s012, dark2.s3456, dark2.s7, dark1.s012);
| CALC_SUM (y2.s4567, y1.s01, y1.s23, dark2.s4567, dark1.s01, dark1.s23);
| CALC_SUM (y2.s567, y1.s0123, y1.s4, dark2.s567, dark1.s0123, dark1.s4);
| CALC_SUM (y2.s67, y1.s0123, y1.s45, dark2.s67, dark1.s0123, dark1.s45);
| CALC_SUM (y2.s7, y1.s0123, y1.s456, dark2.s7, dark1.s0123, dark1.s456);
| }
|
| float8 out_data = data_sum / weight_sum;
| write_imageui(output_dark, (int2)(pos_x, pos_y), convert_uint4(as_ushort4(convert_uchar8(out_data))));
| }
|
|
