/*
|
* function: kernel_3d_denoise
|
* 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 ENABLE_GRADIENT 1
|
|
#ifndef WORKGROUP_WIDTH
|
#define WORKGROUP_WIDTH 2
|
#endif
|
|
#ifndef WORKGROUP_HEIGHT
|
#define WORKGROUP_HEIGHT 32
|
#endif
|
|
#define REF_BLOCK_X_OFFSET 1
|
#define REF_BLOCK_Y_OFFSET 4
|
|
#define REF_BLOCK_WIDTH (WORKGROUP_WIDTH + 2 * REF_BLOCK_X_OFFSET)
|
#define REF_BLOCK_HEIGHT (WORKGROUP_HEIGHT + 2 * REF_BLOCK_Y_OFFSET)
|
|
inline int2 subgroup_pos(const int sg_id, const int sg_lid)
|
{
|
int2 pos;
|
pos.x = mad24(2, sg_id % 2, sg_lid % 2);
|
pos.y = mad24(4, sg_id / 2, sg_lid / 2);
|
|
return pos;
|
}
|
|
inline void average_slice(float8 ref,
|
float8 observe,
|
float8* restore,
|
float2* sum_weight,
|
float gain,
|
float threshold,
|
uint sg_id,
|
uint sg_lid)
|
{
|
float8 grad = 0.0f;
|
float8 gradient = 0.0f;
|
float8 dist = 0.0f;
|
float8 distance = 0.0f;
|
float weight = 0.0f;
|
|
#if ENABLE_GRADIENT
|
// calculate & cumulate gradient
|
if (sg_lid % 2 == 0) {
|
grad = intel_sub_group_shuffle(ref, 4);
|
} else {
|
grad = intel_sub_group_shuffle(ref, 5);
|
}
|
gradient = (float8)(grad.s1, grad.s1, grad.s1, grad.s1, grad.s5, grad.s5, grad.s5, grad.s5);
|
|
// normalize gradient "1/(4*255.0f) = 0.00098039f"
|
grad = fabs(gradient - ref) * 0.00098039f;
|
//grad = mad(-2, gradient, (ref + grad)) * 0.0004902f;
|
|
grad.s0 = (grad.s0 + grad.s1 + grad.s2 + grad.s3);
|
grad.s4 = (grad.s4 + grad.s5 + grad.s6 + grad.s7);
|
#endif
|
// calculate & normalize distance "1/255.0f = 0.00392157f"
|
dist = (observe - ref) * 0.00392157f;
|
dist = dist * dist;
|
|
float8 dist_shuffle[8];
|
dist_shuffle[0] = (intel_sub_group_shuffle(dist, 0));
|
dist_shuffle[1] = (intel_sub_group_shuffle(dist, 1));
|
dist_shuffle[2] = (intel_sub_group_shuffle(dist, 2));
|
dist_shuffle[3] = (intel_sub_group_shuffle(dist, 3));
|
dist_shuffle[4] = (intel_sub_group_shuffle(dist, 4));
|
dist_shuffle[5] = (intel_sub_group_shuffle(dist, 5));
|
dist_shuffle[6] = (intel_sub_group_shuffle(dist, 6));
|
dist_shuffle[7] = (intel_sub_group_shuffle(dist, 7));
|
|
if (sg_lid % 2 == 0) {
|
distance = dist_shuffle[0];
|
distance += dist_shuffle[2];
|
distance += dist_shuffle[4];
|
distance += dist_shuffle[6];
|
}
|
else {
|
distance = dist_shuffle[1];
|
distance += dist_shuffle[3];
|
distance += dist_shuffle[5];
|
distance += dist_shuffle[7];
|
}
|
|
// cumulate distance
|
dist.s0 = (distance.s0 + distance.s1 + distance.s2 + distance.s3);
|
dist.s4 = (distance.s4 + distance.s5 + distance.s6 + distance.s7);
|
gain = (grad.s0 < threshold) ? gain : 2.0f * gain;
|
weight = native_exp(-gain * dist.s0);
|
(*restore).lo = mad(weight, ref.lo, (*restore).lo);
|
(*sum_weight).lo = (*sum_weight).lo + weight;
|
|
gain = (grad.s4 < threshold) ? gain : 2.0f * gain;
|
weight = native_exp(-gain * dist.s4);
|
(*restore).hi = mad(weight, ref.hi, (*restore).hi);
|
(*sum_weight).hi = (*sum_weight).hi + weight;
|
}
|
|
inline void weighted_average (__read_only image2d_t input,
|
__local uchar8* ref_cache,
|
bool load_observe,
|
float8* observe,
|
float8* restore,
|
float2* sum_weight,
|
float gain,
|
float threshold,
|
uint sg_id,
|
uint sg_lid)
|
{
|
sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
|
|
int local_id_x = get_local_id(0);
|
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 start_x = mad24(group_id_x, WORKGROUP_WIDTH, -REF_BLOCK_X_OFFSET);
|
int start_y = mad24(group_id_y, WORKGROUP_HEIGHT, -REF_BLOCK_Y_OFFSET);
|
|
int i = local_id_x + local_id_y * WORKGROUP_WIDTH;
|
for ( int j = i; j < (REF_BLOCK_HEIGHT * REF_BLOCK_WIDTH);
|
j += (WORKGROUP_HEIGHT * WORKGROUP_WIDTH) ) {
|
int corrd_x = start_x + (j % REF_BLOCK_WIDTH);
|
int corrd_y = start_y + (j / REF_BLOCK_WIDTH);
|
|
ref_cache[j] = as_uchar8( convert_ushort4(read_imageui(input,
|
sampler,
|
(int2)(corrd_x, corrd_y))));
|
}
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
#if WORKGROUP_WIDTH == 4
|
int2 pos = subgroup_pos(sg_id, sg_lid);
|
local_id_x = pos.x;
|
local_id_y = pos.y;
|
#endif
|
|
if (load_observe) {
|
(*observe) = convert_float8(
|
ref_cache[mad24(local_id_y + REF_BLOCK_Y_OFFSET,
|
REF_BLOCK_WIDTH,
|
local_id_x + REF_BLOCK_X_OFFSET)]);
|
(*restore) = (*observe);
|
(*sum_weight) = 1.0f;
|
}
|
|
float8 ref[2] = {0.0f, 0.0f};
|
__local uchar4* p_ref = (__local uchar4*)(ref_cache);
|
|
// top-left
|
ref[0] = convert_float8(*(__local uchar8*)(p_ref + mad24(local_id_y,
|
2 * REF_BLOCK_WIDTH,
|
mad24(2, local_id_x, 1))));
|
average_slice(ref[0], *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// top-right
|
ref[1] = convert_float8(*(__local uchar8*)(p_ref + mad24(local_id_y,
|
2 * REF_BLOCK_WIDTH,
|
mad24(2, local_id_x, 3))));
|
average_slice(ref[1], *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// top-mid
|
average_slice((float8)(ref[0].hi, ref[1].lo), *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// mid-left
|
ref[0] = convert_float8(*(__local uchar8*)(p_ref + mad24((local_id_y + 4),
|
2 * REF_BLOCK_WIDTH,
|
mad24(2, local_id_x, 1))));
|
average_slice(ref[0], *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// mid-right
|
ref[1] = convert_float8(*(__local uchar8*)(p_ref + mad24((local_id_y + 4),
|
2 * REF_BLOCK_WIDTH,
|
mad24(2, local_id_x, 3))));
|
average_slice(ref[1], *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// mid-mid
|
if (!load_observe) {
|
average_slice((float8)(ref[0].hi, ref[1].lo), *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
}
|
|
// bottom-left
|
ref[0] = convert_float8(*(__local uchar8*)(p_ref + mad24((local_id_y + 8),
|
2 * REF_BLOCK_WIDTH,
|
mad24(2, local_id_x, 1))));
|
average_slice(ref[0], *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// bottom-right
|
ref[1] = convert_float8(*(__local uchar8*)(p_ref + mad24((local_id_y + 8),
|
2 * REF_BLOCK_WIDTH,
|
mad24(2, local_id_x, 3))));
|
average_slice(ref[1], *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
|
// bottom-mid
|
average_slice((float8)(ref[0].hi, ref[1].lo), *observe, restore, sum_weight, gain, threshold, sg_id, sg_lid);
|
}
|
|
__kernel void kernel_3d_denoise ( 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)
|
{
|
float8 restore = 0.0f;
|
float8 observe = 0.0f;
|
float2 sum_weight = 0.0f;
|
|
const int sg_id = get_sub_group_id();
|
const int sg_lid = (get_local_id(1) * WORKGROUP_WIDTH + get_local_id(0)) % 8;
|
|
__local uchar8 ref_cache[REF_BLOCK_HEIGHT * REF_BLOCK_WIDTH];
|
|
weighted_average (input, ref_cache, true, &observe, &restore, &sum_weight, gain, threshold, sg_id, sg_lid);
|
|
#if ENABLE_IIR_FILERING
|
weighted_average (restoredPrev, ref_cache, false, &observe, &restore, &sum_weight, gain, threshold, sg_id, sg_lid);
|
#else
|
#if REFERENCE_FRAME_COUNT > 1
|
weighted_average (inputPrev1, ref_cache, false, &observe, &restore, &sum_weight, gain, threshold, sg_id, sg_lid);
|
#endif
|
|
#if REFERENCE_FRAME_COUNT > 2
|
weighted_average (inputPrev2, ref_cache, false, &observe, &restore, &sum_weight, gain, threshold, sg_id, sg_lid);
|
#endif
|
#endif
|
|
restore.lo = restore.lo / sum_weight.lo;
|
restore.hi = restore.hi / sum_weight.hi;
|
|
int local_id_x = get_local_id(0);
|
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);
|
|
#if WORKGROUP_WIDTH == 4
|
int2 pos = subgroup_pos(sg_id, sg_lid);
|
local_id_x = pos.x;
|
local_id_y = pos.y;
|
#endif
|
|
int coor_x = mad24(group_id_x, WORKGROUP_WIDTH, local_id_x);
|
int coor_y = mad24(group_id_y, WORKGROUP_HEIGHT, local_id_y);
|
|
write_imageui(output,
|
(int2)(coor_x, coor_y),
|
convert_uint4(as_ushort4(convert_uchar8(restore))));
|
}
|
