#include "cl_viv_vx_ext.h"

_viv_uniform VXC_512Bits uniBilinearTmp1Bgra_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp2Bgra_4x4;
_viv_uniform VXC_512Bits uniDescaleU8_4x4;
_viv_uniform VXC_512Bits uniConvertIntergetoF32_4x4;
_viv_uniform VXC_512Bits uniExtractInt32BgraToU8_2x8;
_viv_uniform VXC_512Bits uniExchangeBgra_2x8;
_viv_uniform VXC_512Bits uniExchangeBgra2_2x8;

_viv_uniform VXC_512Bits uniBilinearTmp1BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp2BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp3BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp4BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp5BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp6BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp7BgraShort_4x4;
_viv_uniform VXC_512Bits uniBilinearTmp8BgraShort_4x4;

_viv_uniform VXC_512Bits uniExtractBfromBgra_4x4;
_viv_uniform VXC_512Bits uniExtractGfromBgra_4x4;
_viv_uniform VXC_512Bits uniExtractRfromBgra_4x4;
_viv_uniform VXC_512Bits uniConvertInt32toUint8_2x8;

_viv_uniform int bOrder;
_viv_uniform int rOrder;
_viv_uniform int zp;
_viv_uniform float outputScale;

__kernel void pre_process_bgra_scale_U8toU8(
         __read_only image2d_array_t  input,
        __write_only image2d_array_t  output,
              global int *            xRatio,
              global int *            yRatio,
              global int *            xOffset,
              global int *            yOffset,
                     float            rMean,
                     float            gMean,
                     float            bMean,
                     float            r_scale,
                     int              reverse_channel,
                     int              trans,
                     float            g_scale,
                     float            b_scale
    )
{
    int4 gidx = get_global_id(0);
    int gidy = get_global_id(1);
    gidx += (int4)(0, 1, 2, 3);

    int4 fx = (gidx * (*xRatio) + ((*xRatio) >> 1)) - (1 << 14);
    int4 sx = fx & 0xffff8000; // Floor
    int fy, sy;
    fx -= sx;
    sx = sx >> 15;
    fx = (fx +(1 << 4)) >> 5;

    // for y
    fy = (gidy * (*yRatio) + ((*yRatio) >> 1)) - (1<< 14);
    sy = fy & 0xffff8000; // Floor
    fy -= sy;
    sy = sy >> 15;

    sy = sy < 0 ? 0 : sy;
    fy = fy < 0 ? 0 : fy;

    fy = (fy + (1<< 4)) >> 5;
    sx = (sx + (*xOffset)) * 4 ;
    sy += (*yOffset);
    int4 srcPos = (int4)(sx.x, sy, sy + 1, sx.y);
    vxc_uchar16 lineBGRA0, lineBGRA1, lineBGRA2, lineBGRA3;
    vxc_uchar16 dataB, dataG, dataR;

    VXC_ReadImage(lineBGRA0, input, srcPos.xy, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    VXC_ReadImage(lineBGRA0, input, srcPos.xz, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(8, 15, 0, VXC_RM_TowardZero, 0));

    VXC_ReadImage(lineBGRA1, input, srcPos.wy, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    VXC_ReadImage(lineBGRA1, input, srcPos.wz, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(8, 15, 0, VXC_RM_TowardZero, 0));

    srcPos.x = sx.z;
    srcPos.w = sx.w;

    VXC_ReadImage(lineBGRA2, input, srcPos.xy, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    VXC_ReadImage(lineBGRA2, input, srcPos.xz, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(8, 15, 0, VXC_RM_TowardZero, 0));

    VXC_ReadImage(lineBGRA3, input, srcPos.wy, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    VXC_ReadImage(lineBGRA3, input, srcPos.wz, VXC_5BITOFFSET_XY(0, 0),
                     VXC_MODIFIER(8, 15, 0, VXC_RM_TowardZero, 0));

    vxc_uchar4 val_u8;
    int4 tmp1, tmp2, result1, result2;
    float4 tmpDst, tmp0;
    float4 mean = (float4)(bMean, gMean, rMean, 0);
    float4 var = (float4)(b_scale, g_scale, r_scale, 0);
    //tmpFx = (int4)(fx.x, fx.x, fx.x, fx.x);
    int tmpV = 1 << 19;
    vxc_short8 tmpFx;
    VXC_DP2x8(tmpFx, fx, fx, VXC_MODIFIER(0, 7, 0, VXC_RM_ToNearestEven, 1), uniConvertInt32toUint8_2x8);
    //tmpFx = fx.xxxx;
    VXC_DP4x4(tmp1, lineBGRA0, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp1BgraShort_4x4);
    VXC_DP4x4(tmp2, lineBGRA0, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp2BgraShort_4x4);
    tmp1 = fy * (tmp2 - tmp1) + (tmp1 << 10);
    VXC_DP4x4(val_u8, tmp1, tmpV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniDescaleU8_4x4);
    VXC_DP4x4(tmp0, val_u8, val_u8, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), \
        uniConvertIntergetoF32_4x4);
    tmpDst = (tmp0 - mean) * var;
    result1 = convert_int4_rte(tmpDst * outputScale + zp);

    //tmpFx = fx.yyyy;
    VXC_DP4x4(tmp1, lineBGRA1, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp3BgraShort_4x4);
    VXC_DP4x4(tmp2, lineBGRA1, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp4BgraShort_4x4);
    tmp1 = fy * (tmp2 - tmp1) + (tmp1 << 10);
    VXC_DP4x4(val_u8, tmp1, tmpV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniDescaleU8_4x4);
    VXC_DP4x4(tmp0, val_u8, val_u8, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), \
        uniConvertIntergetoF32_4x4);
    tmpDst = (tmp0 - mean) * var;
    result2 = convert_int4_rte(tmpDst * outputScale + zp);

    vxc_uchar16 dst, data;
    VXC_DP2x8(dst, result1, result2, VXC_MODIFIER(0, 7, 0, VXC_RM_ToNearestEven, 1), uniExtractInt32BgraToU8_2x8);

    //tmpFx = fx.zzzz;
    VXC_DP4x4(tmp1, lineBGRA2, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp5BgraShort_4x4);
    VXC_DP4x4(tmp2, lineBGRA2, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp6BgraShort_4x4);
    tmp1 = fy * (tmp2 - tmp1) + (tmp1 << 10);
    VXC_DP4x4(val_u8, tmp1, tmpV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniDescaleU8_4x4);
    VXC_DP4x4(tmp0, val_u8, val_u8, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), \
        uniConvertIntergetoF32_4x4);
    tmpDst = (tmp0 - mean) * var;
    result1 = convert_int4_rte(tmpDst * outputScale + zp);

    //tmpFx = fx.wwww;
    VXC_DP4x4(tmp1, lineBGRA3, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp7BgraShort_4x4);
    VXC_DP4x4(tmp2, lineBGRA3, tmpFx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniBilinearTmp8BgraShort_4x4);
    tmp1 = fy * (tmp2 - tmp1) + (tmp1 << 10);
    VXC_DP4x4(val_u8, tmp1, tmpV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniDescaleU8_4x4);
    VXC_DP4x4(tmp0, val_u8, val_u8, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), \
        uniConvertIntergetoF32_4x4);
    tmpDst = (tmp0 - mean) * var;
    result2 = convert_int4_rte(tmpDst * outputScale + zp);

    VXC_DP2x8(dst, result1, result2, VXC_MODIFIER(8, 15, 0, VXC_RM_ToNearestEven, 1), uniExtractInt32BgraToU8_2x8);

    VXC_DP2x8(data, dst, dst, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExchangeBgra_2x8);
    VXC_DP2x8(data, dst, dst, VXC_MODIFIER(8, 15, 0, VXC_RM_TowardZero, 0), uniExchangeBgra2_2x8);

    int4 dstPos = (int4)(get_global_id(0), gidy, 0, 0);
    dstPos.z = bOrder;
    VXC_WriteImage2DArray(output, dstPos, data, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0));
    dstPos.z = 1;
    VXC_WriteImage2DArray(output, dstPos, data.s4567, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0));
    dstPos.z = rOrder;
    VXC_WriteImage2DArray(output, dstPos, data.s89ab, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0));
}

__kernel void pre_process_bgra_copy_U8toU8(
         __read_only image2d_array_t  input,
        __write_only image2d_array_t  output,
              global int *            xRatio,
              global int *            yRatio,
              global int *            xOffset,
              global int *            yOffset,
                     float            rMean,
                     float            gMean,
                     float            bMean,
                     float            r_scale,
                     int              reverse_channel,
                     int              trans,
                     float            g_scale,
                     float            b_scale
)
{
    int2 pos = (int2)((get_global_id(0) + (*xOffset)) << 2, get_global_id(1) + (*yOffset));

    vxc_uchar16 lineBGRA0;
    float4 tmpB, tmpG, tmpR;
    float4 tmpDst;
    int4 result1, result2;
    vxc_uchar16 dst;

    VXC_ReadImage(lineBGRA0, input, pos, VXC_5BITOFFSET_XY(0, 0), VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));
    VXC_DP4x4(tmpB, lineBGRA0, lineBGRA0, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniExtractBfromBgra_4x4);
    VXC_DP4x4(tmpG, lineBGRA0, lineBGRA0, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniExtractGfromBgra_4x4);
    VXC_DP4x4(tmpR, lineBGRA0, lineBGRA0, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniExtractRfromBgra_4x4);

    tmpDst = (tmpB - bMean) * b_scale;
    result1 = convert_int4_rte(tmpDst * outputScale + zp);

    tmpDst = (tmpG - gMean) * g_scale;
    result2 = convert_int4_rte(tmpDst * outputScale + zp);
    VXC_DP2x8(dst, result1, result2, VXC_MODIFIER(0, 7, 0, VXC_RM_ToNearestEven, 1), uniConvertInt32toUint8_2x8);

    int4 dstPos = (int4)(get_global_id(0), get_global_id(1), 0, 0);
    dstPos.z = bOrder;
    VXC_WriteImage2DArray(output, dstPos, dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0));
    dstPos.z = 1;
    VXC_WriteImage2DArray(output, dstPos, dst, VXC_MODIFIER(4, 7, 0, VXC_RM_TowardZero, 0));

    tmpDst = (tmpR - rMean) * r_scale;
    result1 = convert_int4_rte(tmpDst * outputScale + zp);
    VXC_DP2x8(dst, result1, result1, VXC_MODIFIER(0, 7, 0, VXC_RM_ToNearestEven, 1), uniConvertInt32toUint8_2x8);

    dstPos.z = rOrder;
    VXC_WriteImage2DArray(output, dstPos, dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0));
}
