#include "cl_viv_vx_ext.h"

_viv_uniform VXC_512Bits uniVecShift10;
_viv_uniform VXC_512Bits uniAddRShift;
_viv_uniform VXC_512Bits uniGetTempVal;
_viv_uniform VXC_512Bits uniExtractBytes;
_viv_uniform VXC_512Bits uniUnpackToR;
_viv_uniform VXC_512Bits uniUnpackToG;
_viv_uniform VXC_512Bits uniUnpackToB;

_viv_uniform VXC_512Bits uniConvertIntergetoF32_4x4;
_viv_uniform float outputScale;
_viv_uniform VXC_512Bits uniExtract8Data_2x8;
_viv_uniform float outputZP;
_viv_uniform int r_order;
_viv_uniform int b_order;

#define DESCALE(x) (((x) + (1<<19)) >> 20)

#define IMAGE_PRE_PROCESS(dst_name, conv_type, dst_type, copy_type) \
__kernel void pre_process_rgb_scale_U8to##dst_name \
    ( \
__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         f32Var, \
               int           reverse_channel, \
               int           trans \
    ) \
{ \
    int2 ratioXY = (int2)(*xRatio, *yRatio); \
    int4 xPos       = get_global_id(0); \
    int yPos        = get_global_id(1); \
    int2 ratioSufXY = (ratioXY >> 1) - (1 << 14); \
    xPos += (int4)(0, 1, 2, 3); \
 \
    /*x*/ \
    int4 fx0 = xPos * ratioXY.x + ratioSufXY.x; \
    int4 sx = fx0 & 0xffff8000; \
    fx0 -= sx; \
    sx = sx >> 15; \
 \
    vxc_short4 fx; \
    VXC_DP4x4(fx, fx0, 1 << 4, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniAddRShift); \
    /*y*/ \
    int fy = yPos * ratioXY.y + ratioSufXY.y; \
    int sy = fy & 0xffff8000; \
 \
    fy -= sy; \
    sy = sy >> 15; \
 \
    fy = (fy + (1<< 4)) >> 5; \
 \
    vxc_uchar16 line0RGB1, line0RGB2; \
    vxc_uchar16 line1RGB3, line1RGB4; \
    int4 coord; \
    sx = (sx + (*xOffset)) * 3; \
    coord.xyz    = sx.xyz; \
    coord.w        = sy + *yOffset; \
    int2 coord1 = (int2)(sx.w, coord.w); \
    VXC_ReadImage(line0RGB1, input, coord.xw, 0, VXC_MODIFIER(0, 5, 0, VXC_RM_TowardZero, 0)); \
    VXC_ReadImage(line0RGB1, input, coord.yw, 0, VXC_MODIFIER(6, 11, 0, VXC_RM_TowardZero, 0)); \
    VXC_ReadImage(line0RGB2, input, coord.zw, 0, VXC_MODIFIER(0, 5, 0, VXC_RM_TowardZero, 0)); \
    VXC_ReadImage(line0RGB2, input, coord1, 0, VXC_MODIFIER(6, 11, 0, VXC_RM_TowardZero, 0)); \
 \
    VXC_ReadImage(line1RGB3, input, coord.xw, VXC_5BITOFFSET_XY(0, 1), \
        VXC_MODIFIER(0, 5, 0, VXC_RM_TowardZero, 0)); \
    VXC_ReadImage(line1RGB3, input, coord.yw, VXC_5BITOFFSET_XY(0, 1), \
        VXC_MODIFIER(6, 11, 0, VXC_RM_TowardZero, 0)); \
    VXC_ReadImage(line1RGB4, input, coord.zw, VXC_5BITOFFSET_XY(0, 1), \
        VXC_MODIFIER(0, 5, 0, VXC_RM_TowardZero, 0)); \
    VXC_ReadImage(line1RGB4, input, coord1, VXC_5BITOFFSET_XY(0, 1), \
        VXC_MODIFIER(6, 11, 0, VXC_RM_TowardZero, 0)); \
 \
    float4 bgrMean = (float4)(bMean, gMean, rMean, 0); \
 \
    bgrMean *= f32Var; \
 \
    int4 test01, temp1; \
    int4 test02, temp2; \
    int4 tt; \
    vxc_uchar4 val; \
    int4 coord_out = (int4)(xPos.x, yPos, r_order, 0); \
 \
    vxc_uchar8 line1, line2; \
 \
    /*R*/ \
    VXC_DP2x8(line1, line0RGB1, line0RGB2, \
        VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniUnpackToR); \
    VXC_DP2x8(line2, line1RGB3, line1RGB4, \
        VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniUnpackToR); \
 \
    VXC_DP4x4(test01, line1, line1, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniVecShift10); \
    VXC_DP4x4(temp1, line1, fx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniGetTempVal); \
    temp1 = temp1 + test01; \
 \
    VXC_DP4x4(test02, line2, line2, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniVecShift10); \
    VXC_DP4x4(temp2, line2, fx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniGetTempVal); \
    temp2 = temp2 + test02; \
    temp2 = fy * (temp2 - temp1) + (temp1 << 10); \
 \
    vxc_float4 tmp_dst; \
    vxc_uchar4 u8_dst; \
    VXC_DP4x4(u8_dst, temp2, 1 << 19, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 1), uniExtractBytes); \
    VXC_DP4x4(tmp_dst, u8_dst, u8_dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 1), \
        uniConvertIntergetoF32_4x4); \
 \
    /*convert U8 to dst*/ \
    dst_type dst; \
    tmp_dst = tmp_dst * f32Var - bgrMean.zzzz; \
    tmp_dst = tmp_dst * outputScale + outputZP; \
    conv_type dst0; \
    _viv_asm(CONV_RTE, dst0, tmp_dst); \
    VXC_DP2x8(dst, dst0, dst0, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniExtract8Data_2x8); \
    copy_type result; \
    _viv_asm(COPY, result, dst, 16); \
    VXC_WriteImage2DArray(output, coord_out, result, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0)); \
 \
    /*G*/ \
    VXC_DP2x8(line1, line0RGB1, line0RGB2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniUnpackToG); \
    VXC_DP2x8(line2, line1RGB3, line1RGB4, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniUnpackToG); \
 \
    coord_out.z = 1; \
    VXC_DP4x4(test01, line1, line1, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniVecShift10); \
    VXC_DP4x4(temp1, line1, fx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniGetTempVal); \
    temp1 = temp1 + test01; \
 \
    VXC_DP4x4(test02, line2, line2, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniVecShift10); \
    VXC_DP4x4(temp2, line2, fx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniGetTempVal); \
    temp2 = temp2 + test02; \
    temp2 = fy * (temp2 - temp1) + (temp1 << 10); \
 \
    VXC_DP4x4(u8_dst, temp2, 1 << 19, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 1), uniExtractBytes); \
    VXC_DP4x4(tmp_dst, u8_dst, u8_dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 1), \
            uniConvertIntergetoF32_4x4); \
 \
    tmp_dst = tmp_dst * f32Var - bgrMean.y; \
    tmp_dst = tmp_dst * outputScale + outputZP; \
    _viv_asm(CONV_RTE, dst0, tmp_dst); \
    VXC_DP2x8(dst, dst0, dst0, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniExtract8Data_2x8); \
    _viv_asm(COPY, result, dst, 16); \
    VXC_WriteImage2DArray(output, coord_out, result, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0)); \
 \
    /*B*/ \
    VXC_DP2x8(line1, line0RGB1, line0RGB2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniUnpackToB); \
    VXC_DP2x8(line2, line1RGB3, line1RGB4, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniUnpackToB); \
 \
    coord_out.z = b_order; \
    VXC_DP4x4(test01, line1, line1, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniVecShift10); \
    VXC_DP4x4(temp1, line1, fx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniGetTempVal); \
    temp1 = temp1 + test01; \
 \
    VXC_DP4x4(test02, line2, line2, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniVecShift10); \
    VXC_DP4x4(temp2, line2, fx, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniGetTempVal); \
    temp2 = temp2 + test02; \
    temp2 = fy * (temp2 - temp1) + (temp1 << 10); \
 \
    VXC_DP4x4(u8_dst, temp2, 1 << 19, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 1), uniExtractBytes); \
    VXC_DP4x4(tmp_dst, u8_dst, u8_dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 1), \
        uniConvertIntergetoF32_4x4); \
 \
    tmp_dst = tmp_dst * f32Var - bgrMean.x; \
    tmp_dst = tmp_dst * outputScale + outputZP; \
    _viv_asm(CONV_RTE, dst0, tmp_dst); \
    VXC_DP2x8(dst, dst0, dst0, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1), uniExtract8Data_2x8); \
    _viv_asm(COPY, result, dst, 16); \
    VXC_WriteImage2DArray(output, coord_out, result, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0)); \
}
IMAGE_PRE_PROCESS(U8,  uint4, vxc_uchar16, vxc_uchar16)
IMAGE_PRE_PROCESS(I8,  int4,  vxc_char16,  vxc_char16)
IMAGE_PRE_PROCESS(I16, int4,  vxc_short8,  vxc_short8)
IMAGE_PRE_PROCESS(F16, half4, vxc_half8,   vxc_short8)
