/*
* function: kernel_3d_denoise_slm
* 3D Noise Reduction
* gain: The parameter determines the filtering strength for the reference block
* threshold: Noise variances of observed image
* restoredPrev: The previous restored image, image2d_t as read only
* output: restored image, image2d_t as write only
* input: observed image, image2d_t as read only
* inputPrev1: reference image, image2d_t as read only
* inputPrev2: reference image, image2d_t as read only
*/
#ifndef REFERENCE_FRAME_COUNT
#define REFERENCE_FRAME_COUNT 2
#endif
#ifndef ENABLE_IIR_FILERING
#define ENABLE_IIR_FILERING 1
#endif
#define WORK_GROUP_WIDTH 8
#define WORK_GROUP_HEIGHT 1
#define WORK_BLOCK_WIDTH 8
#define WORK_BLOCK_HEIGHT 8
#define REF_BLOCK_X_OFFSET 1
#define REF_BLOCK_Y_OFFSET 4
#define REF_BLOCK_WIDTH (WORK_BLOCK_WIDTH + 2 * REF_BLOCK_X_OFFSET)
#define REF_BLOCK_HEIGHT (WORK_BLOCK_HEIGHT + 2 * REF_BLOCK_Y_OFFSET)
inline void weighted_average (__read_only image2d_t input,
__local float4* ref_cache,
bool load_observe,
__local float4* observe_cache,
float4* restore,
float2* sum_weight,
float gain,
float threshold)
{
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const int local_id_x = get_local_id(0);
const int local_id_y = get_local_id(1);
const int group_id_x = get_group_id(0);
const int group_id_y = get_group_id(1);
int i = local_id_x + local_id_y * WORK_BLOCK_WIDTH;
int start_x = mad24(group_id_x, WORK_BLOCK_WIDTH, -REF_BLOCK_X_OFFSET);
int start_y = mad24(group_id_y, WORK_BLOCK_HEIGHT, -REF_BLOCK_Y_OFFSET);
for (int j = i; j < REF_BLOCK_WIDTH * REF_BLOCK_HEIGHT; j += (WORK_GROUP_WIDTH * WORK_GROUP_HEIGHT)) {
int corrd_x = start_x + (j % REF_BLOCK_WIDTH);
int corrd_y = start_y + (j / REF_BLOCK_WIDTH);
ref_cache[j] = read_imagef(input, sampler, (int2)(corrd_x, corrd_y));
}
barrier(CLK_LOCAL_MEM_FENCE);
if (load_observe) {
for (int i = 0; i < WORK_BLOCK_HEIGHT; i++) {
observe_cache[i * WORK_BLOCK_WIDTH + local_id_x] =
ref_cache[(i + REF_BLOCK_Y_OFFSET) * REF_BLOCK_WIDTH
+ local_id_x + REF_BLOCK_X_OFFSET];
}
}
float4 dist = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
float4 gradient = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
float weight = 0.0f;
#pragma unroll
for (int i = 0; i < 3; i++) {
#pragma unroll
for (int j = 0; j < 3; j++) {
dist = (ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)] -
observe_cache[local_id_x]) *
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)] -
observe_cache[local_id_x]);
dist = mad((ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[WORK_BLOCK_WIDTH + local_id_x]),
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[WORK_BLOCK_WIDTH + local_id_x]),
dist);
dist = mad((ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[2 * WORK_BLOCK_WIDTH + local_id_x]),
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[2 * WORK_BLOCK_WIDTH + local_id_x]),
dist);
dist = mad((ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[3 * WORK_BLOCK_WIDTH + local_id_x]),
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[3 * WORK_BLOCK_WIDTH + local_id_x]),
dist);
gradient = (float4)(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2,
ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2,
ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2,
ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2);
gradient = (gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)]) +
(gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)]) +
(gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)]) +
(gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)]);
gradient.s0 = (gradient.s0 + gradient.s1 + gradient.s2 + gradient.s3) / 15.0f;
gain = (gradient.s0 < threshold) ? gain : 2.0f * gain;
weight = native_exp(-gain * (dist.s0 + dist.s1 + dist.s2 + dist.s3));
weight = (weight < 0) ? 0 : weight;
(*sum_weight).s0 = (*sum_weight).s0 + weight;
restore[0] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)], restore[0]);
restore[1] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)], restore[1]);
restore[2] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)], restore[2]);
restore[3] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)], restore[3]);
}
}
#pragma unroll
for (int i = 1; i < 4; i++) {
#pragma unroll
for (int j = 0; j < 3; j++) {
dist = (ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)] -
observe_cache[4 * WORK_BLOCK_WIDTH + local_id_x]) *
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)] -
observe_cache[4 * WORK_BLOCK_WIDTH + local_id_x]);
dist = mad((ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[5 * WORK_BLOCK_WIDTH + local_id_x]),
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[5 * WORK_BLOCK_WIDTH + local_id_x]),
dist);
dist = mad((ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[6 * WORK_BLOCK_WIDTH + local_id_x]),
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[6 * WORK_BLOCK_WIDTH + local_id_x]),
dist);
dist = mad((ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[7 * WORK_BLOCK_WIDTH + local_id_x]),
(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)] -
observe_cache[7 * WORK_BLOCK_WIDTH + local_id_x]),
dist);
gradient = (float4)(ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2,
ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2,
ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2,
ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)].s2);
gradient = (gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)]) +
(gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)]) +
(gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)]) +
(gradient - ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)]);
gradient.s0 = (gradient.s0 + gradient.s1 + gradient.s2 + gradient.s3) / 15.0f;
gain = (gradient.s0 < threshold) ? gain : 2.0f * gain;
weight = native_exp(-gain * (dist.s0 + dist.s1 + dist.s2 + dist.s3));
weight = (weight < 0) ? 0 : weight;
(*sum_weight).s1 = (*sum_weight).s1 + weight;
restore[4] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, local_id_x + j)], restore[4]);
restore[5] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, REF_BLOCK_WIDTH + local_id_x + j)], restore[5]);
restore[6] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 2 * REF_BLOCK_WIDTH + local_id_x + j)], restore[6]);
restore[7] = mad(weight, ref_cache[mad24(i, 4 * REF_BLOCK_WIDTH, 3 * REF_BLOCK_WIDTH + local_id_x + j)], restore[7]);
}
}
}
__kernel void kernel_3d_denoise_slm( float gain,
float threshold,
__read_only image2d_t restoredPrev,
__write_only image2d_t output,
__read_only image2d_t input,
__read_only image2d_t inputPrev1,
__read_only image2d_t inputPrev2)
{
float4 restore[8] = {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f};
float2 sum_weight = {0.0f, 0.0f};
__local float4 ref_cache[REF_BLOCK_HEIGHT * REF_BLOCK_WIDTH];
__local float4 observe_cache[WORK_BLOCK_HEIGHT * WORK_BLOCK_WIDTH];
weighted_average (input, ref_cache, true, observe_cache, restore, &sum_weight, gain, threshold);
#if 1
#if ENABLE_IIR_FILERING
weighted_average (restoredPrev, ref_cache, false, observe_cache, restore, &sum_weight, gain, threshold);
#else
#if REFERENCE_FRAME_COUNT > 1
weighted_average (inputPrev1, ref_cache, false, observe_cache, restore, &sum_weight, gain, threshold);
#endif
#if REFERENCE_FRAME_COUNT > 2
weighted_average (inputPrev2, ref_cache, false, observe_cache, restore, &sum_weight, gain, threshold);
#endif
#endif
#endif
restore[0] = restore[0] / sum_weight.s0;
restore[1] = restore[1] / sum_weight.s0;
restore[2] = restore[2] / sum_weight.s0;
restore[3] = restore[3] / sum_weight.s0;
restore[4] = restore[4] / sum_weight.s1;
restore[5] = restore[5] / sum_weight.s1;
restore[6] = restore[6] / sum_weight.s1;
restore[7] = restore[7] / sum_weight.s1;
const int global_id_x = get_global_id (0);
const int global_id_y = get_global_id (1);
write_imagef(output, (int2)(global_id_x, 8 * global_id_y), restore[0]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 1)), restore[1]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 2)), restore[2]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 3)), restore[3]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 4)), restore[4]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 5)), restore[5]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 6)), restore[6]);
write_imagef(output, (int2)(global_id_x, mad24(8, global_id_y, 7)), restore[7]);
}