#include "cl_viv_vx_ext.h"

#define SIGNAL_FRAME_8BITS_SH_IMPL(type) \
__kernel void signal_frame_##type##to##type \
    ( \
    __read_only  image2d_array_t input, \
    __write_only image2d_array_t output, \
                 int             frame_step \
    ) \
{ \
    int inner = get_global_id(0); \
    int length_k = get_global_id(1); \
    int frames_id = get_global_id(2); \
 \
    int4 coord = (int4)(inner, length_k, frames_id, frames_id); \
    int2 coord_in = (int2)(inner, frames_id * frame_step + length_k); \
 \
    vxc_uchar16 src; \
    VXC_ReadImage(src, input, coord_in, 0, VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0)); \
    VXC_WriteImage2DArray(output, coord, src, VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0)); \
}
SIGNAL_FRAME_8BITS_SH_IMPL(U8)
SIGNAL_FRAME_8BITS_SH_IMPL(I8)

#define SIGNAL_FRAME_16BITS_SH_IMPL(type) \
__kernel void signal_frame_##type##to##type \
    ( \
    __read_only  image2d_array_t input, \
    __write_only image2d_array_t output, \
                 int             frame_step \
    ) \
{ \
    int inner = get_global_id(0); \
    int length_k = get_global_id(1); \
    int frames_id = get_global_id(2); \
 \
    int4 coord = (int4)(inner, length_k, frames_id, frames_id); \
    int2 coord_in = (int2)(inner, frames_id * frame_step + length_k); \
 \
    vxc_short8 src; \
    VXC_ReadImage(src, input, coord_in, 0, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
    VXC_WriteImage2DArray(output, coord, src, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); \
}
SIGNAL_FRAME_16BITS_SH_IMPL(I16)
SIGNAL_FRAME_16BITS_SH_IMPL(F16)
SIGNAL_FRAME_16BITS_SH_IMPL(BF16)