#include "cl_viv_vx_ext.h"

_viv_uniform int bOrder;
_viv_uniform int rOrder;

_viv_uniform float outputScaleVar_b;
_viv_uniform float outputScaleVar_g;
_viv_uniform float outputScaleVar_r;
_viv_uniform float bMeanScaleVarZp;
_viv_uniform float gMeanScaleVarZp;
_viv_uniform float rMeanScaleVarZp;

_viv_uniform VXC_512Bits uniConvertYUV422toB_4x4;
_viv_uniform VXC_512Bits uniConvertYUV422toG_4x4;
_viv_uniform VXC_512Bits uniConvertYUV422toR_4x4;

_viv_uniform VXC_512Bits uniExtract8Data_2x8;
_viv_uniform VXC_512Bits uniExtractYUVtoShortSub_2x8;
_viv_uniform VXC_512Bits uniConvertUchartoFp32_4x4;

#define YUV422_COPY_SH_IMPL(name, dst_type, conv_type, save_type, copy_bytes) \
__kernel void pre_process_yuv422_copy_##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           r_scale, \
                 int             reverse_channel, \
                 int             trans, \
                 int             yuv422_type, \
                 float           g_scale, \
                 float           b_scale \
    ) \
{ \
    int gidx = get_global_id(0); \
    int gidy = get_global_id(1); \
 \
    int sy = gidy + (*yOffset); \
    int sx = gidx * 2 + (*xOffset * 2); \
 \
    vxc_uchar8 YUV; \
    vxc_short8 tmpYUV; \
 \
    VXC_ReadImage(YUV, input, (int2)(sx,sy), 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
 \
    if (yuv422_type == 1) \
    { \
        YUV.s01234567 = YUV.s10325476; \
    } \
 \
    float4 tmpDstB, tmpDstG, tmpDstR; \
    vxc_uchar4 DstB_uchar, DstG_uchar, DstR_uchar; \
    vxc_short2 value = (vxc_short2)(128,16); \
    VXC_DP2x8(tmpYUV, YUV, value, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtractYUVtoShortSub_2x8); \
    VXC_DP4x4(DstB_uchar, tmpYUV, tmpYUV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1),\
                                          uniConvertYUV422toB_4x4); \
    VXC_DP4x4(DstG_uchar, tmpYUV, tmpYUV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1),\
                                          uniConvertYUV422toG_4x4); \
    VXC_DP4x4(DstR_uchar, tmpYUV, tmpYUV, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 1),\
                                          uniConvertYUV422toR_4x4); \
    VXC_DP4x4(tmpDstB, DstB_uchar, DstB_uchar, \
              VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertUchartoFp32_4x4); \
    VXC_DP4x4(tmpDstG, DstG_uchar, DstG_uchar, \
              VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertUchartoFp32_4x4); \
    VXC_DP4x4(tmpDstR, DstR_uchar, DstR_uchar, \
              VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertUchartoFp32_4x4); \
 \
    conv_type result; \
    dst_type dst0; \
    save_type dst; \
    int4 dstPos = (int4)(gidx, gidy, 0, 0); \
    tmpDstB = tmpDstB * outputScaleVar_b + bMeanScaleVarZp; \
    _viv_asm(CONV_RTE, result, tmpDstB); \
    dstPos.z = bOrder; \
    VXC_DP2x8(dst0, result, result, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 0), uniExtract8Data_2x8); \
    _viv_asm(COPY, dst, dst0, copy_bytes); \
    VXC_WriteImage2DArray(output, dstPos, dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0)); \
 \
    tmpDstG = tmpDstG * outputScaleVar_g + gMeanScaleVarZp; \
    _viv_asm(CONV_RTE, result, tmpDstG); \
    dstPos.z = 1; \
    VXC_DP2x8(dst0, result, result, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 0), uniExtract8Data_2x8); \
    _viv_asm(COPY, dst, dst0, copy_bytes); \
    VXC_WriteImage2DArray(output, dstPos, dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0)); \
 \
    tmpDstR = tmpDstR * outputScaleVar_r + rMeanScaleVarZp; \
    _viv_asm(CONV_RTE, result, tmpDstR); \
    dstPos.z = rOrder; \
    VXC_DP2x8(dst0, result, result, VXC_MODIFIER(0, 3, 0, VXC_RM_ToNearestEven, 0), uniExtract8Data_2x8); \
    _viv_asm(COPY, dst, dst0, copy_bytes); \
    VXC_WriteImage2DArray(output, dstPos, dst, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0)); \
}
YUV422_COPY_SH_IMPL(U8toU8,  vxc_uchar4, int4,  vxc_uchar4, 4)
YUV422_COPY_SH_IMPL(U8toI8,  vxc_char4,  int4,  vxc_char4,  4)
YUV422_COPY_SH_IMPL(U8toI16, vxc_short4, int4,  vxc_short4, 8)
YUV422_COPY_SH_IMPL(U8toF16, vxc_half4,  half4, vxc_short4, 8)
