#include "cl_viv_vx_ext.h"

_viv_uniform VXC_512Bits uniA_Minus_B_2x8;
_viv_uniform VXC_512Bits uniA_Times_B_2x8;
_viv_uniform VXC_512Bits uniA_Plus_B_2x8;
__kernel void grucell_activation_sma_F16_F16_F16toF16
    (
    __read_only image2d_array_t   input0,
    __read_only image2d_array_t   input1,
    __read_only image2d_array_t   input2,
    __write_only image2d_array_t  output,
    __write_only image2d_array_t  h_status
    )
{
    int4 coord = (int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0);

    vxc_half8   src0, src1, src2, minus, dst;
    vxc_ushort8 vec0, vec1, vec2;

    VXC_ReadImage2DArray(vec0, input0, coord, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    _viv_asm(COPY, src0, vec0, 16);
    VXC_ReadImage2DArray(vec1, input1, coord, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    _viv_asm(COPY, src1, vec1, 16);
    VXC_ReadImage2DArray(vec2, input2, coord, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    _viv_asm(COPY, src2, vec2, 16);

    VXC_DP2x8(minus, src0, src1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), uniA_Minus_B_2x8);
    VXC_DP2x8(dst, minus, src2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), uniA_Times_B_2x8);
    VXC_DP2x8(dst, dst, src1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), uniA_Plus_B_2x8);

    VXC_WriteImage2DArray(output, coord, dst, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
    VXC_WriteImage(h_status, coord, dst, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
}

__kernel void grucell_activation_sma_F16_F16_F16toF16_2D
    (
    __read_only image2d_array_t   input0,
    __read_only image2d_array_t   input1,
    __read_only image2d_array_t   input2,
    __write_only image2d_array_t  output,
    __write_only image2d_array_t  h_status
    )
{
    int2 coord = (int2)(get_global_id(0), get_global_id(1));

    vxc_half8   src0, src1, src2, minus, dst;
    vxc_ushort8 vec0, vec1, vec2;

    VXC_ReadImage(vec0, input0, coord, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    _viv_asm(COPY, src0, vec0, 16);
    VXC_ReadImage(vec1, input1, coord, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    _viv_asm(COPY, src1, vec1, 16);
    VXC_ReadImage(vec2, input2, coord, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0));
    _viv_asm(COPY, src2, vec2, 16);

    VXC_DP2x8(minus, src0, src1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), uniA_Minus_B_2x8);
    VXC_DP2x8(dst, minus, src2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), uniA_Times_B_2x8);
    VXC_DP2x8(dst, dst, src1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), uniA_Plus_B_2x8);

    VXC_WriteImage(output, coord, dst, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
    VXC_WriteImage(h_status, coord, dst, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
}

