diff options
Diffstat (limited to 'shaders/cl/kernel_tnr.cl')
| -rw-r--r-- | shaders/cl/kernel_tnr.cl | 159 |
1 files changed, 159 insertions, 0 deletions
diff --git a/shaders/cl/kernel_tnr.cl b/shaders/cl/kernel_tnr.cl new file mode 100644 index 0000000..e5c1c36 --- /dev/null +++ b/shaders/cl/kernel_tnr.cl @@ -0,0 +1,159 @@ +/* + * function: kernel_tnr_yuv + * Temporal Noise Reduction + * inputFrame: image2d_t as read only + * inputFrame0: image2d_t as read only + * outputFrame: image2d_t as write only + * vertical_offset: vertical offset from y to uv + * gain: Blending ratio of previous and current frame + * thr_y: Motion sensitivity for Y, higher value can cause more motion blur + * thr_uv: Motion sensitivity for UV, higher value can cause more motion blur + */ + +__kernel void kernel_tnr_yuv( + __read_only image2d_t inputFrame, __read_only image2d_t inputFrame0, + __write_only image2d_t outputFrame, uint vertical_offset, float gain, float thr_y, float thr_uv) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; + float4 pixel_t0_Y1 = read_imagef(inputFrame0, sampler, (int2)(2 * x, 2 * y)); + float4 pixel_t0_Y2 = read_imagef(inputFrame0, sampler, (int2)(2 * x + 1, 2 * y)); + float4 pixel_t0_Y3 = read_imagef(inputFrame0, sampler, (int2)(2 * x, 2 * y + 1)); + float4 pixel_t0_Y4 = read_imagef(inputFrame0, sampler, (int2)(2 * x + 1, 2 * y + 1)); + + float4 pixel_t0_U = read_imagef(inputFrame0, sampler, (int2)(2 * x, y + vertical_offset)); + float4 pixel_t0_V = read_imagef(inputFrame0, sampler, (int2)(2 * x + 1, y + vertical_offset)); + + float4 pixel_Y1 = read_imagef(inputFrame, sampler, (int2)(2 * x, 2 * y)); + float4 pixel_Y2 = read_imagef(inputFrame, sampler, (int2)(2 * x + 1, 2 * y)); + float4 pixel_Y3 = read_imagef(inputFrame, sampler, (int2)(2 * x, 2 * y + 1)); + float4 pixel_Y4 = read_imagef(inputFrame, sampler, (int2)(2 * x + 1, 2 * y + 1)); + + float4 pixel_U = read_imagef(inputFrame, sampler, (int2)(2 * x, y + vertical_offset)); + float4 pixel_V = read_imagef(inputFrame, sampler, (int2)(2 * x + 1, y + vertical_offset)); + + float diff_max = 0.8f; + + float diff_Y = 0.25f * (fabs(pixel_Y1.x - pixel_t0_Y1.x) + fabs(pixel_Y2.x - pixel_t0_Y2.x) + + fabs(pixel_Y3.x - pixel_t0_Y3.x) + fabs(pixel_Y4.x - pixel_t0_Y4.x)); + + float coeff_Y = (diff_Y < thr_y) ? gain : + (diff_Y * (1 - gain) + diff_max * gain - thr_y) / (diff_max - thr_y); + coeff_Y = (coeff_Y < 1.0f) ? coeff_Y : 1.0f; + + float4 pixel_outY1; + float4 pixel_outY2; + float4 pixel_outY3; + float4 pixel_outY4; + // X'(K) = (1 - gain) * X'(k-1) + gain * X(k) + pixel_outY1.x = pixel_t0_Y1.x + (pixel_Y1.x - pixel_t0_Y1.x) * coeff_Y; + pixel_outY2.x = pixel_t0_Y2.x + (pixel_Y2.x - pixel_t0_Y2.x) * coeff_Y; + pixel_outY3.x = pixel_t0_Y3.x + (pixel_Y3.x - pixel_t0_Y3.x) * coeff_Y; + pixel_outY4.x = pixel_t0_Y4.x + (pixel_Y4.x - pixel_t0_Y4.x) * coeff_Y; + + float diff_U = fabs(pixel_U.x - pixel_t0_U.x); + float diff_V = fabs(pixel_V.x - pixel_t0_V.x); + + float coeff_U = (diff_U < thr_uv) ? gain : + (diff_U * (1 - gain) + diff_max * gain - thr_uv) / (diff_max - thr_uv); + float coeff_V = (diff_V < thr_uv) ? gain : + (diff_V * (1 - gain) + diff_max * gain - thr_uv) / (diff_max - thr_uv); + coeff_U = (coeff_U < 1.0f) ? coeff_U : 1.0f; + coeff_V = (coeff_V < 1.0f) ? coeff_V : 1.0f; + + float4 pixel_outU; + float4 pixel_outV; + pixel_outU.x = pixel_t0_U.x + (pixel_U.x - pixel_t0_U.x) * coeff_U; + pixel_outV.x = pixel_t0_V.x + (pixel_V.x - pixel_t0_V.x) * coeff_V; + + write_imagef(outputFrame, (int2)(2 * x, 2 * y), pixel_outY1); + write_imagef(outputFrame, (int2)(2 * x + 1, 2 * y), pixel_outY2); + write_imagef(outputFrame, (int2)(2 * x, 2 * y + 1), pixel_outY3); + write_imagef(outputFrame, (int2)(2 * x + 1, 2 * y + 1), pixel_outY4); + write_imagef(outputFrame, (int2)(2 * x, y + vertical_offset), pixel_outU); + write_imagef(outputFrame, (int2)(2 * x + 1, y + vertical_offset), pixel_outV); +} + +/* + * function: kernel_tnr_rgb + * Temporal Noise Reduction + * outputFrame: image2d_t as write only + * thr: Motion sensitivity, higher value can cause more motion blur + * frameCount: input frame count to be processed + * inputFrame: image2d_t as read only + */ + +__kernel void kernel_tnr_rgb( + __write_only image2d_t outputFrame, + float tnr_gain, float thr_r, float thr_g, float thr_b, unsigned char frameCount, + __read_only image2d_t inputFrame0, __read_only image2d_t inputFrame1, + __read_only image2d_t inputFrame2, __read_only image2d_t inputFrame3) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; + + float4 pixel_in0; + float4 pixel_in1; + float4 pixel_in2; + float4 pixel_in3; + + float4 pixel_out; + float4 var; + float gain = 0; + int cond; + + pixel_in0 = read_imagef(inputFrame0, sampler, (int2)(x, y)); + pixel_in1 = read_imagef(inputFrame1, sampler, (int2)(x, y)); + + if(frameCount == 4) { + pixel_in2 = read_imagef(inputFrame2, sampler, (int2)(x, y)); + pixel_in3 = read_imagef(inputFrame3, sampler, (int2)(x, y)); + + var.x = (fabs(pixel_in0.x - pixel_in1.x) + fabs(pixel_in1.x - pixel_in2.x) + + fabs(pixel_in2.x - pixel_in3.x)) / 3.0f; + var.y = (fabs(pixel_in0.y - pixel_in1.y) + fabs(pixel_in1.y - pixel_in2.y) + + fabs(pixel_in2.y - pixel_in3.y)) / 3.0f; + var.z = (fabs(pixel_in0.z - pixel_in1.z) + fabs(pixel_in1.z - pixel_in2.z) + + fabs(pixel_in2.z - pixel_in3.z)) / 3.0f; + + cond = (var.x + var.y + var.z) < (thr_r + thr_g + thr_b); + gain = cond ? 1.0f : 0.0f; + + pixel_out.x = (gain * pixel_in0.x + gain * pixel_in1.x + gain * pixel_in2.x + pixel_in3.x) / (1.0f + 3 * gain); + pixel_out.y = (gain * pixel_in0.y + gain * pixel_in1.y + gain * pixel_in2.y + pixel_in3.y) / (1.0f + 3 * gain); + pixel_out.z = (gain * pixel_in0.z + gain * pixel_in1.z + gain * pixel_in2.z + pixel_in3.z) / (1.0f + 3 * gain); + } + else if(frameCount == 3) { + pixel_in2 = read_imagef(inputFrame2, sampler, (int2)(x, y)); + var.x = (fabs(pixel_in0.x - pixel_in1.x) + fabs(pixel_in1.x - pixel_in2.x)) / 2.0f; + var.y = (fabs(pixel_in0.y - pixel_in1.y) + fabs(pixel_in1.y - pixel_in2.y)) / 2.0f; + var.z = (fabs(pixel_in0.z - pixel_in1.z) + fabs(pixel_in1.z - pixel_in2.z)) / 2.0f; + + cond = (var.x + var.y + var.z) < (thr_r + thr_g + thr_b); + gain = cond ? 1.0f : 0.0f; + + pixel_out.x = (gain * pixel_in0.x + gain * pixel_in1.x + pixel_in2.x) / (1.0f + 2 * gain); + pixel_out.y = (gain * pixel_in0.y + gain * pixel_in1.y + pixel_in2.y) / (1.0f + 2 * gain); + pixel_out.z = (gain * pixel_in0.z + gain * pixel_in1.z + pixel_in2.z) / (1.0f + 2 * gain); + } + else if(frameCount == 2) + { + var.x = fabs(pixel_in0.x - pixel_in1.x); + var.y = fabs(pixel_in0.y - pixel_in1.y); + var.z = fabs(pixel_in0.z - pixel_in1.z); + + cond = (var.x + var.y + var.z) < (thr_r + thr_g + thr_b); + gain = cond ? 1.0f : 0.0f; + + pixel_out.x = (gain * pixel_in0.x + pixel_in1.x) / (1.0f + gain); + pixel_out.y = (gain * pixel_in0.y + pixel_in1.y) / (1.0f + gain); + pixel_out.z = (gain * pixel_in0.z + pixel_in1.z) / (1.0f + gain); + } + + write_imagef(outputFrame, (int2)(x, y), pixel_out); +} + |