#pragma OPENCL EXTENSION cl_viv_vx_extension : enable
#include "cl_viv_vx_ext.h"

#define GET_MATRIX_VALUE \
    float4 matrix0; \
    float2 matrix1; \
    Image  img1 = create_image_from_image2d(matrix, 4); \
    __global float* matrix_ptr = (__global float*)img1.ptr; \
    matrix0 = vload4(0, matrix_ptr); \
    matrix1 = vload2(2, matrix_ptr); \

__kernel void custom_warp_affine_nearest_neighbor_U8toU8_rgb_2D_optional_input
(
    __read_only  image2d_array_t input,
    __read_only  image2d_t       matrix,
    __write_only image2d_array_t output,
                 float           _m0,
                 float           _m1,
                 float           _m2,
                 float           _m3,
                 float           _m4,
                 float           _m5
)
{
    int2   coord = (int2)(get_global_id(0) * 3, get_global_id(1));
    int4   coord_in = (int4)(get_global_id(0), get_global_id(1), get_global_id(0) + 1, get_global_id(1));

    float4 coord_f = convert_float4(coord_in);
    int2   coord_matrix = (int2)(0,0);
    GET_MATRIX_VALUE

    coord_f = coord_f.xxzz * matrix0.xyxy + coord_f.yyww * matrix0.zwzw + matrix1.xyxy;

    coord_in.x = floor(coord_f.x) * 3;
    coord_in.y = floor(coord_f.y);
    coord_in.z = floor(coord_f.z) * 3;
    coord_in.w = floor(coord_f.w);

    vxc_uchar16 dst;
    VXC_ReadImage(dst, input, coord_in.xy, 0, VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    coord_in.x = coord_in.x + 1;
    VXC_ReadImage(dst, input, coord_in.xy, 0, VXC_MODIFIER(1, 1, 0, VXC_RM_TowardZero, 0));
    coord_in.x = coord_in.x + 1;
    VXC_ReadImage(dst, input, coord_in.xy, 0, VXC_MODIFIER(2, 2, 0, VXC_RM_TowardZero, 0));

    VXC_ReadImage(dst, input, coord_in.zw, 0, VXC_MODIFIER(3, 3, 0, VXC_RM_TowardZero, 0));
    coord_in.z = coord_in.z + 1;
    VXC_ReadImage(dst, input, coord_in.zw, 0, VXC_MODIFIER(4, 4, 0, VXC_RM_TowardZero, 0));
    coord_in.z = coord_in.z + 1;
    VXC_ReadImage(dst, input, coord_in.zw, 0, VXC_MODIFIER(5, 5, 0, VXC_RM_TowardZero, 0));

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

__kernel void custom_warp_affine_bilinear_U8toU8_rgb_2D_optional_input
(
    __read_only  image2d_array_t input,
    __read_only  image2d_t       matrix,
    __write_only image2d_array_t output,
                 float           _m0,
                 float           _m1,
                 float           _m2,
                 float           _m3,
                 float           _m4,
                 float           _m5
)
{
    int2   coord = (int2)(get_global_id(0) * 3, get_global_id(1));
    int4   coord_in = (int4)(get_global_id(0), get_global_id(1), get_global_id(0) + 1, get_global_id(1));

    float4 coord_f = convert_float4(coord_in);
    int2   coord_matrix = (int2)(0,0);
    GET_MATRIX_VALUE

    coord_f = coord_f.xxzz * matrix0.xyxy + coord_f.yyww * matrix0.zwzw + matrix1.xyxy;

    coord_in.x = floor(coord_f.x) * 3;
    coord_in.y = floor(coord_f.y);
    coord_in.z = floor(coord_f.z) * 3;
    coord_in.w = floor(coord_f.w);

    vxc_uchar16 src0, src1, src_0, src_1, dst;
    VXC_ReadImage(src_0, input, coord_in.xy, VXC_5BITOFFSET_XY(0, 0), VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));
    VXC_ReadImage(src_1, input, coord_in.xy, VXC_5BITOFFSET_XY(0, 1), VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));

    src0.x = src_0.s0;
    src0.y = src_0.s3;
    src1.x = src_1.s0;
    src1.y = src_1.s3;

#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.xy, VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.y, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.x, VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s1;
    src0.y = src_0.s4;
    src1.x = src_1.s1;
    src1.y = src_1.s4;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.xy, VXC_MODIFIER(1, 1, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.y, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.x, VXC_MODIFIER(1, 1, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s2;
    src0.y = src_0.s5;
    src1.x = src_1.s2;
    src1.y = src_1.s5;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.xy, VXC_MODIFIER(2, 2, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.y, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.x, VXC_MODIFIER(2, 2, 0, VXC_RM_TowardZero, 0));
#endif

    VXC_ReadImage(src_0, input, coord_in.zw, VXC_5BITOFFSET_XY(0, 0), VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));
    VXC_ReadImage(src_1, input, coord_in.zw, VXC_5BITOFFSET_XY(0, 1), VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));

    src0.x = src_0.s0;
    src0.y = src_0.s3;
    src1.x = src_1.s0;
    src1.y = src_1.s3;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.zw, VXC_MODIFIER(3, 3, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.w, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.z, VXC_MODIFIER(3, 3, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s1;
    src0.y = src_0.s4;
    src1.x = src_1.s1;
    src1.y = src_1.s4;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.zw, VXC_MODIFIER(4, 4, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.w, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.z, VXC_MODIFIER(4, 4, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s2;
    src0.y = src_0.s5;
    src1.x = src_1.s2;
    src1.y = src_1.s5;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.zw, VXC_MODIFIER(5, 5, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.w, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.z, VXC_MODIFIER(5, 5, 0, VXC_RM_TowardZero, 0));
#endif

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

__kernel void custom_warp_affine_nearest_neighbor_U8toU8_rgb_optional_input
(
    __read_only  image2d_array_t input,
    __read_only  image2d_t       matrix,
    __write_only image2d_array_t output,
                 float           _m0,
                 float           _m1,
                 float           _m2,
                 float           _m3,
                 float           _m4,
                 float           _m5
)
{
    int4   coord = (int4)(get_global_id(0) * 3, get_global_id(1), get_global_id(2), get_global_id(2));
    int4   coord_in = (int4)(get_global_id(0), get_global_id(1), get_global_id(0) + 1, get_global_id(1));

    float4 coord_f = convert_float4(coord_in);
    int2   coord_matrix = (int2)(0,0);
    GET_MATRIX_VALUE

    coord_f = coord_f.xxzz * matrix0.xyxy + coord_f.yyww * matrix0.zwzw + matrix1.xyxy;

    coord_in.x = floor(coord_f.x) * 3;
    coord_in.y = floor(coord_f.y);
    coord_in.z = floor(coord_f.z) * 3;
    coord_in.w = floor(coord_f.w);

    int4 coord_input = (int4)(coord_in.xy, get_global_id(2), get_global_id(2));
    int8 input_desc;
    _viv_asm(COPY, input_desc, input, sizeof(input_desc));
    int baseAddr = (int)coord_input.z * input_desc.s4 + input_desc.s0;
    _viv_asm(MOV, coord_input.w, baseAddr);

    vxc_uchar16 dst;
    VXC_OP4(img_load_3d, dst, input, coord_input.xywz, 0, VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
    coord_input.x = coord_input.x + 1;
    VXC_OP4(img_load_3d, dst, input, coord_input.xywz, 0, VXC_MODIFIER(1, 1, 0, VXC_RM_TowardZero, 0));
    coord_input.x = coord_input.x + 1;
    VXC_OP4(img_load_3d, dst, input, coord_input.xywz, 0, VXC_MODIFIER(2, 2, 0, VXC_RM_TowardZero, 0));
    coord_input.xy = coord_in.zw;
    VXC_OP4(img_load_3d, dst, input, coord_input.xywz, 0, VXC_MODIFIER(3, 3, 0, VXC_RM_TowardZero, 0));
    coord_input.x = coord_input.x + 1;
    VXC_OP4(img_load_3d, dst, input, coord_input.xywz, 0, VXC_MODIFIER(4, 4, 0, VXC_RM_TowardZero, 0));
    coord_input.x = coord_input.x + 1;
    VXC_OP4(img_load_3d, dst, input, coord_input.xywz, 0, VXC_MODIFIER(5, 5, 0, VXC_RM_TowardZero, 0));

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

__kernel void custom_warp_affine_bilinear_U8toU8_rgb_optional_input
(
    __read_only  image2d_array_t input,
    __read_only  image2d_t       matrix,
    __write_only image2d_array_t output,
                 float           _m0,
                 float           _m1,
                 float           _m2,
                 float           _m3,
                 float           _m4,
                 float           _m5
)
{
    int4   coord = (int4)(get_global_id(0) * 3, get_global_id(1), get_global_id(2), get_global_id(2));
    int4   coord_in = (int4)(get_global_id(0), get_global_id(1), get_global_id(0) + 1, get_global_id(1));

    float4 coord_f = convert_float4(coord_in);
    int2   coord_matrix = (int2)(0,0);
    GET_MATRIX_VALUE

    coord_f = coord_f.xxzz * matrix0.xyxy + coord_f.yyww * matrix0.zwzw + matrix1.xyxy;

    coord_in.x = floor(coord_f.x) * 3;
    coord_in.y = floor(coord_f.y);
    coord_in.z = floor(coord_f.z) * 3;
    coord_in.w = floor(coord_f.w);

    int4 coord_input = (int4)(coord_in.xy, get_global_id(2), get_global_id(2));
    int8 input_desc;
    _viv_asm(COPY, input_desc, input, sizeof(input_desc));
    int baseAddr = (int)coord_input.z * input_desc.s4 + input_desc.s0;
    _viv_asm(MOV, coord_input.w, baseAddr);

    vxc_uchar16 src0, src1, src_0, src_1, dst;
    VXC_OP4(img_load_3d, src_0, input, coord_input.xywz, VXC_5BITOFFSET_XY(0, 0),
        VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src_1, input, coord_input.xywz, VXC_5BITOFFSET_XY(0, 1),
        VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));

    src0.x = src_0.s0;
    src0.y = src_0.s3;
    src1.x = src_1.s0;
    src1.y = src_1.s3;

#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.xy, VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.y, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.x, VXC_MODIFIER(0, 0, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s1;
    src0.y = src_0.s4;
    src1.x = src_1.s1;
    src1.y = src_1.s4;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.xy, VXC_MODIFIER(1, 1, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.y, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.x, VXC_MODIFIER(1, 1, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s2;
    src0.y = src_0.s5;
    src1.x = src_1.s2;
    src1.y = src_1.s5;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.xy, VXC_MODIFIER(2, 2, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.y, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.x, VXC_MODIFIER(2, 2, 0, VXC_RM_TowardZero, 0));
#endif

    coord_input.xy = coord_in.zw;
    VXC_OP4(img_load_3d, src_0, input, coord_input.xywz, VXC_5BITOFFSET_XY(0, 0),
        VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));
    VXC_OP4(img_load_3d, src_1, input, coord_input.xywz, VXC_5BITOFFSET_XY(0, 1),
        VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));

    src0.x = src_0.s0;
    src0.y = src_0.s3;
    src1.x = src_1.s0;
    src1.y = src_1.s3;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.zw, VXC_MODIFIER(3, 3, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.w, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.z, VXC_MODIFIER(3, 3, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s1;
    src0.y = src_0.s4;
    src1.x = src_1.s1;
    src1.y = src_1.s4;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.zw, VXC_MODIFIER(4, 4, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.w, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.z, VXC_MODIFIER(4, 4, 0, VXC_RM_TowardZero, 0));
#endif

    src0.x = src_0.s2;
    src0.y = src_0.s5;
    src1.x = src_1.s2;
    src1.y = src_1.s5;
#if (VX_VERSION==1)
    VXC_BiLinear(dst, src0, src1, coord_f.zw, VXC_MODIFIER(5, 5, 0, VXC_RM_TowardZero, 0));
#else
    VXC_Lerp(src0, src0, src1, coord_f.w, VXC_MODIFIER(0, 1, 0, VXC_RM_TowardZero, 0));
    src1.s0 = src0.s1;
    VXC_Lerp(dst, src0, src1, coord_f.z, VXC_MODIFIER(5, 5, 0, VXC_RM_TowardZero, 0));
#endif

    VXC_WriteImage2DArray(output, coord, dst, VXC_MODIFIER(0, 5, 0, VXC_RM_TowardZero, 0));
}