cl_kernel/kernel_yuv_pipe.cl

/*
 * 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

}