1/* 2 * function: kernel_tnr_yuv 3 * Temporal Noise Reduction 4 * inputFrame: image2d_t as read only 5 * inputFrame0: image2d_t as read only 6 * outputFrame: image2d_t as write only 7 * vertical_offset: vertical offset from y to uv 8 * gain: Blending ratio of previous and current frame 9 * thr_y: Motion sensitivity for Y, higher value can cause more motion blur 10 * thr_uv: Motion sensitivity for UV, higher value can cause more motion blur 11 */ 12 13__kernel void kernel_tnr_yuv( 14 __read_only image2d_t inputFrame, __read_only image2d_t inputFrame0, 15 __write_only image2d_t outputFrame, uint vertical_offset, float gain, float thr_y, float thr_uv) 16{ 17 int x = get_global_id(0); 18 int y = get_global_id(1); 19 20 sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; 21 float4 pixel_t0_Y1 = read_imagef(inputFrame0, sampler, (int2)(2 * x, 2 * y)); 22 float4 pixel_t0_Y2 = read_imagef(inputFrame0, sampler, (int2)(2 * x + 1, 2 * y)); 23 float4 pixel_t0_Y3 = read_imagef(inputFrame0, sampler, (int2)(2 * x, 2 * y + 1)); 24 float4 pixel_t0_Y4 = read_imagef(inputFrame0, sampler, (int2)(2 * x + 1, 2 * y + 1)); 25 26 float4 pixel_t0_U = read_imagef(inputFrame0, sampler, (int2)(2 * x, y + vertical_offset)); 27 float4 pixel_t0_V = read_imagef(inputFrame0, sampler, (int2)(2 * x + 1, y + vertical_offset)); 28 29 float4 pixel_Y1 = read_imagef(inputFrame, sampler, (int2)(2 * x, 2 * y)); 30 float4 pixel_Y2 = read_imagef(inputFrame, sampler, (int2)(2 * x + 1, 2 * y)); 31 float4 pixel_Y3 = read_imagef(inputFrame, sampler, (int2)(2 * x, 2 * y + 1)); 32 float4 pixel_Y4 = read_imagef(inputFrame, sampler, (int2)(2 * x + 1, 2 * y + 1)); 33 34 float4 pixel_U = read_imagef(inputFrame, sampler, (int2)(2 * x, y + vertical_offset)); 35 float4 pixel_V = read_imagef(inputFrame, sampler, (int2)(2 * x + 1, y + vertical_offset)); 36 37 float diff_max = 0.8f; 38 39 float diff_Y = 0.25f * (fabs(pixel_Y1.x - pixel_t0_Y1.x) + fabs(pixel_Y2.x - pixel_t0_Y2.x) + 40 fabs(pixel_Y3.x - pixel_t0_Y3.x) + fabs(pixel_Y4.x - pixel_t0_Y4.x)); 41 42 float coeff_Y = (diff_Y < thr_y) ? gain : 43 (diff_Y * (1 - gain) + diff_max * gain - thr_y) / (diff_max - thr_y); 44 coeff_Y = (coeff_Y < 1.0f) ? coeff_Y : 1.0f; 45 46 float4 pixel_outY1; 47 float4 pixel_outY2; 48 float4 pixel_outY3; 49 float4 pixel_outY4; 50 // X'(K) = (1 - gain) * X'(k-1) + gain * X(k) 51 pixel_outY1.x = pixel_t0_Y1.x + (pixel_Y1.x - pixel_t0_Y1.x) * coeff_Y; 52 pixel_outY2.x = pixel_t0_Y2.x + (pixel_Y2.x - pixel_t0_Y2.x) * coeff_Y; 53 pixel_outY3.x = pixel_t0_Y3.x + (pixel_Y3.x - pixel_t0_Y3.x) * coeff_Y; 54 pixel_outY4.x = pixel_t0_Y4.x + (pixel_Y4.x - pixel_t0_Y4.x) * coeff_Y; 55 56 float diff_U = fabs(pixel_U.x - pixel_t0_U.x); 57 float diff_V = fabs(pixel_V.x - pixel_t0_V.x); 58 59 float coeff_U = (diff_U < thr_uv) ? gain : 60 (diff_U * (1 - gain) + diff_max * gain - thr_uv) / (diff_max - thr_uv); 61 float coeff_V = (diff_V < thr_uv) ? gain : 62 (diff_V * (1 - gain) + diff_max * gain - thr_uv) / (diff_max - thr_uv); 63 coeff_U = (coeff_U < 1.0f) ? coeff_U : 1.0f; 64 coeff_V = (coeff_V < 1.0f) ? coeff_V : 1.0f; 65 66 float4 pixel_outU; 67 float4 pixel_outV; 68 pixel_outU.x = pixel_t0_U.x + (pixel_U.x - pixel_t0_U.x) * coeff_U; 69 pixel_outV.x = pixel_t0_V.x + (pixel_V.x - pixel_t0_V.x) * coeff_V; 70 71 write_imagef(outputFrame, (int2)(2 * x, 2 * y), pixel_outY1); 72 write_imagef(outputFrame, (int2)(2 * x + 1, 2 * y), pixel_outY2); 73 write_imagef(outputFrame, (int2)(2 * x, 2 * y + 1), pixel_outY3); 74 write_imagef(outputFrame, (int2)(2 * x + 1, 2 * y + 1), pixel_outY4); 75 write_imagef(outputFrame, (int2)(2 * x, y + vertical_offset), pixel_outU); 76 write_imagef(outputFrame, (int2)(2 * x + 1, y + vertical_offset), pixel_outV); 77} 78 79/* 80 * function: kernel_tnr_rgb 81 * Temporal Noise Reduction 82 * outputFrame: image2d_t as write only 83 * thr: Motion sensitivity, higher value can cause more motion blur 84 * frameCount: input frame count to be processed 85 * inputFrame: image2d_t as read only 86 */ 87 88__kernel void kernel_tnr_rgb( 89 __write_only image2d_t outputFrame, 90 float tnr_gain, float thr_r, float thr_g, float thr_b, unsigned char frameCount, 91 __read_only image2d_t inputFrame0, __read_only image2d_t inputFrame1, 92 __read_only image2d_t inputFrame2, __read_only image2d_t inputFrame3) 93{ 94 int x = get_global_id(0); 95 int y = get_global_id(1); 96 97 sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; 98 99 float4 pixel_in0; 100 float4 pixel_in1; 101 float4 pixel_in2; 102 float4 pixel_in3; 103 104 float4 pixel_out; 105 float4 var; 106 float gain = 0; 107 int cond; 108 109 pixel_in0 = read_imagef(inputFrame0, sampler, (int2)(x, y)); 110 pixel_in1 = read_imagef(inputFrame1, sampler, (int2)(x, y)); 111 112 if(frameCount == 4) { 113 pixel_in2 = read_imagef(inputFrame2, sampler, (int2)(x, y)); 114 pixel_in3 = read_imagef(inputFrame3, sampler, (int2)(x, y)); 115 116 var.x = (fabs(pixel_in0.x - pixel_in1.x) + fabs(pixel_in1.x - pixel_in2.x) + 117 fabs(pixel_in2.x - pixel_in3.x)) / 3.0f; 118 var.y = (fabs(pixel_in0.y - pixel_in1.y) + fabs(pixel_in1.y - pixel_in2.y) + 119 fabs(pixel_in2.y - pixel_in3.y)) / 3.0f; 120 var.z = (fabs(pixel_in0.z - pixel_in1.z) + fabs(pixel_in1.z - pixel_in2.z) + 121 fabs(pixel_in2.z - pixel_in3.z)) / 3.0f; 122 123 cond = (var.x + var.y + var.z) < (thr_r + thr_g + thr_b); 124 gain = cond ? 1.0f : 0.0f; 125 126 pixel_out.x = (gain * pixel_in0.x + gain * pixel_in1.x + gain * pixel_in2.x + pixel_in3.x) / (1.0f + 3 * gain); 127 pixel_out.y = (gain * pixel_in0.y + gain * pixel_in1.y + gain * pixel_in2.y + pixel_in3.y) / (1.0f + 3 * gain); 128 pixel_out.z = (gain * pixel_in0.z + gain * pixel_in1.z + gain * pixel_in2.z + pixel_in3.z) / (1.0f + 3 * gain); 129 } 130 else if(frameCount == 3) { 131 pixel_in2 = read_imagef(inputFrame2, sampler, (int2)(x, y)); 132 var.x = (fabs(pixel_in0.x - pixel_in1.x) + fabs(pixel_in1.x - pixel_in2.x)) / 2.0f; 133 var.y = (fabs(pixel_in0.y - pixel_in1.y) + fabs(pixel_in1.y - pixel_in2.y)) / 2.0f; 134 var.z = (fabs(pixel_in0.z - pixel_in1.z) + fabs(pixel_in1.z - pixel_in2.z)) / 2.0f; 135 136 cond = (var.x + var.y + var.z) < (thr_r + thr_g + thr_b); 137 gain = cond ? 1.0f : 0.0f; 138 139 pixel_out.x = (gain * pixel_in0.x + gain * pixel_in1.x + pixel_in2.x) / (1.0f + 2 * gain); 140 pixel_out.y = (gain * pixel_in0.y + gain * pixel_in1.y + pixel_in2.y) / (1.0f + 2 * gain); 141 pixel_out.z = (gain * pixel_in0.z + gain * pixel_in1.z + pixel_in2.z) / (1.0f + 2 * gain); 142 } 143 else if(frameCount == 2) 144 { 145 var.x = fabs(pixel_in0.x - pixel_in1.x); 146 var.y = fabs(pixel_in0.y - pixel_in1.y); 147 var.z = fabs(pixel_in0.z - pixel_in1.z); 148 149 cond = (var.x + var.y + var.z) < (thr_r + thr_g + thr_b); 150 gain = cond ? 1.0f : 0.0f; 151 152 pixel_out.x = (gain * pixel_in0.x + pixel_in1.x) / (1.0f + gain); 153 pixel_out.y = (gain * pixel_in0.y + pixel_in1.y) / (1.0f + gain); 154 pixel_out.z = (gain * pixel_in0.z + pixel_in1.z) / (1.0f + gain); 155 } 156 157 write_imagef(outputFrame, (int2)(x, y), pixel_out); 158} 159 160