#include "cl_viv_vx_ext.h"

_viv_uniform VXC_512Bits uniConvertInt32toUint8_2x8;
_viv_uniform VXC_512Bits uniConvertFstToFp32_4x4;
_viv_uniform VXC_512Bits uniConvertSecToFp32_4x4;

_viv_uniform float in_scale0;
_viv_uniform float in_scale1;
_viv_uniform float out_scale;
_viv_uniform float in0Tail;
_viv_uniform float in1Tail;
_viv_uniform float out_zp;

#define FLOORDIV_PROCESS(dst_type, save_type, read_type, copy_type, conv_mode, IN0_SCALE, IN0_TAIL,\
                      IN1_SCALE, IN1_TAIL, OUT_SCALE, OUT_OFFSET, read_fun, write_fun) \
    save_type data; \
    read_type read_data0, read_data1; \
    copy_type tmpData0, tmpData1; \
    vxc_float4 in0Val1, in0Val2, in1Val1, in1Val2; \
    vxc_float4 tmpVal1, tmpVal2; \
    dst_type tmpOut1, tmpOut2; \
    read_fun(read_data0, input0, coord, VXC_5BITOFFSET_XY(0,0), \
    VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
    _viv_asm(COPY, tmpData0, read_data0, 16); \
    read_fun(read_data1, input1, coord, VXC_5BITOFFSET_XY(0,0), \
    VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
    _viv_asm(COPY, tmpData1, read_data1, 16); \
    VXC_DP4x4(in0Val1, tmpData0, tmpData0, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertFstToFp32_4x4); \
    VXC_DP4x4(in0Val2, tmpData0, tmpData0, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertSecToFp32_4x4); \
    VXC_DP4x4(in1Val1, tmpData1, tmpData1, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertFstToFp32_4x4); \
    VXC_DP4x4(in1Val2, tmpData1, tmpData1, VXC_MODIFIER(0, 3, 0, VXC_RM_TowardZero, 0), uniConvertSecToFp32_4x4); \
    in0Val1 = in0Val1 * IN0_SCALE + IN0_TAIL; \
    in0Val2 = in0Val2 * IN0_SCALE + IN0_TAIL; \
    in1Val1 = in1Val1 * IN1_SCALE + IN1_TAIL; \
    in1Val2 = in1Val2 * IN1_SCALE + IN1_TAIL; \
    tmpVal1 = floor(in0Val1 / in1Val1) * OUT_SCALE + OUT_OFFSET; \
    tmpVal2 = floor(in0Val2 / in1Val2) * OUT_SCALE + OUT_OFFSET; \
    _viv_asm(conv_mode, tmpOut1, tmpVal1); \
    _viv_asm(conv_mode, tmpOut2, tmpVal2); \
    VXC_DP2x8(data, tmpOut1, tmpOut2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniConvertInt32toUint8_2x8); \
    write_fun(output, coord, data, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \

#define TENSOR_FLOORDIV(src0_name, src1_name, dst_name, dst_type, save_type, read_type, copy_type, \
                    conv_mode, IN0_SCALE, IN0_TAIL, IN1_SCALE, IN1_TAIL, OUT_SCALE, OUT_OFFSET) \
__kernel void floordiv_##src0_name##src1_name##to##dst_name \
    ( \
    image2d_array_t input0, \
    image2d_array_t input1, \
    image2d_array_t output \
    ) \
{ \
    int4 coord = (int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0); \
    FLOORDIV_PROCESS(dst_type, save_type, read_type, copy_type, conv_mode, IN0_SCALE, IN0_TAIL,\
                  IN1_SCALE, IN1_TAIL, OUT_SCALE, OUT_OFFSET, VXC_ReadImage2DArray, VXC_WriteImage2DArray); \
}


TENSOR_FLOORDIV(F16, F16, F16, half4, vxc_short8, vxc_short8,\
                vxc_half8, CONV, 1, 0, 1, 0, 1, 0)
TENSOR_FLOORDIV(F16, F16, I16, short4, vxc_short8, vxc_short8,\
               vxc_half8, CONV_SAT_RTE, 1, 0, 1, 0, out_scale, out_zp)
TENSOR_FLOORDIV(F16, F16, I8,  char4, vxc_char8, vxc_short8,\
                vxc_half8, CONV_SAT_RTE, 1, 0, 1, 0, out_scale, out_zp)
TENSOR_FLOORDIV(F16, F16, U8,  uchar4, vxc_uchar8, vxc_short8,\
               vxc_half8, CONV_SAT_RTE, 1, 0, 1, 0, out_scale, out_zp)

TENSOR_FLOORDIV(I16, I16, I16, short4, vxc_short8, vxc_short8,\
                vxc_short8, CONV_SAT_RTE, in_scale0, in0Tail, in_scale1, in1Tail, out_scale, out_zp)
TENSOR_FLOORDIV(I16, I16, F16, half4, vxc_short8, vxc_short8,\
                vxc_short8, CONV, in_scale0, in0Tail, in_scale1, in1Tail, 1, 0)

TENSOR_FLOORDIV(I8, I8, I8, char4, vxc_char8, vxc_char16,\
                vxc_char16, CONV_SAT_RTE, in_scale0, in0Tail, in_scale1, in1Tail, out_scale, out_zp)
TENSOR_FLOORDIV(I8, I8, F16, half4, vxc_short8, vxc_char16,\
                vxc_char16, CONV, in_scale0, in0Tail, in_scale1, in1Tail, 1, 0)

TENSOR_FLOORDIV(U8, U8, U8,  uchar4, vxc_uchar8, vxc_uchar16,\
                vxc_uchar16, CONV_SAT_RTE, in_scale0, in0Tail, in_scale1, in1Tail, out_scale, out_zp)
TENSOR_FLOORDIV(U8, U8, F16, half4, vxc_short8, vxc_uchar16,\
                vxc_uchar16, CONV, in_scale0, in0Tail, in_scale1, in1Tail, 1, 0)


#define TENSOR_FLOORDIV_2D(src0_name, src1_name, dst_name, dst_type, save_type, read_type, copy_type, \
    conv_mode, IN0_SCALE, IN0_TAIL, IN1_SCALE, IN1_TAIL, OUT_SCALE, OUT_OFFSET) \
__kernel void floordiv_##src0_name##src1_name##to##dst_name##_2D \
    ( \
    image2d_array_t input0, \
    image2d_array_t input1, \
    image2d_array_t output \
    ) \
{ \
    int2 coord = (int2)(get_global_id(0), get_global_id(1)); \
    FLOORDIV_PROCESS(dst_type, save_type, read_type, copy_type, conv_mode, IN0_SCALE, IN0_TAIL,\
                  IN1_SCALE, IN1_TAIL, OUT_SCALE, OUT_OFFSET, VXC_ReadImage, VXC_WriteImage); \
}


TENSOR_FLOORDIV_2D(F16, F16, F16, half4, vxc_short8, vxc_short8,\
                vxc_half8, CONV, 1, 0, 1, 0, 1, 0)
TENSOR_FLOORDIV_2D(F16, F16, I16, short4, vxc_short8, vxc_short8,\
               vxc_half8, CONV_SAT_RTE, 1, 0, 1, 0, out_scale, out_zp)
TENSOR_FLOORDIV_2D(F16, F16, I8,  char4, vxc_char8, vxc_short8,\
                vxc_half8, CONV_SAT_RTE, 1, 0, 1, 0, out_scale, out_zp)
TENSOR_FLOORDIV_2D(F16, F16, U8,  uchar4, vxc_uchar8, vxc_short8,\
               vxc_half8, CONV_SAT_RTE, 1, 0, 1, 0, out_scale, out_zp)

TENSOR_FLOORDIV_2D(I16, I16, I16, short4, vxc_short8, vxc_short8,\
                vxc_short8, CONV_SAT_RTE, in_scale0, in0Tail, in_scale1, in1Tail, out_scale, out_zp)
TENSOR_FLOORDIV_2D(I16, I16, F16, half4, vxc_short8, vxc_short8,\
                vxc_short8, CONV, in_scale0, in0Tail, in_scale1, in1Tail, 1, 0)

TENSOR_FLOORDIV_2D(I8, I8, I8, char4, vxc_char8, vxc_char16,\
                vxc_char16, CONV_SAT_RTE, in_scale0, in0Tail, in_scale1, in1Tail, out_scale, out_zp)
TENSOR_FLOORDIV_2D(I8, I8, F16, half4, vxc_short8, vxc_char16,\
                vxc_char16, CONV, in_scale0, in0Tail, in_scale1, in1Tail, 1, 0)

TENSOR_FLOORDIV_2D(U8, U8, U8,  uchar4, vxc_uchar8, vxc_uchar16,\
                vxc_uchar16, CONV_SAT_RTE, in_scale0, in0Tail, in_scale1, in1Tail, out_scale, out_zp)
TENSOR_FLOORDIV_2D(U8, U8, F16, half4, vxc_short8, vxc_uchar16,\
                vxc_uchar16, CONV, in_scale0, in0Tail, in_scale1, in1Tail, 1, 0)


_viv_uniform VXC_512Bits uniConvBF16toF32_Part0_2x8;
_viv_uniform VXC_512Bits uniConvBF16toF32_Part1_2x8;
_viv_uniform VXC_512Bits uniExtractOddData_2x8;

#define FLOORDIV_BF16_PROCESS(read_fun, write_fun) \
    vxc_short8 read_data0, read_data1, vec0; \
    vxc_float4 in0Val1, in0Val2, in1Val1, in1Val2; \
    vxc_float4 tmpVal1, tmpVal2; \
    vxc_ushort8 dst0, dst1; \
    vxc_ushort8 vect; \
    vxc_short8 zero = (vxc_short8)(0, 0, 0, 0, 0, 0, 0, 0); \
    read_fun(read_data0, input0, coord, VXC_5BITOFFSET_XY(0,0), \
    VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
    VXC_DP2x8(vec0, read_data0, zero, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniConvBF16toF32_Part0_2x8); \
    _viv_asm(COPY, in0Val1, vec0, 16); \
    VXC_DP2x8(vec0, read_data0, zero, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniConvBF16toF32_Part1_2x8); \
    _viv_asm(COPY, in0Val2, vec0, 16); \
    read_fun(read_data1, input1, coord, VXC_5BITOFFSET_XY(0,0), \
    VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
    VXC_DP2x8(vec0, read_data1, zero, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniConvBF16toF32_Part0_2x8); \
    _viv_asm(COPY, in1Val1, vec0, 16); \
    VXC_DP2x8(vec0, read_data1, zero, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniConvBF16toF32_Part1_2x8); \
    _viv_asm(COPY, in1Val2, vec0, 16); \
    tmpVal1 = floor(in0Val1 / in1Val1); \
    tmpVal2 = floor(in0Val2 / in1Val2); \
    _viv_asm(COPY, dst0, tmpVal1, 16); \
    _viv_asm(COPY, dst1, tmpVal2, 16); \
    VXC_DP2x8(vect, dst0, dst1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtractOddData_2x8); \
    write_fun(output, coord, vect, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));

__kernel void floordiv_BF16BF16toBF16
    (
    image2d_array_t input0,
    image2d_array_t input1,
    image2d_array_t output
    )
{
    int4 coord = (int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0);
    FLOORDIV_BF16_PROCESS(VXC_ReadImage2DArray, VXC_WriteImage2DArray);
}

__kernel void floordiv_BF16BF16toBF16_2D
    (
    image2d_array_t input0,
    image2d_array_t input1,
    image2d_array_t output
    )
{
    int2 coord = (int2)(get_global_id(0), get_global_id(1));
    FLOORDIV_BF16_PROCESS(VXC_ReadImage, VXC_WriteImage);
}