• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1/*
2 * function: kernel_yuv_pipe
3 * input:    image2d_t as read only
4 * output:   image2d_t as write only
5 */
6
7#pragma OPENCL FP_CONTRACT OFF
8
9//#define USE_BUFFER_OBJECT 0
10
11unsigned int get_sector_id (float u, float v)
12{
13    u = fabs(u) > 0.00001f ? u : 0.00001f;
14    float tg = v / u;
15    unsigned int se = tg > 1.f ? (tg > 2.f ? 3 : 2) : (tg > 0.5f ? 1 : 0);
16    unsigned int so = tg > -1.f ? (tg > -0.5f ? 3 : 2) : (tg > -2.f ? 1 : 0);
17    return tg > 0 ? (u > 0 ? se : (se + 8)) : (u > 0 ? (so + 12) : (so + 4));
18}
19
20__inline void cl_csc_rgbatonv12(float8 *R, float8 *G, float8 *B, float8 *out, __global float *matrix)
21{
22    out[0] = mad(matrix[0], R[0], mad(matrix[1], G[0], matrix[2] * B[0]));
23    out[1] = mad(matrix[0], R[1], mad(matrix[1], G[1], matrix[2] * B[1]));
24
25    out[2].s0 = mad(matrix[3], R[0].s0, mad(matrix[4], G[0].s0, matrix[5] * B[0].s0));
26    out[2].s1 = mad(matrix[6], R[0].s0, mad(matrix[7], G[0].s0, matrix[8] * B[0].s0));
27    out[2].s2 = mad(matrix[3], R[0].s2, mad(matrix[4], G[0].s2, matrix[5] * B[0].s2));
28    out[2].s3 = mad(matrix[6], R[0].s2, mad(matrix[7], G[0].s2, matrix[8] * B[0].s2));
29    out[2].s4 = mad(matrix[3], R[0].s4, mad(matrix[4], G[0].s4, matrix[5] * B[0].s4));
30    out[2].s5 = mad(matrix[6], R[0].s4, mad(matrix[7], G[0].s4, matrix[8] * B[0].s4));
31    out[2].s6 = mad(matrix[3], R[0].s6, mad(matrix[4], G[0].s6, matrix[5] * B[0].s6));
32    out[2].s7 = mad(matrix[6], R[0].s6, mad(matrix[7], G[0].s6, matrix[8] * B[0].s6));
33
34}
35
36__inline void cl_macc(float8 *in, __global float *table)
37{
38    unsigned int table_id[4];
39    float8 out;
40
41    table_id[0] = get_sector_id(in[0].s0, in[0].s1);
42    table_id[1] = get_sector_id(in[0].s2, in[0].s3);
43    table_id[2] = get_sector_id(in[0].s4, in[0].s5);
44    table_id[3] = get_sector_id(in[0].s6, in[0].s7);
45
46    out.s0 = mad(in[0].s0, table[4 * table_id[0]], in[0].s1 * table[4 * table_id[0] + 1]) + 0.5f;
47    out.s1 = mad(in[0].s0, table[4 * table_id[0] + 2], in[0].s1 * table[4 * table_id[0] + 3]) + 0.5f;
48    out.s2 = mad(in[0].s2, table[4 * table_id[1]], in[0].s3 * table[4 * table_id[1] + 1]) + 0.5f;
49    out.s3 = mad(in[0].s2, table[4 * table_id[1] + 2], in[0].s3 * table[4 * table_id[1] + 3]) + 0.5f;
50    out.s4 = mad(in[0].s4, table[4 * table_id[0]], in[0].s5 * table[4 * table_id[0] + 1]) + 0.5f;
51    out.s5 = mad(in[0].s4, table[4 * table_id[0] + 2], in[0].s5 * table[4 * table_id[0] + 3]) + 0.5f;
52    out.s6 = mad(in[0].s6, table[4 * table_id[1]], in[0].s7 * table[4 * table_id[1] + 1]) + 0.5f;
53    out.s7 = mad(in[0].s6, table[4 * table_id[1] + 2], in[0].s7 * table[4 * table_id[1] + 3]) + 0.5f;
54
55    in[0] = out;
56}
57
58#if USE_BUFFER_OBJECT
59__inline void cl_tnr_yuv(
60    float8 *in, __global uchar8 *inputFramePre,
61    int x, int y,
62    float gain_yuv, float thr_y, float thr_uv,
63    uint vertical_offset, uint x_offset)
64#else
65__inline void cl_tnr_yuv(
66    float8 *in,
67    __read_only image2d_t inputFramePre, __read_only image2d_t inputFramePreUV,
68    int x, int y,
69    float gain_yuv, float thr_y, float thr_uv, uint x_offset)
70#endif
71{
72    float8 in_prev[3];
73    sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
74
75#if USE_BUFFER_OBJECT
76    in_prev[0] = convert_float8(inputFramePre[2 * y * x_offset + x]) / 256.0f;
77    in_prev[1] = convert_float8(inputFramePre[(2 * y + 1) * x_offset + x]) / 256.0f;
78    in_prev[2] = convert_float8(inputFramePre[(y + vertical_offset) * x_offset + x]) / 256.0f;
79#else
80    in_prev[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(inputFramePre, sampler, (int2)(x, 2 * y))))) / 256.0f;
81    in_prev[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(inputFramePre, sampler, (int2)(x, 2 * y + 1))))) / 256.0f;
82    in_prev[2] = convert_float8(as_uchar8(convert_ushort4(read_imageui(inputFramePreUV, sampler, (int2)(x, y))))) / 256.0f;
83#endif
84
85    float diff_max = 0.8f;
86    float diff_Y[4], coeff_Y[4];
87
88    diff_Y[0] = 0.25f * (fabs(in[0].s0 - in_prev[0].s0) + fabs(in[0].s1 - in_prev[0].s1) + fabs(in[1].s0 - in_prev[1].s0) + fabs(in[1].s1 - in_prev[1].s1));
89    diff_Y[1] = 0.25f * (fabs(in[0].s2 - in_prev[0].s2) + fabs(in[0].s3 - in_prev[0].s3) + fabs(in[1].s2 - in_prev[1].s2) + fabs(in[1].s3 - in_prev[1].s3));
90    diff_Y[2] = 0.25f * (fabs(in[0].s4 - in_prev[0].s4) + fabs(in[0].s5 - in_prev[0].s5) + fabs(in[1].s4 - in_prev[1].s4) + fabs(in[1].s5 - in_prev[1].s5));
91    diff_Y[3] = 0.25f * (fabs(in[0].s6 - in_prev[0].s6) + fabs(in[0].s7 - in_prev[0].s7) + fabs(in[1].s6 - in_prev[1].s6) + fabs(in[1].s7 - in_prev[1].s7));
92
93    coeff_Y[0] = (diff_Y[0] < thr_y) ? gain_yuv : (mad(diff_Y[0], 1 - gain_yuv, diff_max * gain_yuv - thr_y) / (diff_max - thr_y));
94    coeff_Y[1] = (diff_Y[1] < thr_y) ? gain_yuv : (mad(diff_Y[1], 1 - gain_yuv, diff_max * gain_yuv - thr_y) / (diff_max - thr_y));
95    coeff_Y[2] = (diff_Y[2] < thr_y) ? gain_yuv : (mad(diff_Y[2], 1 - gain_yuv, diff_max * gain_yuv - thr_y) / (diff_max - thr_y));
96    coeff_Y[3] = (diff_Y[3] < thr_y) ? gain_yuv : (mad(diff_Y[3], 1 - gain_yuv, diff_max * gain_yuv - thr_y) / (diff_max - thr_y));
97
98    coeff_Y[0] = (coeff_Y[0] < 1.0f) ? coeff_Y[0] : 1.0f;
99    coeff_Y[1] = (coeff_Y[1] < 1.0f) ? coeff_Y[1] : 1.0f;
100    coeff_Y[2] = (coeff_Y[2] < 1.0f) ? coeff_Y[2] : 1.0f;
101    coeff_Y[3] = (coeff_Y[3] < 1.0f) ? coeff_Y[3] : 1.0f;
102
103    in[0].s01 = mad(in[0].s01 - in_prev[0].s01, coeff_Y[0], in_prev[0].s01);
104    in[1].s01 = mad(in[1].s01 - in_prev[1].s01, coeff_Y[0], in_prev[1].s01);
105    in[0].s23 = mad(in[0].s23 - in_prev[0].s23, coeff_Y[1], in_prev[0].s23);
106    in[1].s23 = mad(in[1].s23 - in_prev[1].s23, coeff_Y[1], in_prev[1].s23);
107    in[0].s45 = mad(in[0].s45 - in_prev[0].s45, coeff_Y[2], in_prev[0].s45);
108    in[1].s45 = mad(in[1].s45 - in_prev[1].s45, coeff_Y[2], in_prev[1].s45);
109    in[0].s67 = mad(in[0].s67 - in_prev[0].s67, coeff_Y[3], in_prev[0].s67);
110    in[1].s67 = mad(in[1].s67 - in_prev[1].s67, coeff_Y[3], in_prev[1].s67);
111
112    float diff_U[4], diff_V[4], coeff_U[4], coeff_V[4];
113
114    diff_U[0] = fabs(in[3].s0 - in_prev[3].s0);
115    diff_U[1] = fabs(in[3].s2 - in_prev[3].s2);
116    diff_U[2] = fabs(in[3].s4 - in_prev[3].s4);
117    diff_U[3] = fabs(in[3].s6 - in_prev[3].s6);
118
119    diff_V[0] = fabs(in[3].s1 - in_prev[3].s1);
120    diff_V[1] = fabs(in[3].s3 - in_prev[3].s3);
121    diff_V[2] = fabs(in[3].s5 - in_prev[3].s5);
122    diff_V[3] = fabs(in[3].s7 - in_prev[3].s7);
123
124    coeff_U[0] = (diff_U[0] < thr_uv) ? gain_yuv : (mad(diff_U[0], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
125    coeff_U[1] = (diff_U[1] < thr_uv) ? gain_yuv : (mad(diff_U[1], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
126    coeff_U[2] = (diff_U[2] < thr_uv) ? gain_yuv : (mad(diff_U[2], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
127    coeff_U[3] = (diff_U[3] < thr_uv) ? gain_yuv : (mad(diff_U[3], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
128
129    coeff_V[0] = (diff_V[0] < thr_uv) ? gain_yuv : (mad(diff_V[0], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
130    coeff_V[1] = (diff_V[1] < thr_uv) ? gain_yuv : (mad(diff_V[1], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
131    coeff_V[2] = (diff_V[2] < thr_uv) ? gain_yuv : (mad(diff_V[2], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
132    coeff_V[3] = (diff_V[3] < thr_uv) ? gain_yuv : (mad(diff_V[3], 1 - gain_yuv, diff_max * gain_yuv - thr_uv) / (diff_max - thr_uv));
133
134    coeff_U[0] = (coeff_U[0] < 1.0f) ? coeff_U[0] : 1.0f;
135    coeff_U[1] = (coeff_U[1] < 1.0f) ? coeff_U[1] : 1.0f;
136    coeff_U[2] = (coeff_U[2] < 1.0f) ? coeff_U[2] : 1.0f;
137    coeff_U[3] = (coeff_U[3] < 1.0f) ? coeff_U[3] : 1.0f;
138
139    coeff_V[0] = (coeff_V[0] < 1.0f) ? coeff_V[0] : 1.0f;
140    coeff_V[1] = (coeff_V[1] < 1.0f) ? coeff_V[1] : 1.0f;
141    coeff_V[2] = (coeff_V[2] < 1.0f) ? coeff_V[2] : 1.0f;
142    coeff_V[3] = (coeff_V[3] < 1.0f) ? coeff_V[3] : 1.0f;
143
144    in[2].s0 = mad(in[2].s0 - in_prev[2].s0, coeff_U[0], in_prev[2].s0);
145    in[2].s1 = mad(in[2].s1 - in_prev[2].s1, coeff_V[0], in_prev[2].s1);
146    in[2].s2 = mad(in[2].s2 - in_prev[2].s2, coeff_U[1], in_prev[2].s2);
147    in[2].s3 = mad(in[2].s3 - in_prev[2].s3, coeff_V[1], in_prev[2].s3);
148    in[2].s4 = mad(in[2].s4 - in_prev[2].s4, coeff_U[2], in_prev[2].s4);
149    in[2].s5 = mad(in[2].s5 - in_prev[2].s5, coeff_V[2], in_prev[2].s5);
150    in[2].s6 = mad(in[2].s6 - in_prev[2].s6, coeff_U[3], in_prev[2].s6);
151    in[2].s7 = mad(in[2].s7 - in_prev[2].s7, coeff_V[3], in_prev[2].s7);
152
153}
154
155#if USE_BUFFER_OBJECT
156__kernel void kernel_yuv_pipe (
157    __global uchar8 *output,
158    __global uchar8 *inputFramePre, uint vertical_offset,
159    uint plannar_offset,
160    __global float *matrix, __global float *table,
161    float yuv_gain, float thr_y, float thr_uv, uint tnr_yuv_enable,
162    __global ushort8 *inputFrame0)
163
164#else
165
166__kernel void kernel_yuv_pipe (
167    __write_only image2d_t output, __write_only image2d_t output_uv,
168    __read_only image2d_t inputFramePre, __read_only image2d_t inputFramePreUV,
169    uint plannar_offset,
170    __global float *matrix, __global float *table,
171    float yuv_gain, float thr_y, float thr_uv, uint tnr_yuv_enable,
172    __read_only image2d_t inputFrame0)
173
174#endif
175{
176    int x = get_global_id (0);
177    int y = get_global_id (1);
178    int offsetX = get_global_size(0);
179    sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
180    float8 inR[2], inG[2], inB[2];
181    float8 out[3];
182
183#if USE_BUFFER_OBJECT
184    // x [0, 240]
185    // y [0, 540]
186    uint offsetE = 2 * y * offsetX + x;
187    uint offsetO = (2 * y + 1) * offsetX + x;
188    uint offsetUV = (y + vertical_offset) * offsetX + x;
189    uint offsetG = offsetX * plannar_offset;
190    uint offsetB = offsetX * plannar_offset * 2;
191
192    inR[0] = convert_float8(inputFrame0[offsetE]) / 65536.0f;
193    inR[1] = convert_float8(inputFrame0[offsetO]) / 65536.0f;
194    inG[0] = convert_float8(inputFrame0[offsetE + offsetG]) / 65536.0f;
195    inG[1] = convert_float8(inputFrame0[offsetO + offsetG]) / 65536.0f;
196    inB[0] = convert_float8(inputFrame0[offsetE + offsetB]) / 65536.0f;
197    inB[1] = convert_float8(inputFrame0[offsetO + offsetB]) / 65536.0f;
198#else
199    inR[0] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y)))) / 65536.0f;
200    inR[1] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + 1)))) / 65536.0f;
201    inG[0] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + plannar_offset)))) / 65536.0f;
202    inG[1] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + 1 + plannar_offset)))) / 65536.0f;
203    inB[0] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + plannar_offset * 2)))) / 65536.0f;
204    inB[1] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + 1 + plannar_offset * 2)))) / 65536.0f;
205#endif
206
207    cl_csc_rgbatonv12(&inR[0], &inG[0], &inB[0], &out[0], matrix);
208    cl_macc(&out[2], table);
209
210    if (tnr_yuv_enable) {
211#if USE_BUFFER_OBJECT
212        cl_tnr_yuv (&out[0], inputFramePre, x, y, yuv_gain, thr_y, thr_uv, vertical_offset, offsetX);
213#else
214        cl_tnr_yuv (&out[0], inputFramePre, inputFramePreUV, x, y, yuv_gain, thr_y, thr_uv, offsetX);
215#endif
216
217    }
218
219#if USE_BUFFER_OBJECT
220    output[offsetE] = convert_uchar8(out[0] * 255.0f);
221    output[offsetO] = convert_uchar8(out[1] * 255.0f);
222    output[offsetUV] = convert_uchar8(out[2] * 255.0f);
223#else
224    write_imageui(output, (int2)(x, 2 * y), convert_uint4(as_ushort4(convert_uchar8_sat(out[0] * 255.0f))));
225    write_imageui(output, (int2)(x, 2 * y + 1), convert_uint4(as_ushort4(convert_uchar8_sat(out[1] * 255.0f))));
226    write_imageui(output_uv, (int2)(x, y), convert_uint4(as_ushort4(convert_uchar8_sat(out[2] * 255.0f))));
227#endif
228
229}
230
231