/*
 * Copyright (c) 2016-2018 ARM Limited.
 *
 * SPDX-License-Identifier: MIT
 *
 * Permission is hereby granted, free of charge, to any person obtaining a copy
 * of this software and associated documentation files (the "Software"), to
 * deal in the Software without restriction, including without limitation the
 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
 * sell copies of the Software, and to permit persons to whom the Software is
 * furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice shall be included in all
 * copies or substantial portions of the Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 * SOFTWARE.
 */
#include "helpers.h"

/** This function combines three planes to a single RGB image.
 *
 * @param[in] plane0_ptr                           Pointer to the first plane. Supported Format: U8
 * @param[in] plane0_stride_x                      Stride of the first plane in X dimension (in bytes)
 * @param[in] plane0_step_x                        plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane0_stride_y                      Stride of the first plane in Y dimension (in bytes)
 * @param[in] plane0_step_y                        plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane
 * @param[in] plane1_ptr                           Pointer to the second plane. Supported Format: U8
 * @param[in] plane1_stride_x                      Stride of the second plane in X dimension (in bytes)
 * @param[in] plane1_step_x                        plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane1_stride_y                      Stride of the second plane in Y dimension (in bytes)
 * @param[in] plane1_step_y                        plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane
 * @param[in] plane2_ptr                           Pointer to the third plane. Supported Format: U8
 * @param[in] plane2_stride_x                      Stride of the third plane in X dimension (in bytes)
 * @param[in] plane2_step_x                        plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane2_stride_y                      Stride of the third plane in Y dimension (in bytes)
 * @param[in] plane2_step_y                        plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane
 * @param[in] dst_ptr                              Pointer to the destination image. Supported Format: RGB
 * @param[in] dst_stride_x                         Stride of the destination image in X dimension (in bytes)
 * @param[in] dst_step_x                           dst_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_stride_y                         Stride of the destination image in Y dimension (in bytes)
 * @param[in] dst_step_y                           dst_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_offset_first_element_in_bytes    The offset of the first element in the destination image
 */
__kernel void channel_combine_RGB888(
    IMAGE_DECLARATION(plane0),
    IMAGE_DECLARATION(plane1),
    IMAGE_DECLARATION(plane2),
    IMAGE_DECLARATION(dst))
{
    // Get pixels pointer
    Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);
    Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);
    Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);
    Image dst    = CONVERT_TO_IMAGE_STRUCT(dst);

    uchar16 data0 = vload16(0, plane0.ptr);
    uchar16 data1 = vload16(0, plane1.ptr);
    uchar16 data2 = vload16(0, plane2.ptr);

    uchar16 out0 = (uchar16)(data0.s0, data1.s0, data2.s0,
                             data0.s1, data1.s1, data2.s1,
                             data0.s2, data1.s2, data2.s2,
                             data0.s3, data1.s3, data2.s3,
                             data0.s4, data1.s4, data2.s4,
                             data0.s5);
    vstore16(out0, 0, dst.ptr);

    uchar16 out1 = (uchar16)(data1.s5, data2.s5, data0.s6,
                             data1.s6, data2.s6, data0.s7,
                             data1.s7, data2.s7, data0.s8,
                             data1.s8, data2.s8, data0.s9,
                             data1.s9, data2.s9, data0.sA,
                             data1.sA);
    vstore16(out1, 0, dst.ptr + 16);

    uchar16 out2 = (uchar16)(data2.sA, data0.sB, data1.sB,
                             data2.sB, data0.sC, data1.sC,
                             data2.sC, data0.sD, data1.sD,
                             data2.sD, data0.sE, data1.sE,
                             data2.sE, data0.sF, data1.sF,
                             data2.sF);
    vstore16(out2, 0, dst.ptr + 32);
}

/** This function combines three planes to a single RGBA image.
 *
 * @param[in] plane0_ptr                           Pointer to the first plane. Supported Format: U8
 * @param[in] plane0_stride_x                      Stride of the first plane in X dimension (in bytes)
 * @param[in] plane0_step_x                        plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane0_stride_y                      Stride of the first plane in Y dimension (in bytes)
 * @param[in] plane0_step_y                        plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane
 * @param[in] plane1_ptr                           Pointer to the second plane. Supported Format: U8
 * @param[in] plane1_stride_x                      Stride of the second plane in X dimension (in bytes)
 * @param[in] plane1_step_x                        plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane1_stride_y                      Stride of the second plane in Y dimension (in bytes)
 * @param[in] plane1_step_y                        plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane
 * @param[in] plane2_ptr                           Pointer to the third plane. Supported Format: U8
 * @param[in] plane2_stride_x                      Stride of the third plane in X dimension (in bytes)
 * @param[in] plane2_step_x                        plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane2_stride_y                      Stride of the third plane in Y dimension (in bytes)
 * @param[in] plane2_step_y                        plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane
 * @param[in] plane3_ptr                           Pointer to the fourth plane. Supported Format: U8
 * @param[in] plane3_stride_x                      Stride of the fourth plane in X dimension (in bytes)
 * @param[in] plane3_step_x                        plane3_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane3_stride_y                      Stride of the fourth plane in Y dimension (in bytes)
 * @param[in] plane3_step_y                        plane3_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane3_offset_first_element_in_bytes The offset of the first element in the fourth plane
 * @param[in] dst_ptr                              Pointer to the destination image. Supported Format: RGBA
 * @param[in] dst_stride_x                         Stride of the destination image in X dimension (in bytes)
 * @param[in] dst_step_x                           dst_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_stride_y                         Stride of the destination image in Y dimension (in bytes)
 * @param[in] dst_step_y                           dst_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_offset_first_element_in_bytes    The offset of the first element in the destination image
 */
__kernel void channel_combine_RGBA8888(
    IMAGE_DECLARATION(plane0),
    IMAGE_DECLARATION(plane1),
    IMAGE_DECLARATION(plane2),
    IMAGE_DECLARATION(plane3),
    IMAGE_DECLARATION(dst))
{
    // Get pixels pointer
    Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);
    Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);
    Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);
    Image plane3 = CONVERT_TO_IMAGE_STRUCT(plane3);
    Image dst    = CONVERT_TO_IMAGE_STRUCT(dst);

    uchar16 data0 = vload16(0, plane0.ptr);
    uchar16 data1 = vload16(0, plane1.ptr);
    uchar16 data2 = vload16(0, plane2.ptr);
    uchar16 data3 = vload16(0, plane3.ptr);

    uchar16 out0 = (uchar16)(data0.s0, data1.s0, data2.s0, data3.s0,
                             data0.s1, data1.s1, data2.s1, data3.s1,
                             data0.s2, data1.s2, data2.s2, data3.s2,
                             data0.s3, data1.s3, data2.s3, data3.s3);
    vstore16(out0, 0, dst.ptr);

    uchar16 out1 = (uchar16)(data0.s4, data1.s4, data2.s4, data3.s4,
                             data0.s5, data1.s5, data2.s5, data3.s5,
                             data0.s6, data1.s6, data2.s6, data3.s6,
                             data0.s7, data1.s7, data2.s7, data3.s7);
    vstore16(out1, 0, dst.ptr + 16);

    uchar16 out2 = (uchar16)(data0.s8, data1.s8, data2.s8, data3.s8,
                             data0.s9, data1.s9, data2.s9, data3.s9,
                             data0.sA, data1.sA, data2.sA, data3.sA,
                             data0.sB, data1.sB, data2.sB, data3.sB);
    vstore16(out2, 0, dst.ptr + 32);

    uchar16 out3 = (uchar16)(data0.sC, data1.sC, data2.sC, data3.sC,
                             data0.sD, data1.sD, data2.sD, data3.sD,
                             data0.sE, data1.sE, data2.sE, data3.sE,
                             data0.sF, data1.sF, data2.sF, data3.sF);
    vstore16(out3, 0, dst.ptr + 48);
}

/** This function combines three planes to a single YUYV image.
 *
 * @param[in] plane0_ptr                           Pointer to the first plane. Supported Format: U8
 * @param[in] plane0_stride_x                      Stride of the first plane in X dimension (in bytes)
 * @param[in] plane0_step_x                        plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane0_stride_y                      Stride of the first plane in Y dimension (in bytes)
 * @param[in] plane0_step_y                        plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane
 * @param[in] plane1_ptr                           Pointer to the second plane. Supported Format: U8
 * @param[in] plane1_stride_x                      Stride of the second plane in X dimension (in bytes)
 * @param[in] plane1_step_x                        plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane1_stride_y                      Stride of the second plane in Y dimension (in bytes)
 * @param[in] plane1_step_y                        plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane
 * @param[in] plane2_ptr                           Pointer to the third plane. Supported Format: U8
 * @param[in] plane2_stride_x                      Stride of the third plane in X dimension (in bytes)
 * @param[in] plane2_step_x                        plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane2_stride_y                      Stride of the third plane in Y dimension (in bytes)
 * @param[in] plane2_step_y                        plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane
 * @param[in] dst_ptr                              Pointer to the destination image. Supported Format: YUYV
 * @param[in] dst_stride_x                         Stride of the destination image in X dimension (in bytes)
 * @param[in] dst_step_x                           dst_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_stride_y                         Stride of the destination image in Y dimension (in bytes)
 * @param[in] dst_step_y                           dst_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_offset_first_element_in_bytes    The offset of the first element in the destination image
 */
__kernel void channel_combine_YUYV422(
    IMAGE_DECLARATION(plane0),
    IMAGE_DECLARATION(plane1),
    IMAGE_DECLARATION(plane2),
    IMAGE_DECLARATION(dst))
{
    // Get pixels pointer
    Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);
    Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);
    Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);
    Image dst    = CONVERT_TO_IMAGE_STRUCT(dst);

    uchar16 data0 = vload16(0, plane0.ptr);
    uchar8  data1 = vload8(0, plane1.ptr);
    uchar8  data2 = vload8(0, plane2.ptr);

    uchar16 out0 = (uchar16)(data0.s0, data1.s0, data0.s1, data2.s0,
                             data0.s2, data1.s1, data0.s3, data2.s1,
                             data0.s4, data1.s2, data0.s5, data2.s2,
                             data0.s6, data1.s3, data0.s7, data2.s3);
    vstore16(out0, 0, dst.ptr);
    uchar16 out1 = (uchar16)(data0.s8, data1.s4, data0.s9, data2.s4,
                             data0.sA, data1.s5, data0.sB, data2.s5,
                             data0.sC, data1.s6, data0.sD, data2.s6,
                             data0.sE, data1.s7, data0.sF, data2.s7);
    vstore16(out1, 0, dst.ptr + 16);
}

/** This function combines three planes to a single UYUV image.
 *
 * @param[in] plane0_ptr                           Pointer to the first plane. Supported Format: U8
 * @param[in] plane0_stride_x                      Stride of the first plane in X dimension (in bytes)
 * @param[in] plane0_step_x                        plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane0_stride_y                      Stride of the first plane in Y dimension (in bytes)
 * @param[in] plane0_step_y                        plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane
 * @param[in] plane1_ptr                           Pointer to the second plane. Supported Format: U8
 * @param[in] plane1_stride_x                      Stride of the second plane in X dimension (in bytes)
 * @param[in] plane1_step_x                        plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane1_stride_y                      Stride of the second plane in Y dimension (in bytes)
 * @param[in] plane1_step_y                        plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane
 * @param[in] plane2_ptr                           Pointer to the third plane. Supported Format: U8
 * @param[in] plane2_stride_x                      Stride of the third plane in X dimension (in bytes)
 * @param[in] plane2_step_x                        plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] plane2_stride_y                      Stride of the third plane in Y dimension (in bytes)
 * @param[in] plane2_step_y                        plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane
 * @param[in] dst_ptr                              Pointer to the destination image. Supported Format: UYUV
 * @param[in] dst_stride_x                         Stride of the destination image in X dimension (in bytes)
 * @param[in] dst_step_x                           dst_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_stride_y                         Stride of the destination image in Y dimension (in bytes)
 * @param[in] dst_step_y                           dst_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_offset_first_element_in_bytes    The offset of the first element in the destination image
 */
__kernel void channel_combine_UYVY422(
    IMAGE_DECLARATION(plane0),
    IMAGE_DECLARATION(plane1),
    IMAGE_DECLARATION(plane2),
    IMAGE_DECLARATION(dst))
{
    // Get pixels pointer
    Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);
    Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);
    Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);
    Image dst    = CONVERT_TO_IMAGE_STRUCT(dst);

    uchar16 data0 = vload16(0, plane0.ptr);
    uchar8  data1 = vload8(0, plane1.ptr);
    uchar8  data2 = vload8(0, plane2.ptr);

    uchar16 out0 = (uchar16)(data1.s0, data0.s0, data2.s0, data0.s1,
                             data1.s1, data0.s2, data2.s1, data0.s3,
                             data1.s2, data0.s4, data2.s2, data0.s5,
                             data1.s3, data0.s6, data2.s3, data0.s7);
    vstore16(out0, 0, dst.ptr);
    uchar16 out1 = (uchar16)(data1.s4, data0.s8, data2.s4, data0.s9,
                             data1.s5, data0.sA, data2.s5, data0.sB,
                             data1.s6, data0.sC, data2.s6, data0.sD,
                             data1.s7, data0.sE, data2.s7, data0.sF);
    vstore16(out1, 0, dst.ptr + 16);
}

/** This function combines three planes to a single NV12/NV21 image.
 *
 * @note NV12 or NV21 has to be specified through preprocessor macro. eg. -DNV12 performs NV12 channel combine.
 *
 * @param[in] src_plane0_ptr                           Pointer to the first plane. Supported Format: U8
 * @param[in] src_plane0_stride_x                      Stride of the first plane in X dimension (in bytes)
 * @param[in] src_plane0_step_x                        src_plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] src_plane0_stride_y                      Stride of the first plane in Y dimension (in bytes)
 * @param[in] src_plane0_step_y                        src_plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] src_plane0_offset_first_element_in_bytes The offset of the first element in the first plane
 * @param[in] src_plane1_ptr                           Pointer to the second plane. Supported Format: U8
 * @param[in] src_plane1_stride_x                      Stride of the second plane in X dimension (in bytes)
 * @param[in] src_plane1_step_x                        src_plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] src_plane1_stride_y                      Stride of the second plane in Y dimension (in bytes)
 * @param[in] src_plane1_step_y                        src_plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] src_plane1_offset_first_element_in_bytes The offset of the first element in the second plane
 * @param[in] src_plane2_ptr                           Pointer to the third plane. Supported Format: U8
 * @param[in] src_plane2_stride_x                      Stride of the third plane in X dimension (in bytes)
 * @param[in] src_plane2_step_x                        src_plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] src_plane2_stride_y                      Stride of the third plane in Y dimension (in bytes)
 * @param[in] src_plane2_step_y                        src_plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] src_plane2_offset_first_element_in_bytes The offset of the first element in the third plane
 * @param[in] dst_plane0_ptr                           Pointer to the first plane of the destination image. Supported Format: U8
 * @param[in] dst_plane0_stride_x                      Stride of the first plane of the destination image in X dimension (in bytes)
 * @param[in] dst_plane0_step_x                        dst_plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_plane0_stride_y                      Stride of the first plane of the destination image in Y dimension (in bytes)
 * @param[in] dst_plane0_step_y                        dst_plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_plane0_offset_first_element_in_bytes The offset of the first element in the first plane of the destination image
 * @param[in] dst_plane1_ptr                           Pointer to the second plane of the destination image. Supported Format: UV88
 * @param[in] dst_plane1_stride_x                      Stride of the second plane of the destination image in X dimension (in bytes)
 * @param[in] dst_plane1_step_x                        dst_plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_plane1_stride_y                      Stride of the second plane of the destination image in Y dimension (in bytes)
 * @param[in] dst_plane1_step_y                        dst_plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_plane1_offset_first_element_in_bytes The offset of the first element in the second plane of the destination image
 * @param[in] height                                   Sub-sampled height
 */
__kernel void channel_combine_NV(
    IMAGE_DECLARATION(src_plane0),
    IMAGE_DECLARATION(src_plane1),
    IMAGE_DECLARATION(src_plane2),
    IMAGE_DECLARATION(dst_plane0),
    IMAGE_DECLARATION(dst_plane1),
    uint height)
{
    // Get pixels pointer
    Image src_plane0 = CONVERT_TO_IMAGE_STRUCT(src_plane0);
    Image src_plane1 = CONVERT_TO_IMAGE_STRUCT(src_plane1);
    Image src_plane2 = CONVERT_TO_IMAGE_STRUCT(src_plane2);
    Image dst_plane0 = CONVERT_TO_IMAGE_STRUCT(dst_plane0);
    Image dst_plane1 = CONVERT_TO_IMAGE_STRUCT(dst_plane1);

    // Copy plane data
    vstore16(vload16(0, src_plane0.ptr), 0, dst_plane0.ptr);
    vstore16(vload16(0, offset(&src_plane0, 0, height)), 0, (__global uchar *)offset(&dst_plane0, 0, height));

    // Create UV place
    uchar8 data1 = vload8(0, src_plane1.ptr);
    uchar8 data2 = vload8(0, src_plane2.ptr);

#ifdef NV12
    vstore16(shuffle2(data1, data2, (uchar16)(0, 8, 1, 9, 2, 10, 3, 11, 4, 12, 5, 13, 6, 14, 7, 15)), 0, dst_plane1.ptr);
#elif defined(NV21)
    vstore16(shuffle2(data2, data1, (uchar16)(0, 8, 1, 9, 2, 10, 3, 11, 4, 12, 5, 13, 6, 14, 7, 15)), 0, dst_plane1.ptr);
#endif /* NV12 or NV21 */
}

/** This function combines three planes to a single YUV444 or IYUV image.
 *
 * @note YUV444 or IYUV has to be specified through preprocessor macro. eg. -DIYUV performs IYUV channel combine.
 *
 * @param[in] src_plane0_ptr                           Pointer to the first plane. Supported Format: U8
 * @param[in] src_plane0_stride_x                      Stride of the first plane in X dimension (in bytes)
 * @param[in] src_plane0_step_x                        src_plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] src_plane0_stride_y                      Stride of the first plane in Y dimension (in bytes)
 * @param[in] src_plane0_step_y                        src_plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] src_plane0_offset_first_element_in_bytes The offset of the first element in the first plane
 * @param[in] src_plane1_ptr                           Pointer to the second plane. Supported Format: U8
 * @param[in] src_plane1_stride_x                      Stride of the second plane in X dimension (in bytes)
 * @param[in] src_plane1_step_x                        src_plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] src_plane1_stride_y                      Stride of the second plane in Y dimension (in bytes)
 * @param[in] src_plane1_step_y                        src_plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] src_plane1_offset_first_element_in_bytes The offset of the first element in the second plane
 * @param[in] src_plane2_ptr                           Pointer to the third plane. Supported Format: U8
 * @param[in] src_plane2_stride_x                      Stride of the third plane in X dimension (in bytes)
 * @param[in] src_plane2_step_x                        src_plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] src_plane2_stride_y                      Stride of the third plane in Y dimension (in bytes)
 * @param[in] src_plane2_step_y                        src_plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] src_plane2_offset_first_element_in_bytes The offset of the first element in the third plane
 * @param[in] dst_plane0_ptr                           Pointer to the first plane of the destination image. Supported Format: U8
 * @param[in] dst_plane0_stride_x                      Stride of the first plane of the destination image in X dimension (in bytes)
 * @param[in] dst_plane0_step_x                        dst_plane0_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_plane0_stride_y                      Stride of the first plane of the destination image in Y dimension (in bytes)
 * @param[in] dst_plane0_step_y                        dst_plane0_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_plane0_offset_first_element_in_bytes The offset of the first element in the first plane of the destination image
 * @param[in] dst_plane1_ptr                           Pointer to the second plane of the destination image. Supported Format: U8
 * @param[in] dst_plane1_stride_x                      Stride of the second plane of the destination image in X dimension (in bytes)
 * @param[in] dst_plane1_step_x                        dst_plane1_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_plane1_stride_y                      Stride of the second plane of the destination image in Y dimension (in bytes)
 * @param[in] dst_plane1_step_y                        dst_plane1_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_plane1_offset_first_element_in_bytes The offset of the first element in the second plane of the destination image
 * @param[in] dst_plane2_ptr                           Pointer to the third plane of the destination image. Supported Format: U8
 * @param[in] dst_plane2_stride_x                      Stride of the third plane of the destination image in X dimension (in bytes)
 * @param[in] dst_plane2_step_x                        dst_plane2_stride_x * number of elements along X processed per workitem(in bytes)
 * @param[in] dst_plane2_stride_y                      Stride of the third plane of the destination image in Y dimension (in bytes)
 * @param[in] dst_plane2_step_y                        dst_plane2_stride_y * number of elements along Y processed per workitem(in bytes)
 * @param[in] dst_plane2_offset_first_element_in_bytes The offset of the first element in the third plane of the destination image
 * @param[in] height                                   Sub-sampled height
 */
__kernel void copy_planes_3p(
    IMAGE_DECLARATION(src_plane0),
    IMAGE_DECLARATION(src_plane1),
    IMAGE_DECLARATION(src_plane2),
    IMAGE_DECLARATION(dst_plane0),
    IMAGE_DECLARATION(dst_plane1),
    IMAGE_DECLARATION(dst_plane2),
    uint height)
{
    // Get pixels pointer
    Image src_plane0 = CONVERT_TO_IMAGE_STRUCT(src_plane0);
    Image src_plane1 = CONVERT_TO_IMAGE_STRUCT(src_plane1);
    Image src_plane2 = CONVERT_TO_IMAGE_STRUCT(src_plane2);
    Image dst_plane0 = CONVERT_TO_IMAGE_STRUCT(dst_plane0);
    Image dst_plane1 = CONVERT_TO_IMAGE_STRUCT(dst_plane1);
    Image dst_plane2 = CONVERT_TO_IMAGE_STRUCT(dst_plane2);

    // Copy plane data
    vstore16(vload16(0, src_plane0.ptr), 0, dst_plane0.ptr);
#ifdef YUV444
    vstore16(vload16(0, src_plane1.ptr), 0, dst_plane1.ptr);
    vstore16(vload16(0, src_plane2.ptr), 0, dst_plane2.ptr);
#elif defined(IYUV)
    vstore16(vload16(0, offset(&src_plane0, 0, height)), 0, (__global uchar *)offset(&dst_plane0, 0, height));
    vstore8(vload8(0, src_plane1.ptr), 0, dst_plane1.ptr);
    vstore8(vload8(0, src_plane2.ptr), 0, dst_plane2.ptr);
#endif /* YUV444 or IYUV */
}