/*
|
* 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
|
}
|
