/* * function: kernel_yuv_pipe * input: image2d_t as read only * output: image2d_t as write only */ #pragma OPENCL FP_CONTRACT OFF //#define USE_BUFFER_OBJECT 0 unsigned int get_sector_id (float u, float v) { u = fabs(u) > 0.00001f ? u : 0.00001f; float tg = v / u; unsigned int se = tg > 1.f ? (tg > 2.f ? 3 : 2) : (tg > 0.5f ? 1 : 0); unsigned int so = tg > -1.f ? (tg > -0.5f ? 3 : 2) : (tg > -2.f ? 1 : 0); return tg > 0 ? (u > 0 ? se : (se + 8)) : (u > 0 ? (so + 12) : (so + 4)); } __inline void cl_csc_rgbatonv12(float8 *R, float8 *G, float8 *B, float8 *out, __global float *matrix) { out[0] = mad(matrix[0], R[0], mad(matrix[1], G[0], matrix[2] * B[0])); out[1] = mad(matrix[0], R[1], mad(matrix[1], G[1], matrix[2] * B[1])); out[2].s0 = mad(matrix[3], R[0].s0, mad(matrix[4], G[0].s0, matrix[5] * B[0].s0)); out[2].s1 = mad(matrix[6], R[0].s0, mad(matrix[7], G[0].s0, matrix[8] * B[0].s0)); out[2].s2 = mad(matrix[3], R[0].s2, mad(matrix[4], G[0].s2, matrix[5] * B[0].s2)); out[2].s3 = mad(matrix[6], R[0].s2, mad(matrix[7], G[0].s2, matrix[8] * B[0].s2)); out[2].s4 = mad(matrix[3], R[0].s4, mad(matrix[4], G[0].s4, matrix[5] * B[0].s4)); out[2].s5 = mad(matrix[6], R[0].s4, mad(matrix[7], G[0].s4, matrix[8] * B[0].s4)); out[2].s6 = mad(matrix[3], R[0].s6, mad(matrix[4], G[0].s6, matrix[5] * B[0].s6)); out[2].s7 = mad(matrix[6], R[0].s6, mad(matrix[7], G[0].s6, matrix[8] * B[0].s6)); } __inline void cl_macc(float8 *in, __global float *table) { unsigned int table_id[4]; float8 out; table_id[0] = get_sector_id(in[0].s0, in[0].s1); table_id[1] = get_sector_id(in[0].s2, in[0].s3); table_id[2] = get_sector_id(in[0].s4, in[0].s5); table_id[3] = get_sector_id(in[0].s6, in[0].s7); out.s0 = mad(in[0].s0, table[4 * table_id[0]], in[0].s1 * table[4 * table_id[0] + 1]) + 0.5f; out.s1 = mad(in[0].s0, table[4 * table_id[0] + 2], in[0].s1 * table[4 * table_id[0] + 3]) + 0.5f; out.s2 = mad(in[0].s2, table[4 * table_id[1]], in[0].s3 * table[4 * table_id[1] + 1]) + 0.5f; out.s3 = mad(in[0].s2, table[4 * table_id[1] + 2], in[0].s3 * table[4 * table_id[1] + 3]) + 0.5f; out.s4 = mad(in[0].s4, table[4 * table_id[0]], in[0].s5 * table[4 * table_id[0] + 1]) + 0.5f; out.s5 = mad(in[0].s4, table[4 * table_id[0] + 2], in[0].s5 * table[4 * table_id[0] + 3]) + 0.5f; out.s6 = mad(in[0].s6, table[4 * table_id[1]], in[0].s7 * table[4 * table_id[1] + 1]) + 0.5f; out.s7 = mad(in[0].s6, table[4 * table_id[1] + 2], in[0].s7 * table[4 * table_id[1] + 3]) + 0.5f; in[0] = out; } #if USE_BUFFER_OBJECT __inline void cl_tnr_yuv( float8 *in, __global uchar8 *inputFramePre, int x, int y, float gain_yuv, float thr_y, float thr_uv, uint vertical_offset, uint x_offset) #else __inline void cl_tnr_yuv( float8 *in, __read_only image2d_t inputFramePre, __read_only image2d_t inputFramePreUV, int x, int y, float gain_yuv, float thr_y, float thr_uv, uint x_offset) #endif { float8 in_prev[3]; sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; #if USE_BUFFER_OBJECT in_prev[0] = convert_float8(inputFramePre[2 * y * x_offset + x]) / 256.0f; in_prev[1] = convert_float8(inputFramePre[(2 * y + 1) * x_offset + x]) / 256.0f; in_prev[2] = convert_float8(inputFramePre[(y + vertical_offset) * x_offset + x]) / 256.0f; #else in_prev[0] = convert_float8(as_uchar8(convert_ushort4(read_imageui(inputFramePre, sampler, (int2)(x, 2 * y))))) / 256.0f; in_prev[1] = convert_float8(as_uchar8(convert_ushort4(read_imageui(inputFramePre, sampler, (int2)(x, 2 * y + 1))))) / 256.0f; in_prev[2] = convert_float8(as_uchar8(convert_ushort4(read_imageui(inputFramePreUV, sampler, (int2)(x, y))))) / 256.0f; #endif float diff_max = 0.8f; float diff_Y[4], coeff_Y[4]; 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)); 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)); 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)); 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)); 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)); 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)); 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)); 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)); coeff_Y[0] = (coeff_Y[0] < 1.0f) ? coeff_Y[0] : 1.0f; coeff_Y[1] = (coeff_Y[1] < 1.0f) ? coeff_Y[1] : 1.0f; coeff_Y[2] = (coeff_Y[2] < 1.0f) ? coeff_Y[2] : 1.0f; coeff_Y[3] = (coeff_Y[3] < 1.0f) ? coeff_Y[3] : 1.0f; in[0].s01 = mad(in[0].s01 - in_prev[0].s01, coeff_Y[0], in_prev[0].s01); in[1].s01 = mad(in[1].s01 - in_prev[1].s01, coeff_Y[0], in_prev[1].s01); in[0].s23 = mad(in[0].s23 - in_prev[0].s23, coeff_Y[1], in_prev[0].s23); in[1].s23 = mad(in[1].s23 - in_prev[1].s23, coeff_Y[1], in_prev[1].s23); in[0].s45 = mad(in[0].s45 - in_prev[0].s45, coeff_Y[2], in_prev[0].s45); in[1].s45 = mad(in[1].s45 - in_prev[1].s45, coeff_Y[2], in_prev[1].s45); in[0].s67 = mad(in[0].s67 - in_prev[0].s67, coeff_Y[3], in_prev[0].s67); in[1].s67 = mad(in[1].s67 - in_prev[1].s67, coeff_Y[3], in_prev[1].s67); float diff_U[4], diff_V[4], coeff_U[4], coeff_V[4]; diff_U[0] = fabs(in[3].s0 - in_prev[3].s0); diff_U[1] = fabs(in[3].s2 - in_prev[3].s2); diff_U[2] = fabs(in[3].s4 - in_prev[3].s4); diff_U[3] = fabs(in[3].s6 - in_prev[3].s6); diff_V[0] = fabs(in[3].s1 - in_prev[3].s1); diff_V[1] = fabs(in[3].s3 - in_prev[3].s3); diff_V[2] = fabs(in[3].s5 - in_prev[3].s5); diff_V[3] = fabs(in[3].s7 - in_prev[3].s7); 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)); 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)); 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)); 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)); 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)); 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)); 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)); 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)); coeff_U[0] = (coeff_U[0] < 1.0f) ? coeff_U[0] : 1.0f; coeff_U[1] = (coeff_U[1] < 1.0f) ? coeff_U[1] : 1.0f; coeff_U[2] = (coeff_U[2] < 1.0f) ? coeff_U[2] : 1.0f; coeff_U[3] = (coeff_U[3] < 1.0f) ? coeff_U[3] : 1.0f; coeff_V[0] = (coeff_V[0] < 1.0f) ? coeff_V[0] : 1.0f; coeff_V[1] = (coeff_V[1] < 1.0f) ? coeff_V[1] : 1.0f; coeff_V[2] = (coeff_V[2] < 1.0f) ? coeff_V[2] : 1.0f; coeff_V[3] = (coeff_V[3] < 1.0f) ? coeff_V[3] : 1.0f; in[2].s0 = mad(in[2].s0 - in_prev[2].s0, coeff_U[0], in_prev[2].s0); in[2].s1 = mad(in[2].s1 - in_prev[2].s1, coeff_V[0], in_prev[2].s1); in[2].s2 = mad(in[2].s2 - in_prev[2].s2, coeff_U[1], in_prev[2].s2); in[2].s3 = mad(in[2].s3 - in_prev[2].s3, coeff_V[1], in_prev[2].s3); in[2].s4 = mad(in[2].s4 - in_prev[2].s4, coeff_U[2], in_prev[2].s4); in[2].s5 = mad(in[2].s5 - in_prev[2].s5, coeff_V[2], in_prev[2].s5); in[2].s6 = mad(in[2].s6 - in_prev[2].s6, coeff_U[3], in_prev[2].s6); in[2].s7 = mad(in[2].s7 - in_prev[2].s7, coeff_V[3], in_prev[2].s7); } #if USE_BUFFER_OBJECT __kernel void kernel_yuv_pipe ( __global uchar8 *output, __global uchar8 *inputFramePre, uint vertical_offset, uint plannar_offset, __global float *matrix, __global float *table, float yuv_gain, float thr_y, float thr_uv, uint tnr_yuv_enable, __global ushort8 *inputFrame0) #else __kernel void kernel_yuv_pipe ( __write_only image2d_t output, __write_only image2d_t output_uv, __read_only image2d_t inputFramePre, __read_only image2d_t inputFramePreUV, uint plannar_offset, __global float *matrix, __global float *table, float yuv_gain, float thr_y, float thr_uv, uint tnr_yuv_enable, __read_only image2d_t inputFrame0) #endif { int x = get_global_id (0); int y = get_global_id (1); int offsetX = get_global_size(0); sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; float8 inR[2], inG[2], inB[2]; float8 out[3]; #if USE_BUFFER_OBJECT // x [0, 240] // y [0, 540] uint offsetE = 2 * y * offsetX + x; uint offsetO = (2 * y + 1) * offsetX + x; uint offsetUV = (y + vertical_offset) * offsetX + x; uint offsetG = offsetX * plannar_offset; uint offsetB = offsetX * plannar_offset * 2; inR[0] = convert_float8(inputFrame0[offsetE]) / 65536.0f; inR[1] = convert_float8(inputFrame0[offsetO]) / 65536.0f; inG[0] = convert_float8(inputFrame0[offsetE + offsetG]) / 65536.0f; inG[1] = convert_float8(inputFrame0[offsetO + offsetG]) / 65536.0f; inB[0] = convert_float8(inputFrame0[offsetE + offsetB]) / 65536.0f; inB[1] = convert_float8(inputFrame0[offsetO + offsetB]) / 65536.0f; #else inR[0] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y)))) / 65536.0f; inR[1] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + 1)))) / 65536.0f; inG[0] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + plannar_offset)))) / 65536.0f; inG[1] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + 1 + plannar_offset)))) / 65536.0f; inB[0] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + plannar_offset * 2)))) / 65536.0f; inB[1] = convert_float8(as_ushort8(read_imageui(inputFrame0, sampler, (int2)(x, 2 * y + 1 + plannar_offset * 2)))) / 65536.0f; #endif cl_csc_rgbatonv12(&inR[0], &inG[0], &inB[0], &out[0], matrix); cl_macc(&out[2], table); if (tnr_yuv_enable) { #if USE_BUFFER_OBJECT cl_tnr_yuv (&out[0], inputFramePre, x, y, yuv_gain, thr_y, thr_uv, vertical_offset, offsetX); #else cl_tnr_yuv (&out[0], inputFramePre, inputFramePreUV, x, y, yuv_gain, thr_y, thr_uv, offsetX); #endif } #if USE_BUFFER_OBJECT output[offsetE] = convert_uchar8(out[0] * 255.0f); output[offsetO] = convert_uchar8(out[1] * 255.0f); output[offsetUV] = convert_uchar8(out[2] * 255.0f); #else write_imageui(output, (int2)(x, 2 * y), convert_uint4(as_ushort4(convert_uchar8_sat(out[0] * 255.0f)))); write_imageui(output, (int2)(x, 2 * y + 1), convert_uint4(as_ushort4(convert_uchar8_sat(out[1] * 255.0f)))); write_imageui(output_uv, (int2)(x, y), convert_uint4(as_ushort4(convert_uchar8_sat(out[2] * 255.0f)))); #endif }