#include "cl_viv_vx_ext.h"

_viv_uniform VXC_512Bits uniExtract1to8Short_2x8;

#define REPEAT_RES(end_pos) \
coord.y = gidy; \
VXC_OP4_NoDest(img_store_3d, output, coord, src0, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src1, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src2, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src3, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src4, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src5, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src6, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0)); \
coord.y++; \
VXC_OP4_NoDest(img_store_3d, output, coord, src7, VXC_MODIFIER(0, end_pos, 0,VXC_RM_TowardZero, 0));

__kernel void repeat_I16_axis1(
    image2d_array_t  input0, image2d_t  input1, image2d_t  input2,
    image2d_array_t  output, int axis)
{
    int gidy = get_global_id(1);
    int4 coord = (int4)(get_global_id(0), gidy, get_global_id(2), get_global_id(2));
    vxc_short8 src0, src1, src2, src3, src4, src5, src6, src7;

    int8 input_desc, output_desc;
    _viv_asm(COPY, input_desc, input0, sizeof(input_desc));
    int baseAddr_a = (int)get_global_id(2) * input_desc.s4 + input_desc.s0;
    _viv_asm(MOV, coord.z, baseAddr_a);
    _viv_asm(COPY, output_desc, output, sizeof(output_desc));
    int baseAddr = (int)get_global_id(2) * output_desc.s4 + output_desc.s0;

    VXC_OP4(img_load_3d, src0, input0, coord, VXC_5BITOFFSET_XY(0, 0),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src1, input0, coord, VXC_5BITOFFSET_XY(0, 1),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src2, input0, coord, VXC_5BITOFFSET_XY(0, 2),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src3, input0, coord, VXC_5BITOFFSET_XY(0, 3),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src4, input0, coord, VXC_5BITOFFSET_XY(0, 4),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src5, input0, coord, VXC_5BITOFFSET_XY(0, 5),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src6, input0, coord, VXC_5BITOFFSET_XY(0, 6),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src7, input0, coord, VXC_5BITOFFSET_XY(0, 7),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));

    Image img1 = create_image_from_image2d(input1, 4);
    Image img2 = create_image_from_image2d(input2, 4);
    __global int* len_ptr = (__global int*)img1.ptr;
    __global int* index_ptr = (__global int*)img2.ptr;

    int len = len_ptr[get_global_id(0)];
    int start = index_ptr[get_global_id(0)];

    _viv_asm(MOV, coord.z, baseAddr);
    int iter = len >> 3;
    int res = len & 7;
    coord.x = start;

    VXC_DP2x8(src0, src0, src0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src1, src1, src1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src2, src2, src2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src3, src3, src3, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src4, src4, src4, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src5, src5, src5, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src6, src6, src6, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src7, src7, src7, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);

    for(int i = 0; i < iter; i++)
    {
        coord.y = gidy;
        VXC_OP4_NoDest(img_store_3d, output, coord, src0, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src1, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src2, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src3, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src4, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src5, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src6, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src7, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.x += 8;
    }

    if(res == 7)
    {
        REPEAT_RES(6)
    }
    else if(res == 6)
    {
        REPEAT_RES(5)
    }
    else if(res == 5)
    {
        REPEAT_RES(4)
    }
    else if(res == 4)
    {
        REPEAT_RES(3)
    }
    else if(res == 3)
    {
        REPEAT_RES(2)
    }
    else if(res == 2)
    {
        REPEAT_RES(1)
    }
    else if(res == 1)
    {
        REPEAT_RES(0)
    }
}

__kernel void repeat_U8_axis1(
    image2d_array_t  input0, image2d_t  input1, image2d_t  input2,
    image2d_array_t  output, int axis)
{
    int gidy = get_global_id(1);
    int4 coord = (int4)(get_global_id(0), gidy, get_global_id(2), get_global_id(2));
    vxc_uchar16 src0, src1, src2, src3, src4, src5, src6, src7;

    int8 input_desc, output_desc;
    _viv_asm(COPY, input_desc, input0, sizeof(input_desc));
    int baseAddr_a = (int)get_global_id(2) * input_desc.s4 + input_desc.s0;
    _viv_asm(MOV, coord.z, baseAddr_a);
    _viv_asm(COPY, output_desc, output, sizeof(output_desc));
    int baseAddr = (int)get_global_id(2) * output_desc.s4 + output_desc.s0;

    VXC_OP4(img_load_3d, src0, input0, coord, VXC_5BITOFFSET_XY(0, 0),
             VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src1, input0, coord, VXC_5BITOFFSET_XY(0, 1),
             VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src2, input0, coord, VXC_5BITOFFSET_XY(0, 2),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src3, input0, coord, VXC_5BITOFFSET_XY(0, 3),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src4, input0, coord, VXC_5BITOFFSET_XY(0, 4),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src5, input0, coord, VXC_5BITOFFSET_XY(0, 5),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src6, input0, coord, VXC_5BITOFFSET_XY(0, 6),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src7, input0, coord, VXC_5BITOFFSET_XY(0, 7),
            VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));

    Image img1 = create_image_from_image2d(input1, 4);
    Image img2 = create_image_from_image2d(input2, 4);
    __global int* len_ptr = (__global int*)img1.ptr;
    __global int* index_ptr = (__global int*)img2.ptr;

    int len = len_ptr[get_global_id(0)];
    int start = index_ptr[get_global_id(0)];

    _viv_asm(MOV, coord.z, baseAddr);
    int iter = len >> 3;
    int res = len & 7;
    coord.x = start;

    VXC_DP2x8(src0, src0, src0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src1, src1, src1, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src2, src2, src2, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src3, src3, src3, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src4, src4, src4, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src5, src5, src5, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src6, src6, src6, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);
    VXC_DP2x8(src7, src7, src7, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0), uniExtract1to8Short_2x8);

    for(int i = 0; i < iter; i++)
    {
        coord.y = gidy;
        VXC_OP4_NoDest(img_store_3d, output, coord, src0, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src1, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src2, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src3, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src4, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src5, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src6, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.y++;
        VXC_OP4_NoDest(img_store_3d, output, coord, src7, VXC_MODIFIER(0, 7, 0,VXC_RM_TowardZero, 0));
        coord.x += 8;
    }

    if(res == 7)
    {
        REPEAT_RES(6)
    }
    else if(res == 6)
    {
        REPEAT_RES(5)
    }
    else if(res == 5)
    {
        REPEAT_RES(4)
    }
    else if(res == 4)
    {
        REPEAT_RES(3)
    }
    else if(res == 3)
    {
        REPEAT_RES(2)
    }
    else if(res == 2)
    {
        REPEAT_RES(1)
    }
    else if(res == 1)
    {
        REPEAT_RES(0)
    }
}

