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
|
/*
* function: kernel_wavelet_haar_decomposition
* wavelet haar decomposition kernel
* input: input image data as read only
* ll/hl/lh/hh: wavelet decomposition image
* layer: wavelet decomposition layer
* decomLevels: wavelet decomposition levels
*/
#ifndef WAVELET_DENOISE_Y
#define WAVELET_DENOISE_Y 1
#endif
#ifndef WAVELET_DENOISE_UV
#define WAVELET_DENOISE_UV 0
#endif
#ifndef WAVELET_BAYES_SHRINK
#define WAVELET_BAYES_SHRINK 1
#endif
__kernel void kernel_wavelet_haar_decomposition (
__read_only image2d_t input,
__write_only image2d_t ll, __write_only image2d_t hl,
__write_only image2d_t lh, __write_only image2d_t hh,
int layer, int decomLevels,
float hardThresh, float softThresh)
{
int x = get_global_id (0);
int y = get_global_id (1);
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
float8 line[2];
line[0].lo = read_imagef(input, sampler, (int2)(2 * x, 2 * y));
line[0].hi = read_imagef(input, sampler, (int2)(2 * x + 1, 2 * y));
line[1].lo = read_imagef(input, sampler, (int2)(2 * x, 2 * y + 1));
line[1].hi = read_imagef(input, sampler, (int2)(2 * x + 1, 2 * y + 1));
// row transform
float8 row_l;
float8 row_h;
row_l = (float8)(line[0].lo + line[1].lo, line[0].hi + line[1].hi) / 2.0f;
row_h = (float8)(line[0].lo - line[1].lo, line[0].hi - line[1].hi) / 2.0f;
float4 line_ll;
float4 line_hl;
float4 line_lh;
float4 line_hh;
#if WAVELET_DENOISE_Y
// column transform
line_ll = (row_l.odd + row_l.even) / 2.0f;
line_hl = (row_l.odd - row_l.even) / 2.0f;
line_lh = (row_h.odd + row_h.even) / 2.0f;
line_hh = (row_h.odd - row_h.even) / 2.0f;
#endif
#if WAVELET_DENOISE_UV
// U column transform
line_ll.odd = (row_l.odd.odd + row_l.odd.even) / 2.0f;
line_hl.odd = (row_l.odd.odd - row_l.odd.even) / 2.0f;
line_lh.odd = (row_h.odd.odd + row_h.odd.even) / 2.0f;
line_hh.odd = (row_h.odd.odd - row_h.odd.even) / 2.0f;
// V column transform
line_ll.even = (row_l.even.odd + row_l.even.even) / 2.0f;
line_hl.even = (row_l.even.odd - row_l.even.even) / 2.0f;
line_lh.even = (row_h.even.odd + row_h.even.even) / 2.0f;
line_hh.even = (row_h.even.odd - row_h.even.even) / 2.0f;
#endif
write_imagef(ll, (int2)(x, y), line_ll);
write_imagef(hl, (int2)(x, y), line_hl + 0.5f);
write_imagef(lh, (int2)(x, y), line_lh + 0.5f);
write_imagef(hh, (int2)(x, y), line_hh + 0.5f);
}
/*
* function: kernel_wavelet_haar_reconstruction
* wavelet haar reconstruction kernel
* output: output wavelet reconstruction image
* ll/hl/lh/hh: input wavelet transform data as read only
* layer: wavelet reconstruction layer
* decomLevels: wavelet decomposition levels
* threshold: hard/soft denoise thresholding
*/
__constant float uv_threshConst[5] = { 0.1659f, 0.06719f, 0.03343f, 0.01713f, 0.01043f };
__constant float y_threshConst[5] = { 0.06129f, 0.027319f, 0.012643f, 0.006513f, 0.003443f };
__kernel void kernel_wavelet_haar_reconstruction (
__write_only image2d_t output,
__read_only image2d_t ll, __read_only image2d_t hl,
__read_only image2d_t lh, __read_only image2d_t hh,
int layer, int decomLevels,
float hardThresh, float softThresh)
{
int x = get_global_id (0);
int y = get_global_id (1);
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
float thresh = 0.0f;
float4 line_ll;
float4 line_hl;
float4 line_lh;
float4 line_hh;
line_ll = read_imagef(ll, sampler, (int2)(x, y));
line_hl = read_imagef(hl, sampler, (int2)(x, y)) - 0.5f;
line_lh = read_imagef(lh, sampler, (int2)(x, y)) - 0.5f;
line_hh = read_imagef(hh, sampler, (int2)(x, y)) - 0.5f;
#if WAVELET_DENOISE_Y
thresh = hardThresh * y_threshConst[layer - 1];
#endif
#if WAVELET_DENOISE_UV
thresh = hardThresh * uv_threshConst[layer - 1];
#endif
#if !WAVELET_BAYES_SHRINK
// thresholding
line_hl = (line_hl < -thresh) ? line_hl + (thresh - thresh * softThresh) : line_hl;
line_hl = (line_hl > thresh) ? line_hl - (thresh - thresh * softThresh) : line_hl;
line_hl = (line_hl > -thresh && line_hl < thresh) ? line_hl * softThresh : line_hl;
line_lh = (line_lh < -thresh) ? line_lh + (thresh - thresh * softThresh) : line_lh;
line_lh = (line_lh > thresh) ? line_lh - (thresh - thresh * softThresh) : line_lh;
line_lh = (line_lh > -thresh && line_lh < thresh) ? line_lh * softThresh : line_lh;
line_hh = (line_hh < -thresh) ? line_hh + (thresh - thresh * softThresh) : line_hh;
line_hh = (line_hh > thresh) ? line_hh - (thresh - thresh * softThresh) : line_hh;
line_hh = (line_hh > -thresh && line_hh < thresh) ? line_hh * softThresh : line_hh;
#endif
#if WAVELET_DENOISE_Y
// row reconstruction
float8 row_l;
float8 row_h;
row_l = (float8)(line_ll + line_lh, line_hl + line_hh);
row_h = (float8)(line_ll - line_lh, line_hl - line_hh);
// column reconstruction
float8 line[2];
line[0].odd = row_l.lo + row_l.hi;
line[0].even = row_l.lo - row_l.hi;
line[1].odd = row_h.lo + row_h.hi;
line[1].even = row_h.lo - row_h.hi;
write_imagef(output, (int2)(2 * x, 2 * y), line[0].lo);
write_imagef(output, (int2)(2 * x + 1, 2 * y), line[0].hi);
write_imagef(output, (int2)(2 * x, 2 * y + 1), line[1].lo);
write_imagef(output, (int2)(2 * x + 1, 2 * y + 1), line[1].hi);
#endif
#if WAVELET_DENOISE_UV
// row reconstruction
float8 row_l;
float8 row_h;
row_l = (float8)(line_ll + line_lh, line_hl + line_hh);
row_h = (float8)(line_ll - line_lh, line_hl - line_hh);
float8 line[2];
// U column reconstruction
line[0].odd.odd = row_l.lo.odd + row_l.hi.odd;
line[0].odd.even = row_l.lo.odd - row_l.hi.odd;
line[1].odd.odd = row_h.lo.odd + row_h.hi.odd;
line[1].odd.even = row_h.lo.odd - row_h.hi.odd;
// V column reconstruction
line[0].even.odd = row_l.lo.even + row_l.hi.even;
line[0].even.even = row_l.lo.even - row_l.hi.even;
line[1].even.odd = row_h.lo.even + row_h.hi.even;
line[1].even.even = row_h.lo.even - row_h.hi.even;
write_imagef(output, (int2)(2 * x, 2 * y), line[0].lo);
write_imagef(output, (int2)(2 * x + 1, 2 * y), line[0].hi);
write_imagef(output, (int2)(2 * x, 2 * y + 1), line[1].lo);
write_imagef(output, (int2)(2 * x + 1, 2 * y + 1), line[1].hi);
#endif
}
|
